Chiebot-Mirror
/
opencv
mirror of https://github.com/opencv/opencv.git
8
0
Fork 0
Code Issues Projects Releases Wiki Activity

added reflect, reflect101, replicate, wrap border types to ocl::remap

Browse Source
pull/1633/head
Ilya Lavrenov 11 years ago
parent b3e1a4598c
commit 4f68f35a78
4 changed files with 132 additions and 101 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Split View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 4
      modules/ocl/src/imgproc.cpp
  2. 169
      modules/ocl/src/opencl/imgproc_remap.cl
  3. 11
      modules/ocl/test/test_filters.cpp
  4. 49
      modules/ocl/test/test_warp.cpp

4
modules/ocl/src/imgproc.cpp
Unescape View File

@ -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));

169
modules/ocl/src/opencl/imgproc_remap.cl
Unescape View File

@ -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);
}
}

11
modules/ocl/test/test_filters.cpp
Unescape View File

@ -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);
}
};

49
modules/ocl/test/test_warp.cpp
Unescape View File

@ -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<MatType, MatType>, int, bool)
PARAM_TEST_CASE(Remap, MatDepth, Channels, pair<MatType, MatType>, Border, bool)
{
int srcType, map1Type, map2Type;
int borderType;
@ -202,14 +202,15 @@ PARAM_TEST_CASE(Remap, int, int, pair<MatType, MatType>, 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<MatType, MatType>(CV_32FC1, CV_32FC1),
make_pair<MatType, MatType>(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, MatType>((MatType)CV_32FC1, (MatType)CV_32FC1),
pair<MatType, MatType>((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<MatType, MatType>(CV_32FC1, CV_32FC1),
make_pair<MatType, MatType>(CV_32FC2, noType),
make_pair<MatType, MatType>(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, MatType>((MatType)CV_32FC1, (MatType)CV_32FC1),
pair<MatType, MatType>((MatType)CV_32FC2, noType),
pair<MatType, MatType>((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

Write Preview
Loading…
Cancel
Save