diff --git a/modules/imgproc/src/morph.cpp b/modules/imgproc/src/morph.cpp index 3ab495d75..5eec87cb2 100644 --- a/modules/imgproc/src/morph.cpp +++ b/modules/imgproc/src/morph.cpp @@ -1333,69 +1333,98 @@ static bool IPPMorphOp(int op, InputArray _src, OutputArray _dst, if( iterations > 1 ) return false; - if (IPPMorphReplicate( op, src, dst, kernel, ksize, anchor, rectKernel )) - return true; - - return false; + return IPPMorphReplicate( op, src, dst, kernel, ksize, anchor, rectKernel ); } #endif #ifdef HAVE_OPENCL -static bool ocl_morphology_op(InputArray _src, OutputArray _dst, Mat kernel, - const Size & ksize, const Point & anchor, int iterations, int op) +static bool ocl_morphOp(InputArray _src, OutputArray _dst, InputArray _kernel, + Point anchor, int iterations, int op, int borderType, + const Scalar &, int actual_op = -1, InputArray _extraMat = noArray()) { - CV_Assert(op == MORPH_ERODE || op == MORPH_DILATE); - + const ocl::Device & dev = ocl::Device::getDefault(); int type = _src.type(), depth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); - bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0; + bool doubleSupport = dev.doubleFPConfig() > 0; - if (depth == CV_64F && !doubleSupport) + if ((depth == CV_64F && !doubleSupport) || borderType != BORDER_CONSTANT) return false; - UMat kernel8U; - kernel.convertTo(kernel8U, CV_8U); - kernel8U = kernel8U.reshape(1, 1); + Mat kernel = _kernel.getMat(); + bool haveExtraMat = !_extraMat.empty(); + Size ksize = kernel.data ? kernel.size() : Size(3, 3), ssize = _src.size(); + CV_Assert(actual_op <= 3 || haveExtraMat); - bool rectKernel = true; + if (iterations == 0 || kernel.rows*kernel.cols == 1) { - Mat m = kernel.reshape(1, 1); - for (int i = 0; i < m.size().area(); ++i) - if (m.at(i) != 1) - { - rectKernel = false; - break; - } + _src.copyTo(_dst); + return true; } - UMat src = _src.getUMat(); + if (!kernel.data) + { + kernel = getStructuringElement(MORPH_RECT, Size(1+iterations*2,1+iterations*2)); + anchor = Point(iterations, iterations); + iterations = 1; + } + else if( iterations > 1 && countNonZero(kernel) == kernel.rows*kernel.cols ) + { + anchor = Point(anchor.x*iterations, anchor.y*iterations); + kernel = getStructuringElement(MORPH_RECT, + Size(ksize.width + (iterations-1)*(ksize.width-1), + ksize.height + (iterations-1)*(ksize.height-1)), + anchor); + iterations = 1; + } #ifdef ANDROID - size_t localThreads[3] = {16, 8, 1}; + size_t localThreads[2] = { 16, 8 }; #else - size_t localThreads[3] = {16, 16, 1}; + size_t localThreads[2] = { 16, 16 }; #endif - size_t globalThreads[3] = {(src.cols + localThreads[0] - 1) / localThreads[0] *localThreads[0], (src.rows + localThreads[1] - 1) / localThreads[1] *localThreads[1], 1}; + size_t globalThreads[2] = { ssize.width, ssize.height }; if (localThreads[0]*localThreads[1] * 2 < (localThreads[0] + ksize.width - 1) * (localThreads[1] + ksize.height - 1)) return false; - static const char * const op2str[] = { "ERODE", "DILATE" }; - String buildOptions = format("-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D %s%s%s" - " -D T=%s -D DEPTH_%d -D cn=%d -D T1=%s", anchor.x, anchor.y, - (int)localThreads[0], (int)localThreads[1], op2str[op], - doubleSupport ? " -D DOUBLE_SUPPORT" : "", rectKernel ? " -D RECTKERNEL" : "", - ocl::typeToStr(_src.type()), _src.depth(), cn, ocl::typeToStr(depth)); + // build processing + String processing; + Mat kernel8u; + kernel.convertTo(kernel8u, CV_8U); + for (int y = 0; y < kernel8u.rows; ++y) + for (int x = 0; x < kernel8u.cols; ++x) + if (kernel8u.at(y, x) != 0) + processing += format("PROCESS(%d,%d)", y, x); - std::vector kernels; + static const char * const op2str[] = { "OP_ERODE", "OP_DILATE", NULL, NULL, "OP_GRADIENT", "OP_TOPHAT", "OP_BLACKHAT" }; + + char cvt[2][50]; + int wdepth = std::max(depth, CV_32F), scalarcn = cn == 3 ? 4 : cn; + + if (actual_op < 0) + actual_op = op; + + std::vector kernels(iterations); for (int i = 0; i < iterations; i++) { - ocl::Kernel k("morph", ocl::imgproc::morph_oclsrc, buildOptions); - if (k.empty()) + int current_op = iterations == i + 1 ? actual_op : op; + String buildOptions = format("-D RADIUSX=%d -D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D %s%s" + " -D PROCESS_ELEMS=%s -D T=%s -D DEPTH_%d -D cn=%d -D T1=%s" + " -D convertToWT=%s -D convertToT=%s -D ST=%s%s", + anchor.x, anchor.y, (int)localThreads[0], (int)localThreads[1], op2str[op], + doubleSupport ? " -D DOUBLE_SUPPORT" : "", processing.c_str(), + ocl::typeToStr(type), depth, cn, ocl::typeToStr(depth), + ocl::convertTypeStr(depth, wdepth, cn, cvt[0]), + ocl::convertTypeStr(wdepth, depth, cn, cvt[1]), + ocl::typeToStr(CV_MAKE_TYPE(depth, scalarcn)), + current_op == op ? "" : cv::format(" -D %s", op2str[current_op]).c_str()); + + kernels[i].create("morph", ocl::imgproc::morph_oclsrc, buildOptions); + if (kernels[i].empty()) return false; - kernels.push_back(k); } + UMat src = _src.getUMat(), extraMat = _extraMat.getUMat(); _dst.create(src.size(), src.type()); UMat dst = _dst.getUMat(); @@ -1406,9 +1435,13 @@ static bool ocl_morphology_op(InputArray _src, OutputArray _dst, Mat kernel, src.locateROI(wholesize, ofs); int wholecols = wholesize.width, wholerows = wholesize.height; - kernels[0].args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnlyNoSize(dst), - ofs.x, ofs.y, src.cols, src.rows, ocl::KernelArg::PtrReadOnly(kernel8U), - wholecols, wholerows); + if (haveExtraMat) + kernels[0].args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnlyNoSize(dst), + ofs.x, ofs.y, src.cols, src.rows, wholecols, wholerows, + ocl::KernelArg::ReadOnlyNoSize(extraMat)); + else + kernels[0].args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnlyNoSize(dst), + ofs.x, ofs.y, src.cols, src.rows, wholecols, wholerows); return kernels[0].run(2, globalThreads, localThreads, false); } @@ -1422,19 +1455,20 @@ static bool ocl_morphology_op(InputArray _src, OutputArray _dst, Mat kernel, if (i == 0) { int cols = src.cols, rows = src.rows; - src.locateROI(wholesize,ofs); + src.locateROI(wholesize, ofs); src.adjustROI(ofs.y, wholesize.height - rows - ofs.y, ofs.x, wholesize.width - cols - ofs.x); if(src.u != dst.u) source = src; else src.copyTo(source); + src.adjustROI(-ofs.y, -wholesize.height + rows + ofs.y, -ofs.x, -wholesize.width + cols + ofs.x); source.adjustROI(-ofs.y, -wholesize.height + rows + ofs.y, -ofs.x, -wholesize.width + cols + ofs.x); } else { int cols = dst.cols, rows = dst.rows; - dst.locateROI(wholesize,ofs); + dst.locateROI(wholesize, ofs); dst.adjustROI(ofs.y, wholesize.height - rows - ofs.y, ofs.x, wholesize.width - cols - ofs.x); dst.copyTo(source); dst.adjustROI(-ofs.y, -wholesize.height + rows + ofs.y, -ofs.x, -wholesize.width + cols + ofs.x); @@ -1442,13 +1476,18 @@ static bool ocl_morphology_op(InputArray _src, OutputArray _dst, Mat kernel, } source.locateROI(wholesize, ofs); - kernels[i].args(ocl::KernelArg::ReadOnlyNoSize(source), ocl::KernelArg::WriteOnlyNoSize(dst), - ofs.x, ofs.y, source.cols, source.rows, ocl::KernelArg::PtrReadOnly(kernel8U), - wholesize.width, wholesize.height); + if (haveExtraMat && iterations == i + 1) + kernels[i].args(ocl::KernelArg::ReadOnlyNoSize(source), ocl::KernelArg::WriteOnlyNoSize(dst), + ofs.x, ofs.y, source.cols, source.rows, wholesize.width, wholesize.height, + ocl::KernelArg::ReadOnlyNoSize(extraMat)); + else + kernels[i].args(ocl::KernelArg::ReadOnlyNoSize(source), ocl::KernelArg::WriteOnlyNoSize(dst), + ofs.x, ofs.y, source.cols, source.rows, wholesize.width, wholesize.height); if (!kernels[i].run(2, globalThreads, localThreads, false)) return false; } + return true; } @@ -1459,15 +1498,16 @@ static void morphOp( int op, InputArray _src, OutputArray _dst, Point anchor, int iterations, int borderType, const Scalar& borderValue ) { -#ifdef HAVE_OPENCL - int src_type = _src.type(), - src_cn = CV_MAT_CN(src_type), src_depth = CV_MAT_DEPTH(src_type); -#endif - Mat kernel = _kernel.getMat(); Size ksize = kernel.data ? kernel.size() : Size(3,3); anchor = normalizeAnchor(anchor, ksize); + CV_OCL_RUN(_dst.isUMat() && _src.dims() <= 2 && _src.channels() <= 4 && + borderType == cv::BORDER_CONSTANT && borderValue == morphologyDefaultBorderValue() && + (op == MORPH_ERODE || op == MORPH_DILATE) && + anchor.x == ksize.width >> 1 && anchor.y == ksize.height >> 1, + ocl_morphOp(_src, _dst, kernel, anchor, iterations, op, borderType, borderValue) ) + if (iterations == 0 || kernel.rows*kernel.cols == 1) { _src.copyTo(_dst); @@ -1490,12 +1530,6 @@ static void morphOp( int op, InputArray _src, OutputArray _dst, iterations = 1; } - CV_OCL_RUN(_dst.isUMat() && _src.dims() <= 2 && src_cn <= 4 && - (src_depth == CV_8U || src_depth == CV_32F || src_depth == CV_64F ) && - borderType == cv::BORDER_CONSTANT && borderValue == morphologyDefaultBorderValue() && - (op == MORPH_ERODE || op == MORPH_DILATE), - ocl_morphology_op(_src, _dst, kernel, ksize, anchor, iterations, op) ) - #if IPP_VERSION_X100 >= 801 if( IPPMorphOp(op, _src, _dst, kernel, anchor, iterations, borderType, borderValue) ) return; @@ -1515,13 +1549,6 @@ static void morphOp( int op, InputArray _src, OutputArray _dst, parallel_for_(Range(0, nStripes), MorphologyRunner(src, dst, nStripes, iterations, op, kernel, anchor, borderType, borderType, borderValue)); - - //Ptr f = createMorphologyFilter(op, src.type(), - // kernel, anchor, borderType, borderType, borderValue ); - - //f->apply( src, dst ); - //for( int i = 1; i < iterations; i++ ) - // f->apply( dst, dst ); } } @@ -1541,97 +1568,122 @@ void cv::dilate( InputArray src, OutputArray dst, InputArray kernel, morphOp( MORPH_DILATE, src, dst, kernel, anchor, iterations, borderType, borderValue ); } -void cv::morphologyEx( InputArray _src, OutputArray _dst, int op, - InputArray kernel, Point anchor, int iterations, - int borderType, const Scalar& borderValue ) +#ifdef HAVE_OPENCL + +namespace cv { + +static bool ocl_morphologyEx(InputArray _src, OutputArray _dst, int op, + InputArray kernel, Point anchor, int iterations, + int borderType, const Scalar& borderValue) { - int src_type = _src.type(), dst_type = _dst.type(), - src_cn = CV_MAT_CN(src_type), src_depth = CV_MAT_DEPTH(src_type); - - bool use_opencl = cv::ocl::useOpenCL() && _src.isUMat() && _src.size() == _dst.size() && src_type == dst_type && - _src.dims()<=2 && (src_cn == 1 || src_cn == 4) && (anchor.x == -1) && (anchor.y == -1) && - (src_depth == CV_8U || src_depth == CV_32F || src_depth == CV_64F ) && - (borderType == cv::BORDER_CONSTANT) && (borderValue == morphologyDefaultBorderValue()); - - _dst.create(_src.size(), _src.type()); - Mat src, dst, temp; - UMat usrc, udst, utemp; + _dst.createSameSize(_src, _src.type()); + bool submat = _dst.isSubmatrix(); + UMat temp; + _OutputArray _temp = submat ? _dst : _OutputArray(temp); switch( op ) { case MORPH_ERODE: - erode( _src, _dst, kernel, anchor, iterations, borderType, borderValue ); + if (!ocl_morphOp( _src, _dst, kernel, anchor, iterations, MORPH_ERODE, borderType, borderValue )) + return false; break; case MORPH_DILATE: - dilate( _src, _dst, kernel, anchor, iterations, borderType, borderValue ); + if (!ocl_morphOp( _src, _dst, kernel, anchor, iterations, MORPH_DILATE, borderType, borderValue )) + return false; break; case MORPH_OPEN: - erode( _src, _dst, kernel, anchor, iterations, borderType, borderValue ); - dilate( _dst, _dst, kernel, anchor, iterations, borderType, borderValue ); + if (!ocl_morphOp( _src, _temp, kernel, anchor, iterations, MORPH_ERODE, borderType, borderValue )) + return false; + if (!ocl_morphOp( _temp, _dst, kernel, anchor, iterations, MORPH_DILATE, borderType, borderValue )) + return false; + break; + case MORPH_CLOSE: + if (!ocl_morphOp( _src, _temp, kernel, anchor, iterations, MORPH_DILATE, borderType, borderValue )) + return false; + if (!ocl_morphOp( _temp, _dst, kernel, anchor, iterations, MORPH_ERODE, borderType, borderValue )) + return false; + break; + case MORPH_GRADIENT: + if (!ocl_morphOp( _src, temp, kernel, anchor, iterations, MORPH_ERODE, borderType, borderValue )) + return false; + if (!ocl_morphOp( _src, _dst, kernel, anchor, iterations, MORPH_DILATE, borderType, borderValue, MORPH_GRADIENT, temp )) + return false; + break; + case MORPH_TOPHAT: + if (!ocl_morphOp( _src, _temp, kernel, anchor, iterations, MORPH_ERODE, borderType, borderValue )) + return false; + if (!ocl_morphOp( _temp, _dst, kernel, anchor, iterations, MORPH_DILATE, borderType, borderValue, MORPH_TOPHAT, _src )) + return false; + break; + case MORPH_BLACKHAT: + if (!ocl_morphOp( _src, _temp, kernel, anchor, iterations, MORPH_DILATE, borderType, borderValue )) + return false; + if (!ocl_morphOp( _temp, _dst, kernel, anchor, iterations, MORPH_ERODE, borderType, borderValue, MORPH_BLACKHAT, _src )) + return false; + break; + default: + CV_Error( CV_StsBadArg, "unknown morphological operation" ); + } + + return true; +} + +} + +#endif + +void cv::morphologyEx( InputArray _src, OutputArray _dst, int op, + InputArray kernel, Point anchor, int iterations, + int borderType, const Scalar& borderValue ) +{ +#ifdef HAVE_OPENCL + Size ksize = kernel.size(); + anchor = normalizeAnchor(anchor, ksize); + + CV_OCL_RUN(_dst.isUMat() && _src.dims() <= 2 && _src.channels() <= 4 && + anchor.x == ksize.width >> 1 && anchor.y == ksize.height >> 1 && + borderType == cv::BORDER_CONSTANT && borderValue == morphologyDefaultBorderValue(), + ocl_morphologyEx(_src, _dst, op, kernel, anchor, iterations, borderType, borderValue)) +#endif + + Mat src = _src.getMat(), temp; + _dst.create(src.size(), src.type()); + Mat dst = _dst.getMat(); + + switch( op ) + { + case MORPH_ERODE: + erode( src, dst, kernel, anchor, iterations, borderType, borderValue ); + break; + case MORPH_DILATE: + dilate( src, dst, kernel, anchor, iterations, borderType, borderValue ); + break; + case MORPH_OPEN: + erode( src, dst, kernel, anchor, iterations, borderType, borderValue ); + dilate( dst, dst, kernel, anchor, iterations, borderType, borderValue ); break; case CV_MOP_CLOSE: - dilate( _src, _dst, kernel, anchor, iterations, borderType, borderValue ); - erode( _dst, _dst, kernel, anchor, iterations, borderType, borderValue ); + dilate( src, dst, kernel, anchor, iterations, borderType, borderValue ); + erode( dst, dst, kernel, anchor, iterations, borderType, borderValue ); break; case CV_MOP_GRADIENT: - erode( _src, use_opencl ? (cv::OutputArray)utemp : (cv::OutputArray)temp, kernel, anchor, iterations, borderType, borderValue ); - dilate( _src, _dst, kernel, anchor, iterations, borderType, borderValue ); - if(use_opencl) - { - udst = _dst.getUMat(); - subtract(udst, utemp, udst); - } - else - { - dst = _dst.getMat(); - dst -= temp; - } + erode( src, temp, kernel, anchor, iterations, borderType, borderValue ); + dilate( src, dst, kernel, anchor, iterations, borderType, borderValue ); + dst -= temp; break; case CV_MOP_TOPHAT: - if(use_opencl) - { - usrc = _src.getUMat(); - udst = _dst.getUMat(); - if( usrc.u != udst.u ) - utemp = udst; - } - else - { - src = _src.getMat(); - dst = _dst.getMat(); - if( src.data != dst.data ) - temp = dst; - } - erode( _src, use_opencl ? (cv::OutputArray)utemp : (cv::OutputArray)temp, kernel, anchor, iterations, borderType, borderValue ); - dilate( use_opencl ? (cv::OutputArray)utemp : (cv::OutputArray)temp, use_opencl ? (cv::OutputArray)utemp : (cv::OutputArray)temp, kernel, - anchor, iterations, borderType, borderValue ); - if(use_opencl) - subtract(usrc, utemp, udst); - else - dst = src - temp; + if( src.data != dst.data ) + temp = dst; + erode( src, temp, kernel, anchor, iterations, borderType, borderValue ); + dilate( temp, temp, kernel, anchor, iterations, borderType, borderValue ); + dst = src - temp; break; case CV_MOP_BLACKHAT: - if(use_opencl) - { - usrc = _src.getUMat(); - udst = _dst.getUMat(); - if( usrc.u != udst.u ) - utemp = udst; - } - else - { - src = _src.getMat(); - dst = _dst.getMat(); - if( src.data != dst.data ) - temp = dst; - } - dilate( _src, use_opencl ? (cv::OutputArray)utemp : (cv::OutputArray)temp, kernel, anchor, iterations, borderType, borderValue ); - erode( use_opencl ? (cv::OutputArray)utemp : (cv::OutputArray)temp, use_opencl ? (cv::OutputArray)utemp : (cv::OutputArray)temp, kernel, - anchor, iterations, borderType, borderValue ); - if(use_opencl) - subtract(utemp, usrc, udst); - else - dst = temp - src; + if( src.data != dst.data ) + temp = dst; + dilate( src, temp, kernel, anchor, iterations, borderType, borderValue ); + erode( temp, temp, kernel, anchor, iterations, borderType, borderValue ); + dst = temp - src; break; default: CV_Error( CV_StsBadArg, "unknown morphological operation" ); diff --git a/modules/imgproc/src/opencl/morph.cl b/modules/imgproc/src/opencl/morph.cl index a7611c50f..f78af89c9 100644 --- a/modules/imgproc/src/opencl/morph.cl +++ b/modules/imgproc/src/opencl/morph.cl @@ -43,6 +43,8 @@ #endif #endif +#define noconvert + #if cn != 3 #define loadpix(addr) *(__global const T *)(addr) #define storepix(val, addr) *(__global T *)(addr) = val @@ -54,59 +56,75 @@ #endif #ifdef DEPTH_0 -#ifdef ERODE -#define VAL 255 -#endif -#ifdef DILATE -#define VAL 0 -#endif +#define MIN_VAL 0 +#define MAX_VAL UCHAR_MAX +#elif defined DEPTH_1 +#define MIN_VAL SCHAR_MIN +#define MAX_VAL SCHAR_MAX +#elif defined DEPTH_2 +#define MIN_VAL 0 +#define MAX_VAL USHRT_MAX +#elif defined DEPTH_3 +#define MIN_VAL SHRT_MIN +#define MAX_VAL SHRT_MAX +#elif defined DEPTH_4 +#define MIN_VAL INT_MIN +#define MAX_VAL INT_MAX #elif defined DEPTH_5 -#ifdef ERODE -#define VAL FLT_MAX -#endif -#ifdef DILATE -#define VAL -FLT_MAX -#endif +#define MIN_VAL (-FLT_MAX) +#define MAX_VAL FLT_MAX #elif defined DEPTH_6 -#ifdef ERODE -#define VAL DBL_MAX -#endif -#ifdef DILATE -#define VAL -DBL_MAX -#endif +#define MIN_VAL (-DBL_MAX) +#define MAX_VAL DBL_MAX #endif -#ifdef ERODE -#if defined(INTEL_DEVICE) && (DEPTH_0) +#ifdef OP_ERODE +#define VAL MAX_VAL +#elif defined OP_DILATE +#define VAL MIN_VAL +#else +#error "Unknown operation" +#endif + +#ifdef OP_ERODE +#if defined INTEL_DEVICE && defined DEPTH_0 // workaround for bug in Intel HD graphics drivers (10.18.10.3496 or older) #define __CAT(x, y) x##y #define CAT(x, y) __CAT(x, y) #define WA_CONVERT_1 CAT(convert_uint, cn) #define WA_CONVERT_2 CAT(convert_, T) #define convert_uint1 convert_uint -#define MORPH_OP(A,B) WA_CONVERT_2(min(WA_CONVERT_1(A),WA_CONVERT_1(B))) +#define MORPH_OP(A, B) WA_CONVERT_2(min(WA_CONVERT_1(A), WA_CONVERT_1(B))) #else -#define MORPH_OP(A,B) min((A),(B)) +#define MORPH_OP(A, B) min((A), (B)) #endif #endif -#ifdef DILATE -#define MORPH_OP(A,B) max((A),(B)) +#ifdef OP_DILATE +#define MORPH_OP(A, B) max((A), (B)) #endif +#define PROCESS(y, x) \ + res = MORPH_OP(res, LDS_DAT[mad24(l_y + y, width, l_x + x)]); + // BORDER_CONSTANT: iiiiii|abcdefgh|iiiiiii #define ELEM(i, l_edge, r_edge, elem1, elem2) (i) < (l_edge) | (i) >= (r_edge) ? (elem1) : (elem2) +#if defined OP_GRADIENT || defined OP_TOPHAT || defined OP_BLACKHAT +#define EXTRA_PARAMS , __global const uchar * matptr, int mat_step, int mat_offset +#else +#define EXTRA_PARAMS +#endif + __kernel void morph(__global const uchar * srcptr, int src_step, int src_offset, __global uchar * dstptr, int dst_step, int dst_offset, int src_offset_x, int src_offset_y, int cols, int rows, - __constant uchar * mat_kernel, int src_whole_cols, int src_whole_rows) + int src_whole_cols, int src_whole_rows EXTRA_PARAMS) { int gidx = get_global_id(0), gidy = get_global_id(1); int l_x = get_local_id(0), l_y = get_local_id(1); int x = get_group_id(0) * LSIZE0, y = get_group_id(1) * LSIZE1; int start_x = x + src_offset_x - RADIUSX; - int end_x = x + src_offset_x + LSIZE0 + RADIUSX; - int width = end_x - (x + src_offset_x - RADIUSX) + 1; + int width = mad24(RADIUSX, 2, LSIZE0 + 1); int start_y = y + src_offset_y - RADIUSY; int point1 = mad24(l_y, LSIZE0, l_x); int point2 = point1 + LSIZE0 * LSIZE1; @@ -117,7 +135,7 @@ __kernel void morph(__global const uchar * srcptr, int src_step, int src_offset, int start_addr = mad24(cur_y, src_step, cur_x * TSIZE); int start_addr2 = mad24(cur_y2, src_step, cur_x2 * TSIZE); - __local T LDS_DAT[2*LSIZE1*LSIZE0]; + __local T LDS_DAT[2 * LSIZE1 * LSIZE0]; // read pixels from src int end_addr = mad24(src_whole_rows - 1, src_step, src_whole_cols * TSIZE); @@ -128,8 +146,8 @@ __kernel void morph(__global const uchar * srcptr, int src_step, int src_offset, T temp1 = loadpix(srcptr + start_addr2); // judge if read out of boundary - temp0 = ELEM(cur_x, 0, src_whole_cols, (T)(VAL),temp0); - temp0 = ELEM(cur_y, 0, src_whole_rows, (T)(VAL),temp0); + temp0 = ELEM(cur_x, 0, src_whole_cols, (T)(VAL), temp0); + temp0 = ELEM(cur_y, 0, src_whole_rows, (T)(VAL), temp0); temp1 = ELEM(cur_x2, 0, src_whole_cols, (T)(VAL), temp1); temp1 = ELEM(cur_y2, 0, src_whole_rows, (T)(VAL), temp1); @@ -138,24 +156,26 @@ __kernel void morph(__global const uchar * srcptr, int src_step, int src_offset, LDS_DAT[point2] = temp1; barrier(CLK_LOCAL_MEM_FENCE); - T res = (T)(VAL); - for (int i = 0, sizey = 2 * RADIUSY + 1; i < sizey; i++) - for (int j = 0, sizex = 2 * RADIUSX + 1; j < sizex; j++) - { - res = -#ifndef RECTKERNEL - mat_kernel[i*(2*RADIUSX+1)+j] ? -#endif - MORPH_OP(res, LDS_DAT[mad24(l_y + i, width, l_x + j)]) -#ifndef RECTKERNEL - : res -#endif - ; - } - if (gidx < cols && gidy < rows) { + T res = (T)(VAL); + PROCESS_ELEMS; + int dst_index = mad24(gidy, dst_step, mad24(gidx, TSIZE, dst_offset)); + +#if defined OP_GRADIENT || defined OP_TOPHAT || defined OP_BLACKHAT + int mat_index = mad24(gidy, mat_step, mad24(gidx, TSIZE, mat_offset)); + T value = loadpix(matptr + mat_index); + +#ifdef OP_GRADIENT + storepix(convertToT(convertToWT(res) - convertToWT(value)), dstptr + dst_index); +#elif defined OP_TOPHAT + storepix(convertToT(convertToWT(value) - convertToWT(res)), dstptr + dst_index); +#elif defined OP_BLACKHAT + storepix(convertToT(convertToWT(res) - convertToWT(value)), dstptr + dst_index); +#endif +#else // erode or dilate storepix(res, dstptr + dst_index); +#endif } } diff --git a/modules/imgproc/test/ocl/test_filters.cpp b/modules/imgproc/test/ocl/test_filters.cpp index 46d77285d..1fe292788 100644 --- a/modules/imgproc/test/ocl/test_filters.cpp +++ b/modules/imgproc/test/ocl/test_filters.cpp @@ -63,7 +63,7 @@ PARAM_TEST_CASE(FilterTestBase, MatType, BorderType, // border type double, // optional parameter bool, // roi or not - int) //width multiplier + int) // width multiplier { int type, borderType, ksize; Size size; @@ -244,8 +244,8 @@ OCL_TEST_P(Erode, Mat) random_roi(); Mat kernel = randomMat(kernelSize, CV_8UC1, 0, 3); - OCL_OFF(cv::erode(src_roi, dst_roi, kernel, Point(-1,-1), iterations) ); - OCL_ON(cv::erode(usrc_roi, udst_roi, kernel, Point(-1,-1), iterations) ); + OCL_OFF(cv::erode(src_roi, dst_roi, kernel, Point(-1, -1), iterations) ); + OCL_ON(cv::erode(usrc_roi, udst_roi, kernel, Point(-1, -1), iterations) ); Near(); } @@ -266,8 +266,8 @@ OCL_TEST_P(Dilate, Mat) random_roi(); Mat kernel = randomMat(kernelSize, CV_8UC1, 0, 3); - OCL_OFF(cv::dilate(src_roi, dst_roi, kernel, Point(-1,-1), iterations) ); - OCL_ON(cv::dilate(usrc_roi, udst_roi, kernel, Point(-1,-1), iterations) ); + OCL_OFF(cv::dilate(src_roi, dst_roi, kernel, Point(-1, -1), iterations) ); + OCL_ON(cv::dilate(usrc_roi, udst_roi, kernel, Point(-1, -1), iterations) ); Near(); } @@ -289,8 +289,8 @@ OCL_TEST_P(MorphologyEx, Mat) random_roi(); Mat kernel = randomMat(kernelSize, CV_8UC1, 0, 3); - OCL_OFF(cv::morphologyEx(src_roi, dst_roi, op, kernel, Point(-1,-1), iterations) ); - OCL_ON(cv::morphologyEx(usrc_roi, udst_roi, op, kernel, Point(-1,-1), iterations) ); + OCL_OFF(cv::morphologyEx(src_roi, dst_roi, op, kernel, Point(-1, -1), iterations) ); + OCL_ON(cv::morphologyEx(usrc_roi, udst_roi, op, kernel, Point(-1, -1), iterations) ); Near(); } @@ -360,8 +360,8 @@ OCL_INSTANTIATE_TEST_CASE_P(Filter, GaussianBlurTest, Combine( OCL_INSTANTIATE_TEST_CASE_P(Filter, Erode, Combine( Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4, CV_64FC1, CV_64FC4), Values(3, 5, 7), - Values(Size(0,0)),//not used - Values((BorderType)BORDER_CONSTANT),//not used + Values(Size(0, 0)), //not used + Values((BorderType)BORDER_CONSTANT), Values(1.0, 2.0, 3.0), Bool(), Values(1))); // not used @@ -369,20 +369,20 @@ OCL_INSTANTIATE_TEST_CASE_P(Filter, Erode, Combine( OCL_INSTANTIATE_TEST_CASE_P(Filter, Dilate, Combine( Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4, CV_64FC1, CV_64FC4), Values(3, 5, 7), - Values(Size(0,0)),//not used - Values((BorderType)BORDER_CONSTANT),//not used + Values(Size(0, 0)), // not used + Values((BorderType)BORDER_CONSTANT), Values(1.0, 2.0, 3.0), Bool(), - Values(1))); //not used + Values(1))); // not used OCL_INSTANTIATE_TEST_CASE_P(Filter, MorphologyEx, Combine( - Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4, CV_64FC1, CV_64FC4), + Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4), Values(3, 5, 7), - Values(Size(0, 0), Size(0, 1), Size(0, 2), Size(0, 3), Size(0, 4), Size(0, 5), Size(0, 6)), // used as generator of operations - Values((BorderType)BORDER_CONSTANT),// not used + Values(Size(0, 2), Size(0, 3), Size(0, 4), Size(0, 5), Size(0, 6)), // used as generator of operations + Values((BorderType)BORDER_CONSTANT), Values(1.0, 2.0, 3.0), Bool(), - Values(1))); //not used + Values(1))); // not used } } // namespace cvtest::ocl