diff --git a/modules/gpu/CMakeLists.txt b/modules/gpu/CMakeLists.txt index 578957037..dec43263a 100644 --- a/modules/gpu/CMakeLists.txt +++ b/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) diff --git a/modules/gpu/doc/calib3d.rst b/modules/gpu/doc/calib3d.rst new file mode 100644 index 000000000..faa6c0fec --- /dev/null +++ b/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* 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` diff --git a/modules/gpu/include/opencv2/gpu.hpp b/modules/gpu/include/opencv2/gpu.hpp index 0466ac472..72b3f49a2 100644 --- a/modules/gpu/include/opencv2/gpu.hpp +++ b/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* inliers=NULL); + //! removes points (CV_32FC2, single row matrix) with zero mask value CV_EXPORTS void compactPoints(GpuMat &points0, GpuMat &points1, const GpuMat &mask); diff --git a/modules/gpu/perf/perf_calib3d.cpp b/modules/gpu/perf/perf_calib3d.cpp new file mode 100644 index 000000000..185d9cd68 --- /dev/null +++ b/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(0, 1) = 0.f; + camera_mat.at(1, 0) = 0.f; + camera_mat.at(2, 0) = 0.f; + camera_mat.at(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 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); + } +} diff --git a/modules/gpu/perf/perf_precomp.hpp b/modules/gpu/perf/perf_precomp.hpp index 9c75581d9..0329b3a98 100644 --- a/modules/gpu/perf/perf_precomp.hpp +++ b/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 diff --git a/modules/gpucalib3d/src/calib3d.cpp b/modules/gpu/src/calib3d.cpp similarity index 82% rename from modules/gpucalib3d/src/calib3d.cpp rename to modules/gpu/src/calib3d.cpp index 135859094..cb3b1464b 100644 --- a/modules/gpucalib3d/src/calib3d.cpp +++ b/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*) { 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 - 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, 0, 0, reprojectImageTo3D_gpu}, - {reprojectImageTo3D_gpu, 0, 0, reprojectImageTo3D_gpu} - }; - - 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(), 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& src, const PtrStepSzb& dst, int ndisp, const cudaStream_t& stream); -}}} - -namespace -{ - template - 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)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, 0, 0, drawColorDisp_caller, 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 diff --git a/modules/gpucalib3d/src/cuda/calib3d.cu b/modules/gpu/src/cuda/calib3d.cu similarity index 58% rename from modules/gpucalib3d/src/cuda/calib3d.cu rename to modules/gpu/src/cuda/calib3d.cu index d1d59ce23..6085e716d 100644 --- a/modules/gpucalib3d/src/cuda/calib3d.cu +++ b/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 - __global__ void reprojectImageTo3D(const PtrStepSz disp, PtrStep 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::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 - 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<<>>((PtrStepSz)disp, (PtrStepSz)xyz); - cudaSafeCall( cudaGetLastError() ); - - if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); - } - - template void reprojectImageTo3D_gpu(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream); - template void reprojectImageTo3D_gpu(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream); - template void reprojectImageTo3D_gpu(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream); - template void reprojectImageTo3D_gpu(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream); - - /////////////////////////////////// drawColorDisp /////////////////////////////////////////////// - - template - __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<<>>(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& 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<<>>(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 diff --git a/modules/gpu/src/precomp.hpp b/modules/gpu/src/precomp.hpp index 1b5207b38..c662cb1d7 100644 --- a/modules/gpu/src/precomp.hpp +++ b/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" diff --git a/modules/gpucalib3d/test/test_calib3d.cpp b/modules/gpu/test/test_calib3d.cpp similarity index 61% rename from modules/gpucalib3d/test/test_calib3d.cpp rename to modules/gpu/test/test_calib3d.cpp index 5de3d34df..3ad19dcbe 100644 --- a/modules/gpucalib3d/test/test_calib3d.cpp +++ b/modules/gpu/test/test_calib3d.cpp @@ -46,123 +46,6 @@ using namespace cvtest; -////////////////////////////////////////////////////////////////////////// -// StereoBM - -struct StereoBM : testing::TestWithParam -{ - 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 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 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 diff --git a/modules/gpu/test/test_precomp.hpp b/modules/gpu/test/test_precomp.hpp index 1e4248101..3177e88e7 100644 --- a/modules/gpu/test/test_precomp.hpp +++ b/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 diff --git a/modules/gpucalib3d/CMakeLists.txt b/modules/gpucalib3d/CMakeLists.txt deleted file mode 100644 index bb949c4d3..000000000 --- a/modules/gpucalib3d/CMakeLists.txt +++ /dev/null @@ -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) diff --git a/modules/gpucalib3d/doc/gpucalib3d.rst b/modules/gpucalib3d/doc/gpucalib3d.rst deleted file mode 100644 index 5dffaa048..000000000 --- a/modules/gpucalib3d/doc/gpucalib3d.rst +++ /dev/null @@ -1,8 +0,0 @@ -************************************************************* -gpu. GPU-accelerated Camera Calibration and 3D Reconstruction -************************************************************* - -.. toctree:: - :maxdepth: 1 - - camera_calibration_and_3d_reconstruction diff --git a/modules/gpustereo/CMakeLists.txt b/modules/gpustereo/CMakeLists.txt new file mode 100644 index 000000000..c2a302c8a --- /dev/null +++ b/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) diff --git a/modules/gpustereo/doc/gpustereo.rst b/modules/gpustereo/doc/gpustereo.rst new file mode 100644 index 000000000..d3076794c --- /dev/null +++ b/modules/gpustereo/doc/gpustereo.rst @@ -0,0 +1,8 @@ +************************************************ +gpustereo. GPU-accelerated Stereo Correspondence +************************************************ + +.. toctree:: + :maxdepth: 1 + + stereo diff --git a/modules/gpucalib3d/doc/camera_calibration_and_3d_reconstruction.rst b/modules/gpustereo/doc/stereo.rst similarity index 92% rename from modules/gpucalib3d/doc/camera_calibration_and_3d_reconstruction.rst rename to modules/gpustereo/doc/stereo.rst index 587c253d2..cd2add0b9 100644 --- a/modules/gpucalib3d/doc/camera_calibration_and_3d_reconstruction.rst +++ b/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* 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. diff --git a/modules/gpucalib3d/include/opencv2/gpucalib3d.hpp b/modules/gpustereo/include/opencv2/gpustereo.hpp similarity index 91% rename from modules/gpucalib3d/include/opencv2/gpucalib3d.hpp rename to modules/gpustereo/include/opencv2/gpustereo.hpp index 3496d987b..54ce82d08 100644 --- a/modules/gpucalib3d/include/opencv2/gpucalib3d.hpp +++ b/modules/gpustereo/include/opencv2/gpustereo.hpp @@ -40,8 +40,12 @@ // //M*/ -#ifndef __OPENCV_GPUCALIB3D_HPP__ -#define __OPENCV_GPUCALIB3D_HPP__ +#ifndef __OPENCV_GPUSTEREO_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* 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__ */ diff --git a/modules/gpucalib3d/perf/perf_main.cpp b/modules/gpustereo/perf/perf_main.cpp similarity index 97% rename from modules/gpucalib3d/perf/perf_main.cpp rename to modules/gpustereo/perf/perf_main.cpp index b35791cda..d681cdb4d 100644 --- a/modules/gpucalib3d/perf/perf_main.cpp +++ b/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()) diff --git a/modules/gpucalib3d/perf/perf_precomp.cpp b/modules/gpustereo/perf/perf_precomp.cpp similarity index 100% rename from modules/gpucalib3d/perf/perf_precomp.cpp rename to modules/gpustereo/perf/perf_precomp.cpp diff --git a/modules/gpucalib3d/perf/perf_precomp.hpp b/modules/gpustereo/perf/perf_precomp.hpp similarity index 98% rename from modules/gpucalib3d/perf/perf_precomp.hpp rename to modules/gpustereo/perf/perf_precomp.hpp index dc244a72a..f61de9fae 100644 --- a/modules/gpucalib3d/perf/perf_precomp.hpp +++ b/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 diff --git a/modules/gpucalib3d/perf/perf_calib3d.cpp b/modules/gpustereo/perf/perf_stereo.cpp similarity index 68% rename from modules/gpucalib3d/perf/perf_calib3d.cpp rename to modules/gpustereo/perf/perf_stereo.cpp index 725f49c53..e0438c0ae 100644 --- a/modules/gpucalib3d/perf/perf_calib3d.cpp +++ b/modules/gpustereo/perf/perf_stereo.cpp @@ -52,7 +52,7 @@ using namespace perf; typedef std::tr1::tuple 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(0, 1) = 0.f; - camera_mat.at(1, 0) = 0.f; - camera_mat.at(2, 0) = 0.f; - camera_mat.at(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 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))) { diff --git a/modules/gpucalib3d/src/cuda/disp_bilateral_filter.cu b/modules/gpustereo/src/cuda/disparity_bilateral_filter.cu similarity index 100% rename from modules/gpucalib3d/src/cuda/disp_bilateral_filter.cu rename to modules/gpustereo/src/cuda/disparity_bilateral_filter.cu diff --git a/modules/gpucalib3d/src/cuda/stereobm.cu b/modules/gpustereo/src/cuda/stereobm.cu similarity index 100% rename from modules/gpucalib3d/src/cuda/stereobm.cu rename to modules/gpustereo/src/cuda/stereobm.cu diff --git a/modules/gpucalib3d/src/cuda/stereobp.cu b/modules/gpustereo/src/cuda/stereobp.cu similarity index 100% rename from modules/gpucalib3d/src/cuda/stereobp.cu rename to modules/gpustereo/src/cuda/stereobp.cu diff --git a/modules/gpucalib3d/src/cuda/stereocsbp.cu b/modules/gpustereo/src/cuda/stereocsbp.cu similarity index 100% rename from modules/gpucalib3d/src/cuda/stereocsbp.cu rename to modules/gpustereo/src/cuda/stereocsbp.cu diff --git a/modules/gpustereo/src/cuda/util.cu b/modules/gpustereo/src/cuda/util.cu new file mode 100644 index 000000000..1945d2463 --- /dev/null +++ b/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 + __global__ void reprojectImageTo3D(const PtrStepSz disp, PtrStep 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::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 + 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<<>>((PtrStepSz)disp, (PtrStepSz)xyz); + cudaSafeCall( cudaGetLastError() ); + + if (stream == 0) + cudaSafeCall( cudaDeviceSynchronize() ); + } + + template void reprojectImageTo3D_gpu(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream); + template void reprojectImageTo3D_gpu(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream); + template void reprojectImageTo3D_gpu(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream); + template void reprojectImageTo3D_gpu(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream); + + /////////////////////////////////// drawColorDisp /////////////////////////////////////////////// + + template + __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<<>>(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& 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<<>>(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 */ diff --git a/modules/gpucalib3d/src/disparity_bilateral_filter.cpp b/modules/gpustereo/src/disparity_bilateral_filter.cpp similarity index 100% rename from modules/gpucalib3d/src/disparity_bilateral_filter.cpp rename to modules/gpustereo/src/disparity_bilateral_filter.cpp diff --git a/modules/gpucalib3d/src/precomp.cpp b/modules/gpustereo/src/precomp.cpp similarity index 100% rename from modules/gpucalib3d/src/precomp.cpp rename to modules/gpustereo/src/precomp.cpp diff --git a/modules/gpucalib3d/src/precomp.hpp b/modules/gpustereo/src/precomp.hpp similarity index 94% rename from modules/gpucalib3d/src/precomp.hpp rename to modules/gpustereo/src/precomp.hpp index 89396fd57..a552d5ff3 100644 --- a/modules/gpucalib3d/src/precomp.hpp +++ b/modules/gpustereo/src/precomp.hpp @@ -45,11 +45,7 @@ #include -#include "opencv2/gpucalib3d.hpp" -#include "opencv2/gpuarithm.hpp" - -#include "opencv2/calib3d.hpp" -#include "opencv2/imgproc.hpp" +#include "opencv2/gpustereo.hpp" #include "opencv2/core/gpu_private.hpp" diff --git a/modules/gpucalib3d/src/stereobm.cpp b/modules/gpustereo/src/stereobm.cpp similarity index 100% rename from modules/gpucalib3d/src/stereobm.cpp rename to modules/gpustereo/src/stereobm.cpp diff --git a/modules/gpucalib3d/src/stereobp.cpp b/modules/gpustereo/src/stereobp.cpp similarity index 100% rename from modules/gpucalib3d/src/stereobp.cpp rename to modules/gpustereo/src/stereobp.cpp diff --git a/modules/gpucalib3d/src/stereocsbp.cpp b/modules/gpustereo/src/stereocsbp.cpp similarity index 99% rename from modules/gpucalib3d/src/stereocsbp.cpp rename to modules/gpustereo/src/stereocsbp.cpp index 431dfd38c..bd5ef4be9 100644 --- a/modules/gpucalib3d/src/stereocsbp.cpp +++ b/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); } }; diff --git a/modules/gpustereo/src/util.cpp b/modules/gpustereo/src/util.cpp new file mode 100644 index 000000000..9bff6fff2 --- /dev/null +++ b/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 + 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, 0, 0, reprojectImageTo3D_gpu}, + {reprojectImageTo3D_gpu, 0, 0, reprojectImageTo3D_gpu} + }; + + 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(), 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& src, const PtrStepSzb& dst, int ndisp, const cudaStream_t& stream); +}}} + +namespace +{ + template + 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)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, 0, 0, drawColorDisp_caller, 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 diff --git a/modules/gpucalib3d/test/test_main.cpp b/modules/gpustereo/test/test_main.cpp similarity index 100% rename from modules/gpucalib3d/test/test_main.cpp rename to modules/gpustereo/test/test_main.cpp diff --git a/modules/gpucalib3d/test/test_precomp.cpp b/modules/gpustereo/test/test_precomp.cpp similarity index 100% rename from modules/gpucalib3d/test/test_precomp.cpp rename to modules/gpustereo/test/test_precomp.cpp diff --git a/modules/gpucalib3d/test/test_precomp.hpp b/modules/gpustereo/test/test_precomp.hpp similarity index 98% rename from modules/gpucalib3d/test/test_precomp.hpp rename to modules/gpustereo/test/test_precomp.hpp index 8c53f4786..d55b1ec0d 100644 --- a/modules/gpucalib3d/test/test_precomp.hpp +++ b/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 diff --git a/modules/gpustereo/test/test_stereo.cpp b/modules/gpustereo/test/test_stereo.cpp new file mode 100644 index 000000000..0ead03dc5 --- /dev/null +++ b/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 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 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 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 diff --git a/samples/cpp/CMakeLists.txt b/samples/cpp/CMakeLists.txt index d1807e929..ea1b3e102 100644 --- a/samples/cpp/CMakeLists.txt +++ b/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() diff --git a/samples/gpu/CMakeLists.txt b/samples/gpu/CMakeLists.txt index 670c71c9e..f03af5605 100644 --- a/samples/gpu/CMakeLists.txt +++ b/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})