From 48d82fd911d379574cca55fec1286813ed2b8eb1 Mon Sep 17 00:00:00 2001 From: vbystricky Date: Mon, 26 May 2014 16:31:42 +0400 Subject: [PATCH] Fix some errors --- modules/core/src/opencl/lut.cl | 58 +++++++++++++++++----------------- 1 file changed, 29 insertions(+), 29 deletions(-) diff --git a/modules/core/src/opencl/lut.cl b/modules/core/src/opencl/lut.cl index 27428ed2b3..9646809e3b 100644 --- a/modules/core/src/opencl/lut.cl +++ b/modules/core/src/opencl/lut.cl @@ -37,71 +37,71 @@ #if lcn == 1 #if dcn == 4 #define LUT_OP(num)\ - uchar4 idx = vload4(0, (__global const uchar *)(srcptr + src_index + num * src_step));\ - dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\ - dst[0] = lut_l[idx.x];\ - dst[1] = lut_l[idx.y];\ - dst[2] = lut_l[idx.z];\ - dst[3] = lut_l[idx.w]; + __global const uchar4 *idx = (__global const uchar4 *)(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];\ + dst[2] = lut_l[idx->z];\ + dst[3] = lut_l[idx->w]; #elif dcn == 3 #define LUT_OP(num)\ - uchar3 idx = vload3(0, (__global const uchar *)(srcptr + src_index + num * src_step));\ - dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\ + uchar3 idx = vload3(0, (__global const uchar *)(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];\ dst[2] = lut_l[idx.z]; #elif dcn == 2 #define LUT_OP(num)\ - uchar2 idx = vload2(0, (__global const uchar *)(srcptr + src_index + num * src_step));\ - dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\ - dst[0] = lut_l[idx.x];\ - dst[1] = lut_l[idx.y]; + __global const uchar2 * idx = (__global const uchar *)(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]; #elif dcn == 1 #define LUT_OP(num)\ - uchar idx = (__global const uchar *)(srcptr + src_index + num * src_step)[0];\ - dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\ + uchar idx = (__global const uchar *)(srcptr + mad24(num, src_step, src_index))[0];\ + dst = (__global dstT *)(dstptr + mad24(num, dst_step, dst_index));\ dst[0] = lut_l[idx]; #else #define LUT_OP(num)\ - src = (__global const srcT *)(srcptr + src_index + num * src_step);\ - dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\ + 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)\ dst[cn] = lut_l[src[cn]]; #endif #else #if dcn == 4 #define LUT_OP(num)\ - uchar4 src_pixel = vload4(0, (__global const uchar *)(srcptr + src_index + num * src_step));\ - int4 idx = convert_int4(src_pixel) * lcn + (int4)(0, 1, 2, 3);\ - dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\ + __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];\ dst[3] = lut_l[idx.w]; #elif dcn == 3 #define LUT_OP(num)\ - uchar3 src_pixel = vload3(0, (__global const uchar *)(srcptr + src_index + num * src_step));\ + uchar3 src_pixel = vload3(0, (__global const uchar *)(srcptr + mad24(num, src_step, src_index)));\ int3 idx = convert_int3(src_pixel) * lcn + (int3)(0, 1, 2);\ - dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\ + 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]; #elif dcn == 2 #define LUT_OP(num)\ - uchar2 src_pixel = vload2(0, (__global const uchar *)(srcptr + src_index + num * src_step));\ - int2 idx = convert_int2(src_pixel) * lcn + (int2)(0, 1);\ - dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\ + __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];\ dst[1] = lut_l[idx.y]; #elif dcn == 1 //error case (1 < lcn) ==> lcn == scn == dcn #define LUT_OP(num)\ - uchar idx = (__global const uchar *)(srcptr + src_index + num * src_step)[0];\ - dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\ + uchar idx = (__global const uchar *)(srcptr + mad24(num, src_step, src_index))[0];\ + dst = (__global dstT *)(dstptr + mad24(num, dst_step, dst_index));\ dst[0] = lut_l[idx]; #else #define LUT_OP(num)\ - src = (__global const srcT *)(srcptr + src_index + num * src_step);\ - dst = (__global dstT *)(dstptr + dst_index + num * dst_step);\ + 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)\ dst[cn] = lut_l[mad24(src[cn], lcn, cn)]; #endif @@ -134,7 +134,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 dst_index = mad24(y, dst_step, mad24(x, (int)sizeof(dstT) * dcn, dst_offset)); __global const srcT * src; __global dstT * dst; - + int tmp_idx; LUT_OP(0); if (y < rows - 1) {