Fixed support of translation in the GPU-based plane warper
authorAlexey Spizhevoy <no@email>
Mon, 3 Oct 2011 07:42:16 +0000 (07:42 +0000)
committerAlexey Spizhevoy <no@email>
Mon, 3 Oct 2011 07:42:16 +0000 (07:42 +0000)
modules/gpu/include/opencv2/gpu/gpu.hpp
modules/gpu/perf/perf_imgproc.cpp
modules/gpu/src/cuda/imgproc.cu
modules/gpu/src/imgproc.cpp
modules/stitching/src/warpers.cpp

index dba9b30..51aadeb 100644 (file)
@@ -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());\r
 \r
         //! builds plane warping maps\r
-        CV_EXPORTS void buildWarpPlaneMaps(Size src_size, Rect dst_roi, const Mat &K, const Mat& R, float scale,\r
+        CV_EXPORTS void buildWarpPlaneMaps(Size src_size, Rect dst_roi, const Mat &K, const Mat& R, const Mat &T, float scale,\r
                                            GpuMat& map_x, GpuMat& map_y, Stream& stream = Stream::Null());\r
 \r
         //! builds cylindrical warping maps\r
index 98e7edd..f239edb 100644 (file)
@@ -367,7 +367,7 @@ PERF_TEST_P(DevInfo_Size, buildWarpPlaneMaps, testing::Combine(testing::ValuesIn
     SIMPLE_TEST_CYCLE()\r
     {\r
         buildWarpPlaneMaps(size, Rect(0, 0, size.width, size.height), Mat::eye(3, 3, CV_32FC1), \r
-                           Mat::ones(3, 3, CV_32FC1), 1.0, map_x, map_y);\r
+                           Mat::ones(3, 3, CV_32FC1), Mat::zeros(1, 3, CV_32F), 1.0, map_x, map_y);\r
     }\r
 \r
     Mat map_x_host(map_x);\r
index 927c383..79b5e10 100644 (file)
@@ -794,6 +794,7 @@ namespace cv { namespace gpu { namespace imgproc
 \r
         __constant__ float ck_rinv[9];\r
         __constant__ float cr_kinv[9];\r
+        __constant__ float ct[3];\r
         __constant__ float cscale;\r
     }\r
 \r
@@ -805,13 +806,13 @@ namespace cv { namespace gpu { namespace imgproc
         {\r
             using namespace build_warp_maps;\r
 \r
-            float x_ = u / cscale;\r
-            float y_ = v / cscale;\r
+            float x_ = u / cscale - ct[0];\r
+            float y_ = v / cscale - ct[1];\r
 \r
             float z;\r
-            x = ck_rinv[0] * x_ + ck_rinv[1] * y_ + ck_rinv[2];\r
-            y = ck_rinv[3] * x_ + ck_rinv[4] * y_ + ck_rinv[5];\r
-            z = ck_rinv[6] * x_ + ck_rinv[7] * y_ + ck_rinv[8];\r
+            x = ck_rinv[0] * x_ + ck_rinv[1] * y_ + ck_rinv[2] * (1 - ct[2]);\r
+            y = ck_rinv[3] * x_ + ck_rinv[4] * y_ + ck_rinv[5] * (1 - ct[2]);\r
+            z = ck_rinv[6] * x_ + ck_rinv[7] * y_ + ck_rinv[8] * (1 - ct[2]);\r
 \r
             x /= z;\r
             y /= z;\r
@@ -887,11 +888,12 @@ 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 k_rinv[9], const float r_kinv[9], float scale,\r
-                            cudaStream_t stream)\r
+                            const float k_rinv[9], const float r_kinv[9], const float t[3], \r
+                            float scale, cudaStream_t stream)\r
     {\r
         cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ck_rinv, k_rinv, 9*sizeof(float)));\r
         cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr_kinv, r_kinv, 9*sizeof(float)));\r
+        cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ct, t, 3*sizeof(float)));\r
         cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cscale, &scale, sizeof(float)));\r
 \r
         int cols = map_x.cols;\r
@@ -908,8 +910,8 @@ 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 k_rinv[9], const float r_kinv[9], float scale,\r
-                            cudaStream_t stream)\r
+                                  const float k_rinv[9], const float r_kinv[9], float scale,\r
+                                  cudaStream_t stream)\r
     {\r
         cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ck_rinv, k_rinv, 9*sizeof(float)));\r
         cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr_kinv, r_kinv, 9*sizeof(float)));\r
@@ -929,8 +931,8 @@ namespace cv { namespace gpu { namespace imgproc
 \r
 \r
     void buildWarpSphericalMaps(int tl_u, int tl_v, DevMem2Df map_x, DevMem2Df map_y,\r
-                            const float k_rinv[9], const float r_kinv[9], float scale,\r
-                            cudaStream_t stream)\r
+                                const float k_rinv[9], const float r_kinv[9], float scale,\r
+                                cudaStream_t stream)\r
     {\r
         cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::ck_rinv, k_rinv, 9*sizeof(float)));\r
         cudaSafeCall(cudaMemcpyToSymbol(build_warp_maps::cr_kinv, r_kinv, 9*sizeof(float)));\r
index 87d3b5d..8b86ce6 100644 (file)
@@ -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(); }\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::buildWarpPlaneMaps(Size, Rect, const Mat&, const Mat&, float, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
+void cv::gpu::buildWarpPlaneMaps(Size, Rect, const Mat&, const Mat&, const Mat&, float, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
 void cv::gpu::buildWarpCylindricalMaps(Size, Rect, const Mat&, const Mat&, float, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
 void cv::gpu::buildWarpSphericalMaps(Size, Rect, const Mat&, const Mat&, float, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
 void cv::gpu::rotate(const GpuMat&, GpuMat&, Size, double, double, double, int, Stream&) { throw_nogpu(); }\r
@@ -623,15 +623,16 @@ void cv::gpu::warpPerspective(const GpuMat& src, GpuMat& dst, const Mat& M, Size
 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 k_rinv[9], const float r_kinv[9], float scale,\r
+                            const float k_rinv[9], const float r_kinv[9], const float t[3], float scale,\r
                             cudaStream_t stream);\r
 }}}\r
 \r
-void cv::gpu::buildWarpPlaneMaps(Size src_size, Rect dst_roi, const Mat &K, const Mat& R, float scale,\r
-                                 GpuMat& map_x, GpuMat& map_y, Stream& stream)\r
+void cv::gpu::buildWarpPlaneMaps(Size src_size, Rect dst_roi, const Mat &K, const Mat& R, const Mat &T, \r
+                                 float scale, GpuMat& map_x, GpuMat& map_y, Stream& stream)\r
 {\r
     CV_Assert(K.size() == Size(3,3) && K.type() == CV_32F);\r
     CV_Assert(R.size() == Size(3,3) && R.type() == CV_32F);\r
+    CV_Assert((T.size() == Size(3,1) || T.size() == Size(1,3)) && T.type() == CV_32F && T.isContinuous());\r
 \r
     Mat K_Rinv = K * R.t();\r
     Mat R_Kinv = R * K.inv();\r
@@ -640,8 +641,8 @@ void cv::gpu::buildWarpPlaneMaps(Size src_size, Rect dst_roi, const Mat &K, cons
 \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, K_Rinv.ptr<float>(), R_Kinv.ptr<float>(),\r
-                                scale, StreamAccessor::getStream(stream));\r
+    imgproc::buildWarpPlaneMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, K_Rinv.ptr<float>(), R_Kinv.ptr<float>(), \r
+                                T.ptr<float>(), scale, StreamAccessor::getStream(stream));\r
 }\r
 \r
 //////////////////////////////////////////////////////////////////////////////\r
index b96b14a..3bd0497 100644 (file)
@@ -216,8 +216,8 @@ Rect PlaneWarperGpu::buildMaps(Size src_size, const Mat &K, const Mat &R, const
     Point dst_tl, dst_br;\r
     detectResultRoi(src_size, dst_tl, dst_br);\r
 \r
-    gpu::buildWarpPlaneMaps(src_size, Rect(dst_tl, Point(dst_br.x+1, dst_br.y+1)),\r
-                            K, R, projector_.scale, xmap, ymap);\r
+    gpu::buildWarpPlaneMaps(src_size, Rect(dst_tl, Point(dst_br.x + 1, dst_br.y + 1)),\r
+                            K, R, T, projector_.scale, xmap, ymap);\r
 \r
     return Rect(dst_tl, dst_br);\r
 }\r