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

Read 4 pixel for aligned data with 1 channel

Browse Source
This commit is contained in:
VBystricky 2014-06-04 23:50:23 +04:00
parent 8a5f2781fc
commit 0c0ebca855
2 changed files with 24 additions and 6 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -1548,10 +1548,12 @@ 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));
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", dcn, lcn, format("-D dcn=%d -D lcn=%d -D srcT=%s -D dstT=%s%s", 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;
@ -1560,6 +1562,8 @@ static bool ocl_LUT(InputArray _src, InputArray _lut, OutputArray _dst)
ocl::KernelArg::WriteOnly(dst)); ocl::KernelArg::WriteOnly(dst));
size_t globalSize[2] = { dst.cols, (dst.rows + 3) / 4}; size_t globalSize[2] = { dst.cols, (dst.rows + 3) / 4};
if (bAligned)
globalSize[0] = (dst.cols + 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,10 +57,20 @@
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
#define LUT_OP(num)\ #ifdef USE_ALIGNED
uchar idx = (srcptr + mad24(num, src_step, src_index))[0];\ #define LUT_OP(num)\
dst = (__global dstT *)(dstptr + mad24(num, dst_step, dst_index));\ int idx = *(__global const int *)(srcptr + mad24(num, src_step, src_index));\
dst[0] = lut_l[idx]; 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];\
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));\
@ -126,7 +136,11 @@ __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)