generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

If lut table has one channel and src aligned to 4, work with src as with one channel matrix

Browse Source
This commit is contained in:
VBystricky 2014-06-05 19:31:31 +04:00
parent 0c0ebca855
commit 5d924b7a75
2 changed files with 13 additions and 27 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -1548,22 +1548,22 @@ static bool ocl_LUT(InputArray _src, InputArray _lut, OutputArray _dst)
UMat src = _src.getUMat(), lut = _lut.getUMat(); UMat src = _src.getUMat(), lut = _lut.getUMat();
_dst.create(src.size(), CV_MAKETYPE(ddepth, dcn)); _dst.create(src.size(), CV_MAKETYPE(ddepth, dcn));
UMat dst = _dst.getUMat(); 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, ocl::Kernel k("LUT", ocl::core::lut_oclsrc,
format("-D dcn=%d -D lcn=%d -D srcT=%s -D dstT=%s%s", dcn, lcn, format("-D dcn=%d -D lcn=%d -D srcT=%s -D dstT=%s", bAligned ? 4 : dcn, lcn,
ocl::typeToStr(src.depth()), ocl::memopTypeToStr(ddepth), ocl::typeToStr(src.depth()), ocl::memopTypeToStr(ddepth)
bAligned ? " -D USE_ALIGNED" : ""
)); ));
if (k.empty()) if (k.empty())
return false; return false;
k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::ReadOnlyNoSize(lut), int cols = bAligned ? dcn * dst.cols / 4 : dst.cols;
ocl::KernelArg::WriteOnly(dst));
size_t globalSize[2] = { dst.cols, (dst.rows + 3) / 4}; k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::ReadOnlyNoSize(lut),
if (bAligned) ocl::KernelArg::WriteOnlyNoSize(dst), dst.rows, cols);
globalSize[0] = (dst.cols + 3) / 4;
size_t globalSize[2] = { cols, (dst.rows + 3) / 4 };
return k.run(2, globalSize, NULL, false); return k.run(2, globalSize, NULL, false);
} }

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

@ -57,20 +57,10 @@
dst[0] = lut_l[idx->x];\ dst[0] = lut_l[idx->x];\
dst[1] = lut_l[idx->y]; dst[1] = lut_l[idx->y];
#elif dcn == 1 #elif dcn == 1
#ifdef USE_ALIGNED #define LUT_OP(num)\
#define LUT_OP(num)\ uchar idx = (srcptr + mad24(num, src_step, src_index))[0];\
int idx = *(__global const int *)(srcptr + mad24(num, src_step, src_index));\ dst = (__global dstT *)(dstptr + mad24(num, dst_step, dst_index));\
dst = (__global dstT *)(dstptr + mad24(num, dst_step, dst_index));\ dst[0] = lut_l[idx];
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
#else #else
#define LUT_OP(num)\ #define LUT_OP(num)\
src = (__global const srcT *)(srcptr + mad24(num, src_step, src_index));\ 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 dstT lut_l[256 * lcn];
LOCAL_LUT_INIT; LOCAL_LUT_INIT;
#ifdef USE_ALIGNED
int x = 4 * get_global_id(0);
#else
int x = get_global_id(0); int x = get_global_id(0);
#endif
int y = 4 * get_global_id(1); int y = 4 * get_global_id(1);
if (x < cols && y < rows) if (x < cols && y < rows)