From e75cd74f5ab182487465c9a260d7780fae1b14d4 Mon Sep 17 00:00:00 2001 From: vbystricky Date: Mon, 30 Jun 2014 10:37:41 +0400 Subject: [PATCH] Optimize OpenCL version of Laplacian filter for kernel size great than 3 --- modules/imgproc/src/deriv.cpp | 90 ++++++++++- modules/imgproc/src/opencl/laplacian5.cl | 185 +++++++++++++++++++++++ 2 files changed, 271 insertions(+), 4 deletions(-) diff --git a/modules/imgproc/src/deriv.cpp b/modules/imgproc/src/deriv.cpp index c3579fee18..0207a3d469 100644 --- a/modules/imgproc/src/deriv.cpp +++ b/modules/imgproc/src/deriv.cpp @@ -643,21 +643,103 @@ void cv::Scharr( InputArray _src, OutputArray _dst, int ddepth, int dx, int dy, namespace cv { +#define LAPLACIAN_LOCAL_MEM(tileX, tileY, ksize, elsize) (((tileX) + 2 * (int)((ksize) / 2)) * (3 * (tileY) + 2 * (int)((ksize) / 2)) * elsize) + static bool ocl_Laplacian5(InputArray _src, OutputArray _dst, const Mat & kd, const Mat & ks, double scale, double delta, int borderType, int depth, int ddepth) { + const size_t tileSizeX = 16; + const size_t tileSizeYmin = 8; + + const ocl::Device dev = ocl::Device::getDefault(); + + int stype = _src.type(); + int sdepth = CV_MAT_DEPTH(stype), cn = CV_MAT_CN(stype), esz = CV_ELEM_SIZE(stype); + + bool doubleSupport = dev.doubleFPConfig() > 0; + if (!doubleSupport && (sdepth == CV_64F || ddepth == CV_64F)) + return false; + + Mat kernelX = kd.reshape(1, 1); + if (kernelX.cols % 2 != 1) + return false; + Mat kernelY = ks.reshape(1, 1); + if (kernelY.cols % 2 != 1) + return false; + CV_Assert(kernelX.cols == kernelY.cols); + + size_t wgs = dev.maxWorkGroupSize(); + size_t lmsz = dev.localMemSize(); + + size_t src_step = _src.step(), src_offset = _src.offset(); + if (((src_offset % src_step) % esz == 0) && + ( + (borderType == BORDER_CONSTANT || borderType == BORDER_REPLICATE) || + (borderType == BORDER_REFLECT || borderType == BORDER_WRAP || borderType == BORDER_REFLECT_101) && + (_src.cols() >= kernelX.cols && _src.rows() >= kernelY.cols) + ) && + (tileSizeX * tileSizeYmin <= wgs) && + (LAPLACIAN_LOCAL_MEM(tileSizeX, tileSizeYmin, kernelX.cols, cn * 4) <= lmsz) + ) + { + Size size = _src.size(), wholeSize; + Point origin; + int dtype = CV_MAKE_TYPE(ddepth, cn); + int wdepth = CV_32F; + + size_t tileSizeY = wgs / tileSizeX; + while ((tileSizeX * tileSizeY > wgs) || (LAPLACIAN_LOCAL_MEM(tileSizeX, tileSizeY, kernelX.cols, cn * 4) > lmsz)) + { + tileSizeY /= 2; + } + size_t lt2[2] = { tileSizeX, tileSizeY}; + size_t gt2[2] = { lt2[0] * (1 + (size.width - 1) / lt2[0]), 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 RADIUS=%d%s%s" + " -D convertToWT=%s -D convertToDT=%s" + " -D %s -D srcT1=%s -D dstT1=%s -D WT1=%s -D CN=%d ", + (int)lt2[0], (int)lt2[1], kernelX.cols / 2, + ocl::kernelToStr(kernelX, wdepth, "KERNEL_MATRIX_X").c_str(), + ocl::kernelToStr(kernelY, wdepth, "KERNEL_MATRIX_Y").c_str(), + ocl::convertTypeStr(sdepth, wdepth, cn, cvt[0]), + ocl::convertTypeStr(wdepth, ddepth, cn, cvt[1]), + borderMap[borderType], ocl::typeToStr(sdepth), ocl::typeToStr(ddepth), ocl::typeToStr(wdepth), + cn); + + ocl::Kernel k("laplacian", ocl::imgproc::laplacian5_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), + static_cast(scale), static_cast(delta)); + + return k.run(2, gt2, lt2, false); + } int iscale = cvRound(scale), idelta = cvRound(delta); - bool doubleSupport = ocl::Device::getDefault().doubleFPConfig() > 0, - floatCoeff = std::fabs(delta - idelta) > DBL_EPSILON || std::fabs(scale - iscale) > DBL_EPSILON; - int cn = _src.channels(), wdepth = std::max(depth, floatCoeff ? CV_32F : CV_32S), kercn = 1; + bool floatCoeff = std::fabs(delta - idelta) > DBL_EPSILON || std::fabs(scale - iscale) > DBL_EPSILON; + int wdepth = std::max(depth, floatCoeff ? CV_32F : CV_32S), kercn = 1; if (!doubleSupport && wdepth == CV_64F) return false; char cvt[2][40]; ocl::Kernel k("sumConvert", ocl::imgproc::laplacian5_oclsrc, - format("-D srcT=%s -D WT=%s -D dstT=%s -D coeffT=%s -D wdepth=%d " + format("-D ONLY_SUM_CONVERT " + "-D srcT=%s -D WT=%s -D dstT=%s -D coeffT=%s -D wdepth=%d " "-D convertToWT=%s -D convertToDT=%s%s", ocl::typeToStr(CV_MAKE_TYPE(depth, kercn)), ocl::typeToStr(CV_MAKE_TYPE(wdepth, kercn)), diff --git a/modules/imgproc/src/opencl/laplacian5.cl b/modules/imgproc/src/opencl/laplacian5.cl index 3e15e097cb..808c414575 100644 --- a/modules/imgproc/src/opencl/laplacian5.cl +++ b/modules/imgproc/src/opencl/laplacian5.cl @@ -5,8 +5,11 @@ // Copyright (C) 2014, Itseez, Inc., all rights reserved. // Third party copyrights are property of their respective owners. + #define noconvert +#ifdef ONLY_SUM_CONVERT + __kernel void sumConvert(__global const uchar * src1ptr, int src1_step, int src1_offset, __global const uchar * src2ptr, int src2_step, int src2_offset, __global uchar * dstptr, int dst_step, int dst_offset, int dst_rows, int dst_cols, @@ -32,3 +35,185 @@ __kernel void sumConvert(__global const uchar * src1ptr, int src1_step, int src1 #endif } } + +#else + +#if CN==1 +#define srcT srcT1 +#define WT WT1 +#define dstT dstT1 +#else +#define __CAT(x, y) x##y +#define CAT(x, y) __CAT(x, y) + +#define srcT CAT(srcT1, CN) +#define WT CAT(WT1, CN) +#define dstT CAT(dstT1, CN) +#endif + +/////////////////////////////////////////////////////////////////////////////////////////////////// +/////////////////////////////////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 +// gfedcb|abcdefgh|gfedcba +#define EXTRAPOLATE(x, maxV) \ + { \ + (x) = min(((maxV)-1)*2-(x), max((x),-(x)) ); \ + } +#else +#error No extrapolation method +#endif + +#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 +#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 + +// horizontal and vertical filter kernels +// should be defined on host during compile time to avoid overhead +#define DIG(a) a, +__constant WT1 mat_kernelX[] = { KERNEL_MATRIX_X }; +__constant WT1 mat_kernelY[] = { KERNEL_MATRIX_Y }; + +__kernel void laplacian(__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, + WT1 scale, WT1 delta) +{ + __local WT lsmem[BLK_Y + 2 * RADIUS][BLK_X + 2 * RADIUS]; + __local WT lsmemDy1[BLK_Y][BLK_X + 2 * RADIUS]; + __local WT lsmemDy2[BLK_Y][BLK_X + 2 * RADIUS]; + + int lix = get_local_id(0); + int liy = get_local_id(1); + + int x = get_global_id(0); + + int srcX = x + srcOffsetX - RADIUS; + + int clocY = liy; + do + { + int yb = clocY + srcOffsetY - RADIUS; + 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+(RADIUS*2)); + + clocY += BLK_Y; + } + while (clocY < BLK_Y+(RADIUS*2)); + barrier(CLK_LOCAL_MEM_FENCE); + + WT scale_v = (WT)scale; + WT delta_v = (WT)delta; + for (int y = 0; y < dst_rows; y+=BLK_Y) + { + int i, clocX = lix; + WT sum1 = (WT) 0; + WT sum2 = (WT) 0; + do + { + sum1 = (WT) 0; + sum2 = (WT) 0; + for (i=0; i<=2*RADIUS; i++) + { + sum1 = mad(lsmem[liy + i][clocX], mat_kernelY[i], sum1); + sum2 = mad(lsmem[liy + i][clocX], mat_kernelX[i], sum2); + } + lsmemDy1[liy][clocX] = sum1; + lsmemDy2[liy][clocX] = sum2; + clocX += BLK_X; + } + while(clocX < BLK_X+(RADIUS*2)); + barrier(CLK_LOCAL_MEM_FENCE); + + if ((x < dst_cols) && (y + liy < dst_rows)) + { + sum1 = (WT) 0; + sum2 = (WT) 0; + for (i=0; i<=2*RADIUS; i++) + { + sum1 = mad(lsmemDy1[liy][lix+i], mat_kernelX[i], sum1); + sum2 = mad(lsmemDy2[liy][lix+i], mat_kernelY[i], sum2); + } + + WT sum = mad(scale_v, (sum1 + sum2), delta_v); + storepix(convertToDT(sum), Dst + mad24(y + liy, dst_step, mad24(x, DSTSIZE, dst_offset))); + } + + for (int i = liy * BLK_X + lix; i < (RADIUS*2) * (BLK_X+(RADIUS*2)); i += BLK_X * BLK_Y) + { + int clocX = i % (BLK_X+(RADIUS*2)); + int clocY = i / (BLK_X+(RADIUS*2)); + lsmem[clocY][clocX] = lsmem[clocY + BLK_Y][clocX]; + } + barrier(CLK_LOCAL_MEM_FENCE); + + int yb = y + liy + BLK_Y + srcOffsetY + RADIUS; + EXTRAPOLATE(yb, (height)); + + clocX = lix; + int cSrcX = x + srcOffsetX - RADIUS; + do + { + int xb = cSrcX; + EXTRAPOLATE(xb,(width)); + lsmem[liy + 2*RADIUS][clocX] = ELEM(xb, yb, (width), (height), 0 ); + + clocX += BLK_X; + cSrcX += BLK_X; + } + while(clocX < BLK_X+(RADIUS*2)); + barrier(CLK_LOCAL_MEM_FENCE); + } +} + +#endif \ No newline at end of file -- GitLab