Fixed support of translation in the GPU-based plane warper

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

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

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

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<float>(), R_Kinv.ptr<float>(),
-                                scale, StreamAccessor::getStream(stream));
+    imgproc::buildWarpPlaneMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, K_Rinv.ptr<float>(), R_Kinv.ptr<float>(), 
+                                T.ptr<float>(), scale, StreamAccessor::getStream(stream));
 }
 
 //////////////////////////////////////////////////////////////////////////////

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