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

Merge pull request #2874 from vbystricky:ocl_LUTAMDfix

Browse Source
This commit is contained in:
Alexander Alekhin
2014-06-23 14:23:26 +00:00
parent 6952b90ed0 ea6da6ee95
commit 4f1d06d14a
1 changed files with 6 additions and 7 deletions
Download Patch File Download Diff File Expand all files Collapse all files

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

@@ -52,10 +52,10 @@
dst[2] = lut_l[idx.z]; dst[2] = lut_l[idx.z];
#elif dcn == 2 #elif dcn == 2
#define LUT_OP(num)\ #define LUT_OP(num)\
__global const uchar2 * idx = (__global const uchar2 *)(srcptr + mad24(num, src_step, src_index));\ short idx = *(__global const short *)(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->x];\ dst[0] = lut_l[idx & 0xff];\
dst[1] = lut_l[idx->y]; dst[1] = lut_l[(idx >> 8) & 0xff];
#elif dcn == 1 #elif dcn == 1
#define LUT_OP(num)\ #define LUT_OP(num)\
uchar idx = (srcptr + mad24(num, src_step, src_index))[0];\ uchar idx = (srcptr + mad24(num, src_step, src_index))[0];\
@@ -63,7 +63,7 @@
dst[0] = lut_l[idx]; dst[0] = lut_l[idx];
#else #else
#define LUT_OP(num)\ #define LUT_OP(num)\
src = (__global const srcT *)(srcptr + mad24(num, src_step, src_index));\ __global const srcT * src = (__global const srcT *)(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));\
for (int cn = 0; cn < dcn; ++cn)\ for (int cn = 0; cn < dcn; ++cn)\
dst[cn] = lut_l[src[cn]]; dst[cn] = lut_l[src[cn]];
@@ -100,7 +100,7 @@
dst[0] = lut_l[idx]; dst[0] = lut_l[idx];
#else #else
#define LUT_OP(num)\ #define LUT_OP(num)\
src = (__global const srcT *)(srcptr + mad24(num, src_step, src_index));\ __global const srcT *src = (__global const srcT *)(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));\
for (int cn = 0; cn < dcn; ++cn)\ for (int cn = 0; cn < dcn; ++cn)\
dst[cn] = lut_l[mad24(src[cn], lcn, cn)]; dst[cn] = lut_l[mad24(src[cn], lcn, cn)];
@@ -133,8 +133,7 @@ __kernel void LUT(__global const uchar * srcptr, int src_step, int src_offset,
{ {
int src_index = mad24(y, src_step, mad24(x, (int)sizeof(srcT) * dcn, src_offset)); 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)); int dst_index = mad24(y, dst_step, mad24(x, (int)sizeof(dstT) * dcn, dst_offset));
__global const srcT * src; __global dstT * dst; __global dstT * dst;
int tmp_idx;
LUT_OP(0); LUT_OP(0);
if (y < rows - 1) if (y < rows - 1)
{ {