added support of 3-channels output to gpu::reprojectImageTo3D

minor refactoring of gpu tests

modules/gpu/include/opencv2/gpu/gpu.hpp

@@ -605,10 +605,10 @@ CV_EXPORTS void drawColorDisp(const GpuMat& src_disp, GpuMat& dst_disp, int ndis
 
 //! Reprojects disparity image to 3D space.
 //! Supports CV_8U and CV_16S types of input disparity.
-//! The output is a 4-channel floating-point (CV_32FC4) matrix.
+//! The output is a 3- or 4-channel floating-point matrix.
 //! Each element of this matrix will contain the 3D coordinates of the point (x,y,z,1), computed from the disparity map.
 //! Q is the 4x4 perspective transformation matrix that can be obtained with cvStereoRectify.
-CV_EXPORTS void reprojectImageTo3D(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, Stream& stream = Stream::Null());
+CV_EXPORTS void reprojectImageTo3D(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, int dst_cn = 4, Stream& stream = Stream::Null());
 
 //! converts image from one color space to another
 CV_EXPORTS void cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn = 0, Stream& stream = Stream::Null());

modules/gpu/src/cuda/imgproc.cu

@@ -316,62 +316,51 @@ namespace cv { namespace gpu { namespace device
 
         __constant__ float cq[16];
 
-        template <typename T>
-        __global__ void reprojectImageTo3D(const T* disp, size_t disp_step, float* xyzw, size_t xyzw_step, int rows, int cols)
+        template <typename T, typename D>
+        __global__ void reprojectImageTo3D(const DevMem2D_<T> disp, PtrStep<D> xyz)
         {
             const int x = blockIdx.x * blockDim.x + threadIdx.x;
             const int y = blockIdx.y * blockDim.y + threadIdx.y;
 
-            if (y < rows && x < cols)
-            {
+            if (y >= disp.rows || x >= disp.cols)
+                return;
 
-                float qx = cq[1] * y + cq[3], qy = cq[5] * y + cq[7];
-                float qz = cq[9] * y + cq[11], qw = cq[13] * y + cq[15];
+            const float qx = x * cq[ 0] + y * cq[ 1] + cq[ 3];
+            const float qy = x * cq[ 4] + y * cq[ 5] + cq[ 7];
+            const float qz = x * cq[ 8] + y * cq[ 9] + cq[11];
+            const float qw = x * cq[12] + y * cq[13] + cq[15];
 
-                qx += x * cq[0];
-                qy += x * cq[4];
-                qz += x * cq[8];
-                qw += x * cq[12];
+            const T d = disp(y, x);
 
-                T d = *(disp + disp_step * y + x);
+            const float iW = 1.f / (qw + cq[14] * d);
 
-                float iW = 1.f / (qw + cq[14] * d);
-                float4 v;
+            D v = VecTraits<D>::all(1.0f);
             v.x = (qx + cq[2] * d) * iW;
             v.y = (qy + cq[6] * d) * iW;
             v.z = (qz + cq[10] * d) * iW;
-                v.w = 1.f;
 
-                *(float4*)(xyzw + xyzw_step * y + (x * 4)) = v;
-            }
+            xyz(y, x) = v;
         }
 
-        template <typename T>
-        inline void reprojectImageTo3D_caller(const DevMem2D_<T>& disp, const DevMem2Df& xyzw, const float* q, const cudaStream_t& stream)
+        template <typename T, typename D>
+        void reprojectImageTo3D_gpu(const DevMem2Db disp, DevMem2Db xyz, const float* q, cudaStream_t stream)
         {
-            dim3 threads(32, 8, 1);
-            dim3 grid(1, 1, 1);
-            grid.x = divUp(disp.cols, threads.x);
-            grid.y = divUp(disp.rows, threads.y);
+            dim3 block(32, 8);
+            dim3 grid(divUp(disp.cols, block.x), divUp(disp.rows, block.y));
 
             cudaSafeCall( cudaMemcpyToSymbol(cq, q, 16 * sizeof(float)) );
 
-            reprojectImageTo3D<<<grid, threads, 0, stream>>>(disp.data, disp.step / sizeof(T), xyzw.data, xyzw.step / sizeof(float), disp.rows, disp.cols);
+            reprojectImageTo3D<T, D><<<grid, block, 0, stream>>>((DevMem2D_<T>)disp, (DevMem2D_<D>)xyz);
             cudaSafeCall( cudaGetLastError() );
 
             if (stream == 0)
                 cudaSafeCall( cudaDeviceSynchronize() );
         }
 
-        void reprojectImageTo3D_gpu(const DevMem2Db& disp, const DevMem2Df& xyzw, const float* q, const cudaStream_t& stream)
-        {
-            reprojectImageTo3D_caller(disp, xyzw, q, stream);
-        }
-
-        void reprojectImageTo3D_gpu(const DevMem2D_<short>& disp, const DevMem2Df& xyzw, const float* q, const cudaStream_t& stream)
-        {
-            reprojectImageTo3D_caller(disp, xyzw, q, stream);
-        }
+        template void reprojectImageTo3D_gpu<uchar, float3>(const DevMem2Db disp, DevMem2Db xyz, const float* q, cudaStream_t stream);
+        template void reprojectImageTo3D_gpu<uchar, float4>(const DevMem2Db disp, DevMem2Db xyz, const float* q, cudaStream_t stream);
+        template void reprojectImageTo3D_gpu<short, float3>(const DevMem2Db disp, DevMem2Db xyz, const float* q, cudaStream_t stream);
+        template void reprojectImageTo3D_gpu<short, float4>(const DevMem2Db disp, DevMem2Db xyz, const float* q, cudaStream_t stream);
 
         /////////////////////////////////////////// Corner Harris /////////////////////////////////////////////////
 

modules/gpu/src/imgproc.cpp

@@ -50,7 +50,7 @@ using namespace cv::gpu;
 void cv::gpu::meanShiftFiltering(const GpuMat&, GpuMat&, int, int, TermCriteria, Stream&) { throw_nogpu(); }
 void cv::gpu::meanShiftProc(const GpuMat&, GpuMat&, GpuMat&, int, int, TermCriteria, Stream&) { throw_nogpu(); }
 void cv::gpu::drawColorDisp(const GpuMat&, GpuMat&, int, Stream&) { throw_nogpu(); }
-void cv::gpu::reprojectImageTo3D(const GpuMat&, GpuMat&, const Mat&, Stream&) { throw_nogpu(); }
+void cv::gpu::reprojectImageTo3D(const GpuMat&, GpuMat&, const Mat&, int, Stream&) { throw_nogpu(); }
 void cv::gpu::copyMakeBorder(const GpuMat&, GpuMat&, int, int, int, int, int, const Scalar&, 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(); }
@@ -213,33 +213,29 @@ namespace cv { namespace gpu { namespace device
 {
     namespace imgproc
     {
-        void reprojectImageTo3D_gpu(const DevMem2Db& disp, const DevMem2Df& xyzw, const float* q, const cudaStream_t& stream);
-        void reprojectImageTo3D_gpu(const DevMem2D_<short>& disp, const DevMem2Df& xyzw, const float* q, const cudaStream_t& stream);
+        template <typename T, typename D>
+        void reprojectImageTo3D_gpu(const DevMem2Db disp, DevMem2Db xyz, const float* q, cudaStream_t stream);
     }
 }}}
 
-namespace
+void cv::gpu::reprojectImageTo3D(const GpuMat& disp, GpuMat& xyz, const Mat& Q, int dst_cn, Stream& stream)
 {
-    template <typename T>
-    void reprojectImageTo3D_caller(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, const cudaStream_t& stream)
-    {
-        using namespace ::cv::gpu::device::imgproc;
-
-        xyzw.create(disp.rows, disp.cols, CV_32FC4);
-
-        reprojectImageTo3D_gpu((DevMem2D_<T>)disp, xyzw, Q.ptr<float>(), stream);
-    }
+    using namespace cv::gpu::device::imgproc;
 
-    typedef void (*reprojectImageTo3D_caller_t)(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, const cudaStream_t& stream);
+    typedef void (*func_t)(const DevMem2Db disp, DevMem2Db xyz, const float* q, cudaStream_t stream);
+    static const func_t funcs[2][4] = 
+    {
+        {reprojectImageTo3D_gpu<uchar, float3>, 0, 0, reprojectImageTo3D_gpu<short, float3>},
+        {reprojectImageTo3D_gpu<uchar, float4>, 0, 0, reprojectImageTo3D_gpu<short, float4>}
+    };
 
-    const reprojectImageTo3D_caller_t reprojectImageTo3D_callers[] = {reprojectImageTo3D_caller<unsigned char>, 0, 0, reprojectImageTo3D_caller<short>, 0, 0, 0, 0};
-}
+    CV_Assert(disp.type() == CV_8U || disp.type() == CV_16S);
+    CV_Assert(Q.type() == CV_32F && Q.rows == 4 && Q.cols == 4 && Q.isContinuous());
+    CV_Assert(dst_cn == 3 || dst_cn == 4);
 
-void cv::gpu::reprojectImageTo3D(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, Stream& stream)
-{
-    CV_Assert((disp.type() == CV_8U || disp.type() == CV_16S) && Q.type() == CV_32F && Q.rows == 4 && Q.cols == 4);
+    xyz.create(disp.size(), CV_MAKE_TYPE(CV_32F, dst_cn));
 
-    reprojectImageTo3D_callers[disp.type()](disp, xyzw, Q, StreamAccessor::getStream(stream));
+    funcs[dst_cn == 4][disp.type()](disp, xyz, Q.ptr<float>(), StreamAccessor::getStream(stream));
 }
 
 ////////////////////////////////////////////////////////////////////////
@@ -1513,9 +1509,11 @@ void cv::gpu::Canny(const GpuMat& src, CannyBuf& buf, GpuMat& dst, double low_th
 {
     using namespace ::cv::gpu::device::canny;
 
-    CV_Assert(TargetArchs::builtWith(SHARED_ATOMICS) && DeviceInfo().supports(SHARED_ATOMICS));
     CV_Assert(src.type() == CV_8UC1);
 
+    if (!TargetArchs::builtWith(SHARED_ATOMICS) || !DeviceInfo().supports(SHARED_ATOMICS))
+        CV_Error(CV_StsNotImplemented, "The device doesn't support shared atomics");
+
     if( low_thresh > high_thresh )
         std::swap( low_thresh, high_thresh);
 

modules/gpu/test/test_calib3d.cpp

@@ -301,4 +301,45 @@ TEST_P(SolvePnPRansac, Accuracy)
 
 INSTANTIATE_TEST_CASE_P(GPU_Calib3D, SolvePnPRansac, ALL_DEVICES);
 
+////////////////////////////////////////////////////////////////////////////////
+// reprojectImageTo3D
+
+PARAM_TEST_CASE(ReprojectImageTo3D, cv::gpu::DeviceInfo, cv::Size, MatDepth, UseRoi)
+{
+    cv::gpu::DeviceInfo devInfo;
+    cv::Size size;
+    int depth;
+    bool useRoi;
+
+    virtual void SetUp()
+    {
+        devInfo = GET_PARAM(0);
+        size = GET_PARAM(1);
+        depth = GET_PARAM(2);
+        useRoi = GET_PARAM(3);
+
+        cv::gpu::setDevice(devInfo.deviceID());
+    }
+};
+
+TEST_P(ReprojectImageTo3D, Accuracy)
+{
+    cv::Mat disp = randomMat(size, depth, 5.0, 30.0);
+    cv::Mat Q = randomMat(cv::Size(4, 4), CV_32FC1, 0.1, 1.0);
+
+    cv::gpu::GpuMat dst;
+    cv::gpu::reprojectImageTo3D(loadMat(disp, useRoi), dst, Q, 3);
+    
+    cv::Mat dst_gold;
+    cv::reprojectImageTo3D(disp, dst_gold, Q, false);
+
+    EXPECT_MAT_NEAR(dst_gold, dst, 1e-5);
+}
+
+INSTANTIATE_TEST_CASE_P(GPU_Calib3D, ReprojectImageTo3D, testing::Combine(
+    ALL_DEVICES,
+    DIFFERENT_SIZES,
+    testing::Values(MatDepth(CV_8U), MatDepth(CV_16S)),
+    WHOLE_SUBMAT));
+
 } // namespace

modules/gpu/test/test_color.cpp

@@ -41,6 +41,8 @@
 
 #include "precomp.hpp"
 
+namespace {
+
 ///////////////////////////////////////////////////////////////////////////////////////////////////////
 // cvtColor
 
@@ -1652,3 +1654,5 @@ INSTANTIATE_TEST_CASE_P(GPU_ImgProc, SwapChannels, testing::Combine(
     ALL_DEVICES,
     DIFFERENT_SIZES,
     WHOLE_SUBMAT));
+
+} // namespace

modules/gpu/test/test_filters.cpp

@@ -41,15 +41,9 @@
 
 #include "precomp.hpp"
 
-struct KSize : cv::Size
-{
-    KSize() {}
-    KSize(int width, int height) : cv::Size(width, height) {}
-};
-void PrintTo(KSize ksize, std::ostream* os)
-{
-    *os << "kernel size " << ksize.width << "x" << ksize.height;
-}
+namespace {
+
+IMPLEMENT_PARAM_CLASS(KSize, cv::Size)
 
 cv::Mat getInnerROI(cv::InputArray m_, cv::Size ksize)
 {
@@ -107,7 +101,7 @@ INSTANTIATE_TEST_CASE_P(GPU_Filter, Blur, testing::Combine(
     ALL_DEVICES,
     DIFFERENT_SIZES,
     testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)),
-    testing::Values(KSize(3, 3), KSize(5, 5), KSize(7, 7)),
+    testing::Values(KSize(cv::Size(3, 3)), KSize(cv::Size(5, 5)), KSize(cv::Size(7, 7))),
     testing::Values(Anchor(cv::Point(-1, -1)), Anchor(cv::Point(0, 0)), Anchor(cv::Point(2, 2))),
     WHOLE_SUBMAT));
 
@@ -163,7 +157,7 @@ INSTANTIATE_TEST_CASE_P(GPU_Filter, Sobel, testing::Combine(
     ALL_DEVICES,
     DIFFERENT_SIZES,
     testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)),
-    testing::Values(KSize(3, 3), KSize(5, 5), KSize(7, 7)),
+    testing::Values(KSize(cv::Size(3, 3)), KSize(cv::Size(5, 5)), KSize(cv::Size(7, 7))),
     testing::Values(Deriv_X(0), Deriv_X(1), Deriv_X(2)),
     testing::Values(Deriv_Y(0), Deriv_Y(1), Deriv_Y(2)),
     testing::Values(BorderType(cv::BORDER_REFLECT101),
@@ -286,21 +280,21 @@ INSTANTIATE_TEST_CASE_P(GPU_Filter, GaussianBlur, testing::Combine(
     ALL_DEVICES,
     DIFFERENT_SIZES,
     testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)),
-    testing::Values(KSize(3, 3),
-                    KSize(5, 5),
-                    KSize(7, 7),
-                    KSize(9, 9),
-                    KSize(11, 11),
-                    KSize(13, 13),
-                    KSize(15, 15),
-                    KSize(17, 17),
-                    KSize(19, 19),
-                    KSize(21, 21),
-                    KSize(23, 23),
-                    KSize(25, 25),
-                    KSize(27, 27),
-                    KSize(29, 29),
-                    KSize(31, 31)),
+    testing::Values(KSize(cv::Size(3, 3)),
+                    KSize(cv::Size(5, 5)),
+                    KSize(cv::Size(7, 7)),
+                    KSize(cv::Size(9, 9)),
+                    KSize(cv::Size(11, 11)),
+                    KSize(cv::Size(13, 13)),
+                    KSize(cv::Size(15, 15)),
+                    KSize(cv::Size(17, 17)),
+                    KSize(cv::Size(19, 19)),
+                    KSize(cv::Size(21, 21)),
+                    KSize(cv::Size(23, 23)),
+                    KSize(cv::Size(25, 25)),
+                    KSize(cv::Size(27, 27)),
+                    KSize(cv::Size(29, 29)),
+                    KSize(cv::Size(31, 31))),
     testing::Values(BorderType(cv::BORDER_REFLECT101),
                     BorderType(cv::BORDER_REPLICATE),
                     BorderType(cv::BORDER_CONSTANT),
@@ -350,7 +344,7 @@ INSTANTIATE_TEST_CASE_P(GPU_Filter, Laplacian, testing::Combine(
     ALL_DEVICES,
     DIFFERENT_SIZES,
     testing::Values(MatType(CV_8UC1), MatType(CV_8UC4), MatType(CV_32FC1)),
-    testing::Values(KSize(1, 1), KSize(3, 3)),
+    testing::Values(KSize(cv::Size(1, 1)), KSize(cv::Size(3, 3))),
     WHOLE_SUBMAT));
 
 /////////////////////////////////////////////////////////////////////////////////////////////////
@@ -557,6 +551,8 @@ INSTANTIATE_TEST_CASE_P(GPU_Filter, Filter2D, testing::Combine(
     ALL_DEVICES,
     DIFFERENT_SIZES,
     testing::Values(MatType(CV_8UC1), MatType(CV_8UC4), MatType(CV_32FC1)),
-    testing::Values(KSize(3, 3), KSize(5, 5), KSize(7, 7), KSize(11, 11), KSize(13, 13), KSize(15, 15)),
+    testing::Values(KSize(cv::Size(3, 3)), KSize(cv::Size(5, 5)), KSize(cv::Size(7, 7)), KSize(cv::Size(11, 11)), KSize(cv::Size(13, 13)), KSize(cv::Size(15, 15))),
     testing::Values(Anchor(cv::Point(-1, -1)), Anchor(cv::Point(0, 0)), Anchor(cv::Point(2, 2))),
     WHOLE_SUBMAT));
+
+} // namespace

modules/gpu/test/utility.hpp

@@ -265,7 +265,7 @@ void PrintTo(const Inverse& useRoi, std::ostream* os);
     }; \
     inline void PrintTo( name param, std::ostream* os) \
     { \
-        *os << #name <<  "(" << static_cast< type >(param) << ")"; \
+        *os << #name <<  "(" << testing::PrintToString(static_cast< type >(param)) << ")"; \
     }
 
 IMPLEMENT_PARAM_CLASS(Channels, int)