Skip to content

O
opencv
Commit 46e423d2 authored by Roman Donchenko's avatar Roman Donchenko Committed by OpenCV Buildbot
Browse files
Options

Merge pull request #1773 from ilya-lavrenov:ocl_buildWarpPerspectiveMaps

parents 67739141 3727168b
Expand all Hide whitespace changes
Inline Side-by-side
Showing with 170 additions and 283 deletions
modules/ocl/perf/perf_imgproc.cpp
View file @ 46e423d2
......@@ -231,139 +231,6 @@ PERF_TEST_P(integralFixture, integral, OCL_TYPICAL_MAT_SIZES)
OCL_PERF_ELSE
}
///////////// WarpAffine ////////////////////////
typedef Size_MatType WarpAffineFixture;
PERF_TEST_P(WarpAffineFixture, WarpAffine,
::testing::Combine(OCL_TYPICAL_MAT_SIZES,
OCL_PERF_ENUM(CV_8UC1, CV_8UC4)))
{
static const double coeffs[2][3] =
{
{ cos(CV_PI / 6), -sin(CV_PI / 6), 100.0 },
{ sin(CV_PI / 6), cos(CV_PI / 6), -100.0 }
};
Mat M(2, 3, CV_64F, (void *)coeffs);
const int interpolation = INTER_NEAREST;
const Size_MatType_t params = GetParam();
const Size srcSize = get<0>(params);
const int type = get<1>(params);
Mat src(srcSize, type), dst(srcSize, type);
declare.in(src, WARMUP_RNG).out(dst);
if (RUN_OCL_IMPL)
{
ocl::oclMat oclSrc(src), oclDst(srcSize, type);
OCL_TEST_CYCLE() cv::ocl::warpAffine(oclSrc, oclDst, M, srcSize, interpolation);
oclDst.download(dst);
SANITY_CHECK(dst);
}
else if (RUN_PLAIN_IMPL)
{
TEST_CYCLE() cv::warpAffine(src, dst, M, srcSize, interpolation);
SANITY_CHECK(dst);
}
else
OCL_PERF_ELSE
}
///////////// WarpPerspective ////////////////////////
typedef Size_MatType WarpPerspectiveFixture;
PERF_TEST_P(WarpPerspectiveFixture, WarpPerspective,
::testing::Combine(OCL_TYPICAL_MAT_SIZES,
OCL_PERF_ENUM(CV_8UC1, CV_8UC4)))
{
static const double coeffs[3][3] =
{
{cos(CV_PI / 6), -sin(CV_PI / 6), 100.0},
{sin(CV_PI / 6), cos(CV_PI / 6), -100.0},
{0.0, 0.0, 1.0}
};
Mat M(3, 3, CV_64F, (void *)coeffs);
const int interpolation = INTER_LINEAR;
const Size_MatType_t params = GetParam();
const Size srcSize = get<0>(params);
const int type = get<1>(params);
Mat src(srcSize, type), dst(srcSize, type);
declare.in(src, WARMUP_RNG).out(dst)
.time(srcSize == OCL_SIZE_4000 ? 18 : srcSize == OCL_SIZE_2000 ? 5 : 2);
if (RUN_OCL_IMPL)
{
ocl::oclMat oclSrc(src), oclDst(srcSize, type);
OCL_TEST_CYCLE() cv::ocl::warpPerspective(oclSrc, oclDst, M, srcSize, interpolation);
oclDst.download(dst);
SANITY_CHECK(dst);
}
else if (RUN_PLAIN_IMPL)
{
TEST_CYCLE() cv::warpPerspective(src, dst, M, srcSize, interpolation);
SANITY_CHECK(dst);
}
else
OCL_PERF_ELSE
}
///////////// resize ////////////////////////
CV_ENUM(resizeInterType, INTER_NEAREST, INTER_LINEAR)
typedef tuple<Size, MatType, resizeInterType, double> resizeParams;
typedef TestBaseWithParam<resizeParams> resizeFixture;
PERF_TEST_P(resizeFixture, resize,
::testing::Combine(OCL_TYPICAL_MAT_SIZES,
OCL_PERF_ENUM(CV_8UC1, CV_8UC4),
resizeInterType::all(),
::testing::Values(0.5, 2.0)))
{
const resizeParams params = GetParam();
const Size srcSize = get<0>(params);
const int type = get<1>(params), interType = get<2>(params);
double scale = get<3>(params);
Mat src(srcSize, type), dst;
const Size dstSize(cvRound(srcSize.width * scale), cvRound(srcSize.height * scale));
dst.create(dstSize, type);
declare.in(src, WARMUP_RNG).out(dst);
if (interType == INTER_LINEAR && type == CV_8UC4 && OCL_SIZE_4000 == srcSize)
declare.time(11);
if (RUN_OCL_IMPL)
{
ocl::oclMat oclSrc(src), oclDst(dstSize, type);
OCL_TEST_CYCLE() cv::ocl::resize(oclSrc, oclDst, Size(), scale, scale, interType);
oclDst.download(dst);
SANITY_CHECK(dst, 1 + DBL_EPSILON);
}
else if (RUN_PLAIN_IMPL)
{
TEST_CYCLE() cv::resize(src, dst, Size(), scale, scale, interType);
SANITY_CHECK(dst, 1 + DBL_EPSILON);
}
else
OCL_PERF_ELSE
}
///////////// threshold////////////////////////
CV_ENUM(ThreshType, THRESH_BINARY, THRESH_TOZERO_INV)
......@@ -727,67 +594,6 @@ PERF_TEST_P(meanShiftProcFixture, meanShiftProc,
OCL_PERF_ELSE
}
///////////// remap////////////////////////
CV_ENUM(RemapInterType, INTER_NEAREST, INTER_LINEAR)
typedef tuple<Size, MatType, RemapInterType> remapParams;
typedef TestBaseWithParam<remapParams> remapFixture;
PERF_TEST_P(remapFixture, remap,
::testing::Combine(OCL_TYPICAL_MAT_SIZES,
OCL_PERF_ENUM(CV_8UC1, CV_8UC4),
RemapInterType::all()))
{
const remapParams params = GetParam();
const Size srcSize = get<0>(params);
const int type = get<1>(params), interpolation = get<2>(params);
Mat src(srcSize, type), dst(srcSize, type);
declare.in(src, WARMUP_RNG).out(dst);
if (srcSize == OCL_SIZE_4000 && interpolation == INTER_LINEAR)
declare.time(9);
Mat xmap, ymap;
xmap.create(srcSize, CV_32FC1);
ymap.create(srcSize, CV_32FC1);
for (int i = 0; i < srcSize.height; ++i)
{
float * const xmap_row = xmap.ptr<float>(i);
float * const ymap_row = ymap.ptr<float>(i);
for (int j = 0; j < srcSize.width; ++j)
{
xmap_row[j] = (j - srcSize.width * 0.5f) * 0.75f + srcSize.width * 0.5f;
ymap_row[j] = (i - srcSize.height * 0.5f) * 0.75f + srcSize.height * 0.5f;
}
}
const int borderMode = BORDER_CONSTANT;
if (RUN_OCL_IMPL)
{
ocl::oclMat oclSrc(src), oclDst(srcSize, type);
ocl::oclMat oclXMap(xmap), oclYMap(ymap);
OCL_TEST_CYCLE() cv::ocl::remap(oclSrc, oclDst, oclXMap, oclYMap, interpolation, borderMode);
oclDst.download(dst);
SANITY_CHECK(dst, 1 + DBL_EPSILON);
}
else if (RUN_PLAIN_IMPL)
{
TEST_CYCLE() cv::remap(src, dst, xmap, ymap, interpolation, borderMode);
SANITY_CHECK(dst, 1 + DBL_EPSILON);
}
else
OCL_PERF_ELSE
}
///////////// CLAHE ////////////////////////
typedef TestBaseWithParam<Size> CLAHEFixture;
......
modules/ocl/perf/perf_imgwarp.cpp 0 → 100644
View file @ 46e423d2
This diff is collapsed.
modules/ocl/src/build_warps.cpp
View file @ 46e423d2
This diff is collapsed.
modules/ocl/src/opencl/build_warps.cl
View file @ 46e423d2
......@@ -43,31 +43,25 @@
//
//M*/
__kernel
void buildWarpPlaneMaps
(
__global float * map_x,
__global float * map_y,
__constant float * KRT,
int tl_u,
int tl_v,
int cols,
int rows,
int step_x,
int step_y,
float scale
)
__kernel void buildWarpPlaneMaps(__global float * xmap, __global float * ymap,
__constant float * KRT,
int tl_u, int tl_v,
int cols, int rows,
int xmap_step, int ymap_step,
int xmap_offset, int ymap_offset,
float scale)
{
int du = get_global_id(0);
int dv = get_global_id(1);
step_x /= sizeof(float);
step_y /= sizeof(float);
__constant float * ck_rinv = KRT;
__constant float * ct = KRT + 9;
if (du < cols && dv < rows)
{
int xmap_index = mad24(dv, xmap_step, xmap_offset + du);
int ymap_index = mad24(dv, ymap_step, ymap_offset + du);
float u = tl_u + du;
float v = tl_v + dv;
float x, y;
......@@ -83,33 +77,27 @@ __kernel
x /= z;
y /= z;
map_x[dv * step_x + du] = x;
map_y[dv * step_y + du] = y;
xmap[xmap_index] = x;
ymap[ymap_index] = y;
}
}
__kernel
void buildWarpCylindricalMaps
(
__global float * map_x,
__global float * map_y,
__constant float * ck_rinv,
int tl_u,
int tl_v,
int cols,
int rows,
int step_x,
int step_y,
float scale
)
__kernel void buildWarpCylindricalMaps(__global float * xmap, __global float * ymap,
__constant float * ck_rinv,
int tl_u, int tl_v,
int cols, int rows,
int xmap_step, int ymap_step,
int xmap_offset, int ymap_offset,
float scale)
{
int du = get_global_id(0);
int dv = get_global_id(1);
step_x /= sizeof(float);
step_y /= sizeof(float);
if (du < cols && dv < rows)
{
int xmap_index = mad24(dv, xmap_step, xmap_offset + du);
int ymap_index = mad24(dv, ymap_step, ymap_offset + du);
float u = tl_u + du;
float v = tl_v + dv;
float x, y;
......@@ -127,33 +115,27 @@ __kernel
if (z > 0) { x /= z; y /= z; }
else x = y = -1;
map_x[dv * step_x + du] = x;
map_y[dv * step_y + du] = y;
xmap[xmap_index] = x;
ymap[ymap_index] = y;
}
}
__kernel
void buildWarpSphericalMaps
(
__global float * map_x,
__global float * map_y,
__constant float * ck_rinv,
int tl_u,
int tl_v,
int cols,
int rows,
int step_x,
int step_y,
float scale
)
__kernel void buildWarpSphericalMaps(__global float * xmap, __global float * ymap,
__constant float * ck_rinv,
int tl_u, int tl_v,
int cols, int rows,
int xmap_step, int ymap_step,
int xmap_offset, int ymap_offset,
float scale)
{
int du = get_global_id(0);
int dv = get_global_id(1);
step_x /= sizeof(float);
step_y /= sizeof(float);
if (du < cols && dv < rows)
{
int xmap_index = mad24(dv, xmap_step, xmap_offset + du);
int ymap_index = mad24(dv, ymap_step, ymap_offset + du);
float u = tl_u + du;
float v = tl_v + dv;
float x, y;
......@@ -174,63 +156,52 @@ __kernel
if (z > 0) { x /= z; y /= z; }
else x = y = -1;
map_x[dv * step_x + du] = x;
map_y[dv * step_y + du] = y;
xmap[xmap_index] = x;
ymap[ymap_index] = y;
}
}
__kernel
void buildWarpAffineMaps
(
__global float * xmap,
__global float * ymap,
__constant float * c_warpMat,
int cols,
int rows,
int step_x,
int step_y
)
__kernel void buildWarpAffineMaps(__global float * xmap, __global float * ymap,
__constant float * c_warpMat,
int cols, int rows,
int xmap_step, int ymap_step,
int xmap_offset, int ymap_offset)
{
int x = get_global_id(0);
int y = get_global_id(1);
step_x /= sizeof(float);
step_y /= sizeof(float);
if (x < cols && y < rows)
{
const float xcoo = c_warpMat[0] * x + c_warpMat[1] * y + c_warpMat[2];
const float ycoo = c_warpMat[3] * x + c_warpMat[4] * y + c_warpMat[5];
int xmap_index = mad24(y, xmap_step, x + xmap_offset);
int ymap_index = mad24(y, ymap_step, x + ymap_offset);
float xcoo = c_warpMat[0] * x + c_warpMat[1] * y + c_warpMat[2];
float ycoo = c_warpMat[3] * x + c_warpMat[4] * y + c_warpMat[5];
map_x[y * step_x + x] = xcoo;
map_y[y * step_y + x] = ycoo;
xmap[xmap_index] = xcoo;
ymap[ymap_index] = ycoo;
}
}
__kernel
void buildWarpPerspectiveMaps
(
__global float * xmap,
__global float * ymap,
__constant float * c_warpMat,
int cols,
int rows,
int step_x,
int step_y
)
__kernel void buildWarpPerspectiveMaps(__global float * xmap, __global float * ymap,
__constant float * c_warpMat,
int cols, int rows,
int xmap_step, int ymap_step,
int xmap_offset, int ymap_offset)
{
int x = get_global_id(0);
int y = get_global_id(1);
step_x /= sizeof(float);
step_y /= sizeof(float);
if (x < cols && y < rows)
{
const float coeff = 1.0f / (c_warpMat[6] * x + c_warpMat[7] * y + c_warpMat[8]);
int xmap_index = mad24(y, xmap_step, x + xmap_offset);
int ymap_index = mad24(y, ymap_step, x + ymap_offset);
const float xcoo = coeff * (c_warpMat[0] * x + c_warpMat[1] * y + c_warpMat[2]);
const float ycoo = coeff * (c_warpMat[3] * x + c_warpMat[4] * y + c_warpMat[5]);
float coeff = 1.0f / (c_warpMat[6] * x + c_warpMat[7] * y + c_warpMat[8]);
float xcoo = coeff * (c_warpMat[0] * x + c_warpMat[1] * y + c_warpMat[2]);
float ycoo = coeff * (c_warpMat[3] * x + c_warpMat[4] * y + c_warpMat[5]);
map_x[y * step_x + x] = xcoo;
map_y[y * step_y + x] = ycoo;
xmap[xmap_index] = xcoo;
ymap[ymap_index] = ycoo;
}
}
modules/ocl/test/test_warp.cpp
View file @ 46e423d2
......@@ -156,6 +156,114 @@ OCL_TEST_P(WarpPerspective, Mat)
}
}
// buildWarpPerspectiveMaps
PARAM_TEST_CASE(BuildWarpPerspectiveMaps, bool, bool)
{
bool useRoi, mapInverse;
Size dsize;
Mat xmap_whole, ymap_whole, xmap_roi, ymap_roi;
ocl::oclMat gxmap_whole, gymap_whole, gxmap_roi, gymap_roi;
void SetUp()
{
mapInverse = GET_PARAM(0);
useRoi = GET_PARAM(1);
}
void random_roi()
{
dsize = randomSize(1, MAX_VALUE);
Border xmapBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
randomSubMat(xmap_whole, xmap_roi, dsize, xmapBorder, CV_32FC1, -MAX_VALUE, MAX_VALUE);
Border ymapBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
randomSubMat(ymap_whole, ymap_roi, dsize, ymapBorder, CV_32FC1, -MAX_VALUE, MAX_VALUE);
generateOclMat(gxmap_whole, gxmap_roi, xmap_whole, dsize, xmapBorder);
generateOclMat(gymap_whole, gymap_roi, ymap_whole, dsize, ymapBorder);
}
void Near(double threshold = 0.0)
{
Mat whole, roi;
gxmap_whole.download(whole);
gxmap_roi.download(roi);
EXPECT_MAT_NEAR(xmap_whole, whole, threshold);
EXPECT_MAT_NEAR(xmap_roi, roi, threshold);
}
void Near1(double threshold = 0.0)
{
Mat whole, roi;
gymap_whole.download(whole);
gymap_roi.download(roi);
EXPECT_MAT_NEAR(ymap_whole, whole, threshold);
EXPECT_MAT_NEAR(ymap_roi, roi, threshold);
}
};
static void buildWarpPerspectiveMaps(const Mat &M, bool inverse, Size dsize, Mat &xmap, Mat &ymap)
{
CV_Assert(M.rows == 3 && M.cols == 3);
CV_Assert(dsize.area() > 0);
xmap.create(dsize, CV_32FC1);
ymap.create(dsize, CV_32FC1);
float coeffs[3 * 3];
Mat coeffsMat(3, 3, CV_32F, (void *)coeffs);
if (inverse)
M.convertTo(coeffsMat, coeffsMat.type());
else
{
cv::Mat iM;
invert(M, iM);
iM.convertTo(coeffsMat, coeffsMat.type());
}
for (int y = 0; y < dsize.height; ++y)
{
float * const xmap_ptr = xmap.ptr<float>(y);
float * const ymap_ptr = ymap.ptr<float>(y);
for (int x = 0; x < dsize.width; ++x)
{
float coeff = 1.0f / (x * coeffs[6] + y * coeffs[7] + coeffs[8]);
xmap_ptr[x] = (x * coeffs[0] + y * coeffs[1] + coeffs[2]) * coeff;
ymap_ptr[x] = (x * coeffs[3] + y * coeffs[4] + coeffs[5]) * coeff;
}
}
}
OCL_TEST_P(BuildWarpPerspectiveMaps, Mat)
{
for (int j = 0; j < LOOP_TIMES; j++)
{
random_roi();
float cols = static_cast<float>(MAX_VALUE), rows = static_cast<float>(MAX_VALUE);
float cols2 = cols / 2.0f, rows2 = rows / 2.0f;
Point2f sp[] = { Point2f(0.0f, 0.0f), Point2f(cols, 0.0f), Point2f(0.0f, rows), Point2f(cols, rows) };
Point2f dp[] = { Point2f(rng.uniform(0.0f, cols2), rng.uniform(0.0f, rows2)),
Point2f(rng.uniform(cols2, cols), rng.uniform(0.0f, rows2)),
Point2f(rng.uniform(0.0f, cols2), rng.uniform(rows2, rows)),
Point2f(rng.uniform(cols2, cols), rng.uniform(rows2, rows)) };
Mat M = getPerspectiveTransform(sp, dp);
buildWarpPerspectiveMaps(M, mapInverse, dsize, xmap_roi, ymap_roi);
ocl::buildWarpPerspectiveMaps(M, mapInverse, dsize, gxmap_roi, gymap_roi);
Near(5e-3);
Near1(5e-3);
}
}
/////////////////////////////////////////////////////////////////////////////////////////////////
// remap
......@@ -338,6 +446,8 @@ INSTANTIATE_TEST_CASE_P(ImgprocWarp, WarpPerspective, Combine(
Bool(),
Bool()));
INSTANTIATE_TEST_CASE_P(ImgprocWarp, BuildWarpPerspectiveMaps, Combine(Bool(), Bool()));
INSTANTIATE_TEST_CASE_P(ImgprocWarp, Remap_INTER_LINEAR, Combine(
Values(CV_8U, CV_16U, CV_16S, CV_32F, CV_64F),
Values(1, 2, 3, 4),
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment