From 82e6edfba28c91f80fbc20557a3f16906db95bc5 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Mon, 17 Mar 2014 19:59:35 +0400 Subject: [PATCH 1/5] optimized sep filter --- modules/core/include/opencv2/core/ocl.hpp | 2 +- modules/core/src/ocl.cpp | 4 +- modules/imgproc/perf/opencl/perf_filters.cpp | 2 +- modules/imgproc/src/filter.cpp | 99 +++++++--- .../src/opencl/filterSep_singlePass.cl | 177 ++++++++++++++++++ 5 files changed, 251 insertions(+), 33 deletions(-) create mode 100644 modules/imgproc/src/opencl/filterSep_singlePass.cl diff --git a/modules/core/include/opencv2/core/ocl.hpp b/modules/core/include/opencv2/core/ocl.hpp index fb9ec24c56..fdb6f9a0aa 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 7c4f8de9e0..b56f84c16e 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 57b928c289..f7329e3194 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 ea0baf6b09..bb54471c07 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 0000000000..7284da0cbc --- /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); +} From b449b0bf71cf7fb69d8dd7a61ce43ebd5d239f91 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Wed, 19 Mar 2014 15:59:00 +0400 Subject: [PATCH 2/5] simplified cv::sepFilter2D OpenCL part --- modules/imgproc/src/filter.cpp | 133 +++--- modules/imgproc/src/opencl/filterSepCol.cl | 62 +-- modules/imgproc/src/opencl/filterSepRow.cl | 377 ++++-------------- .../src/opencl/filterSep_singlePass.cl | 11 +- modules/imgproc/test/ocl/test_sepfilter2D.cpp | 4 +- 5 files changed, 167 insertions(+), 420 deletions(-) diff --git a/modules/imgproc/src/filter.cpp b/modules/imgproc/src/filter.cpp index bb54471c07..ba2e347af0 100644 --- a/modules/imgproc/src/filter.cpp +++ b/modules/imgproc/src/filter.cpp @@ -41,6 +41,7 @@ //M*/ #include "precomp.hpp" +#define CV_OPENCL_RUN_ASSERT #include "opencl_kernels.hpp" #include @@ -3317,11 +3318,9 @@ static bool ocl_filter2D( InputArray _src, OutputArray _dst, int ddepth, return kernel.run(2, globalsize, localsize, true); } -static bool ocl_sepRowFilter2D( UMat &src, UMat &buf, Mat &kernelX, int anchor, int borderType, bool sync) +static bool ocl_sepRowFilter2D( UMat &src, UMat &buf, Mat &kernelX, int anchor, int borderType) { - int type = src.type(); - int cn = CV_MAT_CN(type); - int sdepth = CV_MAT_DEPTH(type); + int type = src.type(), cn = CV_MAT_CN(type), sdepth = CV_MAT_DEPTH(type); Size bufSize = buf.size(); #ifdef ANDROID @@ -3329,27 +3328,14 @@ static bool ocl_sepRowFilter2D( UMat &src, UMat &buf, Mat &kernelX, int anchor, #else size_t localsize[2] = {16, 16}; #endif + size_t globalsize[2] = {DIVUP(bufSize.width, localsize[0]) * localsize[0], DIVUP(bufSize.height, localsize[1]) * localsize[1]}; - if (CV_8U == sdepth) - { - switch (cn) - { - case 1: - globalsize[0] = DIVUP((bufSize.width + 3) >> 2, localsize[0]) * localsize[0]; - break; - case 2: - globalsize[0] = DIVUP((bufSize.width + 1) >> 1, localsize[0]) * localsize[0]; - break; - case 4: - globalsize[0] = DIVUP(bufSize.width, localsize[0]) * localsize[0]; - break; - } - } + if (type == CV_8UC1) + globalsize[0] = DIVUP((bufSize.width + 3) >> 2, localsize[0]) * localsize[0]; - int radiusX = anchor; - int radiusY = (int)((buf.rows - src.rows) >> 1); + int radiusX = anchor, radiusY = (buf.rows - src.rows) >> 1; - bool isIsolatedBorder = (borderType & BORDER_ISOLATED) != 0; + bool isolated = (borderType & BORDER_ISOLATED) != 0; const char * const borderMap[] = { "BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_REFLECT", "BORDER_WRAP", "BORDER_REFLECT_101" }, * const btype = borderMap[borderType & ~BORDER_ISOLATED]; @@ -3358,49 +3344,38 @@ static bool ocl_sepRowFilter2D( UMat &src, UMat &buf, Mat &kernelX, int anchor, extra_extrapolation |= src.cols < (int)((-radiusX + globalsize[0] + 8 * localsize[0] + 3) >> 1) + 1; extra_extrapolation |= src.cols < radiusX; - cv::String build_options = cv::format("-D RADIUSX=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d -D %s -D %s -D %s", - radiusX, (int)localsize[0], (int)localsize[1], cn, - btype, - extra_extrapolation ? "EXTRA_EXTRAPOLATION" : "NO_EXTRA_EXTRAPOLATION", - isIsolatedBorder ? "BORDER_ISOLATED" : "NO_BORDER_ISOLATED"); + char cvt[40]; + cv::String build_options = cv::format("-D RADIUSX=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d -D %s -D %s -D %s" + " -D srcT=%s -D dstT=%s -D convertToDstT=%s -D srcT1=%s -D dstT1=%s", + radiusX, (int)localsize[0], (int)localsize[1], cn, btype, + extra_extrapolation ? "EXTRA_EXTRAPOLATION" : "NO_EXTRA_EXTRAPOLATION", + isolated ? "BORDER_ISOLATED" : "NO_BORDER_ISOLATED", + ocl::typeToStr(type), ocl::typeToStr(CV_32FC(cn)), + ocl::convertTypeStr(sdepth, CV_32F, cn, cvt), + ocl::typeToStr(sdepth), ocl::typeToStr(CV_32F)); build_options += ocl::kernelToStr(kernelX, CV_32F); Size srcWholeSize; Point srcOffset; src.locateROI(srcWholeSize, srcOffset); - std::stringstream strKernel; - strKernel << "row_filter"; - if (-1 != cn) - strKernel << "_C" << cn; - if (-1 != sdepth) - strKernel << "_D" << sdepth; + String kernelName("row_filter"); + if (type == CV_8UC1) + kernelName += "_C1_D0"; - ocl::Kernel kernelRow; - if (!kernelRow.create(strKernel.str().c_str(), cv::ocl::imgproc::filterSepRow_oclsrc, - build_options)) + ocl::Kernel k(kernelName.c_str(), cv::ocl::imgproc::filterSepRow_oclsrc, + build_options); + if (k.empty()) return false; - int idxArg = 0; - idxArg = kernelRow.set(idxArg, ocl::KernelArg::PtrReadOnly(src)); - idxArg = kernelRow.set(idxArg, (int)(src.step / src.elemSize())); + k.args(ocl::KernelArg::PtrReadOnly(src), (int)(src.step / src.elemSize()), srcOffset.x, + srcOffset.y, src.cols, src.rows, srcWholeSize.width, srcWholeSize.height, + ocl::KernelArg::PtrWriteOnly(buf), (int)(buf.step / buf.elemSize()), + buf.cols, buf.rows, radiusY); - idxArg = kernelRow.set(idxArg, srcOffset.x); - idxArg = kernelRow.set(idxArg, srcOffset.y); - idxArg = kernelRow.set(idxArg, src.cols); - idxArg = kernelRow.set(idxArg, src.rows); - idxArg = kernelRow.set(idxArg, srcWholeSize.width); - idxArg = kernelRow.set(idxArg, srcWholeSize.height); - - idxArg = kernelRow.set(idxArg, ocl::KernelArg::PtrWriteOnly(buf)); - idxArg = kernelRow.set(idxArg, (int)(buf.step / buf.elemSize())); - idxArg = kernelRow.set(idxArg, buf.cols); - idxArg = kernelRow.set(idxArg, buf.rows); - idxArg = kernelRow.set(idxArg, radiusY); - - return kernelRow.run(2, globalsize, localsize, sync); + return k.run(2, globalsize, localsize, false); } -static bool ocl_sepColFilter2D(const UMat &buf, UMat &dst, Mat &kernelY, int anchor, bool sync) +static bool ocl_sepColFilter2D(const UMat &buf, UMat &dst, Mat &kernelY, int anchor) { #ifdef ANDROID size_t localsize[2] = {16, 10}; @@ -3420,28 +3395,23 @@ static bool ocl_sepColFilter2D(const UMat &buf, UMat &dst, Mat &kernelY, int anc globalsize[0] = DIVUP(sz.width, localsize[0]) * localsize[0]; char cvt[40]; - cv::String build_options = cv::format("-D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d -D GENTYPE_SRC=%s -D GENTYPE_DST=%s -D convert_to_DST=%s", - anchor, (int)localsize[0], (int)localsize[1], cn, ocl::typeToStr(buf.type()), - ocl::typeToStr(dtype), ocl::convertTypeStr(CV_32F, ddepth, cn, cvt)); + cv::String build_options = cv::format("-D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d" + " -D srcT=%s -D dstT=%s -D convertToDstT=%s", + anchor, (int)localsize[0], (int)localsize[1], cn, + ocl::typeToStr(buf.type()), ocl::typeToStr(dtype), + ocl::convertTypeStr(CV_32F, ddepth, cn, cvt)); build_options += ocl::kernelToStr(kernelY, CV_32F); - ocl::Kernel kernelCol; - if (!kernelCol.create("col_filter", cv::ocl::imgproc::filterSepCol_oclsrc, build_options)) + ocl::Kernel k("col_filter", cv::ocl::imgproc::filterSepCol_oclsrc, + build_options); + if (k.empty()) return false; - int idxArg = 0; - idxArg = kernelCol.set(idxArg, ocl::KernelArg::PtrReadOnly(buf)); - idxArg = kernelCol.set(idxArg, (int)(buf.step / buf.elemSize())); - idxArg = kernelCol.set(idxArg, buf.cols); - idxArg = kernelCol.set(idxArg, buf.rows); + k.args(ocl::KernelArg::PtrReadOnly(buf), (int)(buf.step / buf.elemSize()), buf.cols, + buf.rows, ocl::KernelArg::PtrWriteOnly(dst), (int)(dst.offset / dst.elemSize()), + (int)(dst.step / dst.elemSize()), dst.cols, dst.rows); - idxArg = kernelCol.set(idxArg, ocl::KernelArg::PtrWriteOnly(dst)); - idxArg = kernelCol.set(idxArg, (int)(dst.offset / dst.elemSize())); - idxArg = kernelCol.set(idxArg, (int)(dst.step / dst.elemSize())); - idxArg = kernelCol.set(idxArg, dst.cols); - idxArg = kernelCol.set(idxArg, dst.rows); - - return kernelCol.run(2, globalsize, localsize, sync); + return k.run(2, globalsize, localsize, false); } const int optimizedSepFilterLocalSize = 16; @@ -3473,12 +3443,14 @@ static bool ocl_sepFilter2D_SinglePass(InputArray _src, OutputArray _dst, 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, + " -D %s -D srcT1=%s -D dstT1=%s -D cn=%d", (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::convertTypeStr(wdepth, ddepth, cn, cvt[1]), borderMap[borderType], + ocl::typeToStr(sdepth), ocl::typeToStr(ddepth), cn); ocl::Kernel k("sep_filter", ocl::imgproc::filterSep_singlePass_oclsrc, opts); if (k.empty()) @@ -3529,10 +3501,13 @@ static bool ocl_sepFilter2D( InputArray _src, OutputArray _dst, int ddepth, 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) +// printf("%d %d\n", imgSize.width, optimizedSepFilterLocalSize + (kernelX.rows >> 1)); +// printf("%d %d\n", imgSize.height, optimizedSepFilterLocalSize + (kernelY.rows >> 1)); + +// 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; @@ -3546,12 +3521,12 @@ static bool ocl_sepFilter2D( InputArray _src, OutputArray _dst, int ddepth, Size srcSize = src.size(); Size bufSize(srcSize.width, srcSize.height + kernelY.cols - 1); UMat buf; buf.create(bufSize, CV_MAKETYPE(CV_32F, cn)); - if (!ocl_sepRowFilter2D(src, buf, kernelX, anchor.x, borderType, false)) + if (!ocl_sepRowFilter2D(src, buf, kernelX, anchor.x, borderType)) return false; _dst.create(srcSize, CV_MAKETYPE(ddepth, cn)); UMat dst = _dst.getUMat(); - return ocl_sepColFilter2D(buf, dst, kernelY, anchor.y, false); + return ocl_sepColFilter2D(buf, dst, kernelY, anchor.y); } #endif diff --git a/modules/imgproc/src/opencl/filterSepCol.cl b/modules/imgproc/src/opencl/filterSepCol.cl index 30a2221cf1..05717c6ad2 100644 --- a/modules/imgproc/src/opencl/filterSepCol.cl +++ b/modules/imgproc/src/opencl/filterSepCol.cl @@ -36,16 +36,6 @@ #define READ_TIMES_COL ((2*(RADIUSY+LSIZE1)-1)/LSIZE1) #define RADIUS 1 -#if CN ==1 -#define ALIGN (((RADIUS)+3)>>2<<2) -#elif CN==2 -#define ALIGN (((RADIUS)+1)>>1<<1) -#elif CN==3 -#define ALIGN (((RADIUS)+3)>>2<<2) -#elif CN==4 -#define ALIGN (RADIUS) -#define READ_TIMES_ROW ((2*(RADIUS+LSIZE0)-1)/LSIZE0) -#endif #define noconvert @@ -65,16 +55,8 @@ The info above maybe obsolete. #define DIG(a) a, __constant float mat_kernel[] = { COEFF }; -__kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void col_filter - (__global const GENTYPE_SRC * restrict src, - const int src_step_in_pixel, - const int src_whole_cols, - const int src_whole_rows, - __global GENTYPE_DST * dst, - const int dst_offset_in_pixel, - const int dst_step_in_pixel, - const int dst_cols, - const int dst_rows) +__kernel void col_filter(__global const srcT * src, int src_step_in_pixel, int src_whole_cols, int src_whole_rows, + __global dstT * dst, int dst_offset_in_pixel, int dst_step_in_pixel, int dst_cols, int dst_rows) { int x = get_global_id(0); int y = get_global_id(1); @@ -85,35 +67,35 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void col_filter int start_addr = mad24(y, src_step_in_pixel, x); int end_addr = mad24(src_whole_rows - 1, src_step_in_pixel, src_whole_cols); - int i; - GENTYPE_SRC sum, temp[READ_TIMES_COL]; - __local GENTYPE_SRC LDS_DAT[LSIZE1 * READ_TIMES_COL][LSIZE0 + 1]; + srcT sum, temp[READ_TIMES_COL]; + __local srcT LDS_DAT[LSIZE1 * READ_TIMES_COL][LSIZE0 + 1]; - //read pixels from src - for(i = 0;i>2<<2) -#elif CN==2 -#define ALIGN (((RADIUS)+1)>>1<<1) -#elif CN==3 -#define ALIGN (((RADIUS)+3)>>2<<2) -#elif CN==4 -#define ALIGN (RADIUS) -#endif #ifdef BORDER_REPLICATE -//BORDER_REPLICATE: aaaaaa|abcdefgh|hhhhhhh +// BORDER_REPLICATE: aaaaaa|abcdefgh|hhhhhhh #define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? (l_edge) : (i)) #define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? (r_edge)-1 : (addr)) #endif #ifdef BORDER_REFLECT -//BORDER_REFLECT: fedcba|abcdefgh|hgfedcb +// BORDER_REFLECT: fedcba|abcdefgh|hgfedcb #define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? -(i)-1 : (i)) #define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? -(i)-1+((r_edge)<<1) : (addr)) #endif #ifdef BORDER_REFLECT_101 -//BORDER_REFLECT_101: gfedcb|abcdefgh|gfedcba +// BORDER_REFLECT_101: gfedcb|abcdefgh|gfedcba #define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? -(i) : (i)) #define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? -(i)-2+((r_edge)<<1) : (addr)) #endif -//blur function does not support BORDER_WRAP #ifdef BORDER_WRAP -//BORDER_WRAP: cdefgh|abcdefgh|abcdefg +// BORDER_WRAP: cdefgh|abcdefgh|abcdefg #define ADDR_L(i, l_edge, r_edge) ((i) < (l_edge) ? (i)+(r_edge) : (i)) #define ADDR_R(i, r_edge, addr) ((i) >= (r_edge) ? (i)-(r_edge) : (addr)) #endif @@ -127,65 +115,56 @@ #endif //BORDER_CONSTANT #endif //EXTRA_EXTRAPOLATION -/********************************************************************************** -These kernels are written for separable filters such as Sobel, Scharr, GaussianBlur. -Now(6/29/2011) the kernels only support 8U data type and the anchor of the convovle -kernel must be in the center. ROI is not supported either. -For channels =1,2,4, each kernels read 4 elements(not 4 pixels), and for channels =3, -the kernel read 4 pixels, save them to LDS and read the data needed from LDS to -calculate the result. -The length of the convovle kernel supported is related to the LSIZE0 and the MAX size -of LDS, which is HW related. -For channels = 1,3 the RADIUS is no more than LSIZE0*2 -For channels = 2, the RADIUS is no more than LSIZE0 -For channels = 4, arbitary RADIUS is supported unless the LDS is not enough -Niko -6/29/2011 -The info above maybe obsolete. -***********************************************************************************/ +#define noconvert + +#if cn != 3 +#define loadpix(addr) *(__global const srcT *)(addr) +#define storepix(val, addr) *(__global dstT *)(addr) = val +#define SRCSIZE ((int)sizeof(srcT)) +#define DSTSIZE ((int)sizeof(dstT)) +#else +#define loadpix(addr) vload3(0, (__global const srcT1 *)(addr)) +#define storepix(val, addr) vstore3(val, 0, (__global dstT1 *)(addr)) +#define SRCSIZE ((int)sizeof(srcT1)*3) +#define DSTSIZE ((int)sizeof(dstT1)*3) +#endif #define DIG(a) a, __constant float mat_kernel[] = { COEFF }; -__kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_C1_D0 - (__global uchar * restrict src, - int src_step_in_pixel, - int src_offset_x, int src_offset_y, - int src_cols, int src_rows, - int src_whole_cols, int src_whole_rows, - __global float * dst, - int dst_step_in_pixel, - int dst_cols, int dst_rows, - int radiusy) +__kernel void row_filter_C1_D0(__global const uchar * src, int src_step_in_pixel, int src_offset_x, int src_offset_y, + int src_cols, int src_rows, int src_whole_cols, int src_whole_rows, + __global float * dst, int dst_step_in_pixel, int dst_cols, int dst_rows, + int radiusy) { int x = get_global_id(0)<<2; int y = get_global_id(1); int l_x = get_local_id(0); int l_y = get_local_id(1); - int start_x = x+src_offset_x - RADIUSX & 0xfffffffc; + int start_x = x + src_offset_x - RADIUSX & 0xfffffffc; int offset = src_offset_x - RADIUSX & 3; int start_y = y + src_offset_y - radiusy; int start_addr = mad24(start_y, src_step_in_pixel, start_x); - int i; + float4 sum; uchar4 temp[READ_TIMES_ROW]; - __local uchar4 LDS_DAT[LSIZE1][READ_TIMES_ROW*LSIZE0+1]; + __local uchar4 LDS_DAT[LSIZE1][READ_TIMES_ROW * LSIZE0 + 1]; #ifdef BORDER_CONSTANT int end_addr = mad24(src_whole_rows - 1, src_step_in_pixel, src_whole_cols); // read pixels from src - for (i = 0; i < READ_TIMES_ROW; i++) + for (int i = 0; i < READ_TIMES_ROW; ++i) { - int current_addr = start_addr+i*LSIZE0*4; - current_addr = ((current_addr < end_addr) && (current_addr > 0)) ? current_addr : 0; - temp[i] = *(__global uchar4*)&src[current_addr]; + int current_addr = mad24(i, LSIZE0 << 2, start_addr); + current_addr = current_addr < end_addr && current_addr > 0 ? current_addr : 0; + temp[i] = *(__global const uchar4 *)&src[current_addr]; } // judge if read out of boundary #ifdef BORDER_ISOLATED - for (i = 0; isrc_whole_cols)| (start_y<0) | (start_y >= src_whole_rows); #endif - int4 index[READ_TIMES_ROW]; - int4 addr; + int4 index[READ_TIMES_ROW], addr; int s_y; if (not_all_in_range) { // judge if read out of boundary - for (i = 0; i < READ_TIMES_ROW; i++) + for (int i = 0; i < READ_TIMES_ROW; ++i) { - index[i] = (int4)(start_x+i*LSIZE0*4) + (int4)(0, 1, 2, 3); + index[i] = (int4)(mad24(i, LSIZE0 << 2, start_x)) + (int4)(0, 1, 2, 3); #ifdef BORDER_ISOLATED EXTRAPOLATE(index[i].x, src_offset_x, src_offset_x + src_cols); EXTRAPOLATE(index[i].y, src_offset_x, src_offset_x + src_cols); @@ -231,6 +209,7 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_ EXTRAPOLATE(index[i].w, 0, src_whole_cols); #endif } + s_y = start_y; #ifdef BORDER_ISOLATED EXTRAPOLATE(s_y, src_offset_y, src_offset_y + src_rows); @@ -239,9 +218,9 @@ __kernel __attribute__((reqd_work_group_size(LSIZE0,LSIZE1,1))) void row_filter_ #endif // read pixels from src - for (i = 0; i 0)) ? current_addr : 0; + int current_addr = mad24(i, LSIZE0, start_addr); + current_addr = current_addr < end_addr && current_addr > 0 ? current_addr : 0; temp[i] = src[current_addr]; } - //judge if read out of boundary + // judge if read out of boundary #ifdef BORDER_ISOLATED - for (i = 0; i 0)) ? current_addr : 0; - temp[i] = src[current_addr]; - } - - // judge if read out of boundary -#ifdef BORDER_ISOLATED - for (i = 0; i 0)) ? current_addr : 0; - temp[i] = src[current_addr]; - } - - // judge if read out of boundary -#ifdef BORDER_ISOLATED - for (i = 0; i Date: Wed, 19 Mar 2014 18:49:33 +0400 Subject: [PATCH 3/5] generalized OpenCL version of cv::sepFilter2D; removed some restrictions and added 3-channels support --- modules/core/src/ocl.cpp | 4 +- modules/imgproc/src/filter.cpp | 86 +++++++++++-------- modules/imgproc/src/opencl/filterSepCol.cl | 47 +++++----- modules/imgproc/src/opencl/filterSepRow.cl | 45 ++++++---- modules/imgproc/test/ocl/test_filters.cpp | 2 +- modules/imgproc/test/ocl/test_sepfilter2D.cpp | 17 ++-- 6 files changed, 112 insertions(+), 89 deletions(-) diff --git a/modules/core/src/ocl.cpp b/modules/core/src/ocl.cpp index b56f84c16e..3a7c718d4f 100644 --- a/modules/core/src/ocl.cpp +++ b/modules/core/src/ocl.cpp @@ -4317,8 +4317,8 @@ String kernelToStr(InputArray _kernel, int ddepth, const char * name) if (ddepth != depth) kernel.convertTo(kernel, ddepth); - typedef std::string (*func_t)(const Mat &); - static const func_t funcs[] = { kerToStr, kerToStr, kerToStr,kerToStr, + typedef std::string (* func_t)(const Mat &); + static const func_t funcs[] = { kerToStr, kerToStr, kerToStr, kerToStr, kerToStr, kerToStr, kerToStr, 0 }; const func_t func = funcs[depth]; CV_Assert(func != 0); diff --git a/modules/imgproc/src/filter.cpp b/modules/imgproc/src/filter.cpp index ba2e347af0..c013a9b16c 100644 --- a/modules/imgproc/src/filter.cpp +++ b/modules/imgproc/src/filter.cpp @@ -41,7 +41,6 @@ //M*/ #include "precomp.hpp" -#define CV_OPENCL_RUN_ASSERT #include "opencl_kernels.hpp" #include @@ -3135,7 +3134,7 @@ template struct Filter2D : public BaseFi // b e h b e h 0 0 // c f i c f i 0 0 template -static int _prepareKernelFilter2D(std::vector& data, const Mat &kernel) +static int _prepareKernelFilter2D(std::vector & data, const Mat & kernel) { Mat _kernel; kernel.convertTo(_kernel, DataDepth::value); int size_y_aligned = ROUNDUP(kernel.rows * 2, 4); @@ -3318,11 +3317,16 @@ static bool ocl_filter2D( InputArray _src, OutputArray _dst, int ddepth, return kernel.run(2, globalsize, localsize, true); } -static bool ocl_sepRowFilter2D( UMat &src, UMat &buf, Mat &kernelX, int anchor, int borderType) +static bool ocl_sepRowFilter2D(const UMat & src, UMat & buf, const Mat & kernelX, int anchor, + int borderType, int ddepth, bool fast8uc1) { int type = src.type(), cn = CV_MAT_CN(type), sdepth = CV_MAT_DEPTH(type); + bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0; Size bufSize = buf.size(); + if (!doubleSupport && (sdepth == CV_64F || ddepth == CV_64F)) + return false; + #ifdef ANDROID size_t localsize[2] = {16, 10}; #else @@ -3330,7 +3334,7 @@ static bool ocl_sepRowFilter2D( UMat &src, UMat &buf, Mat &kernelX, int anchor, #endif size_t globalsize[2] = {DIVUP(bufSize.width, localsize[0]) * localsize[0], DIVUP(bufSize.height, localsize[1]) * localsize[1]}; - if (type == CV_8UC1) + if (fast8uc1) globalsize[0] = DIVUP((bufSize.width + 3) >> 2, localsize[0]) * localsize[0]; int radiusX = anchor, radiusY = (buf.rows - src.rows) >> 1; @@ -3346,20 +3350,21 @@ static bool ocl_sepRowFilter2D( UMat &src, UMat &buf, Mat &kernelX, int anchor, char cvt[40]; cv::String build_options = cv::format("-D RADIUSX=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d -D %s -D %s -D %s" - " -D srcT=%s -D dstT=%s -D convertToDstT=%s -D srcT1=%s -D dstT1=%s", + " -D srcT=%s -D dstT=%s -D convertToDstT=%s -D srcT1=%s -D dstT1=%s%s", radiusX, (int)localsize[0], (int)localsize[1], cn, btype, extra_extrapolation ? "EXTRA_EXTRAPOLATION" : "NO_EXTRA_EXTRAPOLATION", isolated ? "BORDER_ISOLATED" : "NO_BORDER_ISOLATED", ocl::typeToStr(type), ocl::typeToStr(CV_32FC(cn)), ocl::convertTypeStr(sdepth, CV_32F, cn, cvt), - ocl::typeToStr(sdepth), ocl::typeToStr(CV_32F)); + ocl::typeToStr(sdepth), ocl::typeToStr(CV_32F), + doubleSupport ? " -D DOUBLE_SUPPORT" : ""); build_options += ocl::kernelToStr(kernelX, CV_32F); Size srcWholeSize; Point srcOffset; src.locateROI(srcWholeSize, srcOffset); String kernelName("row_filter"); - if (type == CV_8UC1) + if (fast8uc1) kernelName += "_C1_D0"; ocl::Kernel k(kernelName.c_str(), cv::ocl::imgproc::filterSepRow_oclsrc, @@ -3367,39 +3372,47 @@ static bool ocl_sepRowFilter2D( UMat &src, UMat &buf, Mat &kernelX, int anchor, if (k.empty()) return false; - k.args(ocl::KernelArg::PtrReadOnly(src), (int)(src.step / src.elemSize()), srcOffset.x, - srcOffset.y, src.cols, src.rows, srcWholeSize.width, srcWholeSize.height, - ocl::KernelArg::PtrWriteOnly(buf), (int)(buf.step / buf.elemSize()), - buf.cols, buf.rows, radiusY); + if (fast8uc1) + k.args(ocl::KernelArg::PtrReadOnly(src), (int)(src.step / src.elemSize()), srcOffset.x, + srcOffset.y, src.cols, src.rows, srcWholeSize.width, srcWholeSize.height, + ocl::KernelArg::PtrWriteOnly(buf), (int)(buf.step / buf.elemSize()), + buf.cols, buf.rows, radiusY); + else + k.args(ocl::KernelArg::PtrReadOnly(src), (int)src.step, srcOffset.x, + srcOffset.y, src.cols, src.rows, srcWholeSize.width, srcWholeSize.height, + ocl::KernelArg::PtrWriteOnly(buf), (int)buf.step, buf.cols, buf.rows, radiusY); return k.run(2, globalsize, localsize, false); } -static bool ocl_sepColFilter2D(const UMat &buf, UMat &dst, Mat &kernelY, int anchor) +static bool ocl_sepColFilter2D(const UMat & buf, UMat & dst, const Mat & kernelY, int anchor) { + bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0; + if (dst.depth() == CV_64F && !doubleSupport) + return false; + #ifdef ANDROID - size_t localsize[2] = {16, 10}; + size_t localsize[2] = { 16, 10 }; #else - size_t localsize[2] = {16, 16}; + size_t localsize[2] = { 16, 16 }; #endif - size_t globalsize[2] = {0, 0}; + size_t globalsize[2] = { 0, 0 }; int dtype = dst.type(), cn = CV_MAT_CN(dtype), ddepth = CV_MAT_DEPTH(dtype); Size sz = dst.size(); globalsize[1] = DIVUP(sz.height, localsize[1]) * localsize[1]; - - if (dtype == CV_8UC2) - globalsize[0] = DIVUP((sz.width + 1) / 2, localsize[0]) * localsize[0]; - else - globalsize[0] = DIVUP(sz.width, localsize[0]) * localsize[0]; + globalsize[0] = DIVUP(sz.width, localsize[0]) * localsize[0]; char cvt[40]; cv::String build_options = cv::format("-D RADIUSY=%d -D LSIZE0=%d -D LSIZE1=%d -D CN=%d" - " -D srcT=%s -D dstT=%s -D convertToDstT=%s", + " -D srcT=%s -D dstT=%s -D convertToDstT=%s" + " -D srcT1=%s -D dstT1=%s%s", anchor, (int)localsize[0], (int)localsize[1], cn, ocl::typeToStr(buf.type()), ocl::typeToStr(dtype), - ocl::convertTypeStr(CV_32F, ddepth, cn, cvt)); + ocl::convertTypeStr(CV_32F, ddepth, cn, cvt), + ocl::typeToStr(CV_32F), ocl::typeToStr(ddepth), + doubleSupport ? " -D DOUBLE_SUPPORT" : ""); build_options += ocl::kernelToStr(kernelY, CV_32F); ocl::Kernel k("col_filter", cv::ocl::imgproc::filterSepCol_oclsrc, @@ -3407,13 +3420,13 @@ static bool ocl_sepColFilter2D(const UMat &buf, UMat &dst, Mat &kernelY, int anc if (k.empty()) return false; - k.args(ocl::KernelArg::PtrReadOnly(buf), (int)(buf.step / buf.elemSize()), buf.cols, - buf.rows, ocl::KernelArg::PtrWriteOnly(dst), (int)(dst.offset / dst.elemSize()), - (int)(dst.step / dst.elemSize()), dst.cols, dst.rows); + k.args(ocl::KernelArg::ReadOnly(buf), ocl::KernelArg::WriteOnly(dst)); return k.run(2, globalsize, localsize, false); } +#if 0 + const int optimizedSepFilterLocalSize = 16; static bool ocl_sepFilter2D_SinglePass(InputArray _src, OutputArray _dst, @@ -3471,18 +3484,19 @@ static bool ocl_sepFilter2D_SinglePass(InputArray _src, OutputArray _dst, return k.run(2, gt2, lt2, false); } +#endif + static bool ocl_sepFilter2D( InputArray _src, OutputArray _dst, int ddepth, InputArray _kernelX, InputArray _kernelY, Point anchor, double delta, int borderType ) { - Size imgSize = _src.size(); +// Size imgSize = _src.size(); if (abs(delta)> FLT_MIN) return false; 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) ) ) + if (cn > 4) return false; Mat kernelX = _kernelX.getMat().reshape(1, 1); @@ -3501,9 +3515,6 @@ static bool ocl_sepFilter2D( InputArray _src, OutputArray _dst, int ddepth, if (ddepth < 0) ddepth = sdepth; -// printf("%d %d\n", imgSize.width, optimizedSepFilterLocalSize + (kernelX.rows >> 1)); -// printf("%d %d\n", imgSize.height, optimizedSepFilterLocalSize + (kernelY.rows >> 1)); - // CV_OCL_RUN_(kernelY.rows <= 21 && kernelX.rows <= 21 && // imgSize.width > optimizedSepFilterLocalSize + (kernelX.rows >> 1) && // imgSize.height > optimizedSepFilterLocalSize + (kernelY.rows >> 1), @@ -3512,20 +3523,19 @@ static bool ocl_sepFilter2D( InputArray _src, OutputArray _dst, int ddepth, UMat src = _src.getUMat(); Size srcWholeSize; Point srcOffset; src.locateROI(srcWholeSize, srcOffset); - if ( (0 != (srcOffset.x % 4)) || - (0 != (src.cols % 4)) || - (0 != ((src.step / src.elemSize()) % 4)) - ) - return false; + + bool fast8uc1 = type == CV_8UC1 && srcOffset.x % 4 == 0 && + src.cols % 4 == 0 && src.step % 4 == 0; Size srcSize = src.size(); Size bufSize(srcSize.width, srcSize.height + kernelY.cols - 1); - UMat buf; buf.create(bufSize, CV_MAKETYPE(CV_32F, cn)); - if (!ocl_sepRowFilter2D(src, buf, kernelX, anchor.x, borderType)) + UMat buf(bufSize, CV_32FC(cn)); + if (!ocl_sepRowFilter2D(src, buf, kernelX, anchor.x, borderType, ddepth, fast8uc1)) return false; _dst.create(srcSize, CV_MAKETYPE(ddepth, cn)); UMat dst = _dst.getUMat(); + return ocl_sepColFilter2D(buf, dst, kernelY, anchor.y); } diff --git a/modules/imgproc/src/opencl/filterSepCol.cl b/modules/imgproc/src/opencl/filterSepCol.cl index 05717c6ad2..f5d270cf42 100644 --- a/modules/imgproc/src/opencl/filterSepCol.cl +++ b/modules/imgproc/src/opencl/filterSepCol.cl @@ -34,29 +34,36 @@ // // +#ifdef DOUBLE_SUPPORT +#ifdef cl_amd_fp64 +#pragma OPENCL EXTENSION cl_amd_fp64:enable +#elif defined (cl_khr_fp64) +#pragma OPENCL EXTENSION cl_khr_fp64:enable +#endif +#endif + #define READ_TIMES_COL ((2*(RADIUSY+LSIZE1)-1)/LSIZE1) #define RADIUS 1 #define noconvert -/********************************************************************************** -These kernels are written for separable filters such as Sobel, Scharr, GaussianBlur. -Now(6/29/2011) the kernels only support 8U data type and the anchor of the convovle -kernel must be in the center. ROI is not supported either. -Each kernels read 4 elements(not 4 pixels), save them to LDS and read the data needed -from LDS to calculate the result. -The length of the convovle kernel supported is only related to the MAX size of LDS, -which is HW related. -Niko -6/29/2011 -The info above maybe obsolete. -***********************************************************************************/ +#if CN != 3 +#define loadpix(addr) *(__global const srcT *)(addr) +#define storepix(val, addr) *(__global dstT *)(addr) = val +#define SRCSIZE (int)sizeof(srcT) +#define DSTSIZE (int)sizeof(dstT) +#else +#define loadpix(addr) vload3(0, (__global const srcT1 *)(addr)) +#define storepix(val, addr) vstore3(val, 0, (__global dstT1 *)(addr)) +#define SRCSIZE (int)sizeof(srcT1)*3 +#define DSTSIZE (int)sizeof(dstT1)*3 +#endif #define DIG(a) a, __constant float mat_kernel[] = { COEFF }; -__kernel void col_filter(__global const srcT * src, int src_step_in_pixel, int src_whole_cols, int src_whole_rows, - __global dstT * dst, int dst_offset_in_pixel, int dst_step_in_pixel, int dst_cols, int dst_rows) +__kernel void col_filter(__global const uchar * src, int src_step, int src_offset, int src_whole_rows, int src_whole_cols, + __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols) { int x = get_global_id(0); int y = get_global_id(1); @@ -64,8 +71,8 @@ __kernel void col_filter(__global const srcT * src, int src_step_in_pixel, int s int l_x = get_local_id(0); int l_y = get_local_id(1); - int start_addr = mad24(y, src_step_in_pixel, x); - int end_addr = mad24(src_whole_rows - 1, src_step_in_pixel, src_whole_cols); + int start_addr = mad24(y, src_step, x * SRCSIZE); + int end_addr = mad24(src_whole_rows - 1, src_step, src_whole_cols * SRCSIZE); srcT sum, temp[READ_TIMES_COL]; __local srcT LDS_DAT[LSIZE1 * READ_TIMES_COL][LSIZE0 + 1]; @@ -73,9 +80,9 @@ __kernel void col_filter(__global const srcT * src, int src_step_in_pixel, int s // read pixels from src for (int i = 0; i < READ_TIMES_COL; ++i) { - int current_addr = mad24(i, LSIZE1 * src_step_in_pixel, start_addr); + int current_addr = mad24(i, LSIZE1 * src_step, start_addr); current_addr = current_addr < end_addr ? current_addr : 0; - temp[i] = src[current_addr]; + temp[i] = loadpix(src + current_addr); } // save pixels to lds @@ -95,7 +102,7 @@ __kernel void col_filter(__global const srcT * src, int src_step_in_pixel, int s // write the result to dst if (x < dst_cols && y < dst_rows) { - start_addr = mad24(y, dst_step_in_pixel, x + dst_offset_in_pixel); - dst[start_addr] = convertToDstT(sum); + start_addr = mad24(y, dst_step, mad24(DSTSIZE, x, dst_offset)); + storepix(convertToDstT(sum), dst + start_addr); } } diff --git a/modules/imgproc/src/opencl/filterSepRow.cl b/modules/imgproc/src/opencl/filterSepRow.cl index 83968dfc1e..726de448e4 100644 --- a/modules/imgproc/src/opencl/filterSepRow.cl +++ b/modules/imgproc/src/opencl/filterSepRow.cl @@ -34,6 +34,14 @@ // // +#ifdef DOUBLE_SUPPORT +#ifdef cl_amd_fp64 +#pragma OPENCL EXTENSION cl_amd_fp64:enable +#elif defined (cl_khr_fp64) +#pragma OPENCL EXTENSION cl_khr_fp64:enable +#endif +#endif + #define READ_TIMES_ROW ((2*(RADIUSX+LSIZE0)-1)/LSIZE0) //for c4 only #define RADIUS 1 @@ -117,16 +125,16 @@ #define noconvert -#if cn != 3 +#if CN != 3 #define loadpix(addr) *(__global const srcT *)(addr) #define storepix(val, addr) *(__global dstT *)(addr) = val -#define SRCSIZE ((int)sizeof(srcT)) -#define DSTSIZE ((int)sizeof(dstT)) +#define SRCSIZE (int)sizeof(srcT) +#define DSTSIZE (int)sizeof(dstT) #else #define loadpix(addr) vload3(0, (__global const srcT1 *)(addr)) #define storepix(val, addr) vstore3(val, 0, (__global dstT1 *)(addr)) -#define SRCSIZE ((int)sizeof(srcT1)*3) -#define DSTSIZE ((int)sizeof(dstT1)*3) +#define SRCSIZE (int)sizeof(srcT1)*3 +#define DSTSIZE (int)sizeof(dstT1)*3 #endif #define DIG(a) a, @@ -269,32 +277,33 @@ __kernel void row_filter_C1_D0(__global const uchar * src, int src_step_in_pixel dst[start_addr] = sum.x; } -__kernel void row_filter(__global const srcT * src, int src_step_in_pixel, int src_offset_x, int src_offset_y, +__kernel void row_filter(__global const uchar * src, int src_step, int src_offset_x, int src_offset_y, int src_cols, int src_rows, int src_whole_cols, int src_whole_rows, - __global dstT * dst, int dst_step_in_pixel, int dst_cols, int dst_rows, + __global uchar * dst, int dst_step, int dst_cols, int dst_rows, int radiusy) { int x = get_global_id(0); int y = get_global_id(1); int l_x = get_local_id(0); int l_y = get_local_id(1); + int start_x = x + src_offset_x - RADIUSX; int start_y = y + src_offset_y - radiusy; - int start_addr = mad24(start_y, src_step_in_pixel, start_x); + int start_addr = mad24(start_y, src_step, start_x * SRCSIZE); dstT sum; srcT temp[READ_TIMES_ROW]; __local srcT LDS_DAT[LSIZE1][READ_TIMES_ROW * LSIZE0 + 1]; #ifdef BORDER_CONSTANT - int end_addr = mad24(src_whole_rows - 1, src_step_in_pixel, src_whole_cols); + int end_addr = mad24(src_whole_rows - 1, src_step, src_whole_cols * SRCSIZE); // read pixels from src for (int i = 0; i < READ_TIMES_ROW; i++) { - int current_addr = mad24(i, LSIZE0, start_addr); - current_addr = current_addr < end_addr && current_addr > 0 ? current_addr : 0; - temp[i] = src[current_addr]; + int current_addr = mad24(i, LSIZE0 * SRCSIZE, start_addr); + current_addr = current_addr < end_addr && current_addr >= 0 ? current_addr : 0; + temp[i] = loadpix(src + current_addr); } // judge if read out of boundary @@ -312,8 +321,7 @@ __kernel void row_filter(__global const srcT * src, int src_step_in_pixel, int s } #endif #else - int index[READ_TIMES_ROW]; - int s_x, s_y; + int index[READ_TIMES_ROW], s_x, s_y; // judge if read out of boundary for (int i = 0; i < READ_TIMES_ROW; ++i) @@ -328,12 +336,12 @@ __kernel void row_filter(__global const srcT * src, int src_step_in_pixel, int s EXTRAPOLATE(s_x, 0, src_whole_cols); EXTRAPOLATE(s_y, 0, src_whole_rows); #endif - index[i] = mad24(s_y, src_step_in_pixel, s_x); + index[i] = mad24(s_y, src_step, s_x * SRCSIZE); } // read pixels from src for (int i = 0; i < READ_TIMES_ROW; ++i) - temp[i] = src[index[i]]; + temp[i] = loadpix(src + index[i]); #endif // BORDER_CONSTANT // save pixels to lds @@ -349,10 +357,11 @@ __kernel void row_filter(__global const srcT * src, int src_step_in_pixel, int s temp[1] = LDS_DAT[l_y][l_x + RADIUSX + i]; sum += mad(convertToDstT(temp[0]), mat_kernel[RADIUSX - i], convertToDstT(temp[1]) * mat_kernel[RADIUSX + i]); } + // write the result to dst if (x < dst_cols && y < dst_rows) { - start_addr = mad24(y, dst_step_in_pixel, x); - dst[start_addr] = sum; + start_addr = mad24(y, dst_step, x * DSTSIZE); + storepix(sum, dst + start_addr); } } diff --git a/modules/imgproc/test/ocl/test_filters.cpp b/modules/imgproc/test/ocl/test_filters.cpp index fe16fe81d5..04b330527f 100644 --- a/modules/imgproc/test/ocl/test_filters.cpp +++ b/modules/imgproc/test/ocl/test_filters.cpp @@ -312,7 +312,7 @@ OCL_TEST_P(MorphologyEx, Mat) (int)BORDER_REFLECT|BORDER_ISOLATED, (int)BORDER_WRAP|BORDER_ISOLATED, \ (int)BORDER_REFLECT_101|BORDER_ISOLATED*/) // WRAP and ISOLATED are not supported by cv:: version -#define FILTER_TYPES Values(CV_8UC1, CV_8UC2, CV_8UC4, CV_32FC1, CV_32FC4, CV_64FC1, CV_64FC4) +#define FILTER_TYPES Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4) OCL_INSTANTIATE_TEST_CASE_P(Filter, Bilateral, Combine( Values((MatType)CV_8UC1), diff --git a/modules/imgproc/test/ocl/test_sepfilter2D.cpp b/modules/imgproc/test/ocl/test_sepfilter2D.cpp index 09d01d157a..05d46cc2b3 100644 --- a/modules/imgproc/test/ocl/test_sepfilter2D.cpp +++ b/modules/imgproc/test/ocl/test_sepfilter2D.cpp @@ -75,9 +75,9 @@ PARAM_TEST_CASE(SepFilter2D, MatDepth, Channels, BorderType, bool, bool) void random_roi() { Size ksize = randomSize(kernelMinSize, kernelMaxSize); - if (1 != (ksize.width % 2)) + if (1 != ksize.width % 2) ksize.width++; - if (1 != (ksize.height % 2)) + if (1 != ksize.height % 2) ksize.height++; Mat temp = randomMat(Size(ksize.width, 1), CV_MAKE_TYPE(CV_32F, 1), -MAX_VALUE, MAX_VALUE); @@ -86,24 +86,22 @@ PARAM_TEST_CASE(SepFilter2D, MatDepth, Channels, BorderType, bool, bool) cv::normalize(temp, kernelY, 1.0, 0.0, NORM_L1); Size roiSize = randomSize(ksize.width + 16, MAX_VALUE, ksize.height + 20, MAX_VALUE); - std::cout << roiSize << std::endl; int rest = roiSize.width % 4; - if (0 != rest) + if (rest != 0) roiSize.width += (4 - rest); Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); rest = srcBorder.lef % 4; - if (0 != rest) + if (rest != 0) srcBorder.lef += (4 - rest); rest = srcBorder.rig % 4; - if (0 != rest) + if (rest != 0) srcBorder.rig += (4 - rest); randomSubMat(src, src_roi, roiSize, srcBorder, type, -MAX_VALUE, MAX_VALUE); Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); randomSubMat(dst, dst_roi, roiSize, dstBorder, type, -MAX_VALUE, MAX_VALUE); - anchor.x = -1; - anchor.y = -1; + anchor.x = anchor.y = -1; UMAT_UPLOAD_INPUT_PARAMETER(src) UMAT_UPLOAD_OUTPUT_PARAMETER(dst) @@ -128,11 +126,10 @@ OCL_TEST_P(SepFilter2D, Mat) } } - OCL_INSTANTIATE_TEST_CASE_P(ImageProc, SepFilter2D, Combine( Values(CV_8U, CV_32F), - Values(1, 4), + OCL_ALL_CHANNELS, Values( (BorderType)BORDER_CONSTANT, (BorderType)BORDER_REPLICATE, From 63d8a61b9b18db161e0ff94fc7bac0166f0b8e3b Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Wed, 19 Mar 2014 19:12:37 +0400 Subject: [PATCH 4/5] enabled 3-channels support for cv::createSuperResolution_BTVL1 --- modules/imgproc/src/filter.cpp | 15 ++++++--------- modules/imgproc/test/ocl/test_sepfilter2D.cpp | 4 ++-- modules/superres/src/btv_l1.cpp | 4 +--- 3 files changed, 9 insertions(+), 14 deletions(-) diff --git a/modules/imgproc/src/filter.cpp b/modules/imgproc/src/filter.cpp index c013a9b16c..2d7c7404f6 100644 --- a/modules/imgproc/src/filter.cpp +++ b/modules/imgproc/src/filter.cpp @@ -3425,8 +3425,6 @@ static bool ocl_sepColFilter2D(const UMat & buf, UMat & dst, const Mat & kernelY return k.run(2, globalsize, localsize, false); } -#if 0 - const int optimizedSepFilterLocalSize = 16; static bool ocl_sepFilter2D_SinglePass(InputArray _src, OutputArray _dst, @@ -3484,13 +3482,11 @@ static bool ocl_sepFilter2D_SinglePass(InputArray _src, OutputArray _dst, return k.run(2, gt2, lt2, false); } -#endif - static bool ocl_sepFilter2D( InputArray _src, OutputArray _dst, int ddepth, InputArray _kernelX, InputArray _kernelY, Point anchor, double delta, int borderType ) { -// Size imgSize = _src.size(); + Size imgSize = _src.size(); if (abs(delta)> FLT_MIN) return false; @@ -3515,10 +3511,11 @@ static bool ocl_sepFilter2D( InputArray _src, OutputArray _dst, int ddepth, 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) + CV_OCL_RUN_(kernelY.rows <= 21 && kernelX.rows <= 21 && + imgSize.width > optimizedSepFilterLocalSize + (kernelX.rows >> 1) && + imgSize.height > optimizedSepFilterLocalSize + (kernelY.rows >> 1) && + (borderType & BORDER_ISOLATED) != 0, + ocl_sepFilter2D_SinglePass(_src, _dst, _kernelX, _kernelY, borderType, ddepth), true) UMat src = _src.getUMat(); Size srcWholeSize; Point srcOffset; diff --git a/modules/imgproc/test/ocl/test_sepfilter2D.cpp b/modules/imgproc/test/ocl/test_sepfilter2D.cpp index 05d46cc2b3..b724641f45 100644 --- a/modules/imgproc/test/ocl/test_sepfilter2D.cpp +++ b/modules/imgproc/test/ocl/test_sepfilter2D.cpp @@ -85,7 +85,7 @@ PARAM_TEST_CASE(SepFilter2D, MatDepth, Channels, BorderType, bool, bool) temp = randomMat(Size(1, ksize.height), CV_MAKE_TYPE(CV_32F, 1), -MAX_VALUE, MAX_VALUE); cv::normalize(temp, kernelY, 1.0, 0.0, NORM_L1); - Size roiSize = randomSize(ksize.width + 16, MAX_VALUE, ksize.height + 20, MAX_VALUE); + Size roiSize = randomSize(ksize.width, MAX_VALUE, ksize.height, MAX_VALUE); int rest = roiSize.width % 4; if (rest != 0) roiSize.width += (4 - rest); @@ -115,7 +115,7 @@ PARAM_TEST_CASE(SepFilter2D, MatDepth, Channels, BorderType, bool, bool) OCL_TEST_P(SepFilter2D, Mat) { - for (int j = 0; j < test_loop_times; j++) + for (int j = 0; j < test_loop_times + 1; j++) { random_roi(); diff --git a/modules/superres/src/btv_l1.cpp b/modules/superres/src/btv_l1.cpp index 1e4aa48a7d..d54b4b398a 100644 --- a/modules/superres/src/btv_l1.cpp +++ b/modules/superres/src/btv_l1.cpp @@ -1014,10 +1014,8 @@ namespace return; #ifdef HAVE_OPENCL - if (isUmat_ && curFrame_.channels() == 1) + if (isUmat_) curFrame_.copyTo(ucurFrame_); - else - isUmat_ = false; #endif ++storePos_; From 2875ce60ea1c10ed72e393e725c5df6c6238e64f Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Thu, 20 Mar 2014 18:38:31 +0400 Subject: [PATCH 5/5] added 3-channels support to optimized version --- modules/core/include/opencv2/core/mat.hpp | 2 ++ modules/core/src/matrix.cpp | 34 +++++++++++------- modules/imgproc/src/filter.cpp | 35 ++++++++++--------- .../src/opencl/filterSep_singlePass.cl | 17 +++++++-- modules/imgproc/test/ocl/test_sepfilter2D.cpp | 11 +----- 5 files changed, 57 insertions(+), 42 deletions(-) diff --git a/modules/core/include/opencv2/core/mat.hpp b/modules/core/include/opencv2/core/mat.hpp index 33167fab20..6b8368fd57 100644 --- a/modules/core/include/opencv2/core/mat.hpp +++ b/modules/core/include/opencv2/core/mat.hpp @@ -118,6 +118,8 @@ public: virtual int kind() const; virtual int dims(int i=-1) const; + virtual int cols(int i=-1) const; + virtual int rows(int i=-1) const; virtual Size size(int i=-1) const; virtual int sizend(int* sz, int i=-1) const; virtual bool sameSize(const _InputArray& arr) const; diff --git a/modules/core/src/matrix.cpp b/modules/core/src/matrix.cpp index db1ce760f3..0a8d87d91a 100644 --- a/modules/core/src/matrix.cpp +++ b/modules/core/src/matrix.cpp @@ -1416,6 +1416,16 @@ int _InputArray::kind() const return flags & KIND_MASK; } +int _InputArray::rows(int i) const +{ + return size(i).height; +} + +int _InputArray::cols(int i) const +{ + return size(i).width; +} + Size _InputArray::size(int i) const { int k = kind(); @@ -2078,45 +2088,45 @@ void _OutputArray::create(Size _sz, int mtype, int i, bool allowTransposed, int create(2, sizes, mtype, i, allowTransposed, fixedDepthMask); } -void _OutputArray::create(int rows, int cols, int mtype, int i, bool allowTransposed, int fixedDepthMask) const +void _OutputArray::create(int _rows, int _cols, int mtype, int i, bool allowTransposed, int fixedDepthMask) const { int k = kind(); if( k == MAT && i < 0 && !allowTransposed && fixedDepthMask == 0 ) { - CV_Assert(!fixedSize() || ((Mat*)obj)->size.operator()() == Size(cols, rows)); + CV_Assert(!fixedSize() || ((Mat*)obj)->size.operator()() == Size(_cols, _rows)); CV_Assert(!fixedType() || ((Mat*)obj)->type() == mtype); - ((Mat*)obj)->create(rows, cols, mtype); + ((Mat*)obj)->create(_rows, _cols, mtype); return; } if( k == UMAT && i < 0 && !allowTransposed && fixedDepthMask == 0 ) { - CV_Assert(!fixedSize() || ((UMat*)obj)->size.operator()() == Size(cols, rows)); + CV_Assert(!fixedSize() || ((UMat*)obj)->size.operator()() == Size(_cols, _rows)); CV_Assert(!fixedType() || ((UMat*)obj)->type() == mtype); - ((UMat*)obj)->create(rows, cols, mtype); + ((UMat*)obj)->create(_rows, _cols, mtype); return; } if( k == GPU_MAT && i < 0 && !allowTransposed && fixedDepthMask == 0 ) { - CV_Assert(!fixedSize() || ((cuda::GpuMat*)obj)->size() == Size(cols, rows)); + CV_Assert(!fixedSize() || ((cuda::GpuMat*)obj)->size() == Size(_cols, _rows)); CV_Assert(!fixedType() || ((cuda::GpuMat*)obj)->type() == mtype); - ((cuda::GpuMat*)obj)->create(rows, cols, mtype); + ((cuda::GpuMat*)obj)->create(_rows, _cols, mtype); return; } if( k == OPENGL_BUFFER && i < 0 && !allowTransposed && fixedDepthMask == 0 ) { - CV_Assert(!fixedSize() || ((ogl::Buffer*)obj)->size() == Size(cols, rows)); + CV_Assert(!fixedSize() || ((ogl::Buffer*)obj)->size() == Size(_cols, _rows)); CV_Assert(!fixedType() || ((ogl::Buffer*)obj)->type() == mtype); - ((ogl::Buffer*)obj)->create(rows, cols, mtype); + ((ogl::Buffer*)obj)->create(_rows, _cols, mtype); return; } if( k == CUDA_MEM && i < 0 && !allowTransposed && fixedDepthMask == 0 ) { - CV_Assert(!fixedSize() || ((cuda::CudaMem*)obj)->size() == Size(cols, rows)); + CV_Assert(!fixedSize() || ((cuda::CudaMem*)obj)->size() == Size(_cols, _rows)); CV_Assert(!fixedType() || ((cuda::CudaMem*)obj)->type() == mtype); - ((cuda::CudaMem*)obj)->create(rows, cols, mtype); + ((cuda::CudaMem*)obj)->create(_rows, _cols, mtype); return; } - int sizes[] = {rows, cols}; + int sizes[] = {_rows, _cols}; create(2, sizes, mtype, i, allowTransposed, fixedDepthMask); } diff --git a/modules/imgproc/src/filter.cpp b/modules/imgproc/src/filter.cpp index 2d7c7404f6..024f611ffb 100644 --- a/modules/imgproc/src/filter.cpp +++ b/modules/imgproc/src/filter.cpp @@ -3428,7 +3428,7 @@ static bool ocl_sepColFilter2D(const UMat & buf, UMat & dst, const Mat & kernelY const int optimizedSepFilterLocalSize = 16; static bool ocl_sepFilter2D_SinglePass(InputArray _src, OutputArray _dst, - InputArray _row_kernel, InputArray _col_kernel, + Mat row_kernel, Mat col_kernel, int borderType, int ddepth) { Size size = _src.size(), wholeSize; @@ -3439,7 +3439,7 @@ static bool ocl_sepFilter2D_SinglePass(InputArray _src, OutputArray _dst, 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) || + if ((src_offset % src_step) % esz != 0 || (!doubleSupport && (sdepth == CV_64F || ddepth == CV_64F)) || !(borderType == BORDER_CONSTANT || borderType == BORDER_REPLICATE || borderType == BORDER_REFLECT || borderType == BORDER_WRAP || borderType == BORDER_REFLECT_101)) @@ -3454,10 +3454,10 @@ static bool ocl_sepFilter2D_SinglePass(InputArray _src, OutputArray _dst, 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 -D srcT1=%s -D dstT1=%s -D cn=%d", (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(), + " -D %s -D srcT1=%s -D dstT1=%s -D CN=%d", (int)lt2[0], (int)lt2[1], + row_kernel.cols / 2, col_kernel.cols / 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], @@ -3486,12 +3486,13 @@ static bool ocl_sepFilter2D( InputArray _src, OutputArray _dst, int ddepth, InputArray _kernelX, InputArray _kernelY, Point anchor, double delta, int borderType ) { + const ocl::Device & d = ocl::Device::getDefault(); Size imgSize = _src.size(); if (abs(delta)> FLT_MIN) return false; - int type = _src.type(), cn = CV_MAT_CN(type); + int type = _src.type(), sdepth = CV_MAT_DEPTH(type), cn = CV_MAT_CN(type); if (cn > 4) return false; @@ -3502,21 +3503,21 @@ static bool ocl_sepFilter2D( InputArray _src, OutputArray _dst, int ddepth, if (kernelY.cols % 2 != 1) return false; - int sdepth = CV_MAT_DEPTH(type); + if (ddepth < 0) + ddepth = sdepth; + + CV_OCL_RUN_(kernelY.cols <= 21 && kernelX.cols <= 21 && + imgSize.width > optimizedSepFilterLocalSize + (kernelX.cols >> 1) && + imgSize.height > optimizedSepFilterLocalSize + (kernelY.cols >> 1) && + (!(borderType & BORDER_ISOLATED) || _src.offset() == 0) && anchor == Point(-1, -1) && + (d.isIntel() || (d.isAMD() && !d.hostUnifiedMemory())), + ocl_sepFilter2D_SinglePass(_src, _dst, kernelX, kernelY, borderType, ddepth), true) + if (anchor.x < 0) anchor.x = kernelX.cols >> 1; if (anchor.y < 0) anchor.y = kernelY.cols >> 1; - 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) && - (borderType & BORDER_ISOLATED) != 0, - 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 index 705150f7d5..82ccd24192 100644 --- a/modules/imgproc/src/opencl/filterSep_singlePass.cl +++ b/modules/imgproc/src/opencl/filterSep_singlePass.cl @@ -74,8 +74,19 @@ #error No extrapolation method #endif -#define SRC(_x,_y) convertToWT(((global srcT*)(Src+(_y)*src_step))[_x]) -#define DST(_x,_y) (((global dstT*)(Dst+dst_offset+(_y)*dst_step))[_x]) +#if CN != 3 +#define loadpix(addr) *(__global const srcT *)(addr) +#define storepix(val, addr) *(__global dstT *)(addr) = val +#define SRCSIZE (int)sizeof(srcT) +#define DSTSIZE (int)sizeof(dstT) +#else +#define loadpix(addr) vload3(0, (__global const srcT1 *)(addr)) +#define storepix(val, addr) vstore3(val, 0, (__global dstT1 *)(addr)) +#define SRCSIZE (int)sizeof(srcT1)*3 +#define DSTSIZE (int)sizeof(dstT1)*3 +#endif + +#define SRC(_x,_y) convertToWT(loadpix(Src + mad24(_y, src_step, SRCSIZE * _x))) #ifdef BORDER_CONSTANT // CCCCCC|abcdefgh|CCCCCCC @@ -172,5 +183,5 @@ __kernel void sep_filter(__global uchar* Src, int src_step, int srcOffsetX, int sum = mad(lsmemDy[liy][lix+i], mat_kernelX[i], sum); //store result into destination image - DST(x,y) = convertToDstT(sum); + storepix(convertToDstT(sum), Dst + mad24(y, dst_step, mad24(x, DSTSIZE, dst_offset))); } diff --git a/modules/imgproc/test/ocl/test_sepfilter2D.cpp b/modules/imgproc/test/ocl/test_sepfilter2D.cpp index b724641f45..0aafb8a0a4 100644 --- a/modules/imgproc/test/ocl/test_sepfilter2D.cpp +++ b/modules/imgproc/test/ocl/test_sepfilter2D.cpp @@ -86,16 +86,7 @@ PARAM_TEST_CASE(SepFilter2D, MatDepth, Channels, BorderType, bool, bool) cv::normalize(temp, kernelY, 1.0, 0.0, NORM_L1); Size roiSize = randomSize(ksize.width, MAX_VALUE, ksize.height, MAX_VALUE); - int rest = roiSize.width % 4; - if (rest != 0) - roiSize.width += (4 - rest); Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); - rest = srcBorder.lef % 4; - if (rest != 0) - srcBorder.lef += (4 - rest); - rest = srcBorder.rig % 4; - if (rest != 0) - srcBorder.rig += (4 - rest); randomSubMat(src, src_roi, roiSize, srcBorder, type, -MAX_VALUE, MAX_VALUE); Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); @@ -115,7 +106,7 @@ PARAM_TEST_CASE(SepFilter2D, MatDepth, Channels, BorderType, bool, bool) OCL_TEST_P(SepFilter2D, Mat) { - for (int j = 0; j < test_loop_times + 1; j++) + for (int j = 0; j < test_loop_times + 3; j++) { random_roi();