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);
}
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));\
__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)