diff --git a/modules/core/include/opencv2/core/ocl.hpp b/modules/core/include/opencv2/core/ocl.hpp index fb9ec24c5605058867908475ce304c64c370e9e6..fdb6f9a0aae522fca8ebc1b814236f0cfebc6e2c 100644 --- a/modules/core/include/opencv2/core/ocl.hpp +++ b/modules/core/include/opencv2/core/ocl.hpp @@ -592,7 +592,7 @@ protected: CV_EXPORTS const char* convertTypeStr(int sdepth, int ddepth, int cn, char* buf); CV_EXPORTS const char* typeToStr(int t); CV_EXPORTS const char* memopTypeToStr(int t); -CV_EXPORTS String kernelToStr(InputArray _kernel, int ddepth = -1); +CV_EXPORTS String kernelToStr(InputArray _kernel, int ddepth = -1, const char * name = NULL); CV_EXPORTS void getPlatfomsInfo(std::vector& platform_info); CV_EXPORTS int predictOptimalVectorWidth(InputArray src1, InputArray src2 = noArray(), InputArray src3 = noArray(), InputArray src4 = noArray(), InputArray src5 = noArray(), InputArray src6 = noArray(), diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp index 7c4f8de9e0dd5026311bad5b07787522f8860b2e..b56f84c16eaa965e07428e832e44d953237fcfe0 100644 --- a/modules/core/src/ocl.cpp +++ b/modules/core/src/ocl.cpp @@ -4306,7 +4306,7 @@ static std::string kerToStr(const Mat & k) return stream.str(); } -String kernelToStr(InputArray _kernel, int ddepth) +String kernelToStr(InputArray _kernel, int ddepth, const char * name) { Mat kernel = _kernel.getMat().reshape(1, 1); @@ -4323,7 +4323,7 @@ String kernelToStr(InputArray _kernel, int ddepth) const func_t func = funcs[depth]; CV_Assert(func != 0); - return cv::format(" -D COEFF=%s", func(kernel).c_str()); + return cv::format(" -D %s=%s", name ? name : "COEFF", func(kernel).c_str()); } #define PROCESS_SRC(src) \ diff --git a/modules/imgproc/perf/opencl/perf_filters.cpp b/modules/imgproc/perf/opencl/perf_filters.cpp index 57b928c289e03ebd144c5565771a183884588654..f7329e31947f9b5d4457c2749d9f88df5ded116f 100644 --- a/modules/imgproc/perf/opencl/perf_filters.cpp +++ b/modules/imgproc/perf/opencl/perf_filters.cpp @@ -211,7 +211,7 @@ OCL_PERF_TEST_P(SobelFixture, Sobel, OCL_TEST_CYCLE() cv::Sobel(src, dst, -1, dx, dy); - SANITY_CHECK(dst); + SANITY_CHECK(dst, 1e-6); } ///////////// Scharr //////////////////////// diff --git a/modules/imgproc/src/filter.cpp b/modules/imgproc/src/filter.cpp index ea0baf6b09028bff3ca299aab87d85d995584a18..bb54471c070fc57a00571391d4c64224fa0053a9 100644 --- a/modules/imgproc/src/filter.cpp +++ b/modules/imgproc/src/filter.cpp @@ -3350,27 +3350,8 @@ static bool ocl_sepRowFilter2D( UMat &src, UMat &buf, Mat &kernelX, int anchor, int radiusY = (int)((buf.rows - src.rows) >> 1); bool isIsolatedBorder = (borderType & BORDER_ISOLATED) != 0; - const char* btype = NULL; - switch (borderType & ~BORDER_ISOLATED) - { - case BORDER_CONSTANT: - btype = "BORDER_CONSTANT"; - break; - case BORDER_REPLICATE: - btype = "BORDER_REPLICATE"; - break; - case BORDER_REFLECT: - btype = "BORDER_REFLECT"; - break; - case BORDER_WRAP: - btype = "BORDER_WRAP"; - break; - case BORDER_REFLECT101: - btype = "BORDER_REFLECT_101"; - break; - default: - return false; - } + const char * const borderMap[] = { "BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", "BORDER_WRAP", "BORDER_REFLECT_101" }, + * const btype = borderMap[borderType & ~BORDER_ISOLATED]; bool extra_extrapolation = src.rows < (int)((-radiusY + globalsize[1]) >> 1) + 1; extra_extrapolation |= src.rows < radiusY; @@ -3463,36 +3444,96 @@ static bool ocl_sepColFilter2D(const UMat &buf, UMat &dst, Mat &kernelY, int anc return kernelCol.run(2, globalsize, localsize, sync); } +const int optimizedSepFilterLocalSize = 16; + +static bool ocl_sepFilter2D_SinglePass(InputArray _src, OutputArray _dst, + InputArray _row_kernel, InputArray _col_kernel, + int borderType, int ddepth) +{ + Size size = _src.size(), wholeSize; + Point origin; + int stype = _src.type(), sdepth = CV_MAT_DEPTH(stype), cn = CV_MAT_CN(stype), + esz = CV_ELEM_SIZE(stype), wdepth = std::max(std::max(sdepth, ddepth), CV_32F), + dtype = CV_MAKE_TYPE(ddepth, cn); + size_t src_step = _src.step(), src_offset = _src.offset(); + bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0; + + if ((src_offset % src_step) % esz != 0 || (!doubleSupport && sdepth == CV_64F) || + !(borderType == BORDER_CONSTANT || borderType == BORDER_REPLICATE || + borderType == BORDER_REFLECT || borderType == BORDER_WRAP || + borderType == BORDER_REFLECT_101)) + return false; + + size_t lt2[2] = { optimizedSepFilterLocalSize, optimizedSepFilterLocalSize }; + size_t gt2[2] = { lt2[0] * (1 + (size.width - 1) / lt2[0]), lt2[1] * (1 + (size.height - 1) / lt2[1]) }; + + char cvt[2][40]; + const char * const borderMap[] = { "BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", "BORDER_WRAP", + "BORDER_REFLECT_101" }; + + String opts = cv::format("-D BLK_X=%d -D BLK_Y=%d -D RADIUSX=%d -D RADIUSY=%d%s%s" + " -D srcT=%s -D convertToWT=%s -D WT=%s -D dstT=%s -D convertToDstT=%s" + " -D %s", (int)lt2[0], (int)lt2[1], _row_kernel.size().height / 2, _col_kernel.size().height / 2, + ocl::kernelToStr(_row_kernel, CV_32F, "KERNEL_MATRIX_X").c_str(), + ocl::kernelToStr(_col_kernel, CV_32F, "KERNEL_MATRIX_Y").c_str(), + ocl::typeToStr(stype), ocl::convertTypeStr(sdepth, wdepth, cn, cvt[0]), + ocl::typeToStr(CV_MAKE_TYPE(wdepth, cn)), ocl::typeToStr(dtype), + ocl::convertTypeStr(wdepth, ddepth, cn, cvt[1]), borderMap[borderType]); + + ocl::Kernel k("sep_filter", ocl::imgproc::filterSep_singlePass_oclsrc, opts); + if (k.empty()) + return false; + + UMat src = _src.getUMat(); + _dst.create(size, dtype); + UMat dst = _dst.getUMat(); + + int src_offset_x = static_cast((src_offset % src_step) / esz); + int src_offset_y = static_cast(src_offset / src_step); + + src.locateROI(wholeSize, origin); + + k.args(ocl::KernelArg::PtrReadOnly(src), (int)src_step, src_offset_x, src_offset_y, + wholeSize.height, wholeSize.width, ocl::KernelArg::WriteOnly(dst)); + + return k.run(2, gt2, lt2, false); +} + static bool ocl_sepFilter2D( InputArray _src, OutputArray _dst, int ddepth, InputArray _kernelX, InputArray _kernelY, Point anchor, double delta, int borderType ) { + Size imgSize = _src.size(); + if (abs(delta)> FLT_MIN) return false; - int type = _src.type(); + int type = _src.type(), cn = CV_MAT_CN(type); if ( !( (type == CV_8UC1 || type == CV_8UC4 || type == CV_32FC1 || type == CV_32FC4) && (ddepth == CV_32F || ddepth == CV_16S || ddepth == CV_8U || ddepth < 0) ) ) return false; - int cn = CV_MAT_CN(type); - Mat kernelX = _kernelX.getMat().reshape(1, 1); - if (1 != (kernelX.cols % 2)) + if (kernelX.cols % 2 != 1) return false; Mat kernelY = _kernelY.getMat().reshape(1, 1); - if (1 != (kernelY.cols % 2)) + if (kernelY.cols % 2 != 1) return false; int sdepth = CV_MAT_DEPTH(type); - if( anchor.x < 0 ) + if (anchor.x < 0) anchor.x = kernelX.cols >> 1; - if( anchor.y < 0 ) + if (anchor.y < 0) anchor.y = kernelY.cols >> 1; - if( ddepth < 0 ) + if (ddepth < 0) ddepth = sdepth; + CV_OCL_RUN_(kernelY.rows <= 21 && kernelX.rows <= 21 && + imgSize.width > optimizedSepFilterLocalSize + (kernelX.rows >> 1) && + imgSize.height > optimizedSepFilterLocalSize + (kernelY.rows >> 1), + ocl_sepFilter2D_SinglePass(_src, _dst, _kernelX, _kernelY, borderType, ddepth), true) + UMat src = _src.getUMat(); Size srcWholeSize; Point srcOffset; src.locateROI(srcWholeSize, srcOffset); diff --git a/modules/imgproc/src/opencl/filterSep_singlePass.cl b/modules/imgproc/src/opencl/filterSep_singlePass.cl new file mode 100644 index 0000000000000000000000000000000000000000..7284da0cbcaa975086dbd666ef8f51440c86a236 --- /dev/null +++ b/modules/imgproc/src/opencl/filterSep_singlePass.cl @@ -0,0 +1,177 @@ +/*M/////////////////////////////////////////////////////////////////////////////////////// +// +// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. +// +// By downloading, copying, installing or using the software you agree to this license. +// If you do not agree to this license, do not download, install, +// copy or use the software. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2014, Intel Corporation, all rights reserved. +// Third party copyrights are property of their respective owners. +// +// Redistribution and use in source and binary forms, with or without modification, +// are permitted provided that the following conditions are met: +// +// * Redistribution's of source code must retain the above copyright notice, +// this list of conditions and the following disclaimer. +// +// * Redistribution's in binary form must reproduce the above copyright notice, +// this list of conditions and the following disclaimer in the documentation +// and/or other materials provided with the distribution. +// +// * The name of the copyright holders may not be used to endorse or promote products +// derived from this software without specific prior written permission. +// +// This software is provided by the copyright holders and contributors "as is" and +// any express or implied warranties, including, but not limited to, the implied +// warranties of merchantability and fitness for a particular purpose are disclaimed. +// In no event shall the Intel Corporation or contributors be liable for any direct, +// indirect, incidental, special, exemplary, or consequential damages +// (including, but not limited to, procurement of substitute goods or services; +// loss of use, data, or profits; or business interruption) however caused +// and on any theory of liability, whether in contract, strict liability, +// or tort (including negligence or otherwise) arising in any way out of +// the use of this software, even if advised of the possibility of such damage. +// +//M*/ + +/////////////////////////////////////////////////////////////////////////////////////////////////// +/////////////////////////////////Macro for border type//////////////////////////////////////////// +///////////////////////////////////////////////////////////////////////////////////////////////// + +#ifdef BORDER_CONSTANT +// CCCCCC|abcdefgh|CCCCCCC +#define EXTRAPOLATE(x, maxV) +#elif defined BORDER_REPLICATE +// aaaaaa|abcdefgh|hhhhhhh +#define EXTRAPOLATE(x, maxV) \ + { \ + (x) = max(min((x), (maxV) - 1), 0); \ + } +#elif defined BORDER_WRAP +// cdefgh|abcdefgh|abcdefg +#define EXTRAPOLATE(x, maxV) \ + { \ + (x) = ( (x) + (maxV) ) % (maxV); \ + } +#elif defined BORDER_REFLECT +// fedcba|abcdefgh|hgfedcb +#define EXTRAPOLATE(x, maxV) \ + { \ + (x) = min(((maxV)-1)*2-(x)+1, max((x),-(x)-1) ); \ + } +#elif defined BORDER_REFLECT_101 || defined BORDER_REFLECT101 +// gfedcb|abcdefgh|gfedcba +#define EXTRAPOLATE(x, maxV) \ + { \ + (x) = min(((maxV)-1)*2-(x), max((x),-(x)) ); \ + } +#else +#error No extrapolation method +#endif + +#define SRC(_x,_y) convertToWT(((global srcT*)(Src+(_y)*src_step))[_x]) + +#ifdef BORDER_CONSTANT +// CCCCCC|abcdefgh|CCCCCCC +#define ELEM(_x,_y,r_edge,t_edge,const_v) (_x)<0 | (_x) >= (r_edge) | (_y)<0 | (_y) >= (t_edge) ? (const_v) : SRC((_x),(_y)) +#else +#define ELEM(_x,_y,r_edge,t_edge,const_v) SRC((_x),(_y)) +#endif + +#define DST(_x,_y) (((global dstT*)(Dst+dst_offset+(_y)*dst_step))[_x]) + +#define noconvert + +// horizontal and vertical filter kernels +// should be defined on host during compile time to avoid overhead +#define DIG(a) a, +__constant float mat_kernelX[] = { KERNEL_MATRIX_X }; +__constant float mat_kernelY[] = { KERNEL_MATRIX_Y }; + +__kernel void sep_filter(__global uchar* Src, int src_step, int srcOffsetX, int srcOffsetY, int height, int width, + __global uchar* Dst, int dst_step, int dst_offset, int dst_rows, int dst_cols) +{ + // RADIUSX, RADIUSY are filter dimensions + // BLK_X, BLK_Y are local wrogroup sizes + // all these should be defined on host during compile time + // first lsmem array for source pixels used in first pass, + // second lsmemDy for storing first pass results + __local WT lsmem[BLK_Y+2*RADIUSY][BLK_X+2*RADIUSX]; + __local WT lsmemDy[BLK_Y][BLK_X+2*RADIUSX]; + + // get local and global ids - used as image and local memory array indexes + int lix = get_local_id(0); + int liy = get_local_id(1); + + int x = (int)get_global_id(0); + int y = (int)get_global_id(1); + + // calculate pixel position in source image taking image offset into account + int srcX = x + srcOffsetX - RADIUSX; + int srcY = y + srcOffsetY - RADIUSY; + int xb = srcX; + int yb = srcY; + + // extrapolate coordinates, if needed + // and read my own source pixel into local memory + // with account for extra border pixels, which will be read by starting workitems + int clocY = liy; + int cSrcY = srcY; + do + { + int yb = cSrcY; + EXTRAPOLATE(yb, (height)); + + int clocX = lix; + int cSrcX = srcX; + do + { + int xb = cSrcX; + EXTRAPOLATE(xb,(width)); + lsmem[clocY][clocX] = ELEM(xb, yb, (width), (height), 0 ); + + clocX += BLK_X; + cSrcX += BLK_X; + } + while(clocX < BLK_X+(RADIUSX*2)); + + clocY += BLK_Y; + cSrcY += BLK_Y; + } + while (clocY < BLK_Y+(RADIUSY*2)); + barrier(CLK_LOCAL_MEM_FENCE); + + // do vertical filter pass + // and store intermediate results to second local memory array + int i, clocX = lix; + WT sum = 0.0f; + do + { + sum = 0.0f; + for (i=0; i<=2*RADIUSY; i++) + sum = mad(lsmem[liy+i][clocX], mat_kernelY[i], sum); + lsmemDy[liy][clocX] = sum; + clocX += BLK_X; + } + while(clocX < BLK_X+(RADIUSX*2)); + barrier(CLK_LOCAL_MEM_FENCE); + + // if this pixel happened to be out of image borders because of global size rounding, + // then just return + if( x >= dst_cols || y >=dst_rows ) + return; + + // do second horizontal filter pass + // and calculate final result + sum = 0.0f; + for (i=0; i<=2*RADIUSX; i++) + sum = mad(lsmemDy[liy][lix+i], mat_kernelX[i], sum); + + //store result into destination image + DST(x,y) = convertToDstT(sum); +}