If lut table has one channel and src aligned to 4, work with src as with one channel...
authorVBystricky <vladimir.bystricky@stdutility.com>
Thu, 5 Jun 2014 15:31:31 +0000 (19:31 +0400)
committerVBystricky <vladimir.bystricky@stdutility.com>
Thu, 5 Jun 2014 15:31:31 +0000 (19:31 +0400)
modules/core/src/convert.cpp
modules/core/src/opencl/lut.cl

index 162eaac..49eb93a 100644 (file)
@@ -1548,22 +1548,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 == dcn) && (0 == (src.offset % 4)) && (0 == (src.cols % 4));
+    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%s", dcn, lcn,
-                         ocl::typeToStr(src.depth()), ocl::memopTypeToStr(ddepth),
-                         bAligned ? " -D USE_ALIGNED" : ""
+                  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 + 3) / 4};
-    if (bAligned)
-        globalSize[0] = (dst.cols + 3) / 4;
+    size_t globalSize[2] = { cols, (dst.rows + 3) / 4 };
     return k.run(2, globalSize, NULL, false);
 }
 
index 295f0ae..a33d50c 100644 (file)
             dst[0] = lut_l[idx->x];\
             dst[1] = lut_l[idx->y];
     #elif dcn == 1
-        #ifdef USE_ALIGNED
-            #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];
-        #else
-            #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];
-        #endif
+        #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));\
@@ -136,11 +126,7 @@ __kernel void LUT(__global const uchar * srcptr, int src_step, int src_offset,
     __local dstT lut_l[256 * lcn];
     LOCAL_LUT_INIT;
 
-#ifdef USE_ALIGNED
-    int x = 4 * get_global_id(0);
-#else
     int x = get_global_id(0);
-#endif
     int y = 4 * get_global_id(1);
 
     if (x < cols && y < rows)