From dce50b67fc3c3cd2b447d2c5769c2b895f609cab Mon Sep 17 00:00:00 2001 From: Alexey Spizhevoy Date: Mon, 3 Oct 2011 07:42:16 +0000 Subject: [PATCH] Fixed support of translation in the GPU-based plane warper --- modules/gpu/include/opencv2/gpu/gpu.hpp | 2 +- modules/gpu/perf/perf_imgproc.cpp | 2 +- modules/gpu/src/cuda/imgproc.cu | 24 +++++++++++++----------- modules/gpu/src/imgproc.cpp | 13 +++++++------ modules/stitching/src/warpers.cpp | 4 ++-- 5 files changed, 24 insertions(+), 21 deletions(-) diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index dba9b30..51aadeb 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -671,7 +671,7 @@ namespace cv CV_EXPORTS void warpPerspective(const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags = INTER_LINEAR, Stream& stream = Stream::Null()); //! builds plane warping maps - CV_EXPORTS void buildWarpPlaneMaps(Size src_size, Rect dst_roi, const Mat &K, const Mat& R, float scale, + CV_EXPORTS void buildWarpPlaneMaps(Size src_size, Rect dst_roi, const Mat &K, const Mat& R, const Mat &T, float scale, GpuMat& map_x, GpuMat& map_y, Stream& stream = Stream::Null()); //! builds cylindrical warping maps diff --git a/modules/gpu/perf/perf_imgproc.cpp b/modules/gpu/perf/perf_imgproc.cpp index 98e7edd..f239edb 100644 --- a/modules/gpu/perf/perf_imgproc.cpp +++ b/modules/gpu/perf/perf_imgproc.cpp @@ -367,7 +367,7 @@ PERF_TEST_P(DevInfo_Size, buildWarpPlaneMaps, testing::Combine(testing::ValuesIn SIMPLE_TEST_CYCLE() { buildWarpPlaneMaps(size, Rect(0, 0, size.width, size.height), Mat::eye(3, 3, CV_32FC1), - Mat::ones(3, 3, CV_32FC1), 1.0, map_x, map_y); + Mat::ones(3, 3, CV_32FC1), Mat::zeros(1, 3, CV_32F), 1.0, map_x, map_y); } Mat map_x_host(map_x); diff --git a/modules/gpu/src/cuda/imgproc.cu b/modules/gpu/src/cuda/imgproc.cu index 927c383..79b5e10 100644 --- a/modules/gpu/src/cuda/imgproc.cu +++ b/modules/gpu/src/cuda/imgproc.cu @@ -794,6 +794,7 @@ namespace cv { namespace gpu { namespace imgproc __constant__ float ck_rinv[9]; __constant__ float cr_kinv[9]; + __constant__ float ct[3]; __constant__ float cscale; } @@ -805,13 +806,13 @@ namespace cv { namespace gpu { namespace imgproc { using namespace build_warp_maps; - float x_ = u / cscale; - float y_ = v / cscale; + float x_ = u / cscale - ct[0]; + float y_ = v / cscale - ct[1]; float z; - x = ck_rinv[0] * x_ + ck_rinv[1] * y_ + ck_rinv[2]; - y = ck_rinv[3] * x_ + ck_rinv[4] * y_ + ck_rinv[5]; - z = ck_rinv[6] * x_ + ck_rinv[7] * y_ + ck_rinv[8]; + x = ck_rinv[0] * x_ + ck_rinv[1] * y_ + ck_rinv[2] * (1 - ct[2]); + y = ck_rinv[3] * x_ + ck_rinv[4] * y_ + ck_rinv[5] * (1 - ct[2]); + z = ck_rinv[6] * x_ + ck_rinv[7] * y_ + ck_rinv[8] * (1 - ct[2]); x /= z; y /= z; @@ -887,11 +888,12 @@ namespace cv { namespace gpu { namespace imgproc void buildWarpPlaneMaps(int tl_u, int tl_v, DevMem2Df map_x, DevMem2Df map_y, - const float k_rinv[9], const float r_kinv[9], float scale, - cudaStream_t stream) + const float k_rinv[9], const float r_kinv[9], const float t[3], + float scale, cudaStream_t stream) { cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ck_rinv, k_rinv, 9*sizeof(float))); cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr_kinv, r_kinv, 9*sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ct, t, 3*sizeof(float))); cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cscale, &scale, sizeof(float))); int cols = map_x.cols; @@ -908,8 +910,8 @@ namespace cv { namespace gpu { namespace imgproc void buildWarpCylindricalMaps(int tl_u, int tl_v, DevMem2Df map_x, DevMem2Df map_y, - const float k_rinv[9], const float r_kinv[9], float scale, - cudaStream_t stream) + const float k_rinv[9], const float r_kinv[9], float scale, + cudaStream_t stream) { cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ck_rinv, k_rinv, 9*sizeof(float))); cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr_kinv, r_kinv, 9*sizeof(float))); @@ -929,8 +931,8 @@ namespace cv { namespace gpu { namespace imgproc void buildWarpSphericalMaps(int tl_u, int tl_v, DevMem2Df map_x, DevMem2Df map_y, - const float k_rinv[9], const float r_kinv[9], float scale, - cudaStream_t stream) + const float k_rinv[9], const float r_kinv[9], float scale, + cudaStream_t stream) { cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ck_rinv, k_rinv, 9*sizeof(float))); cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr_kinv, r_kinv, 9*sizeof(float))); diff --git a/modules/gpu/src/imgproc.cpp b/modules/gpu/src/imgproc.cpp index 87d3b5d..8b86ce6 100644 --- a/modules/gpu/src/imgproc.cpp +++ b/modules/gpu/src/imgproc.cpp @@ -56,7 +56,7 @@ void cv::gpu::resize(const GpuMat&, GpuMat&, Size, double, double, int, Stream&) void cv::gpu::copyMakeBorder(const GpuMat&, GpuMat&, int, int, int, int, int, const Scalar&, Stream&) { throw_nogpu(); } void cv::gpu::warpAffine(const GpuMat&, GpuMat&, const Mat&, Size, int, Stream&) { throw_nogpu(); } void cv::gpu::warpPerspective(const GpuMat&, GpuMat&, const Mat&, Size, int, Stream&) { throw_nogpu(); } -void cv::gpu::buildWarpPlaneMaps(Size, Rect, const Mat&, const Mat&, float, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); } +void cv::gpu::buildWarpPlaneMaps(Size, Rect, const Mat&, const Mat&, const Mat&, float, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); } void cv::gpu::buildWarpCylindricalMaps(Size, Rect, const Mat&, const Mat&, float, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); } void cv::gpu::buildWarpSphericalMaps(Size, Rect, const Mat&, const Mat&, float, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); } void cv::gpu::rotate(const GpuMat&, GpuMat&, Size, double, double, double, int, Stream&) { throw_nogpu(); } @@ -623,15 +623,16 @@ void cv::gpu::warpPerspective(const GpuMat& src, GpuMat& dst, const Mat& M, Size namespace cv { namespace gpu { namespace imgproc { void buildWarpPlaneMaps(int tl_u, int tl_v, DevMem2Df map_x, DevMem2Df map_y, - const float k_rinv[9], const float r_kinv[9], float scale, + const float k_rinv[9], const float r_kinv[9], const float t[3], float scale, cudaStream_t stream); }}} -void cv::gpu::buildWarpPlaneMaps(Size src_size, Rect dst_roi, const Mat &K, const Mat& R, float scale, - GpuMat& map_x, GpuMat& map_y, Stream& stream) +void cv::gpu::buildWarpPlaneMaps(Size src_size, Rect dst_roi, const Mat &K, const Mat& R, const Mat &T, + float scale, GpuMat& map_x, GpuMat& map_y, Stream& stream) { CV_Assert(K.size() == Size(3,3) && K.type() == CV_32F); CV_Assert(R.size() == Size(3,3) && R.type() == CV_32F); + CV_Assert((T.size() == Size(3,1) || T.size() == Size(1,3)) && T.type() == CV_32F && T.isContinuous()); Mat K_Rinv = K * R.t(); Mat R_Kinv = R * K.inv(); @@ -640,8 +641,8 @@ void cv::gpu::buildWarpPlaneMaps(Size src_size, Rect dst_roi, const Mat &K, cons map_x.create(dst_roi.size(), CV_32F); map_y.create(dst_roi.size(), CV_32F); - imgproc::buildWarpPlaneMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, K_Rinv.ptr(), R_Kinv.ptr(), - scale, StreamAccessor::getStream(stream)); + imgproc::buildWarpPlaneMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, K_Rinv.ptr(), R_Kinv.ptr(), + T.ptr(), scale, StreamAccessor::getStream(stream)); } ////////////////////////////////////////////////////////////////////////////// diff --git a/modules/stitching/src/warpers.cpp b/modules/stitching/src/warpers.cpp index b96b14a..3bd0497 100644 --- a/modules/stitching/src/warpers.cpp +++ b/modules/stitching/src/warpers.cpp @@ -216,8 +216,8 @@ Rect PlaneWarperGpu::buildMaps(Size src_size, const Mat &K, const Mat &R, const Point dst_tl, dst_br; detectResultRoi(src_size, dst_tl, dst_br); - gpu::buildWarpPlaneMaps(src_size, Rect(dst_tl, Point(dst_br.x+1, dst_br.y+1)), - K, R, projector_.scale, xmap, ymap); + gpu::buildWarpPlaneMaps(src_size, Rect(dst_tl, Point(dst_br.x + 1, dst_br.y + 1)), + K, R, T, projector_.scale, xmap, ymap); return Rect(dst_tl, dst_br); } -- 2.7.4