added gpu::buildWarpPlaneMaps and gpu::buildWarpCylindricalMaps functions, integrated into stitching

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

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<PlaneMapper><<<grid,threads>>>(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<CylindricalMapper><<<grid,threads>>>(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
 
 }}}
 
+

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<float>(), Rinv.ptr<float>(),
+                                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<float>(), Rinv.ptr<float>(),
+                                      f, s, 0.5f*src_size.width, 0.5f*src_size.height, StreamAccessor::getStream(stream));
+}
+
 
 //////////////////////////////////////////////////////////////////////////////
 // buildWarpSphericalMaps

modules/stitching/blenders.cpp

@@ -233,9 +233,6 @@ void MultiBandBlender::feed(const Mat &img, const Mat &mask, Point tl)
     copyMakeBorder(img, img_with_border, top, bottom, left, right,
                    BORDER_REFLECT);
     vector<Mat> 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);
 
     // Create the weight map Gaussian pyramid
@@ -341,6 +338,7 @@ void createLaplacePyr(const Mat &img, int num_levels, vector<Mat> &pyr)
 }
 
 
+#if 0
 void createLaplacePyrGpu(const Mat &img, int num_levels, vector<Mat> &pyr)
 {
     pyr.resize(num_levels + 1);
@@ -360,7 +358,7 @@ void createLaplacePyrGpu(const Mat &img, int num_levels, vector<Mat> &pyr)
 
     pyr[num_levels] = gpu_pyr[num_levels];
 }
-
+#endif
 
 
 void restoreImageFromLaplacePyr(vector<Mat> &pyr)

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<cv::Mat>& 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<cv::Mat>& pyr);
+#endif
 
 // Restores source image in-place (result will be stored in pyr[0])
 void restoreImageFromLaplacePyr(std::vector<cv::Mat>& pyr);

modules/stitching/warpers.cpp

@@ -48,9 +48,9 @@ Ptr<Warper> 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;
+}

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<CylindricalProjector>
 public:
     CylindricalWarper(float scale = 300.f) { projector_.scale = scale; }
 
-private:
+protected:
     void detectResultRoi(cv::Point &dst_tl, cv::Point &dst_br)
     {
         WarperBase<CylindricalProjector>::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__