Use 4 pixels for one unit. Some ocl code refactoring

diff --git a/modules/core/src/convert.cpp b/modules/core/src/convert.cpp
index 43fdc00..6e474d9 100644 (file)
--- a/modules/core/src/convert.cpp
+++ b/modules/core/src/convert.cpp
@@ -1547,30 +1547,21 @@ static bool ocl_LUT(InputArray _src, InputArray _lut, OutputArray _dst)
     int sdepth = _src.depth();
 
     UMat src = _src.getUMat(), lut = _lut.getUMat();
-    int dtype = CV_MAKETYPE(ddepth, dcn);
-    _dst.create(src.size(), dtype);
+    _dst.create(src.size(), CV_MAKETYPE(ddepth, dcn));
     UMat dst = _dst.getUMat();
 
-    size_t globalSize[2] = { dst.cols, dst.rows / 2};
+    size_t globalSize[2] = { dst.cols, (dst.rows + 3) / 4};
 
-    cv::String build_opt = format("-D dcn=%d -D lcn=%d -D srcT=%s -D dstT=%s", dcn, lcn,
+    ocl::Kernel k("LUT", ocl::core::lut_oclsrc, format("-D dcn=%d -D lcn=%d -D srcT=%s -D dstT=%s", dcn, lcn,
                          ocl::typeToStr(sdepth), ocl::memopTypeToStr(ddepth)
-                         );
-
-    ocl::Kernel kernel;
-    if ((4 == lcn) && (CV_8U == sdepth))
-        kernel.create("LUTC4", ocl::core::lut_oclsrc, build_opt);
-    else if ((3 == lcn) && (CV_8U == sdepth))
-        kernel.create("LUTC3", ocl::core::lut_oclsrc, build_opt);
-    else
-        kernel.create("LUT", ocl::core::lut_oclsrc, build_opt);
-    if (kernel.empty())
+                         ));
+    if (k.empty())
         return false;
 
-    kernel.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::ReadOnlyNoSize(lut),
+    k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::ReadOnlyNoSize(lut),
            ocl::KernelArg::WriteOnly(dst));
 
-    return kernel.run(2, globalSize, NULL, true);
+    return k.run(2, globalSize, NULL, false);
 }
 
 #endif

diff --git a/modules/core/src/opencl/lut.cl b/modules/core/src/opencl/lut.cl
index 9b06061..27428ed 100644 (file)
--- a/modules/core/src/opencl/lut.cl
+++ b/modules/core/src/opencl/lut.cl
@@ -34,149 +34,120 @@
 //
 //
 
-__kernel void LUTC4(__global const uchar * srcptr, int src_step, int src_offset,
-                  __global const uchar * lutptr, int lut_step, int lut_offset,
-                  __global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols)
-{
-    int x = get_global_id(0);
-    int y = 2 * get_global_id(1);
-
-    __global const dstT * lut = (__global const dstT *)(lutptr + lut_offset);
-
-    __local dstT lut_l[256 * lcn];
-    int init = mad24(get_local_id(1), get_local_size(0), get_local_id(0));
-    int step = get_local_size(0) * get_local_size(1);
-
-    for (int i = init; i < 256 * lcn; i += step)
-    {
-        lut_l[i + 0] = lut[i + 0];
-    }
-    barrier(CLK_LOCAL_MEM_FENCE);
-
-    if (x < cols && y < rows)
-    {
-        int src_index = mad24(y, src_step, mad24(x, (int)sizeof(srcT) * dcn, src_offset));
-        int dst_index = mad24(y, dst_step, mad24(x, (int)sizeof(dstT) * dcn, dst_offset));
-
-        __global const uchar4 * src = (__global const uchar4 *)(srcptr + src_index);
-        int4 idx = convert_int4(src[0]) * lcn + (int4)(0, 1, 2, 3);
-        __global dstT * dst = (__global dstT *)(dstptr + dst_index);
-
-        dst[0] = lut_l[idx.x];
-        dst[1] = lut_l[idx.y];
-        dst[2] = lut_l[idx.z];
-        dst[3] = lut_l[idx.w];
-
-        if (y < rows - 1)
-        {
-            src = (__global const uchar4 *)(srcptr + src_index + src_step);
-            idx = convert_int4(src[0]) * lcn + (int4)(0, 1, 2, 3);
-            dst = (__global dstT *)(dstptr + dst_index + dst_step);
-
-            dst[0] = lut_l[idx.x];
-            dst[1] = lut_l[idx.y];
-            dst[2] = lut_l[idx.z];
+#if lcn == 1
+    #if dcn == 4
+        #define LUT_OP(num)\
+            uchar4 idx = vload4(0, (__global const uchar *)(srcptr + src_index  + num * src_step));\
+            dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\
+            dst[0] = lut_l[idx.x];\
+            dst[1] = lut_l[idx.y];\
+            dst[2] = lut_l[idx.z];\
             dst[3] = lut_l[idx.w];
-        }
-    }
-}
-
-__kernel void LUTC3(__global const uchar * srcptr, int src_step, int src_offset,
-                  __global const uchar * lutptr, int lut_step, int lut_offset,
-                  __global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols)
-{
-    int x = get_global_id(0);
-    int y = 2 * get_global_id(1);
-
-    __global const dstT * lut = (__global const dstT *)(lutptr + lut_offset);
-
-    __local dstT lut_l[256 * lcn];
-    int init = mad24(get_local_id(1), get_local_size(0), get_local_id(0));
-    int step = get_local_size(0) * get_local_size(1);
-
-    for (int i = init; i < 256 * lcn; i += step)
-    {
-        lut_l[i + 0] = lut[i + 0];
-    }
-    barrier(CLK_LOCAL_MEM_FENCE);
-
-    if (x < cols && y < rows)
-    {
-        int src_index = mad24(y, src_step, mad24(x, (int)sizeof(srcT) * dcn, src_offset));
-        int dst_index = mad24(y, dst_step, mad24(x, (int)sizeof(dstT) * dcn, dst_offset));
-
-        uchar3 src_pixel = vload3(0, (__global const uchar *)(srcptr + src_index));
-        int3 idx = convert_int3(src_pixel) * lcn + (int3)(0, 1, 2);
-        __global dstT * dst = (__global dstT *)(dstptr + dst_index);
-
-        dst[0] = lut_l[idx.x];
-        dst[1] = lut_l[idx.y];
-        dst[2] = lut_l[idx.z];
-        if (y < rows - 1)
-        {
-            uchar3 src_pixel = vload3(0, (__global const uchar *)(srcptr + src_index + src_step));
-            idx = convert_int3(src_pixel) * lcn + (int3)(0, 1, 2);
-            dst = (__global dstT *)(dstptr + dst_index + dst_step);
-
-            dst[0] = lut_l[idx.x];
+    #elif dcn == 3
+        #define LUT_OP(num)\
+            uchar3 idx = vload3(0, (__global const uchar *)(srcptr + src_index  + num * src_step));\
+            dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\
+            dst[0] = lut_l[idx.x];\
+            dst[1] = lut_l[idx.y];\
+            dst[2] = lut_l[idx.z];
+    #elif dcn == 2
+        #define LUT_OP(num)\
+            uchar2 idx = vload2(0, (__global const uchar *)(srcptr + src_index  + num * src_step));\
+            dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\
+            dst[0] = lut_l[idx.x];\
             dst[1] = lut_l[idx.y];
+    #elif dcn == 1
+        #define LUT_OP(num)\
+            uchar idx = (__global const uchar *)(srcptr + src_index  + num * src_step)[0];\
+            dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\
+            dst[0] = lut_l[idx];
+    #else
+        #define LUT_OP(num)\
+            src = (__global const srcT *)(srcptr + src_index + num * src_step);\
+            dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\
+            for (int cn = 0; cn < dcn; ++cn)\
+                dst[cn] = lut_l[src[cn]];
+    #endif
+#else
+    #if dcn == 4
+        #define LUT_OP(num)\
+            uchar4 src_pixel = vload4(0, (__global const uchar *)(srcptr + src_index  + num * src_step));\
+            int4 idx = convert_int4(src_pixel) * lcn + (int4)(0, 1, 2, 3);\
+            dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\
+            dst[0] = lut_l[idx.x];\
+            dst[1] = lut_l[idx.y];\
+            dst[2] = lut_l[idx.z];\
+            dst[3] = lut_l[idx.w];
+    #elif dcn == 3
+        #define LUT_OP(num)\
+            uchar3 src_pixel = vload3(0, (__global const uchar *)(srcptr + src_index  + num * src_step));\
+            int3 idx = convert_int3(src_pixel) * lcn + (int3)(0, 1, 2);\
+            dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\
+            dst[0] = lut_l[idx.x];\
+            dst[1] = lut_l[idx.y];\
             dst[2] = lut_l[idx.z];
-        }
+    #elif dcn == 2
+        #define LUT_OP(num)\
+            uchar2 src_pixel = vload2(0, (__global const uchar *)(srcptr + src_index  + num * src_step));\
+            int2 idx = convert_int2(src_pixel) * lcn + (int2)(0, 1);\
+            dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\
+            dst[0] = lut_l[idx.x];\
+            dst[1] = lut_l[idx.y];
+    #elif dcn == 1 //error case (1 < lcn) ==> lcn == scn == dcn
+        #define LUT_OP(num)\
+            uchar idx = (__global const uchar *)(srcptr + src_index  + num * src_step)[0];\
+            dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\
+            dst[0] = lut_l[idx];
+    #else
+        #define LUT_OP(num)\
+            src = (__global const srcT *)(srcptr + src_index + num * src_step);\
+            dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\
+            for (int cn = 0; cn < dcn; ++cn)\
+                dst[cn] = lut_l[mad24(src[cn], lcn, cn)];
+    #endif
+#endif
+
+#define LOCAL_LUT_INIT\
+    {\
+        __global const dstT * lut = (__global const dstT *)(lutptr + lut_offset);\
+        int init = mad24((int)get_local_id(1), (int)get_local_size(0), (int)get_local_id(0));\
+        int step = get_local_size(0) * get_local_size(1);\
+        for (int i = init; i < 256 * lcn; i += step)\
+        {\
+            lut_l[i] = lut[i];\
+        }\
+        barrier(CLK_LOCAL_MEM_FENCE);\
     }
-}
 
 __kernel void LUT(__global const uchar * srcptr, int src_step, int src_offset,
                   __global const uchar * lutptr, int lut_step, int lut_offset,
                   __global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols)
 {
-    __global const dstT * lut = (__global const dstT *)(lutptr + lut_offset);
-
     __local dstT lut_l[256 * lcn];
-    int init = mad24(get_local_id(1), get_local_size(0), get_local_id(0));
-    int step = get_local_size(0) * get_local_size(1);
-
-    for (int i = init; i < 256 * lcn; i += step)
-    {
-        lut_l[i + 0] = lut[i + 0];
-    }
-    barrier(CLK_LOCAL_MEM_FENCE);
+    LOCAL_LUT_INIT;
 
     int x = get_global_id(0);
-    int y = 2 * get_global_id(1);
+    int y = 4 * get_global_id(1);
 
     if (x < cols && y < rows)
     {
         int src_index = mad24(y, src_step, mad24(x, (int)sizeof(srcT) * dcn, src_offset));
-
-        __global const srcT * src = (__global const srcT *)(srcptr + src_index);
-        __global const dstT * lut = (__global const dstT *)(lutptr + lut_offset);
-
         int dst_index = mad24(y, dst_step, mad24(x, (int)sizeof(dstT) * dcn, dst_offset));
-        __global dstT * dst = (__global dstT *)(dstptr + dst_index);
+        __global const srcT * src; __global dstT * dst;
 
-#if lcn == 1
-        #pragma unroll
-        for (int cn = 0; cn < dcn; ++cn)
-            dst[cn] = lut_l[src[cn]];
-#else //lcn == scn == dcn
-        #pragma unroll
-        for (int cn = 0; cn < dcn; ++cn)
-            dst[cn] = lut_l[mad24(src[cn], lcn, cn)];
-#endif
+        LUT_OP(0);
         if (y < rows - 1)
         {
-            src = (__global const srcT *)(srcptr + src_index + src_step);
-            dst = (__global dstT *)(dstptr + dst_index + dst_step);
-
-#if lcn == 1
-            #pragma unroll
-            for (int cn = 0; cn < dcn; ++cn)
-                dst[cn] = lut_l[src[cn]];
-#else //lcn == scn == dcn
-            #pragma unroll
-            for (int cn = 0; cn < dcn; ++cn)
-                dst[cn] = lut_l[mad24(src[cn], lcn, cn)];
-#endif
+            LUT_OP(1);
+            if (y < rows - 2)
+            {
+                LUT_OP(2);
+                if (y < rows - 3)
+                {
+                    LUT_OP(3);
+                }
+            }
         }
+
     }
 }