ocl: don't use vload4 for 3 channel images

alalek · alalek · commit ba8a6e353374 · 2017-03-03T19:36:38.000+03:00
diff --git a/modules/imgproc/src/opencl/cvtcolor.cl b/modules/imgproc/src/opencl/cvtcolor.cl
@@ -160,6 +160,7 @@ enum
 #define CAT(x, y) __CAT(x, y)
 
 #define DATA_TYPE_4 CAT(DATA_TYPE, 4)
+#define DATA_TYPE_3 CAT(DATA_TYPE, 3)
 
 ///////////////////////////////////// RGB <-> GRAY //////////////////////////////////////
 
@@ -182,7 +183,7 @@ __kernel void RGB2Gray(__global const uchar * srcptr, int src_step, int src_offs
             {
                 __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + src_index);
                 __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + dst_index);
-                DATA_TYPE_4 src_pix = vload4(0, src);
+                DATA_TYPE_3 src_pix = vload3(0, src);
 #ifdef DEPTH_5
                 dst[0] = fma(src_pix.B_COMP, B2YF, fma(src_pix.G_COMP, G2YF, src_pix.R_COMP * R2YF));
 #else
@@ -256,7 +257,7 @@ __kernel void RGB2YUV(__global const uchar* srcptr, int src_step, int src_offset
             {
                 __global const DATA_TYPE* src = (__global const DATA_TYPE*)(srcptr + src_index);
                 __global DATA_TYPE* dst = (__global DATA_TYPE*)(dstptr + dst_index);
-                DATA_TYPE_4 src_pix = vload4(0, src);
+                DATA_TYPE_3 src_pix = vload3(0, src);
                 DATA_TYPE b = src_pix.B_COMP, g = src_pix.G_COMP, r = src_pix.R_COMP;
 
 #ifdef DEPTH_5
