From fec21239c809ff363bb40235784276e790e634e0 Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Mon, 4 Aug 2014 12:45:00 +0400 Subject: [PATCH 1/4] Revert optimization for warpAffine INTER_NEAREST mode --- modules/imgproc/src/opencl/warp_affine.cl | 15 ++++++--------- modules/imgproc/src/pyramids.cpp | 3 +++ modules/imgproc/test/ocl/test_pyramids.cpp | 3 ++- 3 files changed, 11 insertions(+), 10 deletions(-) diff --git a/modules/imgproc/src/opencl/warp_affine.cl b/modules/imgproc/src/opencl/warp_affine.cl index 649f10db7a..229336ea15 100644 --- a/modules/imgproc/src/opencl/warp_affine.cl +++ b/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 diff --git a/modules/imgproc/src/pyramids.cpp b/modules/imgproc/src/pyramids.cpp index 2714e08f30..f213445993 100644 --- a/modules/imgproc/src/pyramids.cpp +++ b/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 ); diff --git a/modules/imgproc/test/ocl/test_pyramids.cpp b/modules/imgproc/test/ocl/test_pyramids.cpp index a129c7f771..beff48591e 100644 --- a/modules/imgproc/test/ocl/test_pyramids.cpp +++ b/modules/imgproc/test/ocl/test_pyramids.cpp @@ -94,7 +94,8 @@ OCL_TEST_P(PyrDown, Mat) { for (int j = 0; j < test_loop_times; j++) { - Size src_roiSize = randomSize(1, MAX_VALUE); + // minimal src size is set to 4 since size<4 doesn't make sense + Size src_roiSize = randomSize(4, MAX_VALUE); Size dst_roiSize = Size(randomInt((src_roiSize.width - 1) / 2, (src_roiSize.width + 3) / 2), randomInt((src_roiSize.height - 1) / 2, (src_roiSize.height + 3) / 2)); dst_roiSize = dst_roiSize.area() == 0 ? Size((src_roiSize.width + 1) / 2, (src_roiSize.height + 1) / 2) : dst_roiSize; From eb9fdb01642021a8bd56d44cde5b50295e6403bc Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Tue, 5 Aug 2014 15:40:29 +0400 Subject: [PATCH 2/4] Fixed rounding in remap INTER_LINEAR mode --- modules/imgproc/src/opencl/remap.cl | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/modules/imgproc/src/opencl/remap.cl b/modules/imgproc/src/opencl/remap.cl index 4e45b40bd3..41f5fa85d7 100644 --- a/modules/imgproc/src/opencl/remap.cl +++ b/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); From 2a0b39d30a6e324e51f870536420de5f8df6e687 Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Tue, 5 Aug 2014 17:17:14 +0400 Subject: [PATCH 3/4] Fixed calculate_histogram kernel --- modules/imgproc/src/opencl/histogram.cl | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/modules/imgproc/src/opencl/histogram.cl b/modules/imgproc/src/opencl/histogram.cl index ff8023054f..5c99d5424d 100644 --- a/modules/imgproc/src/opencl/histogram.cl +++ b/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 From 44fbfb2cf6b6dd45ab0ad2a8d31bfafca71bc00f Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Wed, 6 Aug 2014 11:09:56 +0400 Subject: [PATCH 4/4] Fixed extrapolation in pyrDown --- modules/imgproc/src/opencl/pyr_down.cl | 6 +++--- modules/imgproc/test/ocl/test_pyramids.cpp | 3 +-- 2 files changed, 4 insertions(+), 5 deletions(-) diff --git a/modules/imgproc/src/opencl/pyr_down.cl b/modules/imgproc/src/opencl/pyr_down.cl index 4db1a8d811..b6927fa879 100644 --- a/modules/imgproc/src/opencl/pyr_down.cl +++ b/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 diff --git a/modules/imgproc/test/ocl/test_pyramids.cpp b/modules/imgproc/test/ocl/test_pyramids.cpp index beff48591e..a129c7f771 100644 --- a/modules/imgproc/test/ocl/test_pyramids.cpp +++ b/modules/imgproc/test/ocl/test_pyramids.cpp @@ -94,8 +94,7 @@ OCL_TEST_P(PyrDown, Mat) { for (int j = 0; j < test_loop_times; j++) { - // minimal src size is set to 4 since size<4 doesn't make sense - Size src_roiSize = randomSize(4, MAX_VALUE); + Size src_roiSize = randomSize(1, MAX_VALUE); Size dst_roiSize = Size(randomInt((src_roiSize.width - 1) / 2, (src_roiSize.width + 3) / 2), randomInt((src_roiSize.height - 1) / 2, (src_roiSize.height + 3) / 2)); dst_roiSize = dst_roiSize.area() == 0 ? Size((src_roiSize.width + 1) / 2, (src_roiSize.height + 1) / 2) : dst_roiSize;