gpustereo module for stereo correspondence

modules/gpu/CMakeLists.txt

@@ -7,7 +7,7 @@ set(the_description "GPU-accelerated Computer Vision")
 ocv_warnings_disable(CMAKE_CXX_FLAGS -Wundef -Wmissing-declarations -Wshadow -Wunused-parameter)
 
 ocv_define_module(gpu opencv_gpuarithm opencv_gpufilters opencv_gpuwarping opencv_gpuimgproc
-                      opencv_gpufeatures2d opencv_gpuvideo opencv_gpucalib3d opencv_gpuobjdetect)
+                      opencv_gpufeatures2d opencv_gpuvideo opencv_gpustereo opencv_gpuobjdetect)
 
 if(HAVE_CUDA)
   add_subdirectory(perf4au)

modules/gpu/doc/calib3d.rst

@@ -0,0 +1,36 @@
+Camera Calibration and 3D Reconstruction
+========================================
+
+.. highlight:: cpp
+
+
+
+gpu::solvePnPRansac
+-------------------
+Finds the object pose from 3D-2D point correspondences.
+
+.. ocv:function:: void gpu::solvePnPRansac(const Mat& object, const Mat& image, const Mat& camera_mat, const Mat& dist_coef, Mat& rvec, Mat& tvec, bool use_extrinsic_guess=false, int num_iters=100, float max_dist=8.0, int min_inlier_count=100, vector<int>* inliers=NULL)
+
+    :param object: Single-row matrix of object points.
+
+    :param image: Single-row matrix of image points.
+
+    :param camera_mat: 3x3 matrix of intrinsic camera parameters.
+
+    :param dist_coef: Distortion coefficients. See :ocv:func:`undistortPoints` for details.
+
+    :param rvec: Output 3D rotation vector.
+
+    :param tvec: Output 3D translation vector.
+
+    :param use_extrinsic_guess: Flag to indicate that the function must use ``rvec`` and ``tvec`` as an initial transformation guess. It is not supported for now.
+
+    :param num_iters: Maximum number of RANSAC iterations.
+
+    :param max_dist: Euclidean distance threshold to detect whether point is inlier or not.
+
+    :param min_inlier_count: Flag to indicate that the function must stop if greater or equal number of inliers is achieved. It is not supported for now.
+
+    :param inliers: Output vector of inlier indices.
+
+.. seealso:: :ocv:func:`solvePnPRansac`

modules/gpu/include/opencv2/gpu.hpp

@@ -50,7 +50,7 @@
 #include "opencv2/gpuimgproc.hpp"
 #include "opencv2/gpufeatures2d.hpp"
 #include "opencv2/gpuvideo.hpp"
-#include "opencv2/gpucalib3d.hpp"
+#include "opencv2/gpustereo.hpp"
 #include "opencv2/gpuobjdetect.hpp"
 
 namespace cv { namespace gpu {
@@ -71,6 +71,18 @@ CV_EXPORTS void connectivityMask(const GpuMat& image, GpuMat& mask, const cv::Sc
 //! performs connected componnents labeling.
 CV_EXPORTS void labelComponents(const GpuMat& mask, GpuMat& components, int flags = 0, Stream& stream = Stream::Null());
 
+CV_EXPORTS void transformPoints(const GpuMat& src, const Mat& rvec, const Mat& tvec,
+                                GpuMat& dst, Stream& stream = Stream::Null());
+
+CV_EXPORTS void projectPoints(const GpuMat& src, const Mat& rvec, const Mat& tvec,
+                              const Mat& camera_mat, const Mat& dist_coef, GpuMat& dst,
+                              Stream& stream = Stream::Null());
+
+CV_EXPORTS void solvePnPRansac(const Mat& object, const Mat& image, const Mat& camera_mat,
+                               const Mat& dist_coef, Mat& rvec, Mat& tvec, bool use_extrinsic_guess=false,
+                               int num_iters=100, float max_dist=8.0, int min_inlier_count=100,
+                               std::vector<int>* inliers=NULL);
+
 //! removes points (CV_32FC2, single row matrix) with zero mask value
 CV_EXPORTS void compactPoints(GpuMat &points0, GpuMat &points1, const GpuMat &mask);
 

modules/gpu/perf/perf_calib3d.cpp

@@ -0,0 +1,135 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include "perf_precomp.hpp"
+
+using namespace std;
+using namespace testing;
+using namespace perf;
+
+DEF_PARAM_TEST_1(Count, int);
+
+//////////////////////////////////////////////////////////////////////
+// ProjectPoints
+
+PERF_TEST_P(Count, Calib3D_ProjectPoints,
+            Values(5000, 10000, 20000))
+{
+    const int count = GetParam();
+
+    cv::Mat src(1, count, CV_32FC3);
+    declare.in(src, WARMUP_RNG);
+
+    const cv::Mat rvec = cv::Mat::ones(1, 3, CV_32FC1);
+    const cv::Mat tvec = cv::Mat::ones(1, 3, CV_32FC1);
+    const cv::Mat camera_mat = cv::Mat::ones(3, 3, CV_32FC1);
+
+    if (PERF_RUN_GPU())
+    {
+        const cv::gpu::GpuMat d_src(src);
+        cv::gpu::GpuMat dst;
+
+        TEST_CYCLE() cv::gpu::projectPoints(d_src, rvec, tvec, camera_mat, cv::Mat(), dst);
+
+        GPU_SANITY_CHECK(dst);
+    }
+    else
+    {
+        cv::Mat dst;
+
+        TEST_CYCLE() cv::projectPoints(src, rvec, tvec, camera_mat, cv::noArray(), dst);
+
+        CPU_SANITY_CHECK(dst);
+    }
+}
+
+//////////////////////////////////////////////////////////////////////
+// SolvePnPRansac
+
+PERF_TEST_P(Count, Calib3D_SolvePnPRansac,
+            Values(5000, 10000, 20000))
+{
+    declare.time(10.0);
+
+    const int count = GetParam();
+
+    cv::Mat object(1, count, CV_32FC3);
+    declare.in(object, WARMUP_RNG);
+
+    cv::Mat camera_mat(3, 3, CV_32FC1);
+    cv::randu(camera_mat, 0.5, 1);
+    camera_mat.at<float>(0, 1) = 0.f;
+    camera_mat.at<float>(1, 0) = 0.f;
+    camera_mat.at<float>(2, 0) = 0.f;
+    camera_mat.at<float>(2, 1) = 0.f;
+
+    const cv::Mat dist_coef(1, 8, CV_32F, cv::Scalar::all(0));
+
+    cv::Mat rvec_gold(1, 3, CV_32FC1);
+    cv::randu(rvec_gold, 0, 1);
+
+    cv::Mat tvec_gold(1, 3, CV_32FC1);
+    cv::randu(tvec_gold, 0, 1);
+
+    std::vector<cv::Point2f> image_vec;
+    cv::projectPoints(object, rvec_gold, tvec_gold, camera_mat, dist_coef, image_vec);
+
+    const cv::Mat image(1, count, CV_32FC2, &image_vec[0]);
+
+    cv::Mat rvec;
+    cv::Mat tvec;
+
+    if (PERF_RUN_GPU())
+    {
+        TEST_CYCLE() cv::gpu::solvePnPRansac(object, image, camera_mat, dist_coef, rvec, tvec);
+
+        GPU_SANITY_CHECK(rvec, 1e-3);
+        GPU_SANITY_CHECK(tvec, 1e-3);
+    }
+    else
+    {
+        TEST_CYCLE() cv::solvePnPRansac(object, image, camera_mat, dist_coef, rvec, tvec);
+
+        CPU_SANITY_CHECK(rvec, 1e-6);
+        CPU_SANITY_CHECK(tvec, 1e-6);
+    }
+}

modules/gpu/perf/perf_precomp.hpp

@@ -55,6 +55,7 @@
 #include "opencv2/ts/gpu_perf.hpp"
 
 #include "opencv2/gpu.hpp"
+#include "opencv2/calib3d.hpp"
 
 #ifdef GTEST_CREATE_SHARED_LIBRARY
 #error no modules except ts should have GTEST_CREATE_SHARED_LIBRARY defined

modules/gpu/src/calib3d.cpp

@@ -50,8 +50,6 @@ using namespace cv::gpu;
 void cv::gpu::transformPoints(const GpuMat&, const Mat&, const Mat&, GpuMat&, Stream&) { throw_no_cuda(); }
 void cv::gpu::projectPoints(const GpuMat&, const Mat&, const Mat&, const Mat&, const Mat&, GpuMat&, Stream&) { throw_no_cuda(); }
 void cv::gpu::solvePnPRansac(const Mat&, const Mat&, const Mat&, const Mat&, Mat&, Mat&, bool, int, float, int, std::vector<int>*) { throw_no_cuda(); }
-void cv::gpu::reprojectImageTo3D(const GpuMat&, GpuMat&, const Mat&, int, Stream&) { throw_no_cuda(); }
-void cv::gpu::drawColorDisp(const GpuMat&, GpuMat&, int, Stream&) { throw_no_cuda(); }
 
 #else
 
@@ -289,66 +287,4 @@ void cv::gpu::solvePnPRansac(const Mat& object, const Mat& image, const Mat& cam
     }
 }
 
-////////////////////////////////////////////////////////////////////////
-// reprojectImageTo3D
-
-namespace cv { namespace gpu { namespace cudev
-{
-    template <typename T, typename D>
-    void reprojectImageTo3D_gpu(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream);
-}}}
-
-void cv::gpu::reprojectImageTo3D(const GpuMat& disp, GpuMat& xyz, const Mat& Q, int dst_cn, Stream& stream)
-{
-    using namespace cv::gpu::cudev;
-
-    typedef void (*func_t)(const PtrStepSzb disp, PtrStepSzb 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>}
-    };
-
-    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);
-
-    xyz.create(disp.size(), CV_MAKE_TYPE(CV_32F, dst_cn));
-
-    funcs[dst_cn == 4][disp.type()](disp, xyz, Q.ptr<float>(), StreamAccessor::getStream(stream));
-}
-
-////////////////////////////////////////////////////////////////////////
-// drawColorDisp
-
-namespace cv { namespace gpu { namespace cudev
-{
-    void drawColorDisp_gpu(const PtrStepSzb& src, const PtrStepSzb& dst, int ndisp, const cudaStream_t& stream);
-    void drawColorDisp_gpu(const PtrStepSz<short>& src, const PtrStepSzb& dst, int ndisp, const cudaStream_t& stream);
-}}}
-
-namespace
-{
-    template <typename T>
-    void drawColorDisp_caller(const GpuMat& src, GpuMat& dst, int ndisp, const cudaStream_t& stream)
-    {
-        using namespace ::cv::gpu::cudev;
-
-        dst.create(src.size(), CV_8UC4);
-
-        drawColorDisp_gpu((PtrStepSz<T>)src, dst, ndisp, stream);
-    }
-
-    typedef void (*drawColorDisp_caller_t)(const GpuMat& src, GpuMat& dst, int ndisp, const cudaStream_t& stream);
-
-    const drawColorDisp_caller_t drawColorDisp_callers[] = {drawColorDisp_caller<unsigned char>, 0, 0, drawColorDisp_caller<short>, 0, 0, 0, 0};
-}
-
-void cv::gpu::drawColorDisp(const GpuMat& src, GpuMat& dst, int ndisp, Stream& stream)
-{
-    CV_Assert(src.type() == CV_8U || src.type() == CV_16S);
-
-    drawColorDisp_callers[src.type()](src, dst, ndisp, StreamAccessor::getStream(stream));
-}
-
 #endif

modules/gpu/src/cuda/calib3d.cu

@@ -187,189 +187,6 @@ namespace cv { namespace gpu { namespace cudev
             cudaSafeCall( cudaDeviceSynchronize() );
         }
     } // namespace solvepnp_ransac
-
-
-
-    /////////////////////////////////// reprojectImageTo3D ///////////////////////////////////////////////
-
-    __constant__ float cq[16];
-
-    template <typename T, typename D>
-    __global__ void reprojectImageTo3D(const PtrStepSz<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 >= disp.rows || x >= disp.cols)
-            return;
-
-        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];
-
-        const T d = disp(y, x);
-
-        const float iW = 1.f / (qw + cq[14] * d);
-
-        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;
-
-        xyz(y, x) = v;
-    }
-
-    template <typename T, typename D>
-    void reprojectImageTo3D_gpu(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream)
-    {
-        dim3 block(32, 8);
-        dim3 grid(divUp(disp.cols, block.x), divUp(disp.rows, block.y));
-
-        cudaSafeCall( cudaMemcpyToSymbol(cq, q, 16 * sizeof(float)) );
-
-        reprojectImageTo3D<T, D><<<grid, block, 0, stream>>>((PtrStepSz<T>)disp, (PtrStepSz<D>)xyz);
-        cudaSafeCall( cudaGetLastError() );
-
-        if (stream == 0)
-            cudaSafeCall( cudaDeviceSynchronize() );
-    }
-
-    template void reprojectImageTo3D_gpu<uchar, float3>(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream);
-    template void reprojectImageTo3D_gpu<uchar, float4>(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream);
-    template void reprojectImageTo3D_gpu<short, float3>(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream);
-    template void reprojectImageTo3D_gpu<short, float4>(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream);
-
-    /////////////////////////////////// drawColorDisp ///////////////////////////////////////////////
-
-    template <typename T>
-    __device__ unsigned int cvtPixel(T d, int ndisp, float S = 1, float V = 1)
-    {
-        unsigned int H = ((ndisp-d) * 240)/ndisp;
-
-        unsigned int hi = (H/60) % 6;
-        float f = H/60.f - H/60;
-        float p = V * (1 - S);
-        float q = V * (1 - f * S);
-        float t = V * (1 - (1 - f) * S);
-
-        float3 res;
-
-        if (hi == 0) //R = V,	G = t,	B = p
-        {
-            res.x = p;
-            res.y = t;
-            res.z = V;
-        }
-
-        if (hi == 1) // R = q,	G = V,	B = p
-        {
-            res.x = p;
-            res.y = V;
-            res.z = q;
-        }
-
-        if (hi == 2) // R = p,	G = V,	B = t
-        {
-            res.x = t;
-            res.y = V;
-            res.z = p;
-        }
-
-        if (hi == 3) // R = p,	G = q,	B = V
-        {
-            res.x = V;
-            res.y = q;
-            res.z = p;
-        }
-
-        if (hi == 4) // R = t,	G = p,	B = V
-        {
-            res.x = V;
-            res.y = p;
-            res.z = t;
-        }
-
-        if (hi == 5) // R = V,	G = p,	B = q
-        {
-            res.x = q;
-            res.y = p;
-            res.z = V;
-        }
-        const unsigned int b = (unsigned int)(::max(0.f, ::min(res.x, 1.f)) * 255.f);
-        const unsigned int g = (unsigned int)(::max(0.f, ::min(res.y, 1.f)) * 255.f);
-        const unsigned int r = (unsigned int)(::max(0.f, ::min(res.z, 1.f)) * 255.f);
-        const unsigned int a = 255U;
-
-        return (a << 24) + (r << 16) + (g << 8) + b;
-    }
-
-    __global__ void drawColorDisp(uchar* disp, size_t disp_step, uchar* out_image, size_t out_step, int width, int height, int ndisp)
-    {
-        const int x = (blockIdx.x * blockDim.x + threadIdx.x) << 2;
-        const int y = blockIdx.y * blockDim.y + threadIdx.y;
-
-        if(x < width && y < height)
-        {
-            uchar4 d4 = *(uchar4*)(disp + y * disp_step + x);
-
-            uint4 res;
-            res.x = cvtPixel(d4.x, ndisp);
-            res.y = cvtPixel(d4.y, ndisp);
-            res.z = cvtPixel(d4.z, ndisp);
-            res.w = cvtPixel(d4.w, ndisp);
-
-            uint4* line = (uint4*)(out_image + y * out_step);
-            line[x >> 2] = res;
-        }
-    }
-
-    __global__ void drawColorDisp(short* disp, size_t disp_step, uchar* out_image, size_t out_step, int width, int height, int ndisp)
-    {
-        const int x = (blockIdx.x * blockDim.x + threadIdx.x) << 1;
-        const int y = blockIdx.y * blockDim.y + threadIdx.y;
-
-        if(x < width && y < height)
-        {
-            short2 d2 = *(short2*)(disp + y * disp_step + x);
-
-            uint2 res;
-            res.x = cvtPixel(d2.x, ndisp);
-            res.y = cvtPixel(d2.y, ndisp);
-
-            uint2* line = (uint2*)(out_image + y * out_step);
-            line[x >> 1] = res;
-        }
-    }
-
-
-    void drawColorDisp_gpu(const PtrStepSzb& src, const PtrStepSzb& dst, int ndisp, const cudaStream_t& stream)
-    {
-        dim3 threads(16, 16, 1);
-        dim3 grid(1, 1, 1);
-        grid.x = divUp(src.cols, threads.x << 2);
-        grid.y = divUp(src.rows, threads.y);
-
-        drawColorDisp<<<grid, threads, 0, stream>>>(src.data, src.step, dst.data, dst.step, src.cols, src.rows, ndisp);
-        cudaSafeCall( cudaGetLastError() );
-
-        if (stream == 0)
-            cudaSafeCall( cudaDeviceSynchronize() );
-    }
-
-    void drawColorDisp_gpu(const PtrStepSz<short>& src, const PtrStepSzb& dst, int ndisp, const cudaStream_t& stream)
-    {
-        dim3 threads(32, 8, 1);
-        dim3 grid(1, 1, 1);
-        grid.x = divUp(src.cols, threads.x << 1);
-        grid.y = divUp(src.rows, threads.y);
-
-        drawColorDisp<<<grid, threads, 0, stream>>>(src.data, src.step / sizeof(short), dst.data, dst.step, src.cols, src.rows, ndisp);
-        cudaSafeCall( cudaGetLastError() );
-
-        if (stream == 0)
-            cudaSafeCall( cudaDeviceSynchronize() );
-    }
 }}} // namespace cv { namespace gpu { namespace cudev
 
 

modules/gpu/src/precomp.hpp

@@ -44,6 +44,7 @@
 #define __OPENCV_PRECOMP_H__
 
 #include "opencv2/gpu.hpp"
+#include "opencv2/calib3d.hpp"
 
 #include "opencv2/core/gpu_private.hpp"
 

modules/gpu/test/test_calib3d.cpp

@@ -46,123 +46,6 @@
 
 using namespace cvtest;
 
-//////////////////////////////////////////////////////////////////////////
-// StereoBM
-
-struct StereoBM : testing::TestWithParam<cv::gpu::DeviceInfo>
-{
-    cv::gpu::DeviceInfo devInfo;
-
-    virtual void SetUp()
-    {
-        devInfo = GetParam();
-
-        cv::gpu::setDevice(devInfo.deviceID());
-    }
-};
-
-GPU_TEST_P(StereoBM, Regression)
-{
-    cv::Mat left_image  = readImage("stereobm/aloe-L.png", cv::IMREAD_GRAYSCALE);
-    cv::Mat right_image = readImage("stereobm/aloe-R.png", cv::IMREAD_GRAYSCALE);
-    cv::Mat disp_gold   = readImage("stereobm/aloe-disp.png", cv::IMREAD_GRAYSCALE);
-
-    ASSERT_FALSE(left_image.empty());
-    ASSERT_FALSE(right_image.empty());
-    ASSERT_FALSE(disp_gold.empty());
-
-    cv::gpu::StereoBM_GPU bm(0, 128, 19);
-    cv::gpu::GpuMat disp;
-
-    bm(loadMat(left_image), loadMat(right_image), disp);
-
-    EXPECT_MAT_NEAR(disp_gold, disp, 0.0);
-}
-
-INSTANTIATE_TEST_CASE_P(GPU_Calib3D, StereoBM, ALL_DEVICES);
-
-//////////////////////////////////////////////////////////////////////////
-// StereoBeliefPropagation
-
-struct StereoBeliefPropagation : testing::TestWithParam<cv::gpu::DeviceInfo>
-{
-    cv::gpu::DeviceInfo devInfo;
-
-    virtual void SetUp()
-    {
-        devInfo = GetParam();
-
-        cv::gpu::setDevice(devInfo.deviceID());
-    }
-};
-
-GPU_TEST_P(StereoBeliefPropagation, Regression)
-{
-    cv::Mat left_image  = readImage("stereobp/aloe-L.png");
-    cv::Mat right_image = readImage("stereobp/aloe-R.png");
-    cv::Mat disp_gold   = readImage("stereobp/aloe-disp.png", cv::IMREAD_GRAYSCALE);
-
-    ASSERT_FALSE(left_image.empty());
-    ASSERT_FALSE(right_image.empty());
-    ASSERT_FALSE(disp_gold.empty());
-
-    cv::gpu::StereoBeliefPropagation bp(64, 8, 2, 25, 0.1f, 15, 1, CV_16S);
-    cv::gpu::GpuMat disp;
-
-    bp(loadMat(left_image), loadMat(right_image), disp);
-
-    cv::Mat h_disp(disp);
-    h_disp.convertTo(h_disp, disp_gold.depth());
-
-    EXPECT_MAT_NEAR(disp_gold, h_disp, 0.0);
-}
-
-INSTANTIATE_TEST_CASE_P(GPU_Calib3D, StereoBeliefPropagation, ALL_DEVICES);
-
-//////////////////////////////////////////////////////////////////////////
-// StereoConstantSpaceBP
-
-struct StereoConstantSpaceBP : testing::TestWithParam<cv::gpu::DeviceInfo>
-{
-    cv::gpu::DeviceInfo devInfo;
-
-    virtual void SetUp()
-    {
-        devInfo = GetParam();
-
-        cv::gpu::setDevice(devInfo.deviceID());
-    }
-};
-
-GPU_TEST_P(StereoConstantSpaceBP, Regression)
-{
-    cv::Mat left_image  = readImage("csstereobp/aloe-L.png");
-    cv::Mat right_image = readImage("csstereobp/aloe-R.png");
-
-    cv::Mat disp_gold;
-
-    if (supportFeature(devInfo, cv::gpu::FEATURE_SET_COMPUTE_20))
-        disp_gold = readImage("csstereobp/aloe-disp.png", cv::IMREAD_GRAYSCALE);
-    else
-        disp_gold = readImage("csstereobp/aloe-disp_CC1X.png", cv::IMREAD_GRAYSCALE);
-
-    ASSERT_FALSE(left_image.empty());
-    ASSERT_FALSE(right_image.empty());
-    ASSERT_FALSE(disp_gold.empty());
-
-    cv::gpu::StereoConstantSpaceBP csbp(128, 16, 4, 4);
-    cv::gpu::GpuMat disp;
-
-    csbp(loadMat(left_image), loadMat(right_image), disp);
-
-    cv::Mat h_disp(disp);
-    h_disp.convertTo(h_disp, disp_gold.depth());
-
-    EXPECT_MAT_NEAR(disp_gold, h_disp, 1.0);
-}
-
-INSTANTIATE_TEST_CASE_P(GPU_Calib3D, StereoConstantSpaceBP, ALL_DEVICES);
-
 ///////////////////////////////////////////////////////////////////////////////////////////////////////
 // transformPoints
 
@@ -304,45 +187,4 @@ GPU_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());
-    }
-};
-
-GPU_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));
-
 #endif // HAVE_CUDA

modules/gpu/test/test_precomp.hpp

@@ -57,5 +57,6 @@
 #include "opencv2/core.hpp"
 #include "opencv2/core/opengl.hpp"
 #include "opencv2/gpu.hpp"
+#include "opencv2/calib3d.hpp"
 
 #endif

modules/gpucalib3d/CMakeLists.txt

@@ -1,9 +0,0 @@
-if(ANDROID OR IOS)
-  ocv_module_disable(gpucalib3d)
-endif()
-
-set(the_description "GPU-accelerated Camera Calibration and 3D Reconstruction")
-
-ocv_warnings_disable(CMAKE_CXX_FLAGS -Wundef -Wmissing-declarations)
-
-ocv_define_module(gpucalib3d opencv_calib3d opencv_gpuarithm)

modules/gpucalib3d/doc/gpucalib3d.rst

@@ -1,8 +0,0 @@
-*************************************************************
-gpu. GPU-accelerated Camera Calibration and 3D Reconstruction
-*************************************************************
-
-.. toctree::
-    :maxdepth: 1
-
-    camera_calibration_and_3d_reconstruction

modules/gpustereo/CMakeLists.txt

@@ -0,0 +1,9 @@
+if(ANDROID OR IOS)
+  ocv_module_disable(gpustereo)
+endif()
+
+set(the_description "GPU-accelerated Stereo Correspondence")
+
+ocv_warnings_disable(CMAKE_CXX_FLAGS /wd4127 /wd4324 /wd4512 -Wundef -Wmissing-declarations)
+
+ocv_define_module(gpustereo opencv_calib3d)

modules/gpustereo/doc/gpustereo.rst

@@ -0,0 +1,8 @@
+************************************************
+gpustereo. GPU-accelerated Stereo Correspondence
+************************************************
+
+.. toctree::
+    :maxdepth: 1
+
+    stereo

modules/gpustereo/doc/stereo.rst

@@ -1,5 +1,5 @@
-Camera Calibration and 3D Reconstruction
+Stereo Correspondence
-========================================
+=====================
 
 .. highlight:: cpp
 
@@ -462,38 +462,6 @@ Reprojects a disparity image to 3D space.
 
 
 
-gpu::solvePnPRansac
--------------------
-Finds the object pose from 3D-2D point correspondences.
-
-.. ocv:function:: void gpu::solvePnPRansac(const Mat& object, const Mat& image, const Mat& camera_mat, const Mat& dist_coef, Mat& rvec, Mat& tvec, bool use_extrinsic_guess=false, int num_iters=100, float max_dist=8.0, int min_inlier_count=100, vector<int>* inliers=NULL)
-
-    :param object: Single-row matrix of object points.
-
-    :param image: Single-row matrix of image points.
-
-    :param camera_mat: 3x3 matrix of intrinsic camera parameters.
-
-    :param dist_coef: Distortion coefficients. See :ocv:func:`undistortPoints` for details.
-
-    :param rvec: Output 3D rotation vector.
-
-    :param tvec: Output 3D translation vector.
-
-    :param use_extrinsic_guess: Flag to indicate that the function must use ``rvec`` and ``tvec`` as an initial transformation guess. It is not supported for now.
-
-    :param num_iters: Maximum number of RANSAC iterations.
-
-    :param max_dist: Euclidean distance threshold to detect whether point is inlier or not.
-
-    :param min_inlier_count: Flag to indicate that the function must stop if greater or equal number of inliers is achieved. It is not supported for now.
-
-    :param inliers: Output vector of inlier indices.
-
-.. seealso:: :ocv:func:`solvePnPRansac`
-
-
-
 .. [Felzenszwalb2006] Pedro F. Felzenszwalb algorithm [Pedro F. Felzenszwalb and Daniel P. Huttenlocher. *Efficient belief propagation for early vision*. International Journal of Computer Vision, 70(1), October 2006
 
 .. [Yang2010] Q. Yang, L. Wang, and N. Ahuja. *A constant-space belief propagation algorithm for stereo matching*. In CVPR, 2010.

modules/gpustereo/include/opencv2/gpustereo.hpp

@@ -40,8 +40,12 @@
 //
 //M*/
 
-#ifndef __OPENCV_GPUCALIB3D_HPP__
+#ifndef __OPENCV_GPUSTEREO_HPP__
-#define __OPENCV_GPUCALIB3D_HPP__
+#define __OPENCV_GPUSTEREO_HPP__
+
+#ifndef __cplusplus
+#  error gpustereo.hpp header must be compiled as C++
+#endif
 
 #include "opencv2/core/gpumat.hpp"
 
@@ -226,18 +230,6 @@ private:
     GpuMat table_space;
 };
 
-CV_EXPORTS void transformPoints(const GpuMat& src, const Mat& rvec, const Mat& tvec,
-                                GpuMat& dst, Stream& stream = Stream::Null());
-
-CV_EXPORTS void projectPoints(const GpuMat& src, const Mat& rvec, const Mat& tvec,
-                              const Mat& camera_mat, const Mat& dist_coef, GpuMat& dst,
-                              Stream& stream = Stream::Null());
-
-CV_EXPORTS void solvePnPRansac(const Mat& object, const Mat& image, const Mat& camera_mat,
-                               const Mat& dist_coef, Mat& rvec, Mat& tvec, bool use_extrinsic_guess=false,
-                               int num_iters=100, float max_dist=8.0, int min_inlier_count=100,
-                               std::vector<int>* inliers=NULL);
-
 //! Reprojects disparity image to 3D space.
 //! Supports CV_8U and CV_16S types of input disparity.
 //! The output is a 3- or 4-channel floating-point matrix.
@@ -252,4 +244,4 @@ CV_EXPORTS void drawColorDisp(const GpuMat& src_disp, GpuMat& dst_disp, int ndis
 
 }} // namespace cv { namespace gpu {
 
-#endif /* __OPENCV_GPUCALIB3D_HPP__ */
+#endif /* __OPENCV_GPUSTEREO_HPP__ */

modules/gpustereo/perf/perf_main.cpp

@@ -44,4 +44,4 @@
 
 using namespace perf;
 
-CV_PERF_TEST_MAIN(gpuarithm, printCudaInfo())
+CV_PERF_TEST_MAIN(gpustereo, printCudaInfo())

modules/gpustereo/perf/perf_precomp.hpp

@@ -54,8 +54,7 @@
 #include "opencv2/ts.hpp"
 #include "opencv2/ts/gpu_perf.hpp"
 
-#include "opencv2/gpucalib3d.hpp"
+#include "opencv2/gpustereo.hpp"
-
 #include "opencv2/calib3d.hpp"
 
 #ifdef GTEST_CREATE_SHARED_LIBRARY

modules/gpustereo/perf/perf_stereo.cpp

@@ -52,7 +52,7 @@ using namespace perf;
 typedef std::tr1::tuple<string, string> pair_string;
 DEF_PARAM_TEST_1(ImagePair, pair_string);
 
-PERF_TEST_P(ImagePair, Calib3D_StereoBM,
+PERF_TEST_P(ImagePair, StereoBM,
             Values(pair_string("gpu/perf/aloe.png", "gpu/perf/aloeR.png")))
 {
     declare.time(300.0);
@@ -93,7 +93,7 @@ PERF_TEST_P(ImagePair, Calib3D_StereoBM,
 //////////////////////////////////////////////////////////////////////
 // StereoBeliefPropagation
 
-PERF_TEST_P(ImagePair, Calib3D_StereoBeliefPropagation,
+PERF_TEST_P(ImagePair, StereoBeliefPropagation,
             Values(pair_string("gpu/stereobp/aloe-L.png", "gpu/stereobp/aloe-R.png")))
 {
     declare.time(300.0);
@@ -127,7 +127,7 @@ PERF_TEST_P(ImagePair, Calib3D_StereoBeliefPropagation,
 //////////////////////////////////////////////////////////////////////
 // StereoConstantSpaceBP
 
-PERF_TEST_P(ImagePair, Calib3D_StereoConstantSpaceBP,
+PERF_TEST_P(ImagePair, StereoConstantSpaceBP,
             Values(pair_string("gpu/stereobm/aloe-L.png", "gpu/stereobm/aloe-R.png")))
 {
     declare.time(300.0);
@@ -161,7 +161,7 @@ PERF_TEST_P(ImagePair, Calib3D_StereoConstantSpaceBP,
 //////////////////////////////////////////////////////////////////////
 // DisparityBilateralFilter
 
-PERF_TEST_P(ImagePair, Calib3D_DisparityBilateralFilter,
+PERF_TEST_P(ImagePair, DisparityBilateralFilter,
             Values(pair_string("gpu/stereobm/aloe-L.png", "gpu/stereobm/aloe-disp.png")))
 {
     const cv::Mat img = readImage(GET_PARAM(0), cv::IMREAD_GRAYSCALE);
@@ -190,127 +190,10 @@ PERF_TEST_P(ImagePair, Calib3D_DisparityBilateralFilter,
     }
 }
 
-//////////////////////////////////////////////////////////////////////
-// TransformPoints
-
-DEF_PARAM_TEST_1(Count, int);
-
-PERF_TEST_P(Count, Calib3D_TransformPoints,
-            Values(5000, 10000, 20000))
-{
-    const int count = GetParam();
-
-    cv::Mat src(1, count, CV_32FC3);
-    declare.in(src, WARMUP_RNG);
-
-    const cv::Mat rvec = cv::Mat::ones(1, 3, CV_32FC1);
-    const cv::Mat tvec = cv::Mat::ones(1, 3, CV_32FC1);
-
-    if (PERF_RUN_GPU())
-    {
-        const cv::gpu::GpuMat d_src(src);
-        cv::gpu::GpuMat dst;
-
-        TEST_CYCLE() cv::gpu::transformPoints(d_src, rvec, tvec, dst);
-
-        GPU_SANITY_CHECK(dst);
-    }
-    else
-    {
-        FAIL_NO_CPU();
-    }
-}
-
-//////////////////////////////////////////////////////////////////////
-// ProjectPoints
-
-PERF_TEST_P(Count, Calib3D_ProjectPoints,
-            Values(5000, 10000, 20000))
-{
-    const int count = GetParam();
-
-    cv::Mat src(1, count, CV_32FC3);
-    declare.in(src, WARMUP_RNG);
-
-    const cv::Mat rvec = cv::Mat::ones(1, 3, CV_32FC1);
-    const cv::Mat tvec = cv::Mat::ones(1, 3, CV_32FC1);
-    const cv::Mat camera_mat = cv::Mat::ones(3, 3, CV_32FC1);
-
-    if (PERF_RUN_GPU())
-    {
-        const cv::gpu::GpuMat d_src(src);
-        cv::gpu::GpuMat dst;
-
-        TEST_CYCLE() cv::gpu::projectPoints(d_src, rvec, tvec, camera_mat, cv::Mat(), dst);
-
-        GPU_SANITY_CHECK(dst);
-    }
-    else
-    {
-        cv::Mat dst;
-
-        TEST_CYCLE() cv::projectPoints(src, rvec, tvec, camera_mat, cv::noArray(), dst);
-
-        CPU_SANITY_CHECK(dst);
-    }
-}
-
-//////////////////////////////////////////////////////////////////////
-// SolvePnPRansac
-
-PERF_TEST_P(Count, Calib3D_SolvePnPRansac,
-            Values(5000, 10000, 20000))
-{
-    declare.time(10.0);
-
-    const int count = GetParam();
-
-    cv::Mat object(1, count, CV_32FC3);
-    declare.in(object, WARMUP_RNG);
-
-    cv::Mat camera_mat(3, 3, CV_32FC1);
-    cv::randu(camera_mat, 0.5, 1);
-    camera_mat.at<float>(0, 1) = 0.f;
-    camera_mat.at<float>(1, 0) = 0.f;
-    camera_mat.at<float>(2, 0) = 0.f;
-    camera_mat.at<float>(2, 1) = 0.f;
-
-    const cv::Mat dist_coef(1, 8, CV_32F, cv::Scalar::all(0));
-
-    cv::Mat rvec_gold(1, 3, CV_32FC1);
-    cv::randu(rvec_gold, 0, 1);
-
-    cv::Mat tvec_gold(1, 3, CV_32FC1);
-    cv::randu(tvec_gold, 0, 1);
-
-    std::vector<cv::Point2f> image_vec;
-    cv::projectPoints(object, rvec_gold, tvec_gold, camera_mat, dist_coef, image_vec);
-
-    const cv::Mat image(1, count, CV_32FC2, &image_vec[0]);
-
-    cv::Mat rvec;
-    cv::Mat tvec;
-
-    if (PERF_RUN_GPU())
-    {
-        TEST_CYCLE() cv::gpu::solvePnPRansac(object, image, camera_mat, dist_coef, rvec, tvec);
-
-        GPU_SANITY_CHECK(rvec, 1e-3);
-        GPU_SANITY_CHECK(tvec, 1e-3);
-    }
-    else
-    {
-        TEST_CYCLE() cv::solvePnPRansac(object, image, camera_mat, dist_coef, rvec, tvec);
-
-        CPU_SANITY_CHECK(rvec, 1e-6);
-        CPU_SANITY_CHECK(tvec, 1e-6);
-    }
-}
-
 //////////////////////////////////////////////////////////////////////
 // ReprojectImageTo3D
 
-PERF_TEST_P(Sz_Depth, Calib3D_ReprojectImageTo3D,
+PERF_TEST_P(Sz_Depth, ReprojectImageTo3D,
             Combine(GPU_TYPICAL_MAT_SIZES,
                     Values(CV_8U, CV_16S)))
 {
@@ -345,7 +228,7 @@ PERF_TEST_P(Sz_Depth, Calib3D_ReprojectImageTo3D,
 //////////////////////////////////////////////////////////////////////
 // DrawColorDisp
 
-PERF_TEST_P(Sz_Depth, Calib3D_DrawColorDisp,
+PERF_TEST_P(Sz_Depth, DrawColorDisp,
             Combine(GPU_TYPICAL_MAT_SIZES,
                     Values(CV_8U, CV_16S)))
 {

modules/gpustereo/src/cuda/util.cu

@@ -0,0 +1,235 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#if !defined CUDA_DISABLER
+
+#include "opencv2/core/cuda/common.hpp"
+#include "opencv2/core/cuda/transform.hpp"
+#include "opencv2/core/cuda/functional.hpp"
+#include "opencv2/core/cuda/reduce.hpp"
+
+namespace cv { namespace gpu { namespace cudev
+{
+    /////////////////////////////////// reprojectImageTo3D ///////////////////////////////////////////////
+
+    __constant__ float cq[16];
+
+    template <typename T, typename D>
+    __global__ void reprojectImageTo3D(const PtrStepSz<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 >= disp.rows || x >= disp.cols)
+            return;
+
+        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];
+
+        const T d = disp(y, x);
+
+        const float iW = 1.f / (qw + cq[14] * d);
+
+        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;
+
+        xyz(y, x) = v;
+    }
+
+    template <typename T, typename D>
+    void reprojectImageTo3D_gpu(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream)
+    {
+        dim3 block(32, 8);
+        dim3 grid(divUp(disp.cols, block.x), divUp(disp.rows, block.y));
+
+        cudaSafeCall( cudaMemcpyToSymbol(cq, q, 16 * sizeof(float)) );
+
+        reprojectImageTo3D<T, D><<<grid, block, 0, stream>>>((PtrStepSz<T>)disp, (PtrStepSz<D>)xyz);
+        cudaSafeCall( cudaGetLastError() );
+
+        if (stream == 0)
+            cudaSafeCall( cudaDeviceSynchronize() );
+    }
+
+    template void reprojectImageTo3D_gpu<uchar, float3>(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream);
+    template void reprojectImageTo3D_gpu<uchar, float4>(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream);
+    template void reprojectImageTo3D_gpu<short, float3>(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream);
+    template void reprojectImageTo3D_gpu<short, float4>(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream);
+
+    /////////////////////////////////// drawColorDisp ///////////////////////////////////////////////
+
+    template <typename T>
+    __device__ unsigned int cvtPixel(T d, int ndisp, float S = 1, float V = 1)
+    {
+        unsigned int H = ((ndisp-d) * 240)/ndisp;
+
+        unsigned int hi = (H/60) % 6;
+        float f = H/60.f - H/60;
+        float p = V * (1 - S);
+        float q = V * (1 - f * S);
+        float t = V * (1 - (1 - f) * S);
+
+        float3 res;
+
+        if (hi == 0) //R = V,	G = t,	B = p
+        {
+            res.x = p;
+            res.y = t;
+            res.z = V;
+        }
+
+        if (hi == 1) // R = q,	G = V,	B = p
+        {
+            res.x = p;
+            res.y = V;
+            res.z = q;
+        }
+
+        if (hi == 2) // R = p,	G = V,	B = t
+        {
+            res.x = t;
+            res.y = V;
+            res.z = p;
+        }
+
+        if (hi == 3) // R = p,	G = q,	B = V
+        {
+            res.x = V;
+            res.y = q;
+            res.z = p;
+        }
+
+        if (hi == 4) // R = t,	G = p,	B = V
+        {
+            res.x = V;
+            res.y = p;
+            res.z = t;
+        }
+
+        if (hi == 5) // R = V,	G = p,	B = q
+        {
+            res.x = q;
+            res.y = p;
+            res.z = V;
+        }
+        const unsigned int b = (unsigned int)(::max(0.f, ::min(res.x, 1.f)) * 255.f);
+        const unsigned int g = (unsigned int)(::max(0.f, ::min(res.y, 1.f)) * 255.f);
+        const unsigned int r = (unsigned int)(::max(0.f, ::min(res.z, 1.f)) * 255.f);
+        const unsigned int a = 255U;
+
+        return (a << 24) + (r << 16) + (g << 8) + b;
+    }
+
+    __global__ void drawColorDisp(uchar* disp, size_t disp_step, uchar* out_image, size_t out_step, int width, int height, int ndisp)
+    {
+        const int x = (blockIdx.x * blockDim.x + threadIdx.x) << 2;
+        const int y = blockIdx.y * blockDim.y + threadIdx.y;
+
+        if(x < width && y < height)
+        {
+            uchar4 d4 = *(uchar4*)(disp + y * disp_step + x);
+
+            uint4 res;
+            res.x = cvtPixel(d4.x, ndisp);
+            res.y = cvtPixel(d4.y, ndisp);
+            res.z = cvtPixel(d4.z, ndisp);
+            res.w = cvtPixel(d4.w, ndisp);
+
+            uint4* line = (uint4*)(out_image + y * out_step);
+            line[x >> 2] = res;
+        }
+    }
+
+    __global__ void drawColorDisp(short* disp, size_t disp_step, uchar* out_image, size_t out_step, int width, int height, int ndisp)
+    {
+        const int x = (blockIdx.x * blockDim.x + threadIdx.x) << 1;
+        const int y = blockIdx.y * blockDim.y + threadIdx.y;
+
+        if(x < width && y < height)
+        {
+            short2 d2 = *(short2*)(disp + y * disp_step + x);
+
+            uint2 res;
+            res.x = cvtPixel(d2.x, ndisp);
+            res.y = cvtPixel(d2.y, ndisp);
+
+            uint2* line = (uint2*)(out_image + y * out_step);
+            line[x >> 1] = res;
+        }
+    }
+
+
+    void drawColorDisp_gpu(const PtrStepSzb& src, const PtrStepSzb& dst, int ndisp, const cudaStream_t& stream)
+    {
+        dim3 threads(16, 16, 1);
+        dim3 grid(1, 1, 1);
+        grid.x = divUp(src.cols, threads.x << 2);
+        grid.y = divUp(src.rows, threads.y);
+
+        drawColorDisp<<<grid, threads, 0, stream>>>(src.data, src.step, dst.data, dst.step, src.cols, src.rows, ndisp);
+        cudaSafeCall( cudaGetLastError() );
+
+        if (stream == 0)
+            cudaSafeCall( cudaDeviceSynchronize() );
+    }
+
+    void drawColorDisp_gpu(const PtrStepSz<short>& src, const PtrStepSzb& dst, int ndisp, const cudaStream_t& stream)
+    {
+        dim3 threads(32, 8, 1);
+        dim3 grid(1, 1, 1);
+        grid.x = divUp(src.cols, threads.x << 1);
+        grid.y = divUp(src.rows, threads.y);
+
+        drawColorDisp<<<grid, threads, 0, stream>>>(src.data, src.step / sizeof(short), dst.data, dst.step, src.cols, src.rows, ndisp);
+        cudaSafeCall( cudaGetLastError() );
+
+        if (stream == 0)
+            cudaSafeCall( cudaDeviceSynchronize() );
+    }
+}}} // namespace cv { namespace gpu { namespace cudev
+
+
+#endif /* CUDA_DISABLER */

modules/gpustereo/src/precomp.hpp

@@ -45,11 +45,7 @@
 
 #include <limits>
 
-#include "opencv2/gpucalib3d.hpp"
+#include "opencv2/gpustereo.hpp"
-#include "opencv2/gpuarithm.hpp"
-
-#include "opencv2/calib3d.hpp"
-#include "opencv2/imgproc.hpp"
 
 #include "opencv2/core/gpu_private.hpp"
 

modules/gpustereo/src/stereocsbp.cpp

@@ -196,7 +196,7 @@ static void csbp_operator(StereoConstantSpaceBP& rthis, GpuMat& mbuf, GpuMat& te
         for(int _r = 0; _r < 5; ++_r)
         {
             *buf_ptrs[_r] = sub2.rowRange(_r * sub2.rows/5, (_r+1) * sub2.rows/5);
-            assert(buf_ptrs[_r]->cols == cols && buf_ptrs[_r]->rows == rows * rthis.nr_plane);
+            CV_DbgAssert(buf_ptrs[_r]->cols == cols && buf_ptrs[_r]->rows == rows * rthis.nr_plane);
         }
     };
 

modules/gpustereo/src/util.cpp

@@ -0,0 +1,117 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include "precomp.hpp"
+
+using namespace cv;
+using namespace cv::gpu;
+
+#if !defined HAVE_CUDA || defined(CUDA_DISABLER)
+
+void cv::gpu::reprojectImageTo3D(const GpuMat&, GpuMat&, const Mat&, int, Stream&) { throw_no_cuda(); }
+void cv::gpu::drawColorDisp(const GpuMat&, GpuMat&, int, Stream&) { throw_no_cuda(); }
+
+#else
+
+////////////////////////////////////////////////////////////////////////
+// reprojectImageTo3D
+
+namespace cv { namespace gpu { namespace cudev
+{
+    template <typename T, typename D>
+    void reprojectImageTo3D_gpu(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream);
+}}}
+
+void cv::gpu::reprojectImageTo3D(const GpuMat& disp, GpuMat& xyz, const Mat& Q, int dst_cn, Stream& stream)
+{
+    using namespace cv::gpu::cudev;
+
+    typedef void (*func_t)(const PtrStepSzb disp, PtrStepSzb 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>}
+    };
+
+    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);
+
+    xyz.create(disp.size(), CV_MAKE_TYPE(CV_32F, dst_cn));
+
+    funcs[dst_cn == 4][disp.type()](disp, xyz, Q.ptr<float>(), StreamAccessor::getStream(stream));
+}
+
+////////////////////////////////////////////////////////////////////////
+// drawColorDisp
+
+namespace cv { namespace gpu { namespace cudev
+{
+    void drawColorDisp_gpu(const PtrStepSzb& src, const PtrStepSzb& dst, int ndisp, const cudaStream_t& stream);
+    void drawColorDisp_gpu(const PtrStepSz<short>& src, const PtrStepSzb& dst, int ndisp, const cudaStream_t& stream);
+}}}
+
+namespace
+{
+    template <typename T>
+    void drawColorDisp_caller(const GpuMat& src, GpuMat& dst, int ndisp, const cudaStream_t& stream)
+    {
+        using namespace ::cv::gpu::cudev;
+
+        dst.create(src.size(), CV_8UC4);
+
+        drawColorDisp_gpu((PtrStepSz<T>)src, dst, ndisp, stream);
+    }
+
+    typedef void (*drawColorDisp_caller_t)(const GpuMat& src, GpuMat& dst, int ndisp, const cudaStream_t& stream);
+
+    const drawColorDisp_caller_t drawColorDisp_callers[] = {drawColorDisp_caller<unsigned char>, 0, 0, drawColorDisp_caller<short>, 0, 0, 0, 0};
+}
+
+void cv::gpu::drawColorDisp(const GpuMat& src, GpuMat& dst, int ndisp, Stream& stream)
+{
+    CV_Assert(src.type() == CV_8U || src.type() == CV_16S);
+
+    drawColorDisp_callers[src.type()](src, dst, ndisp, StreamAccessor::getStream(stream));
+}
+
+#endif

modules/gpustereo/test/test_precomp.hpp

@@ -54,8 +54,7 @@
 #include "opencv2/ts.hpp"
 #include "opencv2/ts/gpu_test.hpp"
 
-#include "opencv2/gpucalib3d.hpp"
+#include "opencv2/gpustereo.hpp"
-
 #include "opencv2/calib3d.hpp"
 
 #endif

modules/gpustereo/test/test_stereo.cpp

@@ -0,0 +1,207 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other materials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+#include "test_precomp.hpp"
+
+#ifdef HAVE_CUDA
+
+using namespace cvtest;
+
+//////////////////////////////////////////////////////////////////////////
+// StereoBM
+
+struct StereoBM : testing::TestWithParam<cv::gpu::DeviceInfo>
+{
+    cv::gpu::DeviceInfo devInfo;
+
+    virtual void SetUp()
+    {
+        devInfo = GetParam();
+
+        cv::gpu::setDevice(devInfo.deviceID());
+    }
+};
+
+GPU_TEST_P(StereoBM, Regression)
+{
+    cv::Mat left_image  = readImage("stereobm/aloe-L.png", cv::IMREAD_GRAYSCALE);
+    cv::Mat right_image = readImage("stereobm/aloe-R.png", cv::IMREAD_GRAYSCALE);
+    cv::Mat disp_gold   = readImage("stereobm/aloe-disp.png", cv::IMREAD_GRAYSCALE);
+
+    ASSERT_FALSE(left_image.empty());
+    ASSERT_FALSE(right_image.empty());
+    ASSERT_FALSE(disp_gold.empty());
+
+    cv::gpu::StereoBM_GPU bm(0, 128, 19);
+    cv::gpu::GpuMat disp;
+
+    bm(loadMat(left_image), loadMat(right_image), disp);
+
+    EXPECT_MAT_NEAR(disp_gold, disp, 0.0);
+}
+
+INSTANTIATE_TEST_CASE_P(GPU_Stereo, StereoBM, ALL_DEVICES);
+
+//////////////////////////////////////////////////////////////////////////
+// StereoBeliefPropagation
+
+struct StereoBeliefPropagation : testing::TestWithParam<cv::gpu::DeviceInfo>
+{
+    cv::gpu::DeviceInfo devInfo;
+
+    virtual void SetUp()
+    {
+        devInfo = GetParam();
+
+        cv::gpu::setDevice(devInfo.deviceID());
+    }
+};
+
+GPU_TEST_P(StereoBeliefPropagation, Regression)
+{
+    cv::Mat left_image  = readImage("stereobp/aloe-L.png");
+    cv::Mat right_image = readImage("stereobp/aloe-R.png");
+    cv::Mat disp_gold   = readImage("stereobp/aloe-disp.png", cv::IMREAD_GRAYSCALE);
+
+    ASSERT_FALSE(left_image.empty());
+    ASSERT_FALSE(right_image.empty());
+    ASSERT_FALSE(disp_gold.empty());
+
+    cv::gpu::StereoBeliefPropagation bp(64, 8, 2, 25, 0.1f, 15, 1, CV_16S);
+    cv::gpu::GpuMat disp;
+
+    bp(loadMat(left_image), loadMat(right_image), disp);
+
+    cv::Mat h_disp(disp);
+    h_disp.convertTo(h_disp, disp_gold.depth());
+
+    EXPECT_MAT_NEAR(disp_gold, h_disp, 0.0);
+}
+
+INSTANTIATE_TEST_CASE_P(GPU_Stereo, StereoBeliefPropagation, ALL_DEVICES);
+
+//////////////////////////////////////////////////////////////////////////
+// StereoConstantSpaceBP
+
+struct StereoConstantSpaceBP : testing::TestWithParam<cv::gpu::DeviceInfo>
+{
+    cv::gpu::DeviceInfo devInfo;
+
+    virtual void SetUp()
+    {
+        devInfo = GetParam();
+
+        cv::gpu::setDevice(devInfo.deviceID());
+    }
+};
+
+GPU_TEST_P(StereoConstantSpaceBP, Regression)
+{
+    cv::Mat left_image  = readImage("csstereobp/aloe-L.png");
+    cv::Mat right_image = readImage("csstereobp/aloe-R.png");
+
+    cv::Mat disp_gold;
+
+    if (supportFeature(devInfo, cv::gpu::FEATURE_SET_COMPUTE_20))
+        disp_gold = readImage("csstereobp/aloe-disp.png", cv::IMREAD_GRAYSCALE);
+    else
+        disp_gold = readImage("csstereobp/aloe-disp_CC1X.png", cv::IMREAD_GRAYSCALE);
+
+    ASSERT_FALSE(left_image.empty());
+    ASSERT_FALSE(right_image.empty());
+    ASSERT_FALSE(disp_gold.empty());
+
+    cv::gpu::StereoConstantSpaceBP csbp(128, 16, 4, 4);
+    cv::gpu::GpuMat disp;
+
+    csbp(loadMat(left_image), loadMat(right_image), disp);
+
+    cv::Mat h_disp(disp);
+    h_disp.convertTo(h_disp, disp_gold.depth());
+
+    EXPECT_MAT_NEAR(disp_gold, h_disp, 1.0);
+}
+
+INSTANTIATE_TEST_CASE_P(GPU_Stereo, StereoConstantSpaceBP, 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());
+    }
+};
+
+GPU_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_Stereo, ReprojectImageTo3D, testing::Combine(
+    ALL_DEVICES,
+    DIFFERENT_SIZES,
+    testing::Values(MatDepth(CV_8U), MatDepth(CV_16S)),
+    WHOLE_SUBMAT));
+
+#endif // HAVE_CUDA

samples/cpp/CMakeLists.txt

@@ -23,7 +23,7 @@ if(BUILD_EXAMPLES AND OCV_DEPENDENCIES_FOUND)
     ocv_include_directories("${OpenCV_SOURCE_DIR}/modules/gpuimgproc/include")
     ocv_include_directories("${OpenCV_SOURCE_DIR}/modules/gpufeatures2d/include")
     ocv_include_directories("${OpenCV_SOURCE_DIR}/modules/gpuvideo/include")
-    ocv_include_directories("${OpenCV_SOURCE_DIR}/modules/gpucalib3d/include")
+    ocv_include_directories("${OpenCV_SOURCE_DIR}/modules/gpustereo/include")
     ocv_include_directories("${OpenCV_SOURCE_DIR}/modules/gpuobjdetect/include")
     ocv_include_directories("${OpenCV_SOURCE_DIR}/modules/gpu/include")
   endif()

samples/gpu/CMakeLists.txt

@@ -4,7 +4,7 @@ SET(OPENCV_GPU_SAMPLES_REQUIRED_DEPS opencv_core opencv_flann opencv_imgproc ope
                                      opencv_nonfree opencv_softcascade opencv_superres
                                      opencv_gpucodec opencv_gpuarithm opencv_gpufilters opencv_gpuwarping opencv_gpuimgproc
                                      opencv_gpufeatures2d opencv_gpuvideo opencv_gpuobjdetect
-                                     opencv_gpucalib3d opencv_gpulegacy)
+                                     opencv_gpustereo opencv_gpulegacy)
 
 ocv_check_dependencies(${OPENCV_GPU_SAMPLES_REQUIRED_DEPS})