Unroll pyrUp kernel

diff --git a/modules/imgproc/src/opencl/pyr_up.cl b/modules/imgproc/src/opencl/pyr_up.cl
index d754a70..f9b5c8f 100644 (file)
--- 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 42464c1..a0a09ec 100644 (file)
--- 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);
 }