From b08b9ab83b4b05e937a5e464b01f6ace057de8dc Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Wed, 17 Apr 2013 17:51:19 +0400 Subject: [PATCH] gpucalib3d module for camera calibration and stereo correspondence --- modules/gpu/CMakeLists.txt | 2 +- modules/gpu/doc/gpu.rst | 1 - modules/gpu/include/opencv2/gpu.hpp | 186 +------------ modules/gpucalib3d/CMakeLists.txt | 9 + ...mera_calibration_and_3d_reconstruction.rst | 0 modules/gpucalib3d/doc/gpucalib3d.rst | 8 + .../gpucalib3d/include/opencv2/gpucalib3d.hpp | 255 ++++++++++++++++++ .../{gpu => gpucalib3d}/perf/perf_calib3d.cpp | 0 modules/gpucalib3d/perf/perf_main.cpp | 47 ++++ modules/gpucalib3d/perf/perf_precomp.cpp | 43 +++ modules/gpucalib3d/perf/perf_precomp.hpp | 65 +++++ modules/{gpu => gpucalib3d}/src/calib3d.cpp | 76 +++++- .../{gpu => gpucalib3d}/src/cuda/calib3d.cu | 183 +++++++++++++ .../src/cuda/disp_bilateral_filter.cu | 0 .../{gpu => gpucalib3d}/src/cuda/stereobm.cu | 0 .../{gpu => gpucalib3d}/src/cuda/stereobp.cu | 0 .../src/cuda/stereocsbp.cu | 0 .../src/disparity_bilateral_filter.cpp | 0 modules/gpucalib3d/src/precomp.cpp | 43 +++ modules/gpucalib3d/src/precomp.hpp | 56 ++++ modules/{gpu => gpucalib3d}/src/stereobm.cpp | 0 modules/{gpu => gpucalib3d}/src/stereobp.cpp | 0 .../{gpu => gpucalib3d}/src/stereocsbp.cpp | 0 .../{gpu => gpucalib3d}/test/test_calib3d.cpp | 0 modules/gpucalib3d/test/test_main.cpp | 45 ++++ modules/gpucalib3d/test/test_precomp.cpp | 43 +++ modules/gpucalib3d/test/test_precomp.hpp | 61 +++++ .../gpuimgproc/include/opencv2/gpuimgproc.hpp | 12 - modules/gpuimgproc/src/cuda/imgproc.cu | 181 ------------- modules/gpuimgproc/src/imgproc.cpp | 70 ----- samples/cpp/CMakeLists.txt | 1 + samples/gpu/CMakeLists.txt | 3 +- 32 files changed, 931 insertions(+), 459 deletions(-) create mode 100644 modules/gpucalib3d/CMakeLists.txt rename modules/{gpu => gpucalib3d}/doc/camera_calibration_and_3d_reconstruction.rst (100%) create mode 100644 modules/gpucalib3d/doc/gpucalib3d.rst create mode 100644 modules/gpucalib3d/include/opencv2/gpucalib3d.hpp rename modules/{gpu => gpucalib3d}/perf/perf_calib3d.cpp (100%) create mode 100644 modules/gpucalib3d/perf/perf_main.cpp create mode 100644 modules/gpucalib3d/perf/perf_precomp.cpp create mode 100644 modules/gpucalib3d/perf/perf_precomp.hpp rename modules/{gpu => gpucalib3d}/src/calib3d.cpp (80%) rename modules/{gpu => gpucalib3d}/src/cuda/calib3d.cu (58%) rename modules/{gpu => gpucalib3d}/src/cuda/disp_bilateral_filter.cu (100%) rename modules/{gpu => gpucalib3d}/src/cuda/stereobm.cu (100%) rename modules/{gpu => gpucalib3d}/src/cuda/stereobp.cu (100%) rename modules/{gpu => gpucalib3d}/src/cuda/stereocsbp.cu (100%) rename modules/{gpu => gpucalib3d}/src/disparity_bilateral_filter.cpp (100%) create mode 100644 modules/gpucalib3d/src/precomp.cpp create mode 100644 modules/gpucalib3d/src/precomp.hpp rename modules/{gpu => gpucalib3d}/src/stereobm.cpp (100%) rename modules/{gpu => gpucalib3d}/src/stereobp.cpp (100%) rename modules/{gpu => gpucalib3d}/src/stereocsbp.cpp (100%) rename modules/{gpu => gpucalib3d}/test/test_calib3d.cpp (100%) create mode 100644 modules/gpucalib3d/test/test_main.cpp create mode 100644 modules/gpucalib3d/test/test_precomp.cpp create mode 100644 modules/gpucalib3d/test/test_precomp.hpp diff --git a/modules/gpu/CMakeLists.txt b/modules/gpu/CMakeLists.txt index 296545addf..95de6789f2 100644 --- a/modules/gpu/CMakeLists.txt +++ b/modules/gpu/CMakeLists.txt @@ -5,7 +5,7 @@ endif() set(the_description "GPU-accelerated Computer Vision") ocv_add_module(gpu opencv_imgproc opencv_calib3d opencv_objdetect opencv_video opencv_photo opencv_legacy - opencv_gpuarithm opencv_gpufilters opencv_gpuimgproc opencv_gpufeatures2d opencv_gpuvideo + opencv_gpuarithm opencv_gpufilters opencv_gpuimgproc opencv_gpufeatures2d opencv_gpuvideo opencv_gpucalib3d OPTIONAL opencv_gpunvidia) ocv_module_include_directories("${CMAKE_CURRENT_SOURCE_DIR}/src/cuda") diff --git a/modules/gpu/doc/gpu.rst b/modules/gpu/doc/gpu.rst index 3803efd199..bc3b9bdb1f 100644 --- a/modules/gpu/doc/gpu.rst +++ b/modules/gpu/doc/gpu.rst @@ -9,4 +9,3 @@ gpu. GPU-accelerated Computer Vision initalization_and_information data_structures object_detection - camera_calibration_and_3d_reconstruction diff --git a/modules/gpu/include/opencv2/gpu.hpp b/modules/gpu/include/opencv2/gpu.hpp index 524a32fc95..b3fea3fbf5 100644 --- a/modules/gpu/include/opencv2/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu.hpp @@ -55,6 +55,7 @@ #include "opencv2/gpuimgproc.hpp" #include "opencv2/gpufeatures2d.hpp" #include "opencv2/gpuvideo.hpp" +#include "opencv2/gpucalib3d.hpp" #include "opencv2/imgproc.hpp" #include "opencv2/objdetect.hpp" @@ -68,18 +69,6 @@ namespace cv { namespace gpu { ///////////////////////////// Calibration 3D ////////////////////////////////// -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); - //////////////////////////////// Image Labeling //////////////////////////////// @@ -90,190 +79,17 @@ CV_EXPORTS void solvePnPRansac(const Mat& object, const Mat& image, const Mat& c //////////////////////////////// StereoBM_GPU //////////////////////////////// -class CV_EXPORTS StereoBM_GPU -{ -public: - enum { BASIC_PRESET = 0, PREFILTER_XSOBEL = 1 }; - - enum { DEFAULT_NDISP = 64, DEFAULT_WINSZ = 19 }; - - //! the default constructor - StereoBM_GPU(); - //! the full constructor taking the camera-specific preset, number of disparities and the SAD window size. ndisparities must be multiple of 8. - StereoBM_GPU(int preset, int ndisparities = DEFAULT_NDISP, int winSize = DEFAULT_WINSZ); - //! the stereo correspondence operator. Finds the disparity for the specified rectified stereo pair - //! Output disparity has CV_8U type. - void operator()(const GpuMat& left, const GpuMat& right, GpuMat& disparity, Stream& stream = Stream::Null()); - - //! Some heuristics that tries to estmate - // if current GPU will be faster than CPU in this algorithm. - // It queries current active device. - static bool checkIfGpuCallReasonable(); - - int preset; - int ndisp; - int winSize; - - // If avergeTexThreshold == 0 => post procesing is disabled - // If avergeTexThreshold != 0 then disparity is set 0 in each point (x,y) where for left image - // SumOfHorizontalGradiensInWindow(x, y, winSize) < (winSize * winSize) * avergeTexThreshold - // i.e. input left image is low textured. - float avergeTexThreshold; - -private: - GpuMat minSSD, leBuf, riBuf; -}; ////////////////////////// StereoBeliefPropagation /////////////////////////// -// "Efficient Belief Propagation for Early Vision" -// P.Felzenszwalb -class CV_EXPORTS StereoBeliefPropagation -{ -public: - enum { DEFAULT_NDISP = 64 }; - enum { DEFAULT_ITERS = 5 }; - enum { DEFAULT_LEVELS = 5 }; - - static void estimateRecommendedParams(int width, int height, int& ndisp, int& iters, int& levels); - - //! the default constructor - explicit StereoBeliefPropagation(int ndisp = DEFAULT_NDISP, - int iters = DEFAULT_ITERS, - int levels = DEFAULT_LEVELS, - int msg_type = CV_32F); - - //! the full constructor taking the number of disparities, number of BP iterations on each level, - //! number of levels, truncation of data cost, data weight, - //! truncation of discontinuity cost and discontinuity single jump - //! DataTerm = data_weight * min(fabs(I2-I1), max_data_term) - //! DiscTerm = min(disc_single_jump * fabs(f1-f2), max_disc_term) - //! please see paper for more details - StereoBeliefPropagation(int ndisp, int iters, int levels, - float max_data_term, float data_weight, - float max_disc_term, float disc_single_jump, - int msg_type = CV_32F); - - //! the stereo correspondence operator. Finds the disparity for the specified rectified stereo pair, - //! if disparity is empty output type will be CV_16S else output type will be disparity.type(). - void operator()(const GpuMat& left, const GpuMat& right, GpuMat& disparity, Stream& stream = Stream::Null()); - - - //! version for user specified data term - void operator()(const GpuMat& data, GpuMat& disparity, Stream& stream = Stream::Null()); - - int ndisp; - - int iters; - int levels; - - float max_data_term; - float data_weight; - float max_disc_term; - float disc_single_jump; - - int msg_type; -private: - GpuMat u, d, l, r, u2, d2, l2, r2; - std::vector datas; - GpuMat out; -}; /////////////////////////// StereoConstantSpaceBP /////////////////////////// -// "A Constant-Space Belief Propagation Algorithm for Stereo Matching" -// Qingxiong Yang, Liang Wang, Narendra Ahuja -// http://vision.ai.uiuc.edu/~qyang6/ - -class CV_EXPORTS StereoConstantSpaceBP -{ -public: - enum { DEFAULT_NDISP = 128 }; - enum { DEFAULT_ITERS = 8 }; - enum { DEFAULT_LEVELS = 4 }; - enum { DEFAULT_NR_PLANE = 4 }; - - static void estimateRecommendedParams(int width, int height, int& ndisp, int& iters, int& levels, int& nr_plane); - - //! the default constructor - explicit StereoConstantSpaceBP(int ndisp = DEFAULT_NDISP, - int iters = DEFAULT_ITERS, - int levels = DEFAULT_LEVELS, - int nr_plane = DEFAULT_NR_PLANE, - int msg_type = CV_32F); - - //! the full constructor taking the number of disparities, number of BP iterations on each level, - //! number of levels, number of active disparity on the first level, truncation of data cost, data weight, - //! truncation of discontinuity cost, discontinuity single jump and minimum disparity threshold - StereoConstantSpaceBP(int ndisp, int iters, int levels, int nr_plane, - float max_data_term, float data_weight, float max_disc_term, float disc_single_jump, - int min_disp_th = 0, - int msg_type = CV_32F); - - //! the stereo correspondence operator. Finds the disparity for the specified rectified stereo pair, - //! if disparity is empty output type will be CV_16S else output type will be disparity.type(). - void operator()(const GpuMat& left, const GpuMat& right, GpuMat& disparity, Stream& stream = Stream::Null()); - - int ndisp; - - int iters; - int levels; - - int nr_plane; - - float max_data_term; - float data_weight; - float max_disc_term; - float disc_single_jump; - - int min_disp_th; - int msg_type; - bool use_local_init_data_cost; -private: - GpuMat messages_buffers; - - GpuMat temp; - GpuMat out; -}; /////////////////////////// DisparityBilateralFilter /////////////////////////// -// Disparity map refinement using joint bilateral filtering given a single color image. -// Qingxiong Yang, Liang Wang, Narendra Ahuja -// http://vision.ai.uiuc.edu/~qyang6/ - -class CV_EXPORTS DisparityBilateralFilter -{ -public: - enum { DEFAULT_NDISP = 64 }; - enum { DEFAULT_RADIUS = 3 }; - enum { DEFAULT_ITERS = 1 }; - - //! the default constructor - explicit DisparityBilateralFilter(int ndisp = DEFAULT_NDISP, int radius = DEFAULT_RADIUS, int iters = DEFAULT_ITERS); - //! the full constructor taking the number of disparities, filter radius, - //! number of iterations, truncation of data continuity, truncation of disparity continuity - //! and filter range sigma - DisparityBilateralFilter(int ndisp, int radius, int iters, float edge_threshold, float max_disc_threshold, float sigma_range); - - //! the disparity map refinement operator. Refine disparity map using joint bilateral filtering given a single color image. - //! disparity must have CV_8U or CV_16S type, image must have CV_8UC1 or CV_8UC3 type. - void operator()(const GpuMat& disparity, const GpuMat& image, GpuMat& dst, Stream& stream = Stream::Null()); - -private: - int ndisp; - int radius; - int iters; - - float edge_threshold; - float max_disc_threshold; - float sigma_range; - - GpuMat table_color; - GpuMat table_space; -}; //////////////// HOG (Histogram-of-Oriented-Gradients) Descriptor and Object Detector ////////////// diff --git a/modules/gpucalib3d/CMakeLists.txt b/modules/gpucalib3d/CMakeLists.txt new file mode 100644 index 0000000000..bb949c4d35 --- /dev/null +++ b/modules/gpucalib3d/CMakeLists.txt @@ -0,0 +1,9 @@ +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/gpu/doc/camera_calibration_and_3d_reconstruction.rst b/modules/gpucalib3d/doc/camera_calibration_and_3d_reconstruction.rst similarity index 100% rename from modules/gpu/doc/camera_calibration_and_3d_reconstruction.rst rename to modules/gpucalib3d/doc/camera_calibration_and_3d_reconstruction.rst diff --git a/modules/gpucalib3d/doc/gpucalib3d.rst b/modules/gpucalib3d/doc/gpucalib3d.rst new file mode 100644 index 0000000000..5dffaa0481 --- /dev/null +++ b/modules/gpucalib3d/doc/gpucalib3d.rst @@ -0,0 +1,8 @@ +************************************************************* +gpu. GPU-accelerated Camera Calibration and 3D Reconstruction +************************************************************* + +.. toctree:: + :maxdepth: 1 + + camera_calibration_and_3d_reconstruction diff --git a/modules/gpucalib3d/include/opencv2/gpucalib3d.hpp b/modules/gpucalib3d/include/opencv2/gpucalib3d.hpp new file mode 100644 index 0000000000..3496d987b4 --- /dev/null +++ b/modules/gpucalib3d/include/opencv2/gpucalib3d.hpp @@ -0,0 +1,255 @@ +/*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*/ + +#ifndef __OPENCV_GPUCALIB3D_HPP__ +#define __OPENCV_GPUCALIB3D_HPP__ + +#include "opencv2/core/gpumat.hpp" + +namespace cv { namespace gpu { + +class CV_EXPORTS StereoBM_GPU +{ +public: + enum { BASIC_PRESET = 0, PREFILTER_XSOBEL = 1 }; + + enum { DEFAULT_NDISP = 64, DEFAULT_WINSZ = 19 }; + + //! the default constructor + StereoBM_GPU(); + //! the full constructor taking the camera-specific preset, number of disparities and the SAD window size. ndisparities must be multiple of 8. + StereoBM_GPU(int preset, int ndisparities = DEFAULT_NDISP, int winSize = DEFAULT_WINSZ); + + //! the stereo correspondence operator. Finds the disparity for the specified rectified stereo pair + //! Output disparity has CV_8U type. + void operator()(const GpuMat& left, const GpuMat& right, GpuMat& disparity, Stream& stream = Stream::Null()); + + //! Some heuristics that tries to estmate + // if current GPU will be faster than CPU in this algorithm. + // It queries current active device. + static bool checkIfGpuCallReasonable(); + + int preset; + int ndisp; + int winSize; + + // If avergeTexThreshold == 0 => post procesing is disabled + // If avergeTexThreshold != 0 then disparity is set 0 in each point (x,y) where for left image + // SumOfHorizontalGradiensInWindow(x, y, winSize) < (winSize * winSize) * avergeTexThreshold + // i.e. input left image is low textured. + float avergeTexThreshold; + +private: + GpuMat minSSD, leBuf, riBuf; +}; + +// "Efficient Belief Propagation for Early Vision" +// P.Felzenszwalb +class CV_EXPORTS StereoBeliefPropagation +{ +public: + enum { DEFAULT_NDISP = 64 }; + enum { DEFAULT_ITERS = 5 }; + enum { DEFAULT_LEVELS = 5 }; + + static void estimateRecommendedParams(int width, int height, int& ndisp, int& iters, int& levels); + + //! the default constructor + explicit StereoBeliefPropagation(int ndisp = DEFAULT_NDISP, + int iters = DEFAULT_ITERS, + int levels = DEFAULT_LEVELS, + int msg_type = CV_32F); + + //! the full constructor taking the number of disparities, number of BP iterations on each level, + //! number of levels, truncation of data cost, data weight, + //! truncation of discontinuity cost and discontinuity single jump + //! DataTerm = data_weight * min(fabs(I2-I1), max_data_term) + //! DiscTerm = min(disc_single_jump * fabs(f1-f2), max_disc_term) + //! please see paper for more details + StereoBeliefPropagation(int ndisp, int iters, int levels, + float max_data_term, float data_weight, + float max_disc_term, float disc_single_jump, + int msg_type = CV_32F); + + //! the stereo correspondence operator. Finds the disparity for the specified rectified stereo pair, + //! if disparity is empty output type will be CV_16S else output type will be disparity.type(). + void operator()(const GpuMat& left, const GpuMat& right, GpuMat& disparity, Stream& stream = Stream::Null()); + + + //! version for user specified data term + void operator()(const GpuMat& data, GpuMat& disparity, Stream& stream = Stream::Null()); + + int ndisp; + + int iters; + int levels; + + float max_data_term; + float data_weight; + float max_disc_term; + float disc_single_jump; + + int msg_type; +private: + GpuMat u, d, l, r, u2, d2, l2, r2; + std::vector datas; + GpuMat out; +}; + +// "A Constant-Space Belief Propagation Algorithm for Stereo Matching" +// Qingxiong Yang, Liang Wang, Narendra Ahuja +// http://vision.ai.uiuc.edu/~qyang6/ +class CV_EXPORTS StereoConstantSpaceBP +{ +public: + enum { DEFAULT_NDISP = 128 }; + enum { DEFAULT_ITERS = 8 }; + enum { DEFAULT_LEVELS = 4 }; + enum { DEFAULT_NR_PLANE = 4 }; + + static void estimateRecommendedParams(int width, int height, int& ndisp, int& iters, int& levels, int& nr_plane); + + //! the default constructor + explicit StereoConstantSpaceBP(int ndisp = DEFAULT_NDISP, + int iters = DEFAULT_ITERS, + int levels = DEFAULT_LEVELS, + int nr_plane = DEFAULT_NR_PLANE, + int msg_type = CV_32F); + + //! the full constructor taking the number of disparities, number of BP iterations on each level, + //! number of levels, number of active disparity on the first level, truncation of data cost, data weight, + //! truncation of discontinuity cost, discontinuity single jump and minimum disparity threshold + StereoConstantSpaceBP(int ndisp, int iters, int levels, int nr_plane, + float max_data_term, float data_weight, float max_disc_term, float disc_single_jump, + int min_disp_th = 0, + int msg_type = CV_32F); + + //! the stereo correspondence operator. Finds the disparity for the specified rectified stereo pair, + //! if disparity is empty output type will be CV_16S else output type will be disparity.type(). + void operator()(const GpuMat& left, const GpuMat& right, GpuMat& disparity, Stream& stream = Stream::Null()); + + int ndisp; + + int iters; + int levels; + + int nr_plane; + + float max_data_term; + float data_weight; + float max_disc_term; + float disc_single_jump; + + int min_disp_th; + + int msg_type; + + bool use_local_init_data_cost; +private: + GpuMat messages_buffers; + + GpuMat temp; + GpuMat out; +}; + +// Disparity map refinement using joint bilateral filtering given a single color image. +// Qingxiong Yang, Liang Wang, Narendra Ahuja +// http://vision.ai.uiuc.edu/~qyang6/ +class CV_EXPORTS DisparityBilateralFilter +{ +public: + enum { DEFAULT_NDISP = 64 }; + enum { DEFAULT_RADIUS = 3 }; + enum { DEFAULT_ITERS = 1 }; + + //! the default constructor + explicit DisparityBilateralFilter(int ndisp = DEFAULT_NDISP, int radius = DEFAULT_RADIUS, int iters = DEFAULT_ITERS); + + //! the full constructor taking the number of disparities, filter radius, + //! number of iterations, truncation of data continuity, truncation of disparity continuity + //! and filter range sigma + DisparityBilateralFilter(int ndisp, int radius, int iters, float edge_threshold, float max_disc_threshold, float sigma_range); + + //! the disparity map refinement operator. Refine disparity map using joint bilateral filtering given a single color image. + //! disparity must have CV_8U or CV_16S type, image must have CV_8UC1 or CV_8UC3 type. + void operator()(const GpuMat& disparity, const GpuMat& image, GpuMat& dst, Stream& stream = Stream::Null()); + +private: + int ndisp; + int radius; + int iters; + + float edge_threshold; + float max_disc_threshold; + float sigma_range; + + GpuMat table_color; + 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. +//! Each element of this matrix will contain the 3D coordinates of the point (x,y,z,1), computed from the disparity map. +//! Q is the 4x4 perspective transformation matrix that can be obtained with cvStereoRectify. +CV_EXPORTS void reprojectImageTo3D(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, int dst_cn = 4, Stream& stream = Stream::Null()); + +//! Does coloring of disparity image: [0..ndisp) -> [0..240, 1, 1] in HSV. +//! Supported types of input disparity: CV_8U, CV_16S. +//! Output disparity has CV_8UC4 type in BGRA format (alpha = 255). +CV_EXPORTS void drawColorDisp(const GpuMat& src_disp, GpuMat& dst_disp, int ndisp, Stream& stream = Stream::Null()); + +}} // namespace cv { namespace gpu { + +#endif /* __OPENCV_GPUCALIB3D_HPP__ */ diff --git a/modules/gpu/perf/perf_calib3d.cpp b/modules/gpucalib3d/perf/perf_calib3d.cpp similarity index 100% rename from modules/gpu/perf/perf_calib3d.cpp rename to modules/gpucalib3d/perf/perf_calib3d.cpp diff --git a/modules/gpucalib3d/perf/perf_main.cpp b/modules/gpucalib3d/perf/perf_main.cpp new file mode 100644 index 0000000000..b35791cda2 --- /dev/null +++ b/modules/gpucalib3d/perf/perf_main.cpp @@ -0,0 +1,47 @@ +/*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 perf; + +CV_PERF_TEST_MAIN(gpuarithm, printCudaInfo()) diff --git a/modules/gpucalib3d/perf/perf_precomp.cpp b/modules/gpucalib3d/perf/perf_precomp.cpp new file mode 100644 index 0000000000..81f16e8f14 --- /dev/null +++ b/modules/gpucalib3d/perf/perf_precomp.cpp @@ -0,0 +1,43 @@ +/*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" diff --git a/modules/gpucalib3d/perf/perf_precomp.hpp b/modules/gpucalib3d/perf/perf_precomp.hpp new file mode 100644 index 0000000000..dc244a72a0 --- /dev/null +++ b/modules/gpucalib3d/perf/perf_precomp.hpp @@ -0,0 +1,65 @@ +/*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*/ + +#ifdef __GNUC__ +# pragma GCC diagnostic ignored "-Wmissing-declarations" +# if defined __clang__ || defined __APPLE__ +# pragma GCC diagnostic ignored "-Wmissing-prototypes" +# pragma GCC diagnostic ignored "-Wextra" +# endif +#endif + +#ifndef __OPENCV_PERF_PRECOMP_HPP__ +#define __OPENCV_PERF_PRECOMP_HPP__ + +#include "opencv2/ts.hpp" +#include "opencv2/ts/gpu_perf.hpp" + +#include "opencv2/gpucalib3d.hpp" + +#include "opencv2/calib3d.hpp" + +#ifdef GTEST_CREATE_SHARED_LIBRARY +#error no modules except ts should have GTEST_CREATE_SHARED_LIBRARY defined +#endif + +#endif diff --git a/modules/gpu/src/calib3d.cpp b/modules/gpucalib3d/src/calib3d.cpp similarity index 80% rename from modules/gpu/src/calib3d.cpp rename to modules/gpucalib3d/src/calib3d.cpp index abcc3423d8..135859094b 100644 --- a/modules/gpu/src/calib3d.cpp +++ b/modules/gpucalib3d/src/calib3d.cpp @@ -48,10 +48,10 @@ using namespace cv::gpu; #if !defined HAVE_CUDA || defined(CUDA_DISABLER) 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 @@ -150,7 +150,7 @@ namespace } // Computes rotation, translation pair for small subsets if the input data - class TransformHypothesesGenerator + class TransformHypothesesGenerator : public cv::ParallelLoopBody { public: TransformHypothesesGenerator(const Mat& object_, const Mat& image_, const Mat& dist_coef_, @@ -160,7 +160,7 @@ namespace num_points(num_points_), subset_size(subset_size_), rot_matrices(rot_matrices_), transl_vectors(transl_vectors_) {} - void operator()(const BlockedRange& range) const + void operator()(const Range& range) const { // Input data for generation of the current hypothesis std::vector subset_indices(subset_size); @@ -172,7 +172,7 @@ namespace Mat rot_mat(3, 3, CV_64F); Mat transl_vec(1, 3, CV_64F); - for (int iter = range.begin(); iter < range.end(); ++iter) + for (int iter = range.start; iter < range.end; ++iter) { selectRandom(subset_size, num_points, subset_indices); for (int i = 0; i < subset_size; ++i) @@ -238,7 +238,7 @@ void cv::gpu::solvePnPRansac(const Mat& object, const Mat& image, const Mat& cam // Generate set of hypotheses using small subsets of the input data TransformHypothesesGenerator body(object, image_normalized, empty_dist_coef, eye_camera_mat, num_points, subset_size, rot_matrices, transl_vectors); - parallel_for(BlockedRange(0, num_iters), body); + parallel_for_(Range(0, num_iters), body); // Compute scores (i.e. number of inliers) for each hypothesis GpuMat d_object(object); @@ -252,7 +252,7 @@ void cv::gpu::solvePnPRansac(const Mat& object, const Mat& image, const Mat& cam // Find the best hypothesis index Point best_idx; double best_score; - minMaxLoc(d_hypothesis_scores, NULL, &best_score, NULL, &best_idx); + gpu::minMaxLoc(d_hypothesis_scores, NULL, &best_score, NULL, &best_idx); int num_inliers = static_cast(best_score); // Extract the best hypothesis data @@ -289,6 +289,66 @@ void cv::gpu::solvePnPRansac(const Mat& object, const Mat& image, const Mat& cam } } -#endif +//////////////////////////////////////////////////////////////////////// +// 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/gpu/src/cuda/calib3d.cu b/modules/gpucalib3d/src/cuda/calib3d.cu similarity index 58% rename from modules/gpu/src/cuda/calib3d.cu rename to modules/gpucalib3d/src/cuda/calib3d.cu index 6085e716de..d1d59ce23d 100644 --- a/modules/gpu/src/cuda/calib3d.cu +++ b/modules/gpucalib3d/src/cuda/calib3d.cu @@ -187,6 +187,189 @@ 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/cuda/disp_bilateral_filter.cu b/modules/gpucalib3d/src/cuda/disp_bilateral_filter.cu similarity index 100% rename from modules/gpu/src/cuda/disp_bilateral_filter.cu rename to modules/gpucalib3d/src/cuda/disp_bilateral_filter.cu diff --git a/modules/gpu/src/cuda/stereobm.cu b/modules/gpucalib3d/src/cuda/stereobm.cu similarity index 100% rename from modules/gpu/src/cuda/stereobm.cu rename to modules/gpucalib3d/src/cuda/stereobm.cu diff --git a/modules/gpu/src/cuda/stereobp.cu b/modules/gpucalib3d/src/cuda/stereobp.cu similarity index 100% rename from modules/gpu/src/cuda/stereobp.cu rename to modules/gpucalib3d/src/cuda/stereobp.cu diff --git a/modules/gpu/src/cuda/stereocsbp.cu b/modules/gpucalib3d/src/cuda/stereocsbp.cu similarity index 100% rename from modules/gpu/src/cuda/stereocsbp.cu rename to modules/gpucalib3d/src/cuda/stereocsbp.cu diff --git a/modules/gpu/src/disparity_bilateral_filter.cpp b/modules/gpucalib3d/src/disparity_bilateral_filter.cpp similarity index 100% rename from modules/gpu/src/disparity_bilateral_filter.cpp rename to modules/gpucalib3d/src/disparity_bilateral_filter.cpp diff --git a/modules/gpucalib3d/src/precomp.cpp b/modules/gpucalib3d/src/precomp.cpp new file mode 100644 index 0000000000..3c01a2596d --- /dev/null +++ b/modules/gpucalib3d/src/precomp.cpp @@ -0,0 +1,43 @@ +/*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" diff --git a/modules/gpucalib3d/src/precomp.hpp b/modules/gpucalib3d/src/precomp.hpp new file mode 100644 index 0000000000..89396fd57d --- /dev/null +++ b/modules/gpucalib3d/src/precomp.hpp @@ -0,0 +1,56 @@ +/*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*/ + +#ifndef __OPENCV_PRECOMP_H__ +#define __OPENCV_PRECOMP_H__ + +#include + +#include "opencv2/gpucalib3d.hpp" +#include "opencv2/gpuarithm.hpp" + +#include "opencv2/calib3d.hpp" +#include "opencv2/imgproc.hpp" + +#include "opencv2/core/gpu_private.hpp" + +#endif /* __OPENCV_PRECOMP_H__ */ diff --git a/modules/gpu/src/stereobm.cpp b/modules/gpucalib3d/src/stereobm.cpp similarity index 100% rename from modules/gpu/src/stereobm.cpp rename to modules/gpucalib3d/src/stereobm.cpp diff --git a/modules/gpu/src/stereobp.cpp b/modules/gpucalib3d/src/stereobp.cpp similarity index 100% rename from modules/gpu/src/stereobp.cpp rename to modules/gpucalib3d/src/stereobp.cpp diff --git a/modules/gpu/src/stereocsbp.cpp b/modules/gpucalib3d/src/stereocsbp.cpp similarity index 100% rename from modules/gpu/src/stereocsbp.cpp rename to modules/gpucalib3d/src/stereocsbp.cpp diff --git a/modules/gpu/test/test_calib3d.cpp b/modules/gpucalib3d/test/test_calib3d.cpp similarity index 100% rename from modules/gpu/test/test_calib3d.cpp rename to modules/gpucalib3d/test/test_calib3d.cpp diff --git a/modules/gpucalib3d/test/test_main.cpp b/modules/gpucalib3d/test/test_main.cpp new file mode 100644 index 0000000000..eea3d7c008 --- /dev/null +++ b/modules/gpucalib3d/test/test_main.cpp @@ -0,0 +1,45 @@ +/*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" + +CV_GPU_TEST_MAIN("gpu") diff --git a/modules/gpucalib3d/test/test_precomp.cpp b/modules/gpucalib3d/test/test_precomp.cpp new file mode 100644 index 0000000000..0fb6521809 --- /dev/null +++ b/modules/gpucalib3d/test/test_precomp.cpp @@ -0,0 +1,43 @@ +/*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" diff --git a/modules/gpucalib3d/test/test_precomp.hpp b/modules/gpucalib3d/test/test_precomp.hpp new file mode 100644 index 0000000000..8c53f4786e --- /dev/null +++ b/modules/gpucalib3d/test/test_precomp.hpp @@ -0,0 +1,61 @@ +/*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*/ + +#ifdef __GNUC__ +# pragma GCC diagnostic ignored "-Wmissing-declarations" +# if defined __clang__ || defined __APPLE__ +# pragma GCC diagnostic ignored "-Wmissing-prototypes" +# pragma GCC diagnostic ignored "-Wextra" +# endif +#endif + +#ifndef __OPENCV_TEST_PRECOMP_HPP__ +#define __OPENCV_TEST_PRECOMP_HPP__ + +#include "opencv2/ts.hpp" +#include "opencv2/ts/gpu_test.hpp" + +#include "opencv2/gpucalib3d.hpp" + +#include "opencv2/calib3d.hpp" + +#endif diff --git a/modules/gpuimgproc/include/opencv2/gpuimgproc.hpp b/modules/gpuimgproc/include/opencv2/gpuimgproc.hpp index d602d0a131..a0b1e3094f 100644 --- a/modules/gpuimgproc/include/opencv2/gpuimgproc.hpp +++ b/modules/gpuimgproc/include/opencv2/gpuimgproc.hpp @@ -80,18 +80,6 @@ CV_EXPORTS void meanShiftProc(const GpuMat& src, GpuMat& dstr, GpuMat& dstsp, in CV_EXPORTS void meanShiftSegmentation(const GpuMat& src, Mat& dst, int sp, int sr, int minsize, TermCriteria criteria = TermCriteria(TermCriteria::MAX_ITER + TermCriteria::EPS, 5, 1)); -//! Does coloring of disparity image: [0..ndisp) -> [0..240, 1, 1] in HSV. -//! Supported types of input disparity: CV_8U, CV_16S. -//! Output disparity has CV_8UC4 type in BGRA format (alpha = 255). -CV_EXPORTS void drawColorDisp(const GpuMat& src_disp, GpuMat& dst_disp, int ndisp, Stream& stream = Stream::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. -//! Each element of this matrix will contain the 3D coordinates of the point (x,y,z,1), computed from the disparity map. -//! Q is the 4x4 perspective transformation matrix that can be obtained with cvStereoRectify. -CV_EXPORTS void reprojectImageTo3D(const GpuMat& disp, GpuMat& xyzw, const Mat& Q, int dst_cn = 4, Stream& stream = Stream::Null()); - //! converts image from one color space to another CV_EXPORTS void cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn = 0, Stream& stream = Stream::Null()); diff --git a/modules/gpuimgproc/src/cuda/imgproc.cu b/modules/gpuimgproc/src/cuda/imgproc.cu index c6dfbb417b..d2d0d0f3c5 100644 --- a/modules/gpuimgproc/src/cuda/imgproc.cu +++ b/modules/gpuimgproc/src/cuda/imgproc.cu @@ -183,187 +183,6 @@ namespace cv { namespace gpu { namespace cudev //cudaSafeCall( cudaUnbindTexture( tex_meanshift ) ); } - /////////////////////////////////// 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() ); - } - - /////////////////////////////////// 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); - /////////////////////////////////////////// Corner Harris ///////////////////////////////////////////////// texture harrisDxTex(0, cudaFilterModePoint, cudaAddressModeClamp); diff --git a/modules/gpuimgproc/src/imgproc.cpp b/modules/gpuimgproc/src/imgproc.cpp index c21a7b837d..dabf054b60 100644 --- a/modules/gpuimgproc/src/imgproc.cpp +++ b/modules/gpuimgproc/src/imgproc.cpp @@ -49,8 +49,6 @@ using namespace cv::gpu; void cv::gpu::meanShiftFiltering(const GpuMat&, GpuMat&, int, int, TermCriteria, Stream&) { throw_no_cuda(); } void cv::gpu::meanShiftProc(const GpuMat&, GpuMat&, GpuMat&, int, int, TermCriteria, Stream&) { throw_no_cuda(); } -void cv::gpu::drawColorDisp(const GpuMat&, GpuMat&, int, Stream&) { throw_no_cuda(); } -void cv::gpu::reprojectImageTo3D(const GpuMat&, GpuMat&, const Mat&, int, Stream&) { throw_no_cuda(); } void cv::gpu::buildWarpPlaneMaps(Size, Rect, const Mat&, const Mat&, const Mat&, float, GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); } void cv::gpu::buildWarpCylindricalMaps(Size, Rect, const Mat&, const Mat&, float, GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); } void cv::gpu::buildWarpSphericalMaps(Size, Rect, const Mat&, const Mat&, float, GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); } @@ -157,74 +155,6 @@ void cv::gpu::meanShiftProc(const GpuMat& src, GpuMat& dstr, GpuMat& dstsp, int meanShiftProc_gpu(src, dstr, dstsp, sp, sr, maxIter, eps, StreamAccessor::getStream(stream)); } -//////////////////////////////////////////////////////////////////////// -// drawColorDisp - -namespace cv { namespace gpu { namespace cudev -{ - namespace imgproc - { - 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::imgproc; - - 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)); -} - -//////////////////////////////////////////////////////////////////////// -// reprojectImageTo3D - -namespace cv { namespace gpu { namespace cudev -{ - namespace imgproc - { - 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::imgproc; - - 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)); -} - ////////////////////////////////////////////////////////////////////////////// // buildWarpPlaneMaps diff --git a/samples/cpp/CMakeLists.txt b/samples/cpp/CMakeLists.txt index 3cf2391e51..be87bebc76 100644 --- a/samples/cpp/CMakeLists.txt +++ b/samples/cpp/CMakeLists.txt @@ -22,6 +22,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/gpu/include") endif() diff --git a/samples/gpu/CMakeLists.txt b/samples/gpu/CMakeLists.txt index 9ee7fe8f2d..85360f51f1 100644 --- a/samples/gpu/CMakeLists.txt +++ b/samples/gpu/CMakeLists.txt @@ -2,7 +2,8 @@ SET(OPENCV_GPU_SAMPLES_REQUIRED_DEPS opencv_core opencv_flann opencv_imgproc ope opencv_ml opencv_video opencv_objdetect opencv_features2d opencv_calib3d opencv_legacy opencv_contrib opencv_gpu opencv_nonfree opencv_softcascade opencv_superres - opencv_gpucodec opencv_gpuarithm opencv_gpufilters opencv_gpunvidia opencv_gpuimgproc opencv_gpufeatures2d opencv_gpuvideo) + opencv_gpucodec opencv_gpuarithm opencv_gpufilters opencv_gpunvidia opencv_gpuimgproc opencv_gpufeatures2d opencv_gpuvideo + opencv_gpucalib3d) ocv_check_dependencies(${OPENCV_GPU_SAMPLES_REQUIRED_DEPS})