fixed kernel compilation in imgproc module

diff --git a/modules/imgproc/src/opencl/cvtcolor.cl b/modules/imgproc/src/opencl/cvtcolor.cl
index 9ca98b0..9e187f5 100644 (file)
--- 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 0805246..81c973b 100644 (file)
--- 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];
     }
 }