Merge pull request #2755 from vbystricky:oclopt_LUT

modules/core/src/convert.cpp

@@ -1559,17 +1559,22 @@ static bool ocl_LUT(InputArray _src, InputArray _lut, OutputArray _dst)
     UMat src = _src.getUMat(), lut = _lut.getUMat();
     _dst.create(src.size(), CV_MAKETYPE(ddepth, dcn));
     UMat dst = _dst.getUMat();
+    bool bAligned = (1 == lcn) && (0 == (src.offset % 4)) && (0 == ((dcn * src.cols) % 4));
+    // dst.cols == src.cols by params of dst.create
 
     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(src.depth()), ocl::memopTypeToStr(ddepth)));
+                  format("-D dcn=%d -D lcn=%d -D srcT=%s -D dstT=%s", bAligned ? 4 : dcn, lcn,
+                         ocl::typeToStr(src.depth()), ocl::memopTypeToStr(ddepth)
+                         ));
     if (k.empty())
         return false;
 
+    int cols = bAligned ? dcn * dst.cols / 4 : dst.cols;
+
     k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::ReadOnlyNoSize(lut),
-           ocl::KernelArg::WriteOnly(dst));
+        ocl::KernelArg::WriteOnlyNoSize(dst), dst.rows, cols);
 
-    size_t globalSize[2] = { dst.cols, dst.rows };
+    size_t globalSize[2] = { cols, (dst.rows + 3) / 4 };
     return k.run(2, globalSize, NULL, false);
 }
 

modules/core/src/opencl/lut.cl

@@ -34,30 +34,120 @@
 //
 //
 
+#if lcn == 1
+    #if dcn == 4
+        #define LUT_OP(num)\
+            int idx = *(__global const int *)(srcptr + mad24(num, src_step, src_index));\
+            dst = (__global dstT *)(dstptr + mad24(num, dst_step, dst_index));\
+            dst[0] = lut_l[idx & 0xff];\
+            dst[1] = lut_l[(idx >> 8) & 0xff];\
+            dst[2] = lut_l[(idx >> 16) & 0xff];\
+            dst[3] = lut_l[(idx >> 24) & 0xff];
+    #elif dcn == 3
+        #define LUT_OP(num)\
+            uchar3 idx = vload3(0, srcptr + mad24(num, src_step, src_index));\
+            dst = (__global dstT *)(dstptr + mad24(num, dst_step, dst_index));\
+            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)\
+            __global const uchar2 * idx = (__global const uchar2 *)(srcptr + mad24(num, src_step, src_index));\
+            dst = (__global dstT *)(dstptr + mad24(num, dst_step, dst_index));\
+            dst[0] = lut_l[idx->x];\
+            dst[1] = lut_l[idx->y];
+    #elif dcn == 1
+        #define LUT_OP(num)\
+            uchar idx = (srcptr + mad24(num, src_step, src_index))[0];\
+            dst = (__global dstT *)(dstptr + mad24(num, dst_step, dst_index));\
+            dst[0] = lut_l[idx];
+    #else
+        #define LUT_OP(num)\
+            src = (__global const srcT *)(srcptr + mad24(num, src_step, src_index));\
+            dst = (__global dstT *)(dstptr + mad24(num, dst_step, dst_index));\
+            for (int cn = 0; cn < dcn; ++cn)\
+                dst[cn] = lut_l[src[cn]];
+    #endif
+#else
+    #if dcn == 4
+        #define LUT_OP(num)\
+            __global const uchar4 *src_pixel = (__global const uchar4 *)(srcptr + mad24(num, src_step, src_index));\
+            int4 idx = convert_int4(src_pixel[0]) * lcn + (int4)(0, 1, 2, 3);\
+            dst = (__global dstT *)(dstptr + mad24(num, dst_step, 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];
+    #elif dcn == 3
+        #define LUT_OP(num)\
+            uchar3 src_pixel = vload3(0, srcptr + mad24(num, src_step, src_index));\
+            int3 idx = convert_int3(src_pixel) * lcn + (int3)(0, 1, 2);\
+            dst = (__global dstT *)(dstptr + mad24(num, dst_step, dst_index));\
+            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)\
+            __global const uchar2 *src_pixel = (__global const uchar2 *)(srcptr + mad24(num, src_step, src_index));\
+            int2 idx = convert_int2(src_pixel[0]) * lcn + (int2)(0, 1);\
+            dst = (__global dstT *)(dstptr + mad24(num, dst_step, dst_index));\
+            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 = (srcptr + mad24(num, src_step, src_index))[0];\
+            dst = (__global dstT *)(dstptr + mad24(num, dst_step, dst_index));\
+            dst[0] = lut_l[idx];
+    #else
+        #define LUT_OP(num)\
+            src = (__global const srcT *)(srcptr + mad24(num, src_step, src_index));\
+            dst = (__global dstT *)(dstptr + mad24(num, dst_step, dst_index));\
+            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)
 {
+    __local dstT lut_l[256 * lcn];
+    LOCAL_LUT_INIT;
+
     int x = get_global_id(0);
-    int y = 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));
         int dst_index = mad24(y, dst_step, mad24(x, (int)sizeof(dstT) * dcn, dst_offset));
+        __global const srcT * src; __global dstT * dst;
+        int tmp_idx;
+        LUT_OP(0);
+        if (y < rows - 1)
+        {
+            LUT_OP(1);
+            if (y < rows - 2)
+            {
+                LUT_OP(2);
+                if (y < rows - 3)
+                {
+                    LUT_OP(3);
+                }
+            }
+        }
 
-        __global const srcT * src = (__global const srcT *)(srcptr + src_index);
-        __global const dstT * lut = (__global const dstT *)(lutptr + lut_offset);
-        __global dstT * dst = (__global dstT *)(dstptr + dst_index);
-
-#if lcn == 1
-        #pragma unroll
-        for (int cn = 0; cn < dcn; ++cn)
-            dst[cn] = lut[src[cn]];
-#else
-        #pragma unroll
-        for (int cn = 0; cn < dcn; ++cn)
-            dst[cn] = lut[mad24(src[cn], dcn, cn)];
-#endif
     }
 }