added gpu::buildWarpPlaneMaps and gpu::buildWarpCylindricalMaps functions, integrated...

diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp
index 16539b2..4a4df8a 100644 (file)
--- 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\r
         CV_EXPORTS void warpPerspective(const GpuMat& src, GpuMat& dst, const Mat& M, Size dsize, int flags = INTER_LINEAR, Stream& stream = Stream::Null());\r
 \r
+        //! builds plane warping maps\r
+        CV_EXPORTS void buildWarpPlaneMaps(Size src_size, Rect dst_roi, const Mat& R, double f, double s, double dist,\r
+                                           GpuMat& map_x, GpuMat& map_y, Stream& stream = Stream::Null());\r
+\r
+        //! builds cylindrical warping maps\r
+        CV_EXPORTS void buildWarpCylindricalMaps(Size src_size, Rect dst_roi, const Mat& R, double f, double s,\r
+                                                 GpuMat& map_x, GpuMat& map_y, Stream& stream = Stream::Null());\r
+\r
         //! builds spherical warping maps\r
         CV_EXPORTS void buildWarpSphericalMaps(Size src_size, Rect dst_roi, const Mat& R, double f, double s,\r
                                                GpuMat& map_x, GpuMat& map_y, Stream& stream = Stream::Null());\r

diff --git a/modules/gpu/src/cuda/imgproc.cu b/modules/gpu/src/cuda/imgproc.cu
index c1e1ef4..e89b055 100644 (file)
--- 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];\r
         __constant__ float cf, cs;\r
         __constant__ float chalf_w, chalf_h;\r
+        __constant__ float cdist;\r
     }\r
 \r
 \r
+    class PlaneMapper\r
+    {\r
+    public:\r
+        static __device__ __forceinline__ void mapBackward(float u, float v, float &x, float &y)\r
+        {\r
+            using namespace build_warp_maps;\r
+\r
+            float x_ = u / cs;\r
+            float y_ = v / cs;\r
+\r
+            float z;\r
+            x = crinv[0]*x_ + crinv[1]*y_ + crinv[2]*cdist;\r
+            y = crinv[3]*x_ + crinv[4]*y_ + crinv[5]*cdist;\r
+            z = crinv[6]*x_ + crinv[7]*y_ + crinv[8]*cdist;\r
+\r
+            x = cf*x/z + chalf_w;\r
+            y = cf*y/z + chalf_h;\r
+        }\r
+    };\r
+\r
+\r
+    class CylindricalMapper\r
+    {\r
+    public:\r
+        static __device__ __forceinline__ void mapBackward(float u, float v, float &x, float &y)\r
+        {\r
+            using namespace build_warp_maps;\r
+\r
+            u /= cs;\r
+            float x_ = sinf(u);\r
+            float y_ = v / cs;\r
+            float z_ = cosf(u);\r
+\r
+            float z;\r
+            x = crinv[0]*x_ + crinv[1]*y_ + crinv[2]*z_;\r
+            y = crinv[3]*x_ + crinv[4]*y_ + crinv[5]*z_;\r
+            z = crinv[6]*x_ + crinv[7]*y_ + crinv[8]*z_;\r
+\r
+            x = cf*x/z + chalf_w;\r
+            y = cf*y/z + chalf_h;\r
+        }\r
+    };\r
+\r
+\r
     class SphericalMapper\r
     {\r
     public:\r
@@ -1033,6 +1078,55 @@ namespace cv { namespace gpu { namespace imgproc
     }\r
 \r
 \r
+    void buildWarpPlaneMaps(int tl_u, int tl_v, DevMem2Df map_x, DevMem2Df map_y,\r
+                            const float r[9], const float rinv[9], float f, float s, float dist,\r
+                            float half_w, float half_h, cudaStream_t stream)\r
+    {\r
+        cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr, r, 9*sizeof(float)));\r
+        cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::crinv, rinv, 9*sizeof(float)));\r
+        cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cf, &f, sizeof(float)));\r
+        cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cs, &s, sizeof(float)));\r
+        cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::chalf_w, &half_w, sizeof(float)));\r
+        cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::chalf_h, &half_h, sizeof(float)));\r
+        cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cdist, &dist, sizeof(float)));\r
+\r
+        int cols = map_x.cols;\r
+        int rows = map_x.rows;\r
+\r
+        dim3 threads(32, 8);\r
+        dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y));\r
+\r
+        buildWarpMapsKernel<PlaneMapper><<<grid,threads>>>(tl_u, tl_v, cols, rows, map_x, map_y);\r
+        cudaSafeCall(cudaGetLastError());\r
+        if (stream == 0)\r
+            cudaSafeCall(cudaDeviceSynchronize());\r
+    }\r
+\r
+\r
+    void buildWarpCylindricalMaps(int tl_u, int tl_v, DevMem2Df map_x, DevMem2Df map_y,\r
+                                  const float r[9], const float rinv[9], float f, float s,\r
+                                  float half_w, float half_h, cudaStream_t stream)\r
+    {\r
+        cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr, r, 9*sizeof(float)));\r
+        cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::crinv, rinv, 9*sizeof(float)));\r
+        cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cf, &f, sizeof(float)));\r
+        cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cs, &s, sizeof(float)));\r
+        cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::chalf_w, &half_w, sizeof(float)));\r
+        cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::chalf_h, &half_h, sizeof(float)));\r
+\r
+        int cols = map_x.cols;\r
+        int rows = map_x.rows;\r
+\r
+        dim3 threads(32, 8);\r
+        dim3 grid(divUp(cols, threads.x), divUp(rows, threads.y));\r
+\r
+        buildWarpMapsKernel<CylindricalMapper><<<grid,threads>>>(tl_u, tl_v, cols, rows, map_x, map_y);\r
+        cudaSafeCall(cudaGetLastError());\r
+        if (stream == 0)\r
+            cudaSafeCall(cudaDeviceSynchronize());\r
+    }\r
+\r
+\r
     void buildWarpSphericalMaps(int tl_u, int tl_v, DevMem2Df map_x, DevMem2Df map_y,\r
                                 const float r[9], const float rinv[9], float f, float s,\r
                                 float half_w, float half_h, cudaStream_t stream)\r
@@ -1059,3 +1153,4 @@ namespace cv { namespace gpu { namespace imgproc
 \r
 }}}\r
 \r
+\r

diff --git a/modules/gpu/src/imgproc_gpu.cpp b/modules/gpu/src/imgproc_gpu.cpp
index 85e99f8..de23877 100644 (file)
--- 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(); }\r
 void cv::gpu::warpAffine(const GpuMat&, GpuMat&, const Mat&, Size, int, Stream&) { throw_nogpu(); }\r
 void cv::gpu::warpPerspective(const GpuMat&, GpuMat&, const Mat&, Size, int, Stream&) { throw_nogpu(); }\r
-void cv::gpu::buildWarpSphericalMaps(Size, Rect, const Mat&, double, double,\r
-                                     GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
+void cv::gpu::buildWarpPlaneMaps(Size, Rect, const Mat&, double, double, double, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
+void cv::gpu::buildWarpCylindricalMaps(Size, Rect, const Mat&, double, double, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
+void cv::gpu::buildWarpSphericalMaps(Size, Rect, const Mat&, double, double, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
 void cv::gpu::rotate(const GpuMat&, GpuMat&, Size, double, double, double, int, Stream&) { throw_nogpu(); }\r
 void cv::gpu::integral(const GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
 void cv::gpu::integralBuffered(const GpuMat&, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
@@ -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));\r
 }\r
 \r
+//////////////////////////////////////////////////////////////////////////////\r
+// buildWarpPlaneMaps\r
+\r
+namespace cv { namespace gpu { namespace imgproc\r
+{\r
+    void buildWarpPlaneMaps(int tl_u, int tl_v, DevMem2Df map_x, DevMem2Df map_y,\r
+                            const float r[9], const float rinv[9], float f, float s, float dist,\r
+                            float half_w, float half_h, cudaStream_t stream);\r
+}}}\r
+\r
+void cv::gpu::buildWarpPlaneMaps(Size src_size, Rect dst_roi, const Mat& R, double f, double s,\r
+                                 double dist, GpuMat& map_x, GpuMat& map_y, Stream& stream)\r
+{\r
+    CV_Assert(R.size() == Size(3,3) && R.isContinuous() && R.type() == CV_32F);\r
+    Mat Rinv = R.inv();\r
+    CV_Assert(Rinv.isContinuous());\r
+\r
+    map_x.create(dst_roi.size(), CV_32F);\r
+    map_y.create(dst_roi.size(), CV_32F);\r
+    imgproc::buildWarpPlaneMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, R.ptr<float>(), Rinv.ptr<float>(),\r
+                                f, s, dist, 0.5f*src_size.width, 0.5f*src_size.height, StreamAccessor::getStream(stream));\r
+}\r
+\r
+//////////////////////////////////////////////////////////////////////////////\r
+// buildWarpCylyndricalMaps\r
+\r
+namespace cv { namespace gpu { namespace imgproc\r
+{\r
+    void buildWarpCylindricalMaps(int tl_u, int tl_v, DevMem2Df map_x, DevMem2Df map_y,\r
+                                  const float r[9], const float rinv[9], float f, float s,\r
+                                  float half_w, float half_h, cudaStream_t stream);\r
+}}}\r
+\r
+void cv::gpu::buildWarpCylindricalMaps(Size src_size, Rect dst_roi, const Mat& R, double f, double s,\r
+                                       GpuMat& map_x, GpuMat& map_y, Stream& stream)\r
+{\r
+    CV_Assert(R.size() == Size(3,3) && R.isContinuous() && R.type() == CV_32F);\r
+    Mat Rinv = R.inv();\r
+    CV_Assert(Rinv.isContinuous());\r
+\r
+    map_x.create(dst_roi.size(), CV_32F);\r
+    map_y.create(dst_roi.size(), CV_32F);\r
+    imgproc::buildWarpCylindricalMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, R.ptr<float>(), Rinv.ptr<float>(),\r
+                                      f, s, 0.5f*src_size.width, 0.5f*src_size.height, StreamAccessor::getStream(stream));\r
+}\r
+\r
 \r
 //////////////////////////////////////////////////////////////////////////////\r
 // buildWarpSphericalMaps\r

diff --git a/modules/stitching/blenders.cpp b/modules/stitching/blenders.cpp
index b41fa45..99e1b15 100644 (file)
--- 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,\r
                    BORDER_REFLECT);\r
     vector<Mat> src_pyr_laplace;\r
-    if (can_use_gpu_)\r
-        createLaplacePyrGpu(img_with_border, num_bands_, src_pyr_laplace);\r
-    else\r
-        createLaplacePyr(img_with_border, num_bands_, src_pyr_laplace);\r
+    createLaplacePyr(img_with_border, num_bands_, src_pyr_laplace);\r
 \r
     // Create the weight map Gaussian pyramid\r
     Mat weight_map;\r
@@ -341,6 +338,7 @@ void createLaplacePyr(const Mat &img, int num_levels, vector<Mat> &pyr)
 }\r
 \r
 \r
+#if 0\r
 void createLaplacePyrGpu(const Mat &img, int num_levels, vector<Mat> &pyr)\r
 {\r
     pyr.resize(num_levels + 1);\r
@@ -360,7 +358,7 @@ void createLaplacePyrGpu(const Mat &img, int num_levels, vector<Mat> &pyr)
 \r
     pyr[num_levels] = gpu_pyr[num_levels];\r
 }\r
-\r
+#endif\r
 \r
 \r
 void restoreImageFromLaplacePyr(vector<Mat> &pyr)\r

diff --git a/modules/stitching/blenders.hpp b/modules/stitching/blenders.hpp
index 04ede3c..af52f8a 100644 (file)
--- 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);\r
 \r
 void createLaplacePyr(const cv::Mat &img, int num_levels, std::vector<cv::Mat>& pyr);\r
+\r
+// TODO Use it after gpu::pyrDown and gpu::pyrUp are updated\r
+#if 0\r
 void createLaplacePyrGpu(const cv::Mat &img, int num_levels, std::vector<cv::Mat>& pyr);\r
+#endif\r
 \r
 // Restores source image in-place (result will be stored in pyr[0])\r
 void restoreImageFromLaplacePyr(std::vector<cv::Mat>& pyr);\r

diff --git a/modules/stitching/main.cpp b/modules/stitching/main.cpp
index defb32b..26226eb 100644 (file)
--- a/modules/stitching/main.cpp
+++ b/modules/stitching/main.cpp
@@ -547,7 +547,7 @@ int main(int argc, char* argv[])
         else\r
             img = full_img;\r
         full_img.release();\r
-        Size img_size = img.size();\r
+        Size img_size = img.size();                \r
 \r
         // Warp the current image\r
         warper->warp(img, static_cast<float>(cameras[img_idx].focal), cameras[img_idx].R,\r

diff --git a/modules/stitching/warpers.cpp b/modules/stitching/warpers.cpp
index 68916ef..aceedac 100644 (file)
--- a/modules/stitching/warpers.cpp
+++ b/modules/stitching/warpers.cpp
@@ -48,9 +48,9 @@ Ptr<Warper> Warper::createByCameraFocal(float focal, int type, bool try_gpu)
 {\r
     bool can_use_gpu = try_gpu && gpu::getCudaEnabledDeviceCount();\r
     if (type == PLANE)\r
-        return new PlaneWarper(focal);\r
+        return !can_use_gpu ? new PlaneWarper(focal) : new PlaneWarperGpu(focal);\r
     if (type == CYLINDRICAL)\r
-        return new CylindricalWarper(focal);\r
+        return !can_use_gpu ? new CylindricalWarper(focal) : new CylindricalWarperGpu(focal);\r
     if (type == SPHERICAL)\r
         return !can_use_gpu ? new SphericalWarper(focal) : new SphericalWarperGpu(focal);\r
     CV_Error(CV_StsBadArg, "unsupported warping type");\r
@@ -105,6 +105,26 @@ void PlaneWarper::detectResultRoi(Point &dst_tl, Point &dst_br)
 }\r
 \r
 \r
+Point PlaneWarperGpu::warp(const Mat &src, float focal, const cv::Mat &R, cv::Mat &dst, int interp_mode, int border_mode)\r
+{\r
+    src_size_ = src.size();\r
+    projector_.size = src.size();\r
+    projector_.focal = focal;\r
+    projector_.setTransformation(R);\r
+\r
+    cv::Point dst_tl, dst_br;\r
+    detectResultRoi(dst_tl, dst_br);\r
+\r
+    gpu::buildWarpPlaneMaps(src.size(), Rect(dst_tl, Point(dst_br.x+1, dst_br.y+1)),\r
+                            R, focal, projector_.scale, projector_.plane_dist, d_xmap_, d_ymap_);\r
+\r
+    dst.create(dst_br.y - dst_tl.y + 1, dst_br.x - dst_tl.x + 1, src.type());\r
+    remap(src, dst, Mat(d_xmap_), Mat(d_ymap_), interp_mode, border_mode);\r
+\r
+    return dst_tl;\r
+}\r
+\r
+\r
 void SphericalWarper::detectResultRoi(Point &dst_tl, Point &dst_br)\r
 {\r
     detectResultRoiByBorder(dst_tl, dst_br);\r
@@ -168,3 +188,24 @@ Point SphericalWarperGpu::warp(const Mat &src, float focal, const Mat &R, Mat &d
 \r
     return dst_tl;\r
 }\r
+\r
+\r
+Point CylindricalWarperGpu::warp(const Mat &src, float focal, const Mat &R, Mat &dst,\r
+                                 int interp_mode, int border_mode)\r
+{\r
+    src_size_ = src.size();\r
+    projector_.size = src.size();\r
+    projector_.focal = focal;\r
+    projector_.setTransformation(R);\r
+\r
+    cv::Point dst_tl, dst_br;\r
+    detectResultRoi(dst_tl, dst_br);\r
+\r
+    gpu::buildWarpCylindricalMaps(src.size(), Rect(dst_tl, Point(dst_br.x+1, dst_br.y+1)),\r
+                                  R, focal, projector_.scale, d_xmap_, d_ymap_);\r
+\r
+    dst.create(dst_br.y - dst_tl.y + 1, dst_br.x - dst_tl.x + 1, src.type());\r
+    remap(src, dst, Mat(d_xmap_), Mat(d_ymap_), interp_mode, border_mode);\r
+\r
+    return dst_tl;\r
+}\r

diff --git a/modules/stitching/warpers.hpp b/modules/stitching/warpers.hpp
index 854c87b..47a162e 100644 (file)
--- a/modules/stitching/warpers.hpp
+++ b/modules/stitching/warpers.hpp
@@ -109,11 +109,23 @@ public:
         projector_.scale = scale;\r
     }\r
 \r
-private:\r
+protected:\r
     void detectResultRoi(cv::Point &dst_tl, cv::Point &dst_br);\r
 };\r
 \r
 \r
+class PlaneWarperGpu : public PlaneWarper\r
+{\r
+public:\r
+    PlaneWarperGpu(float plane_dist = 1.f, float scale = 1.f) : PlaneWarper(plane_dist, scale) {}\r
+    cv::Point warp(const cv::Mat &src, float focal, const cv::Mat &R, cv::Mat &dst,\r
+                   int interp_mode, int border_mode);\r
+\r
+private:\r
+    cv::gpu::GpuMat d_xmap_, d_ymap_, d_dst_;\r
+};\r
+\r
+\r
 struct SphericalProjector : ProjectorBase\r
 {\r
     void mapForward(float x, float y, float &u, float &v);\r
@@ -158,13 +170,25 @@ class CylindricalWarper : public WarperBase<CylindricalProjector>
 public:\r
     CylindricalWarper(float scale = 300.f) { projector_.scale = scale; }\r
 \r
-private:\r
+protected:\r
     void detectResultRoi(cv::Point &dst_tl, cv::Point &dst_br)\r
     {\r
         WarperBase<CylindricalProjector>::detectResultRoiByBorder(dst_tl, dst_br);\r
     }\r
 };\r
 \r
+\r
+class CylindricalWarperGpu : public CylindricalWarper\r
+{\r
+public:\r
+    CylindricalWarperGpu(float scale = 300.f) : CylindricalWarper(scale) {}\r
+    cv::Point warp(const cv::Mat &src, float focal, const cv::Mat &R, cv::Mat &dst,\r
+                   int interp_mode, int border_mode);\r
+\r
+private:\r
+    cv::gpu::GpuMat d_xmap_, d_ymap_, d_dst_;\r
+};\r
+\r
 #include "warpers_inl.hpp"\r
 \r
 #endif // __OPENCV_WARPERS_HPP__\r