From 561a7f5782f55420fc9a77425a4a81d02d714b4b Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Thu, 21 Nov 2013 18:23:57 +0400 Subject: [PATCH] fixed kernel compilation in imgproc module --- modules/imgproc/src/opencl/cvtcolor.cl | 36 +++++++++++++------------- modules/imgproc/src/opencl/resize.cl | 28 ++++++++++---------- 2 files changed, 32 insertions(+), 32 deletions(-) diff --git a/modules/imgproc/src/opencl/cvtcolor.cl b/modules/imgproc/src/opencl/cvtcolor.cl index 9ca98b0b91..9e187f525c 100644 --- a/modules/imgproc/src/opencl/cvtcolor.cl +++ b/modules/imgproc/src/opencl/cvtcolor.cl @@ -98,9 +98,9 @@ __kernel void RGB2Gray(__global const uchar* srcptr, int srcstep, int srcoffset, if (y < rows && x < cols) { - const DATA_TYPE* src = (const DATA_TYPE*)(srcptr + mad24(y, srcstep, srcoffset + x * scnbytes)); - DATA_TYPE* dst = (DATA_TYPE*)(dstptr + mad24(y, dststep, dstoffset + x * dcnbytes)); -#if defined (DEPTH_5) + __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + mad24(y, srcstep, srcoffset + x * scnbytes)); + __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + mad24(y, dststep, dstoffset + x * dcnbytes)); +#ifdef DEPTH_5 dst[0] = src[bidx] * 0.114f + src[1] * 0.587f + src[(bidx^2)] * 0.299f; #else dst[0] = (DATA_TYPE)CV_DESCALE((src[bidx] * B2Y + src[1] * G2Y + src[(bidx^2)] * R2Y), yuv_shift); @@ -117,8 +117,8 @@ __kernel void Gray2RGB(__global const uchar* srcptr, int srcstep, int srcoffset, if (y < rows && x < cols) { - const DATA_TYPE* src = (const DATA_TYPE*)(srcptr + mad24(y, srcstep, srcoffset + x * scnbytes)); - DATA_TYPE* dst = (DATA_TYPE*)(dstptr + mad24(y, dststep, dstoffset + x * dcnbytes)); + __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + mad24(y, srcstep, srcoffset + x * scnbytes)); + __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + mad24(y, dststep, dstoffset + x * dcnbytes)); DATA_TYPE val = src[0]; dst[0] = dst[1] = dst[2] = val; #if dcn == 4 @@ -141,11 +141,11 @@ __kernel void RGB2YUV(__global const uchar* srcptr, int srcstep, int srcoffset, if (y < rows && x < cols) { - const DATA_TYPE* src = (const DATA_TYPE*)(srcptr + mad24(y, srcstep, srcoffset + x * scnbytes)); - DATA_TYPE* dst = (DATA_TYPE*)(dstptr + mad24(y, dststep, dstoffset + x * dcnbytes)); + __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + mad24(y, srcstep, srcoffset + x * scnbytes)); + __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + mad24(y, dststep, dstoffset + x * dcnbytes)); DATA_TYPE b=src[bidx], g=src[1], r=src[bidx^2]; -#if defined (DEPTH_5) +#ifdef DEPTH_5 __constant float * coeffs = c_RGB2YUVCoeffs_f; const DATA_TYPE Y = b * coeffs[0] + g * coeffs[1] + r * coeffs[2]; const DATA_TYPE U = (b - Y) * coeffs[3] + HALF_MAX; @@ -176,11 +176,11 @@ __kernel void YUV2RGB(__global const uchar* srcptr, int srcstep, int srcoffset, if (y < rows && x < cols) { - const DATA_TYPE* src = (const DATA_TYPE*)(srcptr + mad24(y, srcstep, srcoffset + x * scnbytes)); - DATA_TYPE* dst = (DATA_TYPE*)(dstptr + mad24(y, dststep, dstoffset + x * dcnbytes)); + __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + mad24(y, srcstep, srcoffset + x * scnbytes)); + __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + mad24(y, dststep, dstoffset + x * dcnbytes)); DATA_TYPE Y = src[0], U = src[1], V = src[2]; -#if defined (DEPTH_5) +#ifdef DEPTH_5 __constant float * coeffs = c_YUV2RGBCoeffs_f; const float r = Y + (V - HALF_MAX) * coeffs[3]; const float g = Y + (V - HALF_MAX) * coeffs[2] + (U - HALF_MAX) * coeffs[1]; @@ -217,10 +217,10 @@ __kernel void YUV2RGBA_NV12(__global const uchar* srcptr, int srcstep, int srcof if (y < rows / 2 && x < cols / 2 ) { - __global const uchar* ysrc = srcptr + mad24(y << 1, srcstep, (x << 1) + srcoffset); - __global const uchar* usrc = srcptr + mad24(rows + y, srcstep, (x << 1) + srcoffset); - __global uchar* dst1 = dstptr + mad24(y << 1, dststep, x*(dcn*2) + dstoffset); - __global uchar* dst2 = dstptr + mad24((y << 1) + 1, dststep, x*(dcn*2) + dstoffset); + __global const uchar* ysrc = (__global const uchar*)(srcptr + mad24(y << 1, srcstep, (x << 1) + srcoffset)); + __global const uchar* usrc = (__global const uchar*)(srcptr + mad24(rows + y, srcstep, (x << 1) + srcoffset)); + __global uchar* dst1 = (__global uchar*)(dstptr + mad24(y << 1, dststep, x*(dcn*2) + dstoffset)); + __global uchar* dst2 = (__global uchar*)(dstptr + mad24((y << 1) + 1, dststep, x*(dcn*2) + dstoffset)); int Y1 = ysrc[0]; int Y2 = ysrc[1]; @@ -282,11 +282,11 @@ __kernel void RGB2YCrCb(__global const uchar* srcptr, int srcstep, int srcoffset if (y < rows && x < cols) { - const DATA_TYPE* src = (const DATA_TYPE*)(srcptr + mad24(y, srcstep, srcoffset + x * scnbytes)); - DATA_TYPE* dst = (DATA_TYPE*)(dstptr + mad24(y, dststep, dstoffset + x * dcnbytes)); + __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + mad24(y, srcstep, srcoffset + x * scnbytes)); + __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + mad24(y, dststep, dstoffset + x * dcnbytes)); DATA_TYPE b=src[bidx], g=src[1], r=src[bidx^2]; -#if defined (DEPTH_5) +#ifdef DEPTH_5 __constant float * coeffs = c_RGB2YCrCbCoeffs_f; const DATA_TYPE Y = b * coeffs[0] + g * coeffs[1] + r * coeffs[2]; const DATA_TYPE Cr = (r - Y) * coeffs[3] + HALF_MAX; diff --git a/modules/imgproc/src/opencl/resize.cl b/modules/imgproc/src/opencl/resize.cl index 0805246ca3..81c973b4bc 100644 --- a/modules/imgproc/src/opencl/resize.cl +++ b/modules/imgproc/src/opencl/resize.cl @@ -87,7 +87,7 @@ __kernel void resizeLN(__global const uchar* srcptr, int srcstep, int srcoffset, int y_ = INC(y,srcrows); int x_ = INC(x,srccols); - const PIXTYPE* src = (const PIXTYPE*)(srcptr + mad24(y, srcstep, srcoffset + x*PIXSIZE)); + __global const PIXTYPE* src = (__global const PIXTYPE*)(srcptr + mad24(y, srcstep, srcoffset + x*PIXSIZE)); #if depth == 0 u = u * INTER_RESIZE_COEF_SCALE; @@ -98,10 +98,10 @@ __kernel void resizeLN(__global const uchar* srcptr, int srcstep, int srcoffset, int U1 = rint(INTER_RESIZE_COEF_SCALE - u); int V1 = rint(INTER_RESIZE_COEF_SCALE - v); - WORKTYPE data0 = convertToWT(*(const PIXTYPE*)(srcptr + mad24(y, srcstep, srcoffset + x*PIXSIZE))); - WORKTYPE data1 = convertToWT(*(const PIXTYPE*)(srcptr + mad24(y, srcstep, srcoffset + x_*PIXSIZE))); - WORKTYPE data2 = convertToWT(*(const PIXTYPE*)(srcptr + mad24(y_, srcstep, srcoffset + x*PIXSIZE))); - WORKTYPE data3 = convertToWT(*(const PIXTYPE*)(srcptr + mad24(y_, srcstep, srcoffset + x_*PIXSIZE))); + WORKTYPE data0 = convertToWT(*(__global const PIXTYPE*)(srcptr + mad24(y, srcstep, srcoffset + x*PIXSIZE))); + WORKTYPE data1 = convertToWT(*(__global const PIXTYPE*)(srcptr + mad24(y, srcstep, srcoffset + x_*PIXSIZE))); + WORKTYPE data2 = convertToWT(*(__global const PIXTYPE*)(srcptr + mad24(y_, srcstep, srcoffset + x*PIXSIZE))); + WORKTYPE data3 = convertToWT(*(__global const PIXTYPE*)(srcptr + mad24(y_, srcstep, srcoffset + x_*PIXSIZE))); WORKTYPE val = mul24((WORKTYPE)mul24(U1, V1), data0) + mul24((WORKTYPE)mul24(U, V1), data1) + mul24((WORKTYPE)mul24(U1, V), data2) + mul24((WORKTYPE)mul24(U, V), data3); @@ -109,16 +109,16 @@ __kernel void resizeLN(__global const uchar* srcptr, int srcstep, int srcoffset, #else float u1 = 1.f-u; float v1 = 1.f-v; - WORKTYPE data0 = convertToWT(*(const PIXTYPE*)(srcptr + mad24(y, srcstep, srcoffset + x*PIXSIZE))); - WORKTYPE data1 = convertToWT(*(const PIXTYPE*)(srcptr + mad24(y, srcstep, srcoffset + x_*PIXSIZE))); - WORKTYPE data2 = convertToWT(*(const PIXTYPE*)(srcptr + mad24(y_, srcstep, srcoffset + x*PIXSIZE))); - WORKTYPE data3 = convertToWT(*(const PIXTYPE*)(srcptr + mad24(y_, srcstep, srcoffset + x_*PIXSIZE))); + WORKTYPE data0 = convertToWT(*(__global const PIXTYPE*)(srcptr + mad24(y, srcstep, srcoffset + x*PIXSIZE))); + WORKTYPE data1 = convertToWT(*(__global const PIXTYPE*)(srcptr + mad24(y, srcstep, srcoffset + x_*PIXSIZE))); + WORKTYPE data2 = convertToWT(*(__global const PIXTYPE*)(srcptr + mad24(y_, srcstep, srcoffset + x*PIXSIZE))); + WORKTYPE data3 = convertToWT(*(__global const PIXTYPE*)(srcptr + mad24(y_, srcstep, srcoffset + x_*PIXSIZE))); PIXTYPE uval = u1 * v1 * s_data1 + u * v1 * s_data2 + u1 * v *s_data3 + u * v *s_data4; #endif if(dx < dstcols && dy < dstrows) { - PIXTYPE* dst = (PIXTYPE*)(dstptr + mad24(dy, dststep, dstoffset + dx*PIXSIZE)); + __global PIXTYPE* dst = (__global PIXTYPE*)(dstptr + mad24(dy, dststep, dstoffset + dx*PIXSIZE)); dst[0] = uval; } } @@ -140,10 +140,10 @@ __kernel void resizeNN(__global const uchar* srcptr, int srcstep, int srcoffset, F s2 = dy*ify; int sx = min(convert_int_rtz(s1), srccols-1); int sy = min(convert_int_rtz(s2), srcrows-1); - PIXTYPE* dst = (PIXTYPE*)(dstptr + - mad24(dy, dststep, dstoffset + dx*PIXSIZE)); - const PIXTYPE* src = (const PIXTYPE*)(srcptr + - mad24(sy, srcstep, srcoffset + sx*PIXSIZE)); + + __global PIXTYPE* dst = (__global PIXTYPE*)(dstptr + mad24(dy, dststep, dstoffset + dx*PIXSIZE)); + __global const PIXTYPE* src = (__global const PIXTYPE*)(srcptr + mad24(sy, srcstep, srcoffset + sx*PIXSIZE)); + dst[0] = src[0]; } } -- GitLab