diff --git a/modules/core/src/convert.cpp b/modules/core/src/convert.cpp index 43fdc00b2..6e474d951 100644 --- a/modules/core/src/convert.cpp +++ b/modules/core/src/convert.cpp @@ -1547,30 +1547,21 @@ static bool ocl_LUT(InputArray _src, InputArray _lut, OutputArray _dst) int sdepth = _src.depth(); UMat src = _src.getUMat(), lut = _lut.getUMat(); - int dtype = CV_MAKETYPE(ddepth, dcn); - _dst.create(src.size(), dtype); + _dst.create(src.size(), CV_MAKETYPE(ddepth, dcn)); UMat dst = _dst.getUMat(); - size_t globalSize[2] = { dst.cols, dst.rows / 2}; + size_t globalSize[2] = { dst.cols, (dst.rows + 3) / 4}; - cv::String build_opt = format("-D dcn=%d -D lcn=%d -D srcT=%s -D dstT=%s", dcn, lcn, + 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(sdepth), ocl::memopTypeToStr(ddepth) - ); - - ocl::Kernel kernel; - if ((4 == lcn) && (CV_8U == sdepth)) - kernel.create("LUTC4", ocl::core::lut_oclsrc, build_opt); - else if ((3 == lcn) && (CV_8U == sdepth)) - kernel.create("LUTC3", ocl::core::lut_oclsrc, build_opt); - else - kernel.create("LUT", ocl::core::lut_oclsrc, build_opt); - if (kernel.empty()) + )); + if (k.empty()) return false; - kernel.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::ReadOnlyNoSize(lut), + k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::ReadOnlyNoSize(lut), ocl::KernelArg::WriteOnly(dst)); - return kernel.run(2, globalSize, NULL, true); + return k.run(2, globalSize, NULL, false); } #endif diff --git a/modules/core/src/opencl/lut.cl b/modules/core/src/opencl/lut.cl index 9b0606145..27428ed2b 100644 --- a/modules/core/src/opencl/lut.cl +++ b/modules/core/src/opencl/lut.cl @@ -34,149 +34,120 @@ // // -__kernel void LUTC4(__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) -{ - int x = get_global_id(0); - int y = 2 * get_global_id(1); - - __global const dstT * lut = (__global const dstT *)(lutptr + lut_offset); - - __local dstT lut_l[256 * lcn]; - int init = mad24(get_local_id(1), get_local_size(0), 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 + 0] = lut[i + 0]; - } - 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 const uchar4 * src = (__global const uchar4 *)(srcptr + src_index); - int4 idx = convert_int4(src[0]) * lcn + (int4)(0, 1, 2, 3); - __global dstT * 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]; - - if (y < rows - 1) - { - src = (__global const uchar4 *)(srcptr + src_index + src_step); - idx = convert_int4(src[0]) * lcn + (int4)(0, 1, 2, 3); - dst = (__global dstT *)(dstptr + dst_index + dst_step); - - dst[0] = lut_l[idx.x]; - dst[1] = lut_l[idx.y]; - dst[2] = lut_l[idx.z]; +#if lcn == 1 + #if dcn == 4 + #define LUT_OP(num)\ + uchar4 idx = vload4(0, (__global const uchar *)(srcptr + src_index + num * src_step));\ + dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\ + dst[0] = lut_l[idx.x];\ + dst[1] = lut_l[idx.y];\ + dst[2] = lut_l[idx.z];\ dst[3] = lut_l[idx.w]; - } - } -} - -__kernel void LUTC3(__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) -{ - int x = get_global_id(0); - int y = 2 * get_global_id(1); - - __global const dstT * lut = (__global const dstT *)(lutptr + lut_offset); - - __local dstT lut_l[256 * lcn]; - int init = mad24(get_local_id(1), get_local_size(0), 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 + 0] = lut[i + 0]; - } - 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)); - - uchar3 src_pixel = vload3(0, (__global const uchar *)(srcptr + src_index)); - int3 idx = convert_int3(src_pixel) * lcn + (int3)(0, 1, 2); - __global dstT * dst = (__global dstT *)(dstptr + dst_index); - - dst[0] = lut_l[idx.x]; - dst[1] = lut_l[idx.y]; - dst[2] = lut_l[idx.z]; - if (y < rows - 1) - { - uchar3 src_pixel = vload3(0, (__global const uchar *)(srcptr + src_index + src_step)); - idx = convert_int3(src_pixel) * lcn + (int3)(0, 1, 2); - dst = (__global dstT *)(dstptr + dst_index + dst_step); - - dst[0] = lut_l[idx.x]; - dst[1] = lut_l[idx.y]; + #elif dcn == 3 + #define LUT_OP(num)\ + uchar3 idx = vload3(0, (__global const uchar *)(srcptr + src_index + num * src_step));\ + dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\ + 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)\ + uchar2 idx = vload2(0, (__global const uchar *)(srcptr + src_index + num * src_step));\ + dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\ + dst[0] = lut_l[idx.x];\ + dst[1] = lut_l[idx.y]; + #elif dcn == 1 + #define LUT_OP(num)\ + uchar idx = (__global const uchar *)(srcptr + src_index + num * src_step)[0];\ + dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\ + dst[0] = lut_l[idx]; + #else + #define LUT_OP(num)\ + src = (__global const srcT *)(srcptr + src_index + num * src_step);\ + dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\ + for (int cn = 0; cn < dcn; ++cn)\ + dst[cn] = lut_l[src[cn]]; + #endif +#else + #if dcn == 4 + #define LUT_OP(num)\ + uchar4 src_pixel = vload4(0, (__global const uchar *)(srcptr + src_index + num * src_step));\ + int4 idx = convert_int4(src_pixel) * lcn + (int4)(0, 1, 2, 3);\ + dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\ + 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, (__global const uchar *)(srcptr + src_index + num * src_step));\ + int3 idx = convert_int3(src_pixel) * lcn + (int3)(0, 1, 2);\ + dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\ + 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)\ + uchar2 src_pixel = vload2(0, (__global const uchar *)(srcptr + src_index + num * src_step));\ + int2 idx = convert_int2(src_pixel) * lcn + (int2)(0, 1);\ + dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\ + 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 = (__global const uchar *)(srcptr + src_index + num * src_step)[0];\ + dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\ + dst[0] = lut_l[idx]; + #else + #define LUT_OP(num)\ + src = (__global const srcT *)(srcptr + src_index + num * src_step);\ + dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\ + 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) { - __global const dstT * lut = (__global const dstT *)(lutptr + lut_offset); - __local dstT lut_l[256 * lcn]; - int init = mad24(get_local_id(1), get_local_size(0), 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 + 0] = lut[i + 0]; - } - barrier(CLK_LOCAL_MEM_FENCE); + LOCAL_LUT_INIT; int x = get_global_id(0); - int y = 2 * 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)); - - __global const srcT * src = (__global const srcT *)(srcptr + src_index); - __global const dstT * lut = (__global const dstT *)(lutptr + lut_offset); - int dst_index = mad24(y, dst_step, mad24(x, (int)sizeof(dstT) * dcn, dst_offset)); - __global dstT * dst = (__global dstT *)(dstptr + dst_index); + __global const srcT * src; __global dstT * dst; -#if lcn == 1 - #pragma unroll - for (int cn = 0; cn < dcn; ++cn) - dst[cn] = lut_l[src[cn]]; -#else //lcn == scn == dcn - #pragma unroll - for (int cn = 0; cn < dcn; ++cn) - dst[cn] = lut_l[mad24(src[cn], lcn, cn)]; -#endif + LUT_OP(0); if (y < rows - 1) { - src = (__global const srcT *)(srcptr + src_index + src_step); - dst = (__global dstT *)(dstptr + dst_index + dst_step); - -#if lcn == 1 - #pragma unroll - for (int cn = 0; cn < dcn; ++cn) - dst[cn] = lut_l[src[cn]]; -#else //lcn == scn == dcn - #pragma unroll - for (int cn = 0; cn < dcn; ++cn) - dst[cn] = lut_l[mad24(src[cn], lcn, cn)]; -#endif + LUT_OP(1); + if (y < rows - 2) + { + LUT_OP(2); + if (y < rows - 3) + { + LUT_OP(3); + } + } } + } }