From 4f68f35a7868a1e2e8544de7c29f92ee20cb36b3 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Sat, 19 Oct 2013 01:00:22 +0400 Subject: [PATCH] added reflect, reflect101, replicate, wrap border types to ocl::remap --- modules/ocl/src/imgproc.cpp | 4 +- modules/ocl/src/opencl/imgproc_remap.cl | 169 +++++++++++++----------- modules/ocl/test/test_filters.cpp | 11 +- modules/ocl/test/test_warp.cpp | 49 ++++--- 4 files changed, 132 insertions(+), 101 deletions(-) diff --git a/modules/ocl/src/imgproc.cpp b/modules/ocl/src/imgproc.cpp index 35395211fd..10b6804869 100644 --- a/modules/ocl/src/imgproc.cpp +++ b/modules/ocl/src/imgproc.cpp @@ -195,7 +195,8 @@ namespace cv CV_Assert((map1.type() == CV_16SC2 && !map2.data) || (map1.type() == CV_32FC2 && !map2.data) || (map1.type() == CV_32FC1 && map2.type() == CV_32FC1)); CV_Assert(!map2.data || map2.size() == map1.size()); - CV_Assert(borderType == BORDER_CONSTANT); + CV_Assert(borderType == BORDER_CONSTANT || borderType == BORDER_REPLICATE || borderType == BORDER_WRAP + || borderType == BORDER_REFLECT_101 || borderType == BORDER_REFLECT); dst.create(map1.size(), src.type()); @@ -452,6 +453,7 @@ namespace cv } bordertype &= ~cv::BORDER_ISOLATED; + // TODO need to remove this conditions and fix the code if (bordertype == cv::BORDER_REFLECT || bordertype == cv::BORDER_WRAP) { CV_Assert((_src.cols >= left) && (_src.cols >= right) && (_src.rows >= top) && (_src.rows >= bottom)); diff --git a/modules/ocl/src/opencl/imgproc_remap.cl b/modules/ocl/src/opencl/imgproc_remap.cl index 2627f2460f..23899bdbbc 100644 --- a/modules/ocl/src/opencl/imgproc_remap.cl +++ b/modules/ocl/src/opencl/imgproc_remap.cl @@ -51,6 +51,70 @@ #endif #endif +#ifdef INTER_NEAREST +#define convertToWT +#endif + +#ifdef BORDER_CONSTANT +#define EXTRAPOLATE(v2, v) v = scalar; +#elif defined BORDER_REPLICATE +#define EXTRAPOLATE(v2, v) \ + { \ + v2 = max(min(v2, (int2)(src_cols - 1, src_rows - 1)), zero); \ + v = convertToWT(src[mad24(v2.y, src_step, v2.x + src_offset)]); \ + } +#elif defined BORDER_WRAP +#define EXTRAPOLATE(v2, v) \ + { \ + if (v2.x < 0) \ + v2.x -= ((v2.x - src_cols + 1) / src_cols) * src_cols; \ + if (v2.x >= src_cols) \ + v2.x %= src_cols; \ + \ + if (v2.y < 0) \ + v2.y -= ((v2.y - src_rows + 1) / src_rows) * src_rows; \ + if( v2.y >= src_rows ) \ + v2.y %= src_rows; \ + v = convertToWT(src[mad24(v2.y, src_step, v2.x + src_offset)]); \ + } +#elif defined(BORDER_REFLECT) || defined(BORDER_REFLECT_101) +#ifdef BORDER_REFLECT +#define DELTA int delta = 0 +#else +#define DELTA int delta = 1 +#endif +#define EXTRAPOLATE(v2, v) \ + { \ + DELTA; \ + if (src_cols == 1) \ + v2.x = 0; \ + else \ + do \ + { \ + if( v2.x < 0 ) \ + v2.x = -v2.x - 1 + delta; \ + else \ + v2.x = src_cols - 1 - (v2.x - src_cols) - delta; \ + } \ + while (v2.x >= src_cols || v2.x < 0); \ + \ + if (src_rows == 1) \ + v2.y = 0; \ + else \ + do \ + { \ + if( v2.y < 0 ) \ + v2.y = -v2.y - 1 + delta; \ + else \ + v2.y = src_rows - 1 - (v2.y - src_rows) - delta; \ + } \ + while (v2.y >= src_rows || v2.y < 0); \ + v = convertToWT(src[mad24(v2.y, src_step, v2.x + src_offset)]); \ + } +#else +#error No extrapolation method +#endif + #define NEED_EXTRAPOLATION(gx, gy) (gx >= src_cols || gy >= src_rows || gx < 0 || gy < 0) #ifdef INTER_NEAREST @@ -75,11 +139,8 @@ __kernel void remap_2_32FC1(__global const T * restrict src, __global T * dst, if (NEED_EXTRAPOLATION(gx, gy)) { -#ifdef BORDER_CONSTANT - dst[dstIdx] = scalar; -#else -#error No extrapolation method -#endif + int2 gxy = (int2)(gx, gy), zero = (int2)(0); + EXTRAPOLATE(gxy, dst[dstIdx]); } else { @@ -107,11 +168,8 @@ __kernel void remap_32FC2(__global const T * restrict src, __global T * dst, __g if (NEED_EXTRAPOLATION(gx, gy)) { -#ifdef BORDER_CONSTANT - dst[dstIdx] = scalar; -#else -#error No extrapolation method -#endif + int2 zero = (int2)(0); + EXTRAPOLATE(gxy, dst[dstIdx]); } else { @@ -139,11 +197,8 @@ __kernel void remap_16SC2(__global const T * restrict src, __global T * dst, __g if (NEED_EXTRAPOLATION(gx, gy)) { -#ifdef BORDER_CONSTANT - dst[dstIdx] = scalar; -#else -#error No extrapolation method -#endif + int2 zero = (int2)(0); + EXTRAPOLATE(gxy, dst[dstIdx]); } else { @@ -176,56 +231,37 @@ __kernel void remap_2_32FC1(__global T const * restrict src, __global T * dst, int2 map_dataB = (int2)(map_dataA.x + 1, map_dataA.y); int2 map_dataC = (int2)(map_dataA.x, map_dataA.y + 1); int2 map_dataD = (int2)(map_dataA.x + 1, map_dataA.y +1); + int2 zero = (int2)(0); float2 _u = map_data - convert_float2(map_dataA); WT2 u = convertToWT2(convert_int2_rte(convertToWT2(_u) * (WT2)32)) / (WT2)32; - WT nval = convertToWT(nVal); - WT a = nval, b = nval, c = nval, d = nval; + WT scalar = convertToWT(nVal); + WT a = scalar, b = scalar, c = scalar, d = scalar; if (!NEED_EXTRAPOLATION(map_dataA.x, map_dataA.y)) a = convertToWT(src[mad24(map_dataA.y, src_step, map_dataA.x + src_offset)]); else - { -#ifdef BORDER_CONSTANT -#else -#error No extrapolation method -#endif - } + EXTRAPOLATE(map_dataA, a); if (!NEED_EXTRAPOLATION(map_dataB.x, map_dataB.y)) b = convertToWT(src[mad24(map_dataB.y, src_step, map_dataB.x + src_offset)]); else - { -#ifdef BORDER_CONSTANT -#else -#error No extrapolation method -#endif - } + EXTRAPOLATE(map_dataB, b); if (!NEED_EXTRAPOLATION(map_dataC.x, map_dataC.y)) c = convertToWT(src[mad24(map_dataC.y, src_step, map_dataC.x + src_offset)]); else - { -#ifdef BORDER_CONSTANT -#else -#error No extrapolation method -#endif - } + EXTRAPOLATE(map_dataC, c); if (!NEED_EXTRAPOLATION(map_dataD.x, map_dataD.y)) d = convertToWT(src[mad24(map_dataD.y, src_step, map_dataD.x + src_offset)]); else - { -#ifdef BORDER_CONSTANT -#else -#error No extrapolation method -#endif - } + EXTRAPOLATE(map_dataD, d); - WT dst_data = a * (WT)(1.0 - u.x) * (WT)(1.0 - u.y) + - b * (WT)(u.x) * (WT)(1.0 - u.y) + - c * (WT)(1.0 - u.x) * (WT)(u.y) + - d * (WT)(u.x) * (WT)(u.y); + WT dst_data = a * (WT)(1 - u.x) * (WT)(1 - u.y) + + b * (WT)(u.x) * (WT)(1 - u.y) + + c * (WT)(1 - u.x) * (WT)(u.y) + + d * (WT)(u.x) * (WT)(u.y); dst[dstIdx] = convertToT(dst_data); } } @@ -248,57 +284,38 @@ __kernel void remap_32FC2(__global T const * restrict src, __global T * dst, int2 map_dataA = convert_int2_sat_rtn(map_data); int2 map_dataB = (int2)(map_dataA.x + 1, map_dataA.y); int2 map_dataC = (int2)(map_dataA.x, map_dataA.y + 1); - int2 map_dataD = (int2)(map_dataA.x + 1, map_dataA.y +1); + int2 map_dataD = (int2)(map_dataA.x + 1, map_dataA.y + 1); + int2 zero = (int2)(0); float2 _u = map_data - convert_float2(map_dataA); WT2 u = convertToWT2(convert_int2_rte(convertToWT2(_u) * (WT2)32)) / (WT2)32; - WT nval = convertToWT(nVal); - WT a = nval, b = nval, c = nval, d = nval; + WT scalar = convertToWT(nVal); + WT a = scalar, b = scalar, c = scalar, d = scalar; if (!NEED_EXTRAPOLATION(map_dataA.x, map_dataA.y)) a = convertToWT(src[mad24(map_dataA.y, src_step, map_dataA.x + src_offset)]); else - { -#ifdef BORDER_CONSTANT -#else -#error No extrapolation method -#endif - } + EXTRAPOLATE(map_dataA, a); if (!NEED_EXTRAPOLATION(map_dataB.x, map_dataB.y)) b = convertToWT(src[mad24(map_dataB.y, src_step, map_dataB.x + src_offset)]); else - { -#ifdef BORDER_CONSTANT -#else -#error No extrapolation method -#endif - } + EXTRAPOLATE(map_dataB, b); if (!NEED_EXTRAPOLATION(map_dataC.x, map_dataC.y)) c = convertToWT(src[mad24(map_dataC.y, src_step, map_dataC.x + src_offset)]); else - { -#ifdef BORDER_CONSTANT -#else -#error No extrapolation method -#endif - } + EXTRAPOLATE(map_dataC, c); if (!NEED_EXTRAPOLATION(map_dataD.x, map_dataD.y)) d = convertToWT(src[mad24(map_dataD.y, src_step, map_dataD.x + src_offset)]); else - { -#ifdef BORDER_CONSTANT -#else -#error No extrapolation method -#endif - } + EXTRAPOLATE(map_dataD, d); - WT dst_data = a * (WT)(1.0 - u.x) * (WT)(1.0 - u.y) + - b * (WT)(u.x) * (WT)(1.0 - u.y) + - c * (WT)(1.0 - u.x) * (WT)(u.y) + - d * (WT)(u.x) * (WT)(u.y); + WT dst_data = a * (WT)(1 - u.x) * (WT)(1 - u.y) + + b * (WT)(u.x) * (WT)(1 - u.y) + + c * (WT)(1 - u.x) * (WT)(u.y) + + d * (WT)(u.x) * (WT)(u.y); dst[dstIdx] = convertToT(dst_data); } } diff --git a/modules/ocl/test/test_filters.cpp b/modules/ocl/test/test_filters.cpp index 1c7dd21fbf..cf1857479d 100644 --- a/modules/ocl/test/test_filters.cpp +++ b/modules/ocl/test/test_filters.cpp @@ -62,8 +62,7 @@ PARAM_TEST_CASE(FilterTestBase, MatType, int, // border type, or iteration bool) // roi or not { - int type, borderType; - int ksize; + int type, borderType, ksize; bool useRoi; Mat src, dst_whole, src_roi, dst_roi; @@ -92,8 +91,12 @@ PARAM_TEST_CASE(FilterTestBase, MatType, void Near(double threshold = 0.0) { - EXPECT_MAT_NEAR(dst_whole, Mat(gdst_whole), threshold); - EXPECT_MAT_NEAR(dst_roi, Mat(gdst_roi), threshold); + Mat roi, whole; + gdst_whole.download(whole); + gdst_roi.download(roi); + + EXPECT_MAT_NEAR(dst_whole, whole, threshold); + EXPECT_MAT_NEAR(dst_roi, roi, threshold); } }; diff --git a/modules/ocl/test/test_warp.cpp b/modules/ocl/test/test_warp.cpp index 6e153fb1a0..717bbc7a2e 100644 --- a/modules/ocl/test/test_warp.cpp +++ b/modules/ocl/test/test_warp.cpp @@ -64,7 +64,7 @@ static MatType noType = -1; ///////////////////////////////////////////////////////////////////////////////////////////////// // warpAffine & warpPerspective -PARAM_TEST_CASE(WarpTestBase, MatType, int, bool, bool) +PARAM_TEST_CASE(WarpTestBase, MatType, Interpolation, bool, bool) { int type, interpolation; Size dsize; @@ -164,7 +164,7 @@ OCL_TEST_P(WarpPerspective, Mat) ///////////////////////////////////////////////////////////////////////////////////////////////// // remap -PARAM_TEST_CASE(Remap, int, int, pair, int, bool) +PARAM_TEST_CASE(Remap, MatDepth, Channels, pair, Border, bool) { int srcType, map1Type, map2Type; int borderType; @@ -202,14 +202,15 @@ PARAM_TEST_CASE(Remap, int, int, pair, int, bool) randomSubMat(src, src_roi, srcROISize, srcBorder, srcType, 5, 256); Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0); - randomSubMat(dst, dst_roi, dstROISize, dstBorder, srcType, -MAX_VALUE, MAX_VALUE << 1); + randomSubMat(dst, dst_roi, dstROISize, dstBorder, srcType, -MAX_VALUE, MAX_VALUE); + int mapMaxValue = MAX_VALUE << 2; Border map1Border = randomBorder(0, useRoi ? MAX_VALUE : 0); - randomSubMat(map1, map1_roi, dstROISize, map1Border, map1Type, -MAX_VALUE, MAX_VALUE << 1); + randomSubMat(map1, map1_roi, dstROISize, map1Border, map1Type, -mapMaxValue, mapMaxValue); Border map2Border = randomBorder(0, useRoi ? MAX_VALUE : 0); if (map2Type != noType) - randomSubMat(map2, map2_roi, dstROISize, map2Border, map2Type, -MAX_VALUE, MAX_VALUE << 1); + randomSubMat(map2, map2_roi, dstROISize, map2Border, map2Type, -mapMaxValue, mapMaxValue); generateOclMat(gsrc, gsrc_roi, src, srcROISize, srcBorder); generateOclMat(gdst, gdst_roi, dst, dstROISize, dstBorder); @@ -262,7 +263,7 @@ OCL_TEST_P(Remap_INTER_LINEAR, Mat) ///////////////////////////////////////////////////////////////////////////////////////////////// // resize -PARAM_TEST_CASE(Resize, MatType, double, double, int, bool) +PARAM_TEST_CASE(Resize, MatType, double, double, Interpolation, bool) { int type, interpolation; double fx, fy; @@ -325,38 +326,46 @@ OCL_TEST_P(Resize, Mat) INSTANTIATE_TEST_CASE_P(ImgprocWarp, WarpAffine, Combine( Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4), - Values((int)INTER_NEAREST, (int)INTER_LINEAR, (int)INTER_CUBIC), + Values((Interpolation)INTER_NEAREST, (Interpolation)INTER_LINEAR, (Interpolation)INTER_CUBIC), Bool(), Bool())); INSTANTIATE_TEST_CASE_P(ImgprocWarp, WarpPerspective, Combine( Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4), - Values((int)INTER_NEAREST, (int)INTER_LINEAR, (int)INTER_CUBIC), + Values((Interpolation)INTER_NEAREST, (Interpolation)INTER_LINEAR, (Interpolation)INTER_CUBIC), Bool(), Bool())); INSTANTIATE_TEST_CASE_P(ImgprocWarp, Remap_INTER_LINEAR, Combine( - testing::Values(CV_8U, CV_16U, CV_16S, CV_32F, CV_64F), - testing::Range(1, 5), - Values(make_pair(CV_32FC1, CV_32FC1), - make_pair(CV_32FC2, noType)), - Values((int)BORDER_CONSTANT), + Values(CV_8U, CV_16U, CV_16S, CV_32F, CV_64F), + Values(1, 2, 3, 4), + Values(pair((MatType)CV_32FC1, (MatType)CV_32FC1), + pair((MatType)CV_32FC2, noType)), + Values((Border)BORDER_CONSTANT, + (Border)BORDER_REPLICATE, + (Border)BORDER_WRAP, + (Border)BORDER_REFLECT, + (Border)BORDER_REFLECT_101), Bool())); INSTANTIATE_TEST_CASE_P(ImgprocWarp, Remap_INTER_NEAREST, Combine( - testing::Values(CV_8U, CV_16U, CV_16S, CV_32F, CV_64F), - testing::Range(1, 5), - Values(make_pair(CV_32FC1, CV_32FC1), - make_pair(CV_32FC2, noType), - make_pair(CV_16SC2, noType)), - Values((int)BORDER_CONSTANT), + Values(CV_8U, CV_16U, CV_16S, CV_32F, CV_64F), + Values(1, 2, 3, 4), + Values(pair((MatType)CV_32FC1, (MatType)CV_32FC1), + pair((MatType)CV_32FC2, noType), + pair((MatType)CV_16SC2, noType)), + Values((Border)BORDER_CONSTANT, + (Border)BORDER_REPLICATE, + (Border)BORDER_WRAP, + (Border)BORDER_REFLECT, + (Border)BORDER_REFLECT_101), Bool())); INSTANTIATE_TEST_CASE_P(ImgprocWarp, Resize, Combine( Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1, CV_32FC3, CV_32FC4), Values(0.5, 1.5, 2.0), Values(0.5, 1.5, 2.0), - Values((int)INTER_NEAREST, (int)INTER_LINEAR), + Values((Interpolation)INTER_NEAREST, (Interpolation)INTER_LINEAR), Bool())); #endif // HAVE_OPENCL -- 2.34.1