From 4910242732f7ddb2cabb6e209166e996d6515af0 Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Fri, 23 May 2014 14:58:34 +0400 Subject: [PATCH 1/3] Unroll pyrUp kernel --- modules/imgproc/src/opencl/pyr_up.cl | 173 ++++++++++++++++----------- modules/imgproc/src/pyramids.cpp | 9 +- 2 files changed, 110 insertions(+), 72 deletions(-) diff --git a/modules/imgproc/src/opencl/pyr_up.cl b/modules/imgproc/src/opencl/pyr_up.cl index d754a70e08..f9b5c8f961 100644 --- a/modules/imgproc/src/opencl/pyr_up.cl +++ b/modules/imgproc/src/opencl/pyr_up.cl @@ -68,102 +68,139 @@ #define PIXSIZE ((int)sizeof(T1)*3) #endif +#define EXTRAPOLATE(x, maxV) min(maxV - 1, (int) abs(x)) + #define noconvert __kernel void pyrUp(__global const uchar * src, int src_step, int src_offset, int src_rows, int src_cols, __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols) { - const int x = get_global_id(0); - const int y = get_global_id(1); - - const int lsizex = get_local_size(0); - const int lsizey = get_local_size(1); - - const int tidx = get_local_id(0); - const int tidy = get_local_id(1); + const int lx = 2*get_local_id(0); + const int ly = 2*get_local_id(1); - __local FT s_srcPatch[10][10]; - __local FT s_dstPatch[20][16]; + __local FT s_srcPatch[LOCAL_SIZE+2][LOCAL_SIZE+2]; + __local FT s_dstPatch[2*LOCAL_SIZE+4][2*LOCAL_SIZE]; __global uchar * dstData = dst + dst_offset; __global const uchar * srcData = src + src_offset; - if( tidx < 10 && tidy < 10 ) + if( lx < (LOCAL_SIZE+2) && lx < (LOCAL_SIZE+2) ) { - int srcx = mad24((int)get_group_id(0), lsizex>>1, tidx) - 1; - int srcy = mad24((int)get_group_id(1), lsizey>>1, tidy) - 1; - - srcx = abs(srcx); - srcx = min(src_cols - 1, srcx); - - srcy = abs(srcy); - srcy = min(src_rows - 1, srcy); - - s_srcPatch[tidy][tidx] = convertToFT(loadpix(srcData + srcy * src_step + srcx * PIXSIZE)); + int srcx = mad24((int)get_group_id(0), LOCAL_SIZE, lx) - 1; + int srcy = mad24((int)get_group_id(1), LOCAL_SIZE, ly) - 1; + + int srcx1 = EXTRAPOLATE(srcx, src_cols); + int srcx2 = EXTRAPOLATE(srcx+1, src_cols); + int srcy1 = EXTRAPOLATE(srcy, src_rows); + int srcy2 = EXTRAPOLATE(srcy+1, src_rows); + s_srcPatch[ly][lx] = convertToFT(loadpix(srcData + srcy1 * src_step + srcx1 * PIXSIZE)); + s_srcPatch[ly+1][lx] = convertToFT(loadpix(srcData + srcy2 * src_step + srcx1 * PIXSIZE)); + s_srcPatch[ly][lx+1] = convertToFT(loadpix(srcData + srcy1 * src_step + srcx2 * PIXSIZE)); + s_srcPatch[ly+1][lx+1] = convertToFT(loadpix(srcData + srcy2 * src_step + srcx2 * PIXSIZE)); } barrier(CLK_LOCAL_MEM_FENCE); - FT sum = 0.f; - const FT evenFlag = (FT)((tidx & 1) == 0); - const FT oddFlag = (FT)((tidx & 1) != 0); - const bool eveny = ((tidy & 1) == 0); + FT sum; const FT co1 = 0.375f; const FT co2 = 0.25f; const FT co3 = 0.0625f; - if(eveny) - { - sum = ( evenFlag* co3 ) * s_srcPatch[1 + (tidy >> 1)][1 + ((tidx - 2) >> 1)]; - sum = sum + ( oddFlag * co2 ) * s_srcPatch[1 + (tidy >> 1)][1 + ((tidx - 1) >> 1)]; - sum = sum + ( evenFlag* co1 ) * s_srcPatch[1 + (tidy >> 1)][1 + ((tidx ) >> 1)]; - sum = sum + ( oddFlag * co2 ) * s_srcPatch[1 + (tidy >> 1)][1 + ((tidx + 1) >> 1)]; - sum = sum + ( evenFlag* co3 ) * s_srcPatch[1 + (tidy >> 1)][1 + ((tidx + 2) >> 1)]; - } + // (x,y) + sum = co3 * s_srcPatch[1 + (ly >> 1)][1 + ((lx - 2) >> 1)]; + sum = sum + co1 * s_srcPatch[1 + (ly >> 1)][1 + ((lx ) >> 1)]; + sum = sum + co3 * s_srcPatch[1 + (ly >> 1)][1 + ((lx + 2) >> 1)]; + + s_dstPatch[2 + ly][lx] = sum; + + // (x+1,y) + sum = co2 * s_srcPatch[1 + (ly >> 1)][1 + ((lx + 1 - 1) >> 1)]; + sum = sum + co2 * s_srcPatch[1 + (ly >> 1)][1 + ((lx + 1 + 1) >> 1)]; + s_dstPatch[2 + ly][lx+1] = sum; - s_dstPatch[2 + tidy][tidx] = sum; + // (x, y+1) (x+1, y+1) + s_dstPatch[2 + ly+1][lx] = 0.f; + s_dstPatch[2 + ly+1][lx+1] = 0.f; - if (tidy < 2) + if (ly < 1) { - sum = 0; - - if (eveny) - { - sum = (evenFlag * co3 ) * s_srcPatch[lsizey-16][1 + ((tidx - 2) >> 1)]; - sum = sum + ( oddFlag * co2 ) * s_srcPatch[lsizey-16][1 + ((tidx - 1) >> 1)]; - sum = sum + (evenFlag * co1 ) * s_srcPatch[lsizey-16][1 + ((tidx ) >> 1)]; - sum = sum + ( oddFlag * co2 ) * s_srcPatch[lsizey-16][1 + ((tidx + 1) >> 1)]; - sum = sum + (evenFlag * co3 ) * s_srcPatch[lsizey-16][1 + ((tidx + 2) >> 1)]; - } - - s_dstPatch[tidy][tidx] = sum; + // (x,y) + sum = co3 * s_srcPatch[0][1 + ((lx - 2) >> 1)]; + sum = sum + co1 * s_srcPatch[0][1 + ((lx ) >> 1)]; + sum = sum + co3 * s_srcPatch[0][1 + ((lx + 2) >> 1)]; + s_dstPatch[ly][lx] = sum; + + // (x+1,y) + sum = co2 * s_srcPatch[0][1 + ((lx + 1 - 1) >> 1)]; + sum = sum + co2 * s_srcPatch[0][1 + ((lx + 1 + 1) >> 1)]; + s_dstPatch[ly][lx+1] = sum; + + // (x, y+1) (x+1, y+1) + s_dstPatch[ly+1][lx] = 0.f; + s_dstPatch[ly+1][lx+1] = 0.f; } - if (tidy > 13) + if (ly > 2*LOCAL_SIZE-3) { - sum = 0; - - if (eveny) - { - sum = (evenFlag * co3) * s_srcPatch[lsizey-7][1 + ((tidx - 2) >> 1)]; - sum = sum + ( oddFlag * co2) * s_srcPatch[lsizey-7][1 + ((tidx - 1) >> 1)]; - sum = sum + (evenFlag * co1) * s_srcPatch[lsizey-7][1 + ((tidx ) >> 1)]; - sum = sum + ( oddFlag * co2) * s_srcPatch[lsizey-7][1 + ((tidx + 1) >> 1)]; - sum = sum + (evenFlag * co3) * s_srcPatch[lsizey-7][1 + ((tidx + 2) >> 1)]; - } - s_dstPatch[4 + tidy][tidx] = sum; + // (x,y) + sum = co3 * s_srcPatch[LOCAL_SIZE+1][1 + ((lx - 2) >> 1)]; + sum = sum + co1 * s_srcPatch[LOCAL_SIZE+1][1 + ((lx ) >> 1)]; + sum = sum + co3 * s_srcPatch[LOCAL_SIZE+1][1 + ((lx + 2) >> 1)]; + s_dstPatch[4 + ly][lx] = sum; + + // (x+1,y) + sum = co2 * s_srcPatch[LOCAL_SIZE+1][1 + ((lx + 1 - 1) >> 1)]; + sum = sum + co2 * s_srcPatch[LOCAL_SIZE+1][1 + ((lx + 1 + 1) >> 1)]; + s_dstPatch[4 + ly][lx+1] = sum; + + // (x, y+1) (x+1, y+1) + s_dstPatch[4 + ly+1][lx] = 0.f; + s_dstPatch[4 + ly+1][lx+1] = 0.f; } barrier(CLK_LOCAL_MEM_FENCE); - - sum = co3 * s_dstPatch[2 + tidy - 2][tidx]; - sum = sum + co2 * s_dstPatch[2 + tidy - 1][tidx]; - sum = sum + co1 * s_dstPatch[2 + tidy ][tidx]; - sum = sum + co2 * s_dstPatch[2 + tidy + 1][tidx]; - sum = sum + co3 * s_dstPatch[2 + tidy + 2][tidx]; - - if ((x < dst_cols) && (y < dst_rows)) - storepix(convertToT(4.0f * sum), dstData + y * dst_step + x * PIXSIZE); + int dst_x = 2*get_global_id(0); + int dst_y = 2*get_global_id(1); + + // (x,y) + sum = co3 * s_dstPatch[2 + ly - 2][lx]; + sum = sum + co2 * s_dstPatch[2 + ly - 1][lx]; + sum = sum + co1 * s_dstPatch[2 + ly ][lx]; + sum = sum + co2 * s_dstPatch[2 + ly + 1][lx]; + sum = sum + co3 * s_dstPatch[2 + ly + 2][lx]; + + if ((dst_x < dst_cols) && (dst_y < dst_rows)) + storepix(convertToT(4.0f * sum), dstData + dst_y * dst_step + dst_x * PIXSIZE); + + // (x+1,y) + sum = co3 * s_dstPatch[2 + ly - 2][lx+1]; + sum = sum + co2 * s_dstPatch[2 + ly - 1][lx+1]; + sum = sum + co1 * s_dstPatch[2 + ly ][lx+1]; + sum = sum + co2 * s_dstPatch[2 + ly + 1][lx+1]; + sum = sum + co3 * s_dstPatch[2 + ly + 2][lx+1]; + + if ((dst_x+1 < dst_cols) && (dst_y < dst_rows)) + storepix(convertToT(4.0f * sum), dstData + dst_y * dst_step + (dst_x+1) * PIXSIZE); + + // (x,y+1) + sum = co3 * s_dstPatch[2 + ly+1 - 2][lx]; + sum = sum + co2 * s_dstPatch[2 + ly+1 - 1][lx]; + sum = sum + co1 * s_dstPatch[2 + ly+1 ][lx]; + sum = sum + co2 * s_dstPatch[2 + ly+1 + 1][lx]; + sum = sum + co3 * s_dstPatch[2 + ly+1 + 2][lx]; + + if ((dst_x < dst_cols) && (dst_y+1 < dst_rows)) + storepix(convertToT(4.0f * sum), dstData + (dst_y+1) * dst_step + dst_x * PIXSIZE); + + // (x+1,y+1) + sum = co3 * s_dstPatch[2 + ly+1 - 2][lx+1]; + sum = sum + co2 * s_dstPatch[2 + ly+1 - 1][lx+1]; + sum = sum + co1 * s_dstPatch[2 + ly+1 ][lx+1]; + sum = sum + co2 * s_dstPatch[2 + ly+1 + 1][lx+1]; + sum = sum + co3 * s_dstPatch[2 + ly+1 + 2][lx+1]; + + if ((dst_x+1 < dst_cols) && (dst_y+1 < dst_rows)) + storepix(convertToT(4.0f * sum), dstData + (dst_y+1) * dst_step + (dst_x+1) * PIXSIZE); } diff --git a/modules/imgproc/src/pyramids.cpp b/modules/imgproc/src/pyramids.cpp index 42464c1a5d..a0a09ec683 100644 --- a/modules/imgproc/src/pyramids.cpp +++ b/modules/imgproc/src/pyramids.cpp @@ -467,23 +467,24 @@ static bool ocl_pyrUp( InputArray _src, OutputArray _dst, const Size& _dsz, int UMat dst = _dst.getUMat(); int float_depth = depth == CV_64F ? CV_64F : CV_32F; + int local_size = 8; char cvt[2][50]; String buildOptions = format( "-D T=%s -D FT=%s -D convertToT=%s -D convertToFT=%s%s " - "-D T1=%s -D cn=%d", + "-D T1=%s -D cn=%d -D LOCAL_SIZE=%d", ocl::typeToStr(type), ocl::typeToStr(CV_MAKETYPE(float_depth, channels)), ocl::convertTypeStr(float_depth, depth, channels, cvt[0]), ocl::convertTypeStr(depth, float_depth, channels, cvt[1]), doubleSupport ? " -D DOUBLE_SUPPORT" : "", - ocl::typeToStr(depth), channels + ocl::typeToStr(depth), channels, local_size ); ocl::Kernel k("pyrUp", ocl::imgproc::pyr_up_oclsrc, buildOptions); if (k.empty()) return false; k.args(ocl::KernelArg::ReadOnly(src), ocl::KernelArg::WriteOnly(dst)); - size_t globalThreads[2] = {dst.cols, dst.rows}; - size_t localThreads[2] = {16, 16}; + size_t globalThreads[2] = {dst.cols/2, dst.rows/2}; + size_t localThreads[2] = {local_size, local_size}; return k.run(2, globalThreads, localThreads, false); } From 8e548450497915de4867e3d676b00362a53a4d91 Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Tue, 27 May 2014 10:52:20 +0400 Subject: [PATCH 2/3] Removed useless multiplication by 4 --- modules/imgproc/src/opencl/pyr_up.cl | 113 +++++++++++++++++++++++++-- modules/imgproc/src/pyramids.cpp | 15 +++- 2 files changed, 118 insertions(+), 10 deletions(-) diff --git a/modules/imgproc/src/opencl/pyr_up.cl b/modules/imgproc/src/opencl/pyr_up.cl index f9b5c8f961..dc70c8fbed 100644 --- a/modules/imgproc/src/opencl/pyr_up.cl +++ b/modules/imgproc/src/opencl/pyr_up.cl @@ -72,9 +72,106 @@ #define noconvert - __kernel void pyrUp(__global const uchar * src, int src_step, int src_offset, int src_rows, int src_cols, __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols) +{ + const int x = get_global_id(0); + const int y = get_global_id(1); + + const int lsizex = get_local_size(0); + const int lsizey = get_local_size(1); + + const int tidx = get_local_id(0); + const int tidy = get_local_id(1); + + __local FT s_srcPatch[10][10]; + __local FT s_dstPatch[20][16]; + + __global uchar * dstData = dst + dst_offset; + __global const uchar * srcData = src + src_offset; + + if( tidx < 10 && tidy < 10 ) + { + int srcx = mad24((int)get_group_id(0), lsizex>>1, tidx) - 1; + int srcy = mad24((int)get_group_id(1), lsizey>>1, tidy) - 1; + + srcx = abs(srcx); + srcx = min(src_cols - 1, srcx); + + srcy = abs(srcy); + srcy = min(src_rows - 1, srcy); + + s_srcPatch[tidy][tidx] = convertToFT(loadpix(srcData + srcy * src_step + srcx * PIXSIZE)); + } + + barrier(CLK_LOCAL_MEM_FENCE); + + FT sum = 0.f; + const FT evenFlag = (FT)((tidx & 1) == 0); + const FT oddFlag = (FT)((tidx & 1) != 0); + const bool eveny = ((tidy & 1) == 0); + + const FT co1 = 0.75f; + const FT co2 = 0.5f; + const FT co3 = 0.125f; + + if(eveny) + { + sum = ( evenFlag* co3 ) * s_srcPatch[1 + (tidy >> 1)][1 + ((tidx - 2) >> 1)]; + sum = sum + ( oddFlag * co2 ) * s_srcPatch[1 + (tidy >> 1)][1 + ((tidx - 1) >> 1)]; + sum = sum + ( evenFlag* co1 ) * s_srcPatch[1 + (tidy >> 1)][1 + ((tidx ) >> 1)]; + sum = sum + ( oddFlag * co2 ) * s_srcPatch[1 + (tidy >> 1)][1 + ((tidx + 1) >> 1)]; + sum = sum + ( evenFlag* co3 ) * s_srcPatch[1 + (tidy >> 1)][1 + ((tidx + 2) >> 1)]; + } + + s_dstPatch[2 + tidy][tidx] = sum; + + if (tidy < 2) + { + sum = 0; + + if (eveny) + { + sum = (evenFlag * co3 ) * s_srcPatch[lsizey-16][1 + ((tidx - 2) >> 1)]; + sum = sum + ( oddFlag * co2 ) * s_srcPatch[lsizey-16][1 + ((tidx - 1) >> 1)]; + sum = sum + (evenFlag * co1 ) * s_srcPatch[lsizey-16][1 + ((tidx ) >> 1)]; + sum = sum + ( oddFlag * co2 ) * s_srcPatch[lsizey-16][1 + ((tidx + 1) >> 1)]; + sum = sum + (evenFlag * co3 ) * s_srcPatch[lsizey-16][1 + ((tidx + 2) >> 1)]; + } + + s_dstPatch[tidy][tidx] = sum; + } + + if (tidy > 13) + { + sum = 0; + + if (eveny) + { + sum = (evenFlag * co3) * s_srcPatch[lsizey-7][1 + ((tidx - 2) >> 1)]; + sum = sum + ( oddFlag * co2) * s_srcPatch[lsizey-7][1 + ((tidx - 1) >> 1)]; + sum = sum + (evenFlag * co1) * s_srcPatch[lsizey-7][1 + ((tidx ) >> 1)]; + sum = sum + ( oddFlag * co2) * s_srcPatch[lsizey-7][1 + ((tidx + 1) >> 1)]; + sum = sum + (evenFlag * co3) * s_srcPatch[lsizey-7][1 + ((tidx + 2) >> 1)]; + } + s_dstPatch[4 + tidy][tidx] = sum; + } + + barrier(CLK_LOCAL_MEM_FENCE); + + sum = co3 * s_dstPatch[2 + tidy - 2][tidx]; + sum = sum + co2 * s_dstPatch[2 + tidy - 1][tidx]; + sum = sum + co1 * s_dstPatch[2 + tidy ][tidx]; + sum = sum + co2 * s_dstPatch[2 + tidy + 1][tidx]; + sum = sum + co3 * s_dstPatch[2 + tidy + 2][tidx]; + + if ((x < dst_cols) && (y < dst_rows)) + storepix(convertToT(sum), dstData + y * dst_step + x * PIXSIZE); +} + + +__kernel void pyrUp_unrolled(__global const uchar * src, int src_step, int src_offset, int src_rows, int src_cols, + __global uchar * dst, int dst_step, int dst_offset, int dst_rows, int dst_cols) { const int lx = 2*get_local_id(0); const int ly = 2*get_local_id(1); @@ -104,9 +201,9 @@ __kernel void pyrUp(__global const uchar * src, int src_step, int src_offset, in FT sum; - const FT co1 = 0.375f; - const FT co2 = 0.25f; - const FT co3 = 0.0625f; + const FT co1 = 0.75f; + const FT co2 = 0.5f; + const FT co3 = 0.125f; // (x,y) sum = co3 * s_srcPatch[1 + (ly >> 1)][1 + ((lx - 2) >> 1)]; @@ -172,7 +269,7 @@ __kernel void pyrUp(__global const uchar * src, int src_step, int src_offset, in sum = sum + co3 * s_dstPatch[2 + ly + 2][lx]; if ((dst_x < dst_cols) && (dst_y < dst_rows)) - storepix(convertToT(4.0f * sum), dstData + dst_y * dst_step + dst_x * PIXSIZE); + storepix(convertToT(sum), dstData + dst_y * dst_step + dst_x * PIXSIZE); // (x+1,y) sum = co3 * s_dstPatch[2 + ly - 2][lx+1]; @@ -182,7 +279,7 @@ __kernel void pyrUp(__global const uchar * src, int src_step, int src_offset, in sum = sum + co3 * s_dstPatch[2 + ly + 2][lx+1]; if ((dst_x+1 < dst_cols) && (dst_y < dst_rows)) - storepix(convertToT(4.0f * sum), dstData + dst_y * dst_step + (dst_x+1) * PIXSIZE); + storepix(convertToT(sum), dstData + dst_y * dst_step + (dst_x+1) * PIXSIZE); // (x,y+1) sum = co3 * s_dstPatch[2 + ly+1 - 2][lx]; @@ -192,7 +289,7 @@ __kernel void pyrUp(__global const uchar * src, int src_step, int src_offset, in sum = sum + co3 * s_dstPatch[2 + ly+1 + 2][lx]; if ((dst_x < dst_cols) && (dst_y+1 < dst_rows)) - storepix(convertToT(4.0f * sum), dstData + (dst_y+1) * dst_step + dst_x * PIXSIZE); + storepix(convertToT(sum), dstData + (dst_y+1) * dst_step + dst_x * PIXSIZE); // (x+1,y+1) sum = co3 * s_dstPatch[2 + ly+1 - 2][lx+1]; @@ -202,5 +299,5 @@ __kernel void pyrUp(__global const uchar * src, int src_step, int src_offset, in sum = sum + co3 * s_dstPatch[2 + ly+1 + 2][lx+1]; if ((dst_x+1 < dst_cols) && (dst_y+1 < dst_rows)) - storepix(convertToT(4.0f * sum), dstData + (dst_y+1) * dst_step + (dst_x+1) * PIXSIZE); + storepix(convertToT(sum), dstData + (dst_y+1) * dst_step + (dst_x+1) * PIXSIZE); } diff --git a/modules/imgproc/src/pyramids.cpp b/modules/imgproc/src/pyramids.cpp index a0a09ec683..319ff82000 100644 --- a/modules/imgproc/src/pyramids.cpp +++ b/modules/imgproc/src/pyramids.cpp @@ -478,12 +478,23 @@ static bool ocl_pyrUp( InputArray _src, OutputArray _dst, const Size& _dsz, int doubleSupport ? " -D DOUBLE_SUPPORT" : "", ocl::typeToStr(depth), channels, local_size ); - ocl::Kernel k("pyrUp", ocl::imgproc::pyr_up_oclsrc, buildOptions); + size_t globalThreads[2]; + ocl::Kernel k; + if (ocl::Device::getDefault().isIntel() && channels == 1) + { + k.create("pyrUp_unrolled", ocl::imgproc::pyr_up_oclsrc, buildOptions); + globalThreads[0] = dst.cols/2; globalThreads[1] = dst.rows/2; + } + else + { + k.create("pyrUp", ocl::imgproc::pyr_up_oclsrc, buildOptions); + local_size = 16; + globalThreads[0] = dst.cols; globalThreads[1] = dst.rows; + } if (k.empty()) return false; k.args(ocl::KernelArg::ReadOnly(src), ocl::KernelArg::WriteOnly(dst)); - size_t globalThreads[2] = {dst.cols/2, dst.rows/2}; size_t localThreads[2] = {local_size, local_size}; return k.run(2, globalThreads, localThreads, false); From 06fb5da7c805ed33c6b65b54e910dd29f79c5b00 Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Wed, 11 Jun 2014 15:15:15 +0400 Subject: [PATCH 3/3] Removed storing of zeros in local memory --- modules/imgproc/src/opencl/pyr_up.cl | 153 ++++++++------------------- modules/imgproc/src/pyramids.cpp | 14 +-- 2 files changed, 47 insertions(+), 120 deletions(-) diff --git a/modules/imgproc/src/opencl/pyr_up.cl b/modules/imgproc/src/opencl/pyr_up.cl index dc70c8fbed..1fdc58266a 100644 --- a/modules/imgproc/src/opencl/pyr_up.cl +++ b/modules/imgproc/src/opencl/pyr_up.cl @@ -78,28 +78,19 @@ __kernel void pyrUp(__global const uchar * src, int src_step, int src_offset, in const int x = get_global_id(0); const int y = get_global_id(1); - const int lsizex = get_local_size(0); - const int lsizey = get_local_size(1); - const int tidx = get_local_id(0); const int tidy = get_local_id(1); - __local FT s_srcPatch[10][10]; - __local FT s_dstPatch[20][16]; + __local FT s_srcPatch[LOCAL_SIZE/2 + 2][LOCAL_SIZE/2 + 2]; + __local FT s_dstPatch[LOCAL_SIZE/2 + 2][LOCAL_SIZE]; __global uchar * dstData = dst + dst_offset; __global const uchar * srcData = src + src_offset; - if( tidx < 10 && tidy < 10 ) + if( tidx < (LOCAL_SIZE/2 + 2) && tidy < LOCAL_SIZE/2 + 2 ) { - int srcx = mad24((int)get_group_id(0), lsizex>>1, tidx) - 1; - int srcy = mad24((int)get_group_id(1), lsizey>>1, tidy) - 1; - - srcx = abs(srcx); - srcx = min(src_cols - 1, srcx); - - srcy = abs(srcy); - srcy = min(src_rows - 1, srcy); + int srcx = EXTRAPOLATE(mad24((int)get_group_id(0), LOCAL_SIZE/2, tidx) - 1, src_cols); + int srcy = EXTRAPOLATE(mad24((int)get_group_id(1), LOCAL_SIZE/2, tidy) - 1, src_rows); s_srcPatch[tidy][tidx] = convertToFT(loadpix(srcData + srcy * src_step + srcx * PIXSIZE)); } @@ -107,63 +98,30 @@ __kernel void pyrUp(__global const uchar * src, int src_step, int src_offset, in barrier(CLK_LOCAL_MEM_FENCE); FT sum = 0.f; - const FT evenFlag = (FT)((tidx & 1) == 0); - const FT oddFlag = (FT)((tidx & 1) != 0); - const bool eveny = ((tidy & 1) == 0); const FT co1 = 0.75f; const FT co2 = 0.5f; const FT co3 = 0.125f; - if(eveny) - { - sum = ( evenFlag* co3 ) * s_srcPatch[1 + (tidy >> 1)][1 + ((tidx - 2) >> 1)]; - sum = sum + ( oddFlag * co2 ) * s_srcPatch[1 + (tidy >> 1)][1 + ((tidx - 1) >> 1)]; - sum = sum + ( evenFlag* co1 ) * s_srcPatch[1 + (tidy >> 1)][1 + ((tidx ) >> 1)]; - sum = sum + ( oddFlag * co2 ) * s_srcPatch[1 + (tidy >> 1)][1 + ((tidx + 1) >> 1)]; - sum = sum + ( evenFlag* co3 ) * s_srcPatch[1 + (tidy >> 1)][1 + ((tidx + 2) >> 1)]; - } - - s_dstPatch[2 + tidy][tidx] = sum; + const FT coef1 = (tidx & 1) == 0 ? co1 : (FT) 0; + const FT coef2 = (tidx & 1) == 0 ? co3 : co2; + const FT coefy1 = (tidy & 1) == 0 ? co1 : (FT) 0; + const FT coefy2 = (tidy & 1) == 0 ? co3 : co2; - if (tidy < 2) + if(tidy < LOCAL_SIZE/2 + 2) { - sum = 0; - - if (eveny) - { - sum = (evenFlag * co3 ) * s_srcPatch[lsizey-16][1 + ((tidx - 2) >> 1)]; - sum = sum + ( oddFlag * co2 ) * s_srcPatch[lsizey-16][1 + ((tidx - 1) >> 1)]; - sum = sum + (evenFlag * co1 ) * s_srcPatch[lsizey-16][1 + ((tidx ) >> 1)]; - sum = sum + ( oddFlag * co2 ) * s_srcPatch[lsizey-16][1 + ((tidx + 1) >> 1)]; - sum = sum + (evenFlag * co3 ) * s_srcPatch[lsizey-16][1 + ((tidx + 2) >> 1)]; - } + sum = coef2* s_srcPatch[tidy][1 + ((tidx - 1) >> 1)]; + sum = mad(coef1, s_srcPatch[tidy][1 + ((tidx ) >> 1)], sum); + sum = mad(coef2, s_srcPatch[tidy][1 + ((tidx + 2) >> 1)], sum); s_dstPatch[tidy][tidx] = sum; } - if (tidy > 13) - { - sum = 0; - - if (eveny) - { - sum = (evenFlag * co3) * s_srcPatch[lsizey-7][1 + ((tidx - 2) >> 1)]; - sum = sum + ( oddFlag * co2) * s_srcPatch[lsizey-7][1 + ((tidx - 1) >> 1)]; - sum = sum + (evenFlag * co1) * s_srcPatch[lsizey-7][1 + ((tidx ) >> 1)]; - sum = sum + ( oddFlag * co2) * s_srcPatch[lsizey-7][1 + ((tidx + 1) >> 1)]; - sum = sum + (evenFlag * co3) * s_srcPatch[lsizey-7][1 + ((tidx + 2) >> 1)]; - } - s_dstPatch[4 + tidy][tidx] = sum; - } - barrier(CLK_LOCAL_MEM_FENCE); - sum = co3 * s_dstPatch[2 + tidy - 2][tidx]; - sum = sum + co2 * s_dstPatch[2 + tidy - 1][tidx]; - sum = sum + co1 * s_dstPatch[2 + tidy ][tidx]; - sum = sum + co2 * s_dstPatch[2 + tidy + 1][tidx]; - sum = sum + co3 * s_dstPatch[2 + tidy + 2][tidx]; + sum = coefy2* s_dstPatch[1 + ((tidy - 1) >> 1)][tidx]; + sum = mad(coefy1, s_dstPatch[1 + ((tidy ) >> 1)][tidx], sum); + sum = mad(coefy2, s_dstPatch[1 + ((tidy + 2) >> 1)][tidx], sum); if ((x < dst_cols) && (y < dst_rows)) storepix(convertToT(sum), dstData + y * dst_step + x * PIXSIZE); @@ -177,12 +135,12 @@ __kernel void pyrUp_unrolled(__global const uchar * src, int src_step, int src_o const int ly = 2*get_local_id(1); __local FT s_srcPatch[LOCAL_SIZE+2][LOCAL_SIZE+2]; - __local FT s_dstPatch[2*LOCAL_SIZE+4][2*LOCAL_SIZE]; + __local FT s_dstPatch[LOCAL_SIZE+2][2*LOCAL_SIZE]; __global uchar * dstData = dst + dst_offset; __global const uchar * srcData = src + src_offset; - if( lx < (LOCAL_SIZE+2) && lx < (LOCAL_SIZE+2) ) + if( lx < (LOCAL_SIZE+2) && ly < (LOCAL_SIZE+2) ) { int srcx = mad24((int)get_group_id(0), LOCAL_SIZE, lx) - 1; int srcy = mad24((int)get_group_id(1), LOCAL_SIZE, ly) - 1; @@ -209,17 +167,13 @@ __kernel void pyrUp_unrolled(__global const uchar * src, int src_step, int src_o sum = co3 * s_srcPatch[1 + (ly >> 1)][1 + ((lx - 2) >> 1)]; sum = sum + co1 * s_srcPatch[1 + (ly >> 1)][1 + ((lx ) >> 1)]; sum = sum + co3 * s_srcPatch[1 + (ly >> 1)][1 + ((lx + 2) >> 1)]; - - s_dstPatch[2 + ly][lx] = sum; + + s_dstPatch[1 + get_local_id(1)][lx] = sum; // (x+1,y) sum = co2 * s_srcPatch[1 + (ly >> 1)][1 + ((lx + 1 - 1) >> 1)]; sum = sum + co2 * s_srcPatch[1 + (ly >> 1)][1 + ((lx + 1 + 1) >> 1)]; - s_dstPatch[2 + ly][lx+1] = sum; - - // (x, y+1) (x+1, y+1) - s_dstPatch[2 + ly+1][lx] = 0.f; - s_dstPatch[2 + ly+1][lx+1] = 0.f; + s_dstPatch[1 + get_local_id(1)][lx+1] = sum; if (ly < 1) { @@ -227,16 +181,12 @@ __kernel void pyrUp_unrolled(__global const uchar * src, int src_step, int src_o sum = co3 * s_srcPatch[0][1 + ((lx - 2) >> 1)]; sum = sum + co1 * s_srcPatch[0][1 + ((lx ) >> 1)]; sum = sum + co3 * s_srcPatch[0][1 + ((lx + 2) >> 1)]; - s_dstPatch[ly][lx] = sum; - + s_dstPatch[0][lx] = sum; + // (x+1,y) sum = co2 * s_srcPatch[0][1 + ((lx + 1 - 1) >> 1)]; sum = sum + co2 * s_srcPatch[0][1 + ((lx + 1 + 1) >> 1)]; - s_dstPatch[ly][lx+1] = sum; - - // (x, y+1) (x+1, y+1) - s_dstPatch[ly+1][lx] = 0.f; - s_dstPatch[ly+1][lx+1] = 0.f; + s_dstPatch[0][lx+1] = sum; } if (ly > 2*LOCAL_SIZE-3) @@ -245,59 +195,40 @@ __kernel void pyrUp_unrolled(__global const uchar * src, int src_step, int src_o sum = co3 * s_srcPatch[LOCAL_SIZE+1][1 + ((lx - 2) >> 1)]; sum = sum + co1 * s_srcPatch[LOCAL_SIZE+1][1 + ((lx ) >> 1)]; sum = sum + co3 * s_srcPatch[LOCAL_SIZE+1][1 + ((lx + 2) >> 1)]; - s_dstPatch[4 + ly][lx] = sum; + s_dstPatch[LOCAL_SIZE+1][lx] = sum; // (x+1,y) sum = co2 * s_srcPatch[LOCAL_SIZE+1][1 + ((lx + 1 - 1) >> 1)]; sum = sum + co2 * s_srcPatch[LOCAL_SIZE+1][1 + ((lx + 1 + 1) >> 1)]; - s_dstPatch[4 + ly][lx+1] = sum; - - // (x, y+1) (x+1, y+1) - s_dstPatch[4 + ly+1][lx] = 0.f; - s_dstPatch[4 + ly+1][lx+1] = 0.f; + s_dstPatch[LOCAL_SIZE+1][lx+1] = sum; } barrier(CLK_LOCAL_MEM_FENCE); int dst_x = 2*get_global_id(0); int dst_y = 2*get_global_id(1); - - // (x,y) - sum = co3 * s_dstPatch[2 + ly - 2][lx]; - sum = sum + co2 * s_dstPatch[2 + ly - 1][lx]; - sum = sum + co1 * s_dstPatch[2 + ly ][lx]; - sum = sum + co2 * s_dstPatch[2 + ly + 1][lx]; - sum = sum + co3 * s_dstPatch[2 + ly + 2][lx]; if ((dst_x < dst_cols) && (dst_y < dst_rows)) + { + // (x,y) + sum = co3 * s_dstPatch[1 + get_local_id(1) - 1][lx]; + sum = sum + co1 * s_dstPatch[1 + get_local_id(1) ][lx]; + sum = sum + co3 * s_dstPatch[1 + get_local_id(1) + 1][lx]; storepix(convertToT(sum), dstData + dst_y * dst_step + dst_x * PIXSIZE); - // (x+1,y) - sum = co3 * s_dstPatch[2 + ly - 2][lx+1]; - sum = sum + co2 * s_dstPatch[2 + ly - 1][lx+1]; - sum = sum + co1 * s_dstPatch[2 + ly ][lx+1]; - sum = sum + co2 * s_dstPatch[2 + ly + 1][lx+1]; - sum = sum + co3 * s_dstPatch[2 + ly + 2][lx+1]; - - if ((dst_x+1 < dst_cols) && (dst_y < dst_rows)) + // (x+1,y) + sum = co3 * s_dstPatch[1 + get_local_id(1) - 1][lx+1]; + sum = sum + co1 * s_dstPatch[1 + get_local_id(1) ][lx+1]; + sum = sum + co3 * s_dstPatch[1 + get_local_id(1) + 1][lx+1]; storepix(convertToT(sum), dstData + dst_y * dst_step + (dst_x+1) * PIXSIZE); - // (x,y+1) - sum = co3 * s_dstPatch[2 + ly+1 - 2][lx]; - sum = sum + co2 * s_dstPatch[2 + ly+1 - 1][lx]; - sum = sum + co1 * s_dstPatch[2 + ly+1 ][lx]; - sum = sum + co2 * s_dstPatch[2 + ly+1 + 1][lx]; - sum = sum + co3 * s_dstPatch[2 + ly+1 + 2][lx]; - - if ((dst_x < dst_cols) && (dst_y+1 < dst_rows)) + // (x,y+1) + sum = co2 * s_dstPatch[1 + get_local_id(1) ][lx]; + sum = sum + co2 * s_dstPatch[1 + get_local_id(1) + 1][lx]; storepix(convertToT(sum), dstData + (dst_y+1) * dst_step + dst_x * PIXSIZE); - // (x+1,y+1) - sum = co3 * s_dstPatch[2 + ly+1 - 2][lx+1]; - sum = sum + co2 * s_dstPatch[2 + ly+1 - 1][lx+1]; - sum = sum + co1 * s_dstPatch[2 + ly+1 ][lx+1]; - sum = sum + co2 * s_dstPatch[2 + ly+1 + 1][lx+1]; - sum = sum + co3 * s_dstPatch[2 + ly+1 + 2][lx+1]; - - if ((dst_x+1 < dst_cols) && (dst_y+1 < dst_rows)) + // (x+1,y+1) + sum = co2 * s_dstPatch[1 + get_local_id(1) ][lx+1]; + sum = sum + co2 * s_dstPatch[1 + get_local_id(1) + 1][lx+1]; storepix(convertToT(sum), dstData + (dst_y+1) * dst_step + (dst_x+1) * PIXSIZE); + } } diff --git a/modules/imgproc/src/pyramids.cpp b/modules/imgproc/src/pyramids.cpp index 319ff82000..1d51d9412e 100644 --- a/modules/imgproc/src/pyramids.cpp +++ b/modules/imgproc/src/pyramids.cpp @@ -467,7 +467,7 @@ static bool ocl_pyrUp( InputArray _src, OutputArray _dst, const Size& _dsz, int UMat dst = _dst.getUMat(); int float_depth = depth == CV_64F ? CV_64F : CV_32F; - int local_size = 8; + const int local_size = 16; char cvt[2][50]; String buildOptions = format( "-D T=%s -D FT=%s -D convertToT=%s -D convertToFT=%s%s " @@ -478,25 +478,21 @@ static bool ocl_pyrUp( InputArray _src, OutputArray _dst, const Size& _dsz, int doubleSupport ? " -D DOUBLE_SUPPORT" : "", ocl::typeToStr(depth), channels, local_size ); - size_t globalThreads[2]; + size_t globalThreads[2] = { dst.cols, dst.rows }; + size_t localThreads[2] = { local_size, local_size }; ocl::Kernel k; if (ocl::Device::getDefault().isIntel() && channels == 1) { k.create("pyrUp_unrolled", ocl::imgproc::pyr_up_oclsrc, buildOptions); globalThreads[0] = dst.cols/2; globalThreads[1] = dst.rows/2; - } + } else - { k.create("pyrUp", ocl::imgproc::pyr_up_oclsrc, buildOptions); - local_size = 16; - globalThreads[0] = dst.cols; globalThreads[1] = dst.rows; - } + if (k.empty()) return false; k.args(ocl::KernelArg::ReadOnly(src), ocl::KernelArg::WriteOnly(dst)); - size_t localThreads[2] = {local_size, local_size}; - return k.run(2, globalThreads, localThreads, false); }