From ad454d83b9900f6dd640f7775c6dc6dccd352d37 Mon Sep 17 00:00:00 2001 From: Alexey Spizhevoy Date: Fri, 1 Jul 2011 07:07:54 +0000 Subject: [PATCH] added gpu::buildWarpPlaneMaps and gpu::buildWarpCylindricalMaps functions, integrated into stitching --- modules/gpu/include/opencv2/gpu/gpu.hpp | 8 +++ modules/gpu/src/cuda/imgproc.cu | 95 +++++++++++++++++++++++++++++++++ modules/gpu/src/imgproc_gpu.cpp | 51 +++++++++++++++++- modules/stitching/blenders.cpp | 8 ++- modules/stitching/blenders.hpp | 4 ++ modules/stitching/main.cpp | 2 +- modules/stitching/warpers.cpp | 45 +++++++++++++++- modules/stitching/warpers.hpp | 28 +++++++++- 8 files changed, 229 insertions(+), 12 deletions(-) diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 16539b2..4a4df8a 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -622,6 +622,14 @@ namespace cv //! Supports INTER_NEAREST, INTER_LINEAR, INTER_CUBIC 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& R, double f, double s, double dist, + GpuMat& map_x, GpuMat& map_y, Stream& stream = Stream::Null()); + + //! builds cylindrical warping maps + CV_EXPORTS void buildWarpCylindricalMaps(Size src_size, Rect dst_roi, const Mat& R, double f, double s, + GpuMat& map_x, GpuMat& map_y, Stream& stream = Stream::Null()); + //! builds spherical warping maps CV_EXPORTS void buildWarpSphericalMaps(Size src_size, Rect dst_roi, const Mat& R, double f, double s, GpuMat& map_x, GpuMat& map_y, Stream& stream = Stream::Null()); diff --git a/modules/gpu/src/cuda/imgproc.cu b/modules/gpu/src/cuda/imgproc.cu index c1e1ef4..e89b055 100644 --- a/modules/gpu/src/cuda/imgproc.cu +++ b/modules/gpu/src/cuda/imgproc.cu @@ -986,9 +986,54 @@ namespace cv { namespace gpu { namespace imgproc __constant__ float crinv[9]; __constant__ float cf, cs; __constant__ float chalf_w, chalf_h; + __constant__ float cdist; } + class PlaneMapper + { + public: + static __device__ __forceinline__ void mapBackward(float u, float v, float &x, float &y) + { + using namespace build_warp_maps; + + float x_ = u / cs; + float y_ = v / cs; + + float z; + x = crinv[0]*x_ + crinv[1]*y_ + crinv[2]*cdist; + y = crinv[3]*x_ + crinv[4]*y_ + crinv[5]*cdist; + z = crinv[6]*x_ + crinv[7]*y_ + crinv[8]*cdist; + + x = cf*x/z + chalf_w; + y = cf*y/z + chalf_h; + } + }; + + + class CylindricalMapper + { + public: + static __device__ __forceinline__ void mapBackward(float u, float v, float &x, float &y) + { + using namespace build_warp_maps; + + u /= cs; + float x_ = sinf(u); + float y_ = v / cs; + float z_ = cosf(u); + + float z; + x = crinv[0]*x_ + crinv[1]*y_ + crinv[2]*z_; + y = crinv[3]*x_ + crinv[4]*y_ + crinv[5]*z_; + z = crinv[6]*x_ + crinv[7]*y_ + crinv[8]*z_; + + x = cf*x/z + chalf_w; + y = cf*y/z + chalf_h; + } + }; + + class SphericalMapper { public: @@ -1033,6 +1078,55 @@ namespace cv { namespace gpu { namespace imgproc } + void buildWarpPlaneMaps(int tl_u, int tl_v, DevMem2Df map_x, DevMem2Df map_y, + const float r[9], const float rinv[9], float f, float s, float dist, + float half_w, float half_h, cudaStream_t stream) + { + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr, r, 9*sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::crinv, rinv, 9*sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cf, &f, sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cs, &s, sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::chalf_w, &half_w, sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::chalf_h, &half_h, sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cdist, &dist, sizeof(float))); + + int cols = map_x.cols; + int rows = map_x.rows; + + dim3 threads(32, 8); + dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y)); + + buildWarpMapsKernel<<>>(tl_u, tl_v, cols, rows, map_x, map_y); + cudaSafeCall(cudaGetLastError()); + if (stream == 0) + cudaSafeCall(cudaDeviceSynchronize()); + } + + + void buildWarpCylindricalMaps(int tl_u, int tl_v, DevMem2Df map_x, DevMem2Df map_y, + const float r[9], const float rinv[9], float f, float s, + float half_w, float half_h, cudaStream_t stream) + { + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr, r, 9*sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::crinv, rinv, 9*sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cf, &f, sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cs, &s, sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::chalf_w, &half_w, sizeof(float))); + cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::chalf_h, &half_h, sizeof(float))); + + int cols = map_x.cols; + int rows = map_x.rows; + + dim3 threads(32, 8); + dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y)); + + buildWarpMapsKernel<<>>(tl_u, tl_v, cols, rows, map_x, map_y); + cudaSafeCall(cudaGetLastError()); + if (stream == 0) + cudaSafeCall(cudaDeviceSynchronize()); + } + + void buildWarpSphericalMaps(int tl_u, int tl_v, DevMem2Df map_x, DevMem2Df map_y, const float r[9], const float rinv[9], float f, float s, float half_w, float half_h, cudaStream_t stream) @@ -1059,3 +1153,4 @@ namespace cv { namespace gpu { namespace imgproc }}} + diff --git a/modules/gpu/src/imgproc_gpu.cpp b/modules/gpu/src/imgproc_gpu.cpp index 85e99f8..de23877 100644 --- a/modules/gpu/src/imgproc_gpu.cpp +++ b/modules/gpu/src/imgproc_gpu.cpp @@ -56,8 +56,9 @@ void cv::gpu::resize(const GpuMat&, GpuMat&, Size, double, double, int, Stream&) void cv::gpu::copyMakeBorder(const GpuMat&, GpuMat&, 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::buildWarpSphericalMaps(Size, Rect, const Mat&, double, double, - GpuMat&, GpuMat&, Stream&) { throw_nogpu(); } +void cv::gpu::buildWarpPlaneMaps(Size, Rect, const Mat&, double, double, double, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); } +void cv::gpu::buildWarpCylindricalMaps(Size, Rect, const Mat&, double, double, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); } +void cv::gpu::buildWarpSphericalMaps(Size, Rect, const Mat&, double, double, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); } void cv::gpu::rotate(const GpuMat&, GpuMat&, Size, double, double, double, int, Stream&) { throw_nogpu(); } void cv::gpu::integral(const GpuMat&, GpuMat&, Stream&) { throw_nogpu(); } void cv::gpu::integralBuffered(const GpuMat&, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); } @@ -510,6 +511,52 @@ void cv::gpu::warpPerspective(const GpuMat& src, GpuMat& dst, const Mat& M, Size nppWarpCaller(src, dst, coeffs, dsize, flags, npp_warpPerspective_8u, npp_warpPerspective_16u, npp_warpPerspective_32s, npp_warpPerspective_32f, StreamAccessor::getStream(s)); } +////////////////////////////////////////////////////////////////////////////// +// buildWarpPlaneMaps + +namespace cv { namespace gpu { namespace imgproc +{ + void buildWarpPlaneMaps(int tl_u, int tl_v, DevMem2Df map_x, DevMem2Df map_y, + const float r[9], const float rinv[9], float f, float s, float dist, + float half_w, float half_h, cudaStream_t stream); +}}} + +void cv::gpu::buildWarpPlaneMaps(Size src_size, Rect dst_roi, const Mat& R, double f, double s, + double dist, GpuMat& map_x, GpuMat& map_y, Stream& stream) +{ + CV_Assert(R.size() == Size(3,3) && R.isContinuous() && R.type() == CV_32F); + Mat Rinv = R.inv(); + CV_Assert(Rinv.isContinuous()); + + 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, R.ptr(), Rinv.ptr(), + f, s, dist, 0.5f*src_size.width, 0.5f*src_size.height, StreamAccessor::getStream(stream)); +} + +////////////////////////////////////////////////////////////////////////////// +// buildWarpCylyndricalMaps + +namespace cv { namespace gpu { namespace imgproc +{ + void buildWarpCylindricalMaps(int tl_u, int tl_v, DevMem2Df map_x, DevMem2Df map_y, + const float r[9], const float rinv[9], float f, float s, + float half_w, float half_h, cudaStream_t stream); +}}} + +void cv::gpu::buildWarpCylindricalMaps(Size src_size, Rect dst_roi, const Mat& R, double f, double s, + GpuMat& map_x, GpuMat& map_y, Stream& stream) +{ + CV_Assert(R.size() == Size(3,3) && R.isContinuous() && R.type() == CV_32F); + Mat Rinv = R.inv(); + CV_Assert(Rinv.isContinuous()); + + map_x.create(dst_roi.size(), CV_32F); + map_y.create(dst_roi.size(), CV_32F); + imgproc::buildWarpCylindricalMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, R.ptr(), Rinv.ptr(), + f, s, 0.5f*src_size.width, 0.5f*src_size.height, StreamAccessor::getStream(stream)); +} + ////////////////////////////////////////////////////////////////////////////// // buildWarpSphericalMaps diff --git a/modules/stitching/blenders.cpp b/modules/stitching/blenders.cpp index b41fa45..99e1b15 100644 --- a/modules/stitching/blenders.cpp +++ b/modules/stitching/blenders.cpp @@ -233,10 +233,7 @@ void MultiBandBlender::feed(const Mat &img, const Mat &mask, Point tl) copyMakeBorder(img, img_with_border, top, bottom, left, right, BORDER_REFLECT); vector src_pyr_laplace; - if (can_use_gpu_) - createLaplacePyrGpu(img_with_border, num_bands_, src_pyr_laplace); - else - createLaplacePyr(img_with_border, num_bands_, src_pyr_laplace); + createLaplacePyr(img_with_border, num_bands_, src_pyr_laplace); // Create the weight map Gaussian pyramid Mat weight_map; @@ -341,6 +338,7 @@ void createLaplacePyr(const Mat &img, int num_levels, vector &pyr) } +#if 0 void createLaplacePyrGpu(const Mat &img, int num_levels, vector &pyr) { pyr.resize(num_levels + 1); @@ -360,7 +358,7 @@ void createLaplacePyrGpu(const Mat &img, int num_levels, vector &pyr) pyr[num_levels] = gpu_pyr[num_levels]; } - +#endif void restoreImageFromLaplacePyr(vector &pyr) diff --git a/modules/stitching/blenders.hpp b/modules/stitching/blenders.hpp index 04ede3c..af52f8a 100644 --- a/modules/stitching/blenders.hpp +++ b/modules/stitching/blenders.hpp @@ -108,7 +108,11 @@ void normalize(const cv::Mat& weight, cv::Mat& src); void createWeightMap(const cv::Mat& mask, float sharpness, cv::Mat& weight); void createLaplacePyr(const cv::Mat &img, int num_levels, std::vector& pyr); + +// TODO Use it after gpu::pyrDown and gpu::pyrUp are updated +#if 0 void createLaplacePyrGpu(const cv::Mat &img, int num_levels, std::vector& pyr); +#endif // Restores source image in-place (result will be stored in pyr[0]) void restoreImageFromLaplacePyr(std::vector& pyr); diff --git a/modules/stitching/main.cpp b/modules/stitching/main.cpp index defb32b..26226eb 100644 --- a/modules/stitching/main.cpp +++ b/modules/stitching/main.cpp @@ -547,7 +547,7 @@ int main(int argc, char* argv[]) else img = full_img; full_img.release(); - Size img_size = img.size(); + Size img_size = img.size(); // Warp the current image warper->warp(img, static_cast(cameras[img_idx].focal), cameras[img_idx].R, diff --git a/modules/stitching/warpers.cpp b/modules/stitching/warpers.cpp index 68916ef..aceedac 100644 --- a/modules/stitching/warpers.cpp +++ b/modules/stitching/warpers.cpp @@ -48,9 +48,9 @@ Ptr Warper::createByCameraFocal(float focal, int type, bool try_gpu) { bool can_use_gpu = try_gpu && gpu::getCudaEnabledDeviceCount(); if (type == PLANE) - return new PlaneWarper(focal); + return !can_use_gpu ? new PlaneWarper(focal) : new PlaneWarperGpu(focal); if (type == CYLINDRICAL) - return new CylindricalWarper(focal); + return !can_use_gpu ? new CylindricalWarper(focal) : new CylindricalWarperGpu(focal); if (type == SPHERICAL) return !can_use_gpu ? new SphericalWarper(focal) : new SphericalWarperGpu(focal); CV_Error(CV_StsBadArg, "unsupported warping type"); @@ -105,6 +105,26 @@ void PlaneWarper::detectResultRoi(Point &dst_tl, Point &dst_br) } +Point PlaneWarperGpu::warp(const Mat &src, float focal, const cv::Mat &R, cv::Mat &dst, int interp_mode, int border_mode) +{ + src_size_ = src.size(); + projector_.size = src.size(); + projector_.focal = focal; + projector_.setTransformation(R); + + cv::Point dst_tl, dst_br; + detectResultRoi(dst_tl, dst_br); + + gpu::buildWarpPlaneMaps(src.size(), Rect(dst_tl, Point(dst_br.x+1, dst_br.y+1)), + R, focal, projector_.scale, projector_.plane_dist, d_xmap_, d_ymap_); + + dst.create(dst_br.y - dst_tl.y + 1, dst_br.x - dst_tl.x + 1, src.type()); + remap(src, dst, Mat(d_xmap_), Mat(d_ymap_), interp_mode, border_mode); + + return dst_tl; +} + + void SphericalWarper::detectResultRoi(Point &dst_tl, Point &dst_br) { detectResultRoiByBorder(dst_tl, dst_br); @@ -168,3 +188,24 @@ Point SphericalWarperGpu::warp(const Mat &src, float focal, const Mat &R, Mat &d return dst_tl; } + + +Point CylindricalWarperGpu::warp(const Mat &src, float focal, const Mat &R, Mat &dst, + int interp_mode, int border_mode) +{ + src_size_ = src.size(); + projector_.size = src.size(); + projector_.focal = focal; + projector_.setTransformation(R); + + cv::Point dst_tl, dst_br; + detectResultRoi(dst_tl, dst_br); + + gpu::buildWarpCylindricalMaps(src.size(), Rect(dst_tl, Point(dst_br.x+1, dst_br.y+1)), + R, focal, projector_.scale, d_xmap_, d_ymap_); + + dst.create(dst_br.y - dst_tl.y + 1, dst_br.x - dst_tl.x + 1, src.type()); + remap(src, dst, Mat(d_xmap_), Mat(d_ymap_), interp_mode, border_mode); + + return dst_tl; +} diff --git a/modules/stitching/warpers.hpp b/modules/stitching/warpers.hpp index 854c87b..47a162e 100644 --- a/modules/stitching/warpers.hpp +++ b/modules/stitching/warpers.hpp @@ -109,11 +109,23 @@ public: projector_.scale = scale; } -private: +protected: void detectResultRoi(cv::Point &dst_tl, cv::Point &dst_br); }; +class PlaneWarperGpu : public PlaneWarper +{ +public: + PlaneWarperGpu(float plane_dist = 1.f, float scale = 1.f) : PlaneWarper(plane_dist, scale) {} + cv::Point warp(const cv::Mat &src, float focal, const cv::Mat &R, cv::Mat &dst, + int interp_mode, int border_mode); + +private: + cv::gpu::GpuMat d_xmap_, d_ymap_, d_dst_; +}; + + struct SphericalProjector : ProjectorBase { void mapForward(float x, float y, float &u, float &v); @@ -158,13 +170,25 @@ class CylindricalWarper : public WarperBase public: CylindricalWarper(float scale = 300.f) { projector_.scale = scale; } -private: +protected: void detectResultRoi(cv::Point &dst_tl, cv::Point &dst_br) { WarperBase::detectResultRoiByBorder(dst_tl, dst_br); } }; + +class CylindricalWarperGpu : public CylindricalWarper +{ +public: + CylindricalWarperGpu(float scale = 300.f) : CylindricalWarper(scale) {} + cv::Point warp(const cv::Mat &src, float focal, const cv::Mat &R, cv::Mat &dst, + int interp_mode, int border_mode); + +private: + cv::gpu::GpuMat d_xmap_, d_ymap_, d_dst_; +}; + #include "warpers_inl.hpp" #endif // __OPENCV_WARPERS_HPP__ -- 2.7.4