diff --git a/modules/imgproc/src/deriv.cpp b/modules/imgproc/src/deriv.cpp index d7d3e9560d..88fa5733ed 100644 --- a/modules/imgproc/src/deriv.cpp +++ b/modules/imgproc/src/deriv.cpp @@ -517,6 +517,66 @@ static bool ipp_sobel(InputArray _src, OutputArray _dst, int ddepth, int dx, int } #endif +#ifdef HAVE_OPENCL +namespace cv +{ +static bool ocl_sepFilter3x3_8UC1(InputArray _src, OutputArray _dst, int ddepth, + InputArray _kernelX, InputArray _kernelY, double delta, int borderType) +{ + const ocl::Device & dev = ocl::Device::getDefault(); + int type = _src.type(), sdepth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); + + if ( !(dev.isIntel() && (type == CV_8UC1) && (ddepth == CV_8U) && + (_src.offset() == 0) && (_src.step() % 4 == 0) && + (_src.cols() % 16 == 0) && (_src.rows() % 2 == 0)) ) + return false; + + Mat kernelX = _kernelX.getMat().reshape(1, 1); + if (kernelX.cols % 2 != 1) + return false; + Mat kernelY = _kernelY.getMat().reshape(1, 1); + if (kernelY.cols % 2 != 1) + return false; + + if (ddepth < 0) + ddepth = sdepth; + + Size size = _src.size(); + size_t globalsize[2] = { 0, 0 }; + size_t localsize[2] = { 0, 0 }; + + globalsize[0] = size.width / 16; + globalsize[1] = size.height / 2; + + const char * const borderMap[] = { "BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", 0, "BORDER_REFLECT_101" }; + char build_opts[1024]; + sprintf(build_opts, "-D %s %s%s", borderMap[borderType], + ocl::kernelToStr(kernelX, CV_32F, "KERNEL_MATRIX_X").c_str(), + ocl::kernelToStr(kernelY, CV_32F, "KERNEL_MATRIX_Y").c_str()); + + ocl::Kernel kernel("sepFilter3x3_8UC1_cols16_rows2", cv::ocl::imgproc::sepFilter3x3_oclsrc, build_opts); + if (kernel.empty()) + return false; + + UMat src = _src.getUMat(); + _dst.create(size, CV_MAKETYPE(ddepth, cn)); + if (!(_dst.offset() == 0 && _dst.step() % 4 == 0)) + return false; + UMat dst = _dst.getUMat(); + + int idxArg = kernel.set(0, ocl::KernelArg::PtrReadOnly(src)); + idxArg = kernel.set(idxArg, (int)src.step); + idxArg = kernel.set(idxArg, ocl::KernelArg::PtrWriteOnly(dst)); + idxArg = kernel.set(idxArg, (int)dst.step); + idxArg = kernel.set(idxArg, (int)dst.rows); + idxArg = kernel.set(idxArg, (int)dst.cols); + idxArg = kernel.set(idxArg, static_cast(delta)); + + return kernel.run(2, globalsize, (localsize[0] == 0) ? NULL : localsize, false); +} +} +#endif + void cv::Sobel( InputArray _src, OutputArray _dst, int ddepth, int dx, int dy, int ksize, double scale, double delta, int borderType ) { @@ -554,6 +614,11 @@ void cv::Sobel( InputArray _src, OutputArray _dst, int ddepth, int dx, int dy, else ky *= scale; } + + CV_OCL_RUN(_dst.isUMat() && _src.dims() <= 2 && ksize == 3 && + (size_t)_src.rows() > ky.total() && (size_t)_src.cols() > kx.total(), + ocl_sepFilter3x3_8UC1(_src, _dst, ddepth, kx, ky, delta, borderType)); + sepFilter2D( _src, _dst, ddepth, kx, ky, Point(-1, -1), delta, borderType ); } @@ -593,6 +658,11 @@ void cv::Scharr( InputArray _src, OutputArray _dst, int ddepth, int dx, int dy, else ky *= scale; } + + CV_OCL_RUN(_dst.isUMat() && _src.dims() <= 2 && + (size_t)_src.rows() > ky.total() && (size_t)_src.cols() > kx.total(), + ocl_sepFilter3x3_8UC1(_src, _dst, ddepth, kx, ky, delta, borderType)); + sepFilter2D( _src, _dst, ddepth, kx, ky, Point(-1, -1), delta, borderType ); } diff --git a/modules/imgproc/src/opencl/sepFilter3x3.cl b/modules/imgproc/src/opencl/sepFilter3x3.cl new file mode 100644 index 0000000000..fc8bc06f94 --- /dev/null +++ b/modules/imgproc/src/opencl/sepFilter3x3.cl @@ -0,0 +1,135 @@ +// This file is part of OpenCV project. +// It is subject to the license terms in the LICENSE file found in the top-level directory +// of this distribution and at http://opencv.org/license.html. + +#define DIG(a) a, +__constant float kx[] = { KERNEL_MATRIX_X }; +__constant float ky[] = { KERNEL_MATRIX_Y }; + +#define OP(delta, y, x) (convert_float16(arr[(y + delta) * 3 + x]) * ky[y] * kx[x]) + +__kernel void sepFilter3x3_8UC1_cols16_rows2(__global const uint* src, int src_step, + __global uint* dst, int dst_step, + int rows, int cols, float delta) +{ + int block_x = get_global_id(0); + int y = get_global_id(1) * 2; + int ssx, dsx; + + if ((block_x * 16) >= cols || y >= rows) return; + + uint4 line[4]; + uint4 line_out[2]; + uchar a; uchar16 b; uchar c; + uchar d; uchar16 e; uchar f; + uchar g; uchar16 h; uchar i; + uchar j; uchar16 k; uchar l; + + ssx = dsx = 1; + int src_index = block_x * 4 * ssx + (y - 1) * (src_step / 4); + line[1] = vload4(0, src + src_index + (src_step / 4)); + line[2] = vload4(0, src + src_index + 2 * (src_step / 4)); + +#ifdef BORDER_CONSTANT + line[0] = (y == 0) ? (uint4)0 : vload4(0, src + src_index); + line[3] = (y == (rows - 2)) ? (uint4)0 : vload4(0, src + src_index + 3 * (src_step / 4)); +#elif defined BORDER_REFLECT_101 + line[0] = (y == 0) ? line[2] : vload4(0, src + src_index); + line[3] = (y == (rows - 2)) ? line[1] : vload4(0, src + src_index + 3 * (src_step / 4)); +#elif defined (BORDER_REPLICATE) || defined(BORDER_REFLECT) + line[0] = (y == 0) ? line[1] : vload4(0, src + src_index); + line[3] = (y == (rows - 2)) ? line[2] : vload4(0, src + src_index + 3 * (src_step / 4)); +#endif + + __global uchar *src_p = (__global uchar *)src; + + src_index = block_x * 16 * ssx + (y - 1) * src_step; + bool line_end = ((block_x + 1) * 16 == cols); + + b = as_uchar16(line[0]); + e = as_uchar16(line[1]); + h = as_uchar16(line[2]); + k = as_uchar16(line[3]); + +#ifdef BORDER_CONSTANT + a = (block_x == 0 || y == 0) ? 0 : src_p[src_index - 1]; + c = (line_end || y == 0) ? 0 : src_p[src_index + 16]; + + d = (block_x == 0) ? 0 : src_p[src_index + src_step - 1]; + f = line_end ? 0 : src_p[src_index + src_step + 16]; + + g = (block_x == 0) ? 0 : src_p[src_index + 2 * src_step - 1]; + i = line_end ? 0 : src_p[src_index + 2 * src_step + 16]; + + j = (block_x == 0 || y == (rows - 2)) ? 0 : src_p[src_index + 3 * src_step - 1]; + l = (line_end || y == (rows - 2))? 0 : src_p[src_index + 3 * src_step + 16]; + +#elif defined BORDER_REFLECT_101 + int offset; + offset = (y == 0) ? (2 * src_step) : 0; + + a = (block_x == 0) ? src_p[src_index + offset + 1] : src_p[src_index + offset - 1]; + c = line_end ? src_p[src_index + offset + 14] : src_p[src_index + offset + 16]; + + d = (block_x == 0) ? src_p[src_index + src_step + 1] : src_p[src_index + src_step - 1]; + f = line_end ? src_p[src_index + src_step + 14] : src_p[src_index + src_step + 16]; + + g = (block_x == 0) ? src_p[src_index + 2 * src_step + 1] : src_p[src_index + 2 * src_step - 1]; + i = line_end ? src_p[src_index + 2 * src_step + 14] : src_p[src_index + 2 * src_step + 16]; + + offset = (y == (rows - 2)) ? (1 * src_step) : (3 * src_step); + + j = (block_x == 0) ? src_p[src_index + offset + 1] : src_p[src_index + offset - 1]; + l = line_end ? src_p[src_index + offset + 14] : src_p[src_index + offset + 16]; + +#elif defined (BORDER_REPLICATE) || defined(BORDER_REFLECT) + int offset; + offset = (y == 0) ? (1 * src_step) : 0; + + a = (block_x == 0) ? src_p[src_index + offset] : src_p[src_index + offset - 1]; + c = line_end ? src_p[src_index + offset + 15] : src_p[src_index + offset + 16]; + + d = (block_x == 0) ? src_p[src_index + src_step] : src_p[src_index + src_step - 1]; + f = line_end ? src_p[src_index + src_step + 15] : src_p[src_index + src_step + 16]; + + g = (block_x == 0) ? src_p[src_index + 2 * src_step] : src_p[src_index + 2 * src_step - 1]; + i = line_end ? src_p[src_index + 2 * src_step + 15] : src_p[src_index + 2 * src_step + 16]; + + offset = (y == (rows - 2)) ? (2 * src_step) : (3 * src_step); + + j = (block_x == 0) ? src_p[src_index + offset] : src_p[src_index + offset - 1]; + l = line_end ? src_p[src_index + offset + 15] : src_p[src_index + offset + 16]; + +#endif + + uchar16 arr[12]; + float16 sum[2]; + + arr[0] = (uchar16)(a, b.s0123, b.s456789ab, b.scde); + arr[1] = b; + arr[2] = (uchar16)(b.s123, b.s4567, b.s89abcdef, c); + arr[3] = (uchar16)(d, e.s0123, e.s456789ab, e.scde); + arr[4] = e; + arr[5] = (uchar16)(e.s123, e.s4567, e.s89abcdef, f); + arr[6] = (uchar16)(g, h.s0123, h.s456789ab, h.scde); + arr[7] = h; + arr[8] = (uchar16)(h.s123, h.s4567, h.s89abcdef, i); + arr[9] = (uchar16)(j, k.s0123, k.s456789ab, k.scde); + arr[10] = k; + arr[11] = (uchar16)(k.s123, k.s4567, k.s89abcdef, l); + + sum[0] = OP(0, 0, 0) + OP(0, 0, 1) + OP(0, 0, 2) + + OP(0, 1, 0) + OP(0, 1, 1) + OP(0, 1, 2) + + OP(0, 2, 0) + OP(0, 2, 1) + OP(0, 2, 2); + + sum[1] = OP(1, 0, 0) + OP(1, 0, 1) + OP(1, 0, 2) + + OP(1, 1, 0) + OP(1, 1, 1) + OP(1, 1, 2) + + OP(1, 2, 0) + OP(1, 2, 1) + OP(1, 2, 2); + + line_out[0] = as_uint4(convert_uchar16_sat_rte(sum[0] + delta)); + line_out[1] = as_uint4(convert_uchar16_sat_rte(sum[1] + delta)); + + int dst_index = block_x * 4 * dsx + y * (dst_step / 4); + vstore4(line_out[0], 0, dst + dst_index); + vstore4(line_out[1], 0, dst + dst_index + (dst_step / 4)); +} diff --git a/modules/imgproc/test/ocl/test_filters.cpp b/modules/imgproc/test/ocl/test_filters.cpp index 8dd5b95c7c..927cd2210b 100644 --- a/modules/imgproc/test/ocl/test_filters.cpp +++ b/modules/imgproc/test/ocl/test_filters.cpp @@ -187,6 +187,84 @@ OCL_TEST_P(SobelTest, Mat) } } +PARAM_TEST_CASE(Deriv3x3_cols16_rows2_Base, MatType, + int, // kernel size + Size, // dx, dy + BorderType, // border type + double, // optional parameter + bool, // roi or not + int) // width multiplier +{ + int type, borderType, ksize; + Size size; + double param; + bool useRoi; + int widthMultiple; + + TEST_DECLARE_INPUT_PARAMETER(src); + TEST_DECLARE_OUTPUT_PARAMETER(dst); + + virtual void SetUp() + { + type = GET_PARAM(0); + ksize = GET_PARAM(1); + size = GET_PARAM(2); + borderType = GET_PARAM(3); + param = GET_PARAM(4); + useRoi = GET_PARAM(5); + widthMultiple = GET_PARAM(6); + } + + void random_roi() + { + size = Size(3, 3); + + Size roiSize = randomSize(size.width, MAX_VALUE, size.height, MAX_VALUE); + roiSize.width = std::max(size.width + 13, roiSize.width & (~0xf)); + roiSize.height = std::max(size.height + 1, roiSize.height & (~0x1)); + + Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); + randomSubMat(src, src_roi, roiSize, srcBorder, type, 5, 256); + + Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); + randomSubMat(dst, dst_roi, roiSize, dstBorder, type, -60, 70); + + UMAT_UPLOAD_INPUT_PARAMETER(src); + UMAT_UPLOAD_OUTPUT_PARAMETER(dst); + } + + void Near() + { + Near(1, false); + } + + void Near(double threshold, bool relative) + { + if (relative) + OCL_EXPECT_MATS_NEAR_RELATIVE(dst, threshold); + else + OCL_EXPECT_MATS_NEAR(dst, threshold); + } +}; + +typedef Deriv3x3_cols16_rows2_Base Sobel3x3_cols16_rows2; + +OCL_TEST_P(Sobel3x3_cols16_rows2, Mat) +{ + int dx = size.width, dy = size.height; + double scale = param; + + for (int j = 0; j < test_loop_times; j++) + { + random_roi(); + + OCL_OFF(cv::Sobel(src_roi, dst_roi, -1, dx, dy, ksize, scale, /* delta */0, borderType)); + OCL_ON(cv::Sobel(usrc_roi, udst_roi, -1, dx, dy, ksize, scale, /* delta */0, borderType)); + + Near(); + } +} + ///////////////////////////////////////////////////////////////////////////////////////////////// // Scharr @@ -208,6 +286,24 @@ OCL_TEST_P(ScharrTest, Mat) } } +typedef Deriv3x3_cols16_rows2_Base Scharr3x3_cols16_rows2; + +OCL_TEST_P(Scharr3x3_cols16_rows2, Mat) +{ + int dx = size.width, dy = size.height; + double scale = param; + + for (int j = 0; j < test_loop_times; j++) + { + random_roi(); + + OCL_OFF(cv::Scharr(src_roi, dst_roi, -1, dx, dy, scale, /* delta */ 0, borderType)); + OCL_ON(cv::Scharr(usrc_roi, udst_roi, -1, dx, dy, scale, /* delta */ 0, borderType)); + + Near(); + } +} + ///////////////////////////////////////////////////////////////////////////////////////////////// // GaussianBlur @@ -552,6 +648,15 @@ OCL_INSTANTIATE_TEST_CASE_P(Filter, SobelTest, Combine( Bool(), Values(1))); // not used +OCL_INSTANTIATE_TEST_CASE_P(Filter, Sobel3x3_cols16_rows2, Combine( + Values((MatType)CV_8UC1), + Values(3), // kernel size + Values(Size(1, 0), Size(1, 1), Size(2, 0), Size(2, 1)), // dx, dy + FILTER_BORDER_SET_NO_WRAP_NO_ISOLATED, + Values(0.0), // not used + Bool(), + Values(1))); // not used + OCL_INSTANTIATE_TEST_CASE_P(Filter, ScharrTest, Combine( FILTER_TYPES, Values(0), // not used @@ -561,6 +666,15 @@ OCL_INSTANTIATE_TEST_CASE_P(Filter, ScharrTest, Combine( Bool(), Values(1))); // not used +OCL_INSTANTIATE_TEST_CASE_P(Filter, Scharr3x3_cols16_rows2, Combine( + FILTER_TYPES, + Values(0), // not used + Values(Size(0, 1), Size(1, 0)), // dx, dy + FILTER_BORDER_SET_NO_WRAP_NO_ISOLATED, + Values(1.0, 0.2), // kernel scale + Bool(), + Values(1))); // not used + OCL_INSTANTIATE_TEST_CASE_P(Filter, GaussianBlurTest, Combine( FILTER_TYPES, Values(3, 5), // kernel size