From 10c772fa7f899b9139a32f3cf887b83e536f3793 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Wed, 2 Jul 2014 17:06:34 +0400 Subject: [PATCH] minor optimization of cv::LUT --- modules/core/src/convert.cpp | 14 ++-- modules/core/src/opencl/lut.cl | 138 +++++++++++++++++---------------- 2 files changed, 76 insertions(+), 76 deletions(-) diff --git a/modules/core/src/convert.cpp b/modules/core/src/convert.cpp index 7b2684c85..21d5bdaca 100644 --- a/modules/core/src/convert.cpp +++ b/modules/core/src/convert.cpp @@ -1729,22 +1729,18 @@ 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 + int kercn = lcn == 1 ? std::min(4, ocl::predictOptimalVectorWidth(_dst)) : dcn; ocl::Kernel k("LUT", ocl::core::lut_oclsrc, - format("-D dcn=%d -D lcn=%d -D srcT=%s -D dstT=%s", bAligned ? 4 : dcn, lcn, - ocl::typeToStr(src.depth()), ocl::memopTypeToStr(ddepth) - )); + format("-D dcn=%d -D lcn=%d -D srcT=%s -D dstT=%s", kercn, 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::WriteOnlyNoSize(dst), dst.rows, cols); + ocl::KernelArg::WriteOnly(dst, dcn, kercn)); - size_t globalSize[2] = { cols, (dst.rows + 3) / 4 }; + size_t globalSize[2] = { dst.cols * dcn / kercn, (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 9bcd1b66f..81ca4349a 100644 --- a/modules/core/src/opencl/lut.cl +++ b/modules/core/src/opencl/lut.cl @@ -36,114 +36,118 @@ #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];\ + #define LUT_OP \ + int idx = *(__global const int *)(srcptr + src_index); \ + dst = (__global dstT *)(dstptr + 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];\ + #define LUT_OP \ + uchar3 idx = vload3(0, srcptr + src_index); \ + dst = (__global dstT *)(dstptr + 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)\ - short idx = *(__global const short *)(srcptr + mad24(num, src_step, src_index));\ - dst = (__global dstT *)(dstptr + mad24(num, dst_step, dst_index));\ - dst[0] = lut_l[idx & 0xff];\ + #define LUT_OP \ + short idx = *(__global const short *)(srcptr + src_index); \ + dst = (__global dstT *)(dstptr + dst_index); \ + dst[0] = lut_l[idx & 0xff]; \ dst[1] = lut_l[(idx >> 8) & 0xff]; #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));\ + #define LUT_OP \ + uchar idx = (srcptr + src_index)[0]; \ + dst = (__global dstT *)(dstptr + dst_index); \ dst[0] = lut_l[idx]; #else - #define LUT_OP(num)\ - __global const srcT * 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)\ + #define LUT_OP \ + __global const srcT * src = (__global const srcT *)(srcptr + src_index); \ + dst = (__global dstT *)(dstptr + 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];\ + #define LUT_OP \ + __global const uchar4 * src_pixel = (__global const uchar4 *)(srcptr + src_index); \ + int4 idx = mad24(convert_int4(src_pixel[0]), (int4)(lcn), (int4)(0, 1, 2, 3)); \ + 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]; #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];\ + #define LUT_OP \ + uchar3 src_pixel = vload3(0, srcptr + src_index); \ + int3 idx = mad24(convert_int3(src_pixel), (int3)(lcn), (int3)(0, 1, 2)); \ + dst = (__global dstT *)(dstptr + 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];\ + #define LUT_OP \ + __global const uchar2 * src_pixel = (__global const uchar2 *)(srcptr + src_index); \ + int2 idx = mad24(convert_int2(src_pixel[0]), lcn, (int2)(0, 1)); \ + dst = (__global dstT *)(dstptr + 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));\ + #define LUT_OP \ + uchar idx = (srcptr + src_index)[0]; \ + dst = (__global dstT *)(dstptr + dst_index); \ dst[0] = lut_l[idx]; #else - #define LUT_OP(num)\ - __global const srcT *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)\ + #define LUT_OP \ + __global const srcT * src = (__global const srcT *)(srcptr + src_index); \ + dst = (__global dstT *)(dstptr + 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 = 4 * get_global_id(1); + int y = get_global_id(1) << 2; + + __local dstT lut_l[256 * lcn]; + __global const dstT * lut = (__global const dstT *)(lutptr + lut_offset); + + for (int i = mad24((int)get_local_id(1), (int)get_local_size(0), (int)get_local_id(0)), + step = get_local_size(0) * get_local_size(1); i < 256 * lcn; i += step) + lut_l[i] = lut[i]; + 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 dstT * dst; - LUT_OP(0); + + LUT_OP; + if (y < rows - 1) { - LUT_OP(1); + src_index += src_step; + dst_index += dst_step; + LUT_OP; + if (y < rows - 2) { - LUT_OP(2); + src_index += src_step; + dst_index += dst_step; + LUT_OP; + if (y < rows - 3) { - LUT_OP(3); + src_index += src_step; + dst_index += dst_step; + LUT_OP; } } }