Chiebot-Mirror
/
opencv
mirror of https://github.com/opencv/opencv.git
8
0
Fork 0
Code Issues Projects Releases Wiki Activity

Merge pull request #2927 from ilya-lavrenov:tapi_lut

Browse Source
pull/2969/head
Alexander Alekhin 11 years ago
parent 347eaa8685 10c772fa7f
commit a0816c6d15
2 changed files with 75 additions and 75 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Split View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 14
      modules/core/src/convert.cpp
  2. 136
      modules/core/src/opencl/lut.cl

14
modules/core/src/convert.cpp
Unescape 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);
}

136
modules/core/src/opencl/lut.cl
Unescape 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)
{
int x = get_global_id(0);
int y = get_global_id(1) << 2;
__local dstT lut_l[256 * lcn];
LOCAL_LUT_INIT;
__global const dstT * lut = (__global const dstT *)(lutptr + lut_offset);
int x = get_global_id(0);
int y = 4 * get_global_id(1);
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;
}
}
}

Write Preview
Loading…
Cancel
Save