From 5d924b7a75d9f93aa57371474ad6e40035c41ea7 Mon Sep 17 00:00:00 2001 From: VBystricky Date: Thu, 5 Jun 2014 19:31:31 +0400 Subject: [PATCH] If lut table has one channel and src aligned to 4, work with src as with one channel matrix --- modules/core/src/convert.cpp | 16 ++++++++-------- modules/core/src/opencl/lut.cl | 22 ++++------------------ 2 files changed, 12 insertions(+), 26 deletions(-) diff --git a/modules/core/src/convert.cpp b/modules/core/src/convert.cpp index 162eaac..49eb93a 100644 --- a/modules/core/src/convert.cpp +++ b/modules/core/src/convert.cpp @@ -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); } diff --git a/modules/core/src/opencl/lut.cl b/modules/core/src/opencl/lut.cl index 295f0ae..a33d50c 100644 --- a/modules/core/src/opencl/lut.cl +++ b/modules/core/src/opencl/lut.cl @@ -57,20 +57,10 @@ 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) -- 2.7.4