mirror/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2025-01-18 22:44:02 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

minor optimization of cv::LUT

Browse Source
This commit is contained in:
Ilya Lavrenov 2014-07-02 17:06:34 +04:00
parent 15197eb09b
commit 10c772fa7f
2 changed files with 76 additions and 76 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

14
modules/core/src/convert.cpp
View File

@ -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);
}

138
modules/core/src/opencl/lut.cl
View File

@ -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;
}
}
}