Merge pull request #3044 from akarsakov:fix_ocl_tests

modules/imgproc/src/opencl/histogram.cl

@@ -47,7 +47,7 @@
 
 #define noconvert
 
-__kernel void calculate_histogram(__global const uchar * src, int src_step, int src_offset, int src_rows, int src_cols,
+__kernel void calculate_histogram(__global const uchar * src_ptr, int src_step, int src_offset, int src_rows, int src_cols,
                                   __global uchar * histptr, int total)
 {
     int lid = get_local_id(0);
@@ -61,6 +61,7 @@ __kernel void calculate_histogram(__global const uchar * src, int src_step, int
         localhist[i] = 0;
     barrier(CLK_LOCAL_MEM_FENCE);
 
+    __global const uchar * src = src_ptr + src_offset;
     int src_index;
 
     for (int grain = HISTS_COUNT * WGS * kercn; id < total; id += grain)
@@ -68,7 +69,7 @@ __kernel void calculate_histogram(__global const uchar * src, int src_step, int
 #ifdef HAVE_SRC_CONT
         src_index = id;
 #else
-        src_index = mad24(id / src_cols, src_step, src_offset + id % src_cols);
+        src_index = mad24(id / src_cols, src_step, id % src_cols);
 #endif
 
 #if kercn == 1

modules/imgproc/src/opencl/pyr_down.cl

@@ -53,16 +53,16 @@
 
 #if defined BORDER_REPLICATE
 // aaaaaa|abcdefgh|hhhhhhh
-#define EXTRAPOLATE(x, maxV) clamp(x, 0, maxV-1)
+#define EXTRAPOLATE(x, maxV) clamp((x), 0, (maxV)-1)
 #elif defined BORDER_WRAP
 // cdefgh|abcdefgh|abcdefg
 #define EXTRAPOLATE(x, maxV) ( (x) + (maxV) ) % (maxV)
 #elif defined BORDER_REFLECT
 // fedcba|abcdefgh|hgfedcb
-#define EXTRAPOLATE(x, maxV) min(((maxV)-1)*2-(x)+1, max((x),-(x)-1) )
+#define EXTRAPOLATE(x, maxV) clamp(min(((maxV)-1)*2-(x)+1, max((x),-(x)-1) ), 0, (maxV)-1)
 #elif defined BORDER_REFLECT_101 || defined BORDER_REFLECT101
 // gfedcb|abcdefgh|gfedcba
-#define EXTRAPOLATE(x, maxV) min(((maxV)-1)*2-(x), max((x),-(x)) )
+#define EXTRAPOLATE(x, maxV) clamp(min(((maxV)-1)*2-(x), max((x),-(x)) ), 0, (maxV)-1)
 #else
 #error No extrapolation method
 #endif

modules/imgproc/src/opencl/remap.cl

@@ -413,9 +413,9 @@ __kernel void remap_2_32FC1(__global const uchar * srcptr, int src_step, int src
                 __global T * dst = (__global T *)(dstptr + dst_index);
 
 #if defined BORDER_CONSTANT
-
                 float xf = map1[0], yf = map2[0];
-                int sx = convert_int_sat_rtn(xf), sy = convert_int_sat_rtn(yf);
+                int sx = convert_int_sat_rtz(mad(xf, INTER_TAB_SIZE, 0.5f)) >> INTER_BITS;
+                int sy = convert_int_sat_rtz(mad(yf, INTER_TAB_SIZE, 0.5f)) >> INTER_BITS;
 
                 __constant float * coeffs_x = coeffs + ((convert_int_rte(xf * INTER_TAB_SIZE) & (INTER_TAB_SIZE - 1)) << 1);
                 __constant float * coeffs_y = coeffs + ((convert_int_rte(yf * INTER_TAB_SIZE) & (INTER_TAB_SIZE - 1)) << 1);

modules/imgproc/src/opencl/warp_affine.cl

@@ -98,15 +98,15 @@ __kernel void warpAffine(__global const uchar * srcptr, int src_step, int src_of
     {
         int round_delta = (AB_SCALE >> 1);
 
-        int X0 = rint(fma(M[0], dx, fma(M[1], dy0, M[2])) * AB_SCALE) + round_delta;
-        int Y0 = rint(fma(M[3], dx, fma(M[4], dy0, M[5])) * AB_SCALE) + round_delta;
-
-        int XSTEP = (int)(M[1] * AB_SCALE);
-        int YSTEP = (int)(M[4] * AB_SCALE);
+        int X0_ = rint(M[0] * dx * AB_SCALE);
+        int Y0_ = rint(M[3] * dx * AB_SCALE);
         int dst_index = mad24(dy0, dst_step, mad24(dx, pixsize, dst_offset));
 
         for (int dy = dy0, dy1 = min(dst_rows, dy0 + rowsPerWI); dy < dy1; ++dy, dst_index += dst_step)
         {
+            int X0 = X0_ + rint(fma(M[1], dy, M[2]) * AB_SCALE) + round_delta;
+            int Y0 = Y0_ + rint(fma(M[4], dy, M[5]) * AB_SCALE) + round_delta;
+
             short sx = convert_short_sat(X0 >> AB_BITS);
             short sy = convert_short_sat(Y0 >> AB_BITS);
 
@@ -117,9 +117,6 @@ __kernel void warpAffine(__global const uchar * srcptr, int src_step, int src_of
             }
             else
                 storepix(scalar, dstptr + dst_index);
-
-            X0 += XSTEP;
-            Y0 += YSTEP;
         }
     }
 }
@@ -376,4 +373,4 @@ __kernel void warpAffine(__global const uchar * srcptr, int src_step, int src_of
     }
 }
 
-#endif
+#endif
\ No newline at end of file

modules/imgproc/src/pyramids.cpp

@@ -413,6 +413,9 @@ static bool ocl_pyrDown( InputArray _src, OutputArray _dst, const Size& _dsz, in
 
     Size ssize = _src.size();
     Size dsize = _dsz.area() == 0 ? Size((ssize.width + 1) / 2, (ssize.height + 1) / 2) : _dsz;
+    if (dsize.height < 2 || dsize.width < 2)
+        return false;
+
     CV_Assert( ssize.width > 0 && ssize.height > 0 &&
             std::abs(dsize.width*2 - ssize.width) <= 2 &&
             std::abs(dsize.height*2 - ssize.height) <= 2 );