diff --git a/modules/gpu/CMakeLists.txt b/modules/gpu/CMakeLists.txt
index 296545add..95de6789f 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 3803efd19..bc3b9bdb1 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 524a32fc9..b3fea3fbf 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<int>* 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<GpuMat> 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 000000000..bb949c4d3
--- /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 000000000..5dffaa048
--- /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 000000000..3496d987b
--- /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<GpuMat> 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<int>* inliers=NULL);
+
+//! Reprojects disparity image to 3D space.
+//! Supports CV_8U and CV_16S types of input disparity.
+//! The output is a 3- or 4-channel floating-point matrix.
+//! 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 000000000..b35791cda
--- /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 000000000..81f16e8f1
--- /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 000000000..dc244a72a
--- /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 abcc3423d..135859094 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<int>*) { throw_no_cuda(); }
+void cv::gpu::reprojectImageTo3D(const GpuMat&, GpuMat&, const Mat&, int, Stream&) { throw_no_cuda(); }
+void cv::gpu::drawColorDisp(const GpuMat&, GpuMat&, int, Stream&) { throw_no_cuda(); }
#else
@@ -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<int> 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<int>(best_score);
// Extract the best hypothesis data
@@ -289,6 +289,66 @@ void cv::gpu::solvePnPRansac(const Mat& object, const Mat& image, const Mat& cam
}
}
+////////////////////////////////////////////////////////////////////////
+// reprojectImageTo3D
+
+namespace cv { namespace gpu { namespace cudev
+{
+ template <typename T, typename D>
+ void reprojectImageTo3D_gpu(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream);
+}}}
+
+void cv::gpu::reprojectImageTo3D(const GpuMat& disp, GpuMat& xyz, const Mat& Q, int dst_cn, Stream& stream)
+{
+ using namespace cv::gpu::cudev;
+
+ typedef void (*func_t)(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream);
+ static const func_t funcs[2][4] =
+ {
+ {reprojectImageTo3D_gpu<uchar, float3>, 0, 0, reprojectImageTo3D_gpu<short, float3>},
+ {reprojectImageTo3D_gpu<uchar, float4>, 0, 0, reprojectImageTo3D_gpu<short, float4>}
+ };
+
+ CV_Assert(disp.type() == CV_8U || disp.type() == CV_16S);
+ CV_Assert(Q.type() == CV_32F && Q.rows == 4 && Q.cols == 4 && Q.isContinuous());
+ CV_Assert(dst_cn == 3 || dst_cn == 4);
+
+ xyz.create(disp.size(), CV_MAKE_TYPE(CV_32F, dst_cn));
+
+ funcs[dst_cn == 4][disp.type()](disp, xyz, Q.ptr<float>(), StreamAccessor::getStream(stream));
+}
+
+////////////////////////////////////////////////////////////////////////
+// drawColorDisp
+
+namespace cv { namespace gpu { namespace cudev
+{
+ void drawColorDisp_gpu(const PtrStepSzb& src, const PtrStepSzb& dst, int ndisp, const cudaStream_t& stream);
+ void drawColorDisp_gpu(const PtrStepSz<short>& src, const PtrStepSzb& dst, int ndisp, const cudaStream_t& stream);
+}}}
+
+namespace
+{
+ template <typename T>
+ void drawColorDisp_caller(const GpuMat& src, GpuMat& dst, int ndisp, const cudaStream_t& stream)
+ {
+ using namespace ::cv::gpu::cudev;
+
+ dst.create(src.size(), CV_8UC4);
+
+ drawColorDisp_gpu((PtrStepSz<T>)src, dst, ndisp, stream);
+ }
+
+ typedef void (*drawColorDisp_caller_t)(const GpuMat& src, GpuMat& dst, int ndisp, const cudaStream_t& stream);
+
+ const drawColorDisp_caller_t drawColorDisp_callers[] = {drawColorDisp_caller<unsigned char>, 0, 0, drawColorDisp_caller<short>, 0, 0, 0, 0};
+}
+
+void cv::gpu::drawColorDisp(const GpuMat& src, GpuMat& dst, int ndisp, Stream& stream)
+{
+ CV_Assert(src.type() == CV_8U || src.type() == CV_16S);
+
+ drawColorDisp_callers[src.type()](src, dst, ndisp, StreamAccessor::getStream(stream));
+}
+
#endif
-
-
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 6085e716d..d1d59ce23 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 <typename T, typename D>
+ __global__ void reprojectImageTo3D(const PtrStepSz<T> disp, PtrStep<D> xyz)
+ {
+ const int x = blockIdx.x * blockDim.x + threadIdx.x;
+ const int y = blockIdx.y * blockDim.y + threadIdx.y;
+
+ if (y >= disp.rows || x >= disp.cols)
+ return;
+
+ const float qx = x * cq[ 0] + y * cq[ 1] + cq[ 3];
+ const float qy = x * cq[ 4] + y * cq[ 5] + cq[ 7];
+ const float qz = x * cq[ 8] + y * cq[ 9] + cq[11];
+ const float qw = x * cq[12] + y * cq[13] + cq[15];
+
+ const T d = disp(y, x);
+
+ const float iW = 1.f / (qw + cq[14] * d);
+
+ D v = VecTraits<D>::all(1.0f);
+ v.x = (qx + cq[2] * d) * iW;
+ v.y = (qy + cq[6] * d) * iW;
+ v.z = (qz + cq[10] * d) * iW;
+
+ xyz(y, x) = v;
+ }
+
+ template <typename T, typename D>
+ void reprojectImageTo3D_gpu(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream)
+ {
+ dim3 block(32, 8);
+ dim3 grid(divUp(disp.cols, block.x), divUp(disp.rows, block.y));
+
+ cudaSafeCall( cudaMemcpyToSymbol(cq, q, 16 * sizeof(float)) );
+
+ reprojectImageTo3D<T, D><<<grid, block, 0, stream>>>((PtrStepSz<T>)disp, (PtrStepSz<D>)xyz);
+ cudaSafeCall( cudaGetLastError() );
+
+ if (stream == 0)
+ cudaSafeCall( cudaDeviceSynchronize() );
+ }
+
+ template void reprojectImageTo3D_gpu<uchar, float3>(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream);
+ template void reprojectImageTo3D_gpu<uchar, float4>(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream);
+ template void reprojectImageTo3D_gpu<short, float3>(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream);
+ template void reprojectImageTo3D_gpu<short, float4>(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream);
+
+ /////////////////////////////////// drawColorDisp ///////////////////////////////////////////////
+
+ template <typename T>
+ __device__ unsigned int cvtPixel(T d, int ndisp, float S = 1, float V = 1)
+ {
+ unsigned int H = ((ndisp-d) * 240)/ndisp;
+
+ unsigned int hi = (H/60) % 6;
+ float f = H/60.f - H/60;
+ float p = V * (1 - S);
+ float q = V * (1 - f * S);
+ float t = V * (1 - (1 - f) * S);
+
+ float3 res;
+
+ if (hi == 0) //R = V, G = t, B = p
+ {
+ res.x = p;
+ res.y = t;
+ res.z = V;
+ }
+
+ if (hi == 1) // R = q, G = V, B = p
+ {
+ res.x = p;
+ res.y = V;
+ res.z = q;
+ }
+
+ if (hi == 2) // R = p, G = V, B = t
+ {
+ res.x = t;
+ res.y = V;
+ res.z = p;
+ }
+
+ if (hi == 3) // R = p, G = q, B = V
+ {
+ res.x = V;
+ res.y = q;
+ res.z = p;
+ }
+
+ if (hi == 4) // R = t, G = p, B = V
+ {
+ res.x = V;
+ res.y = p;
+ res.z = t;
+ }
+
+ if (hi == 5) // R = V, G = p, B = q
+ {
+ res.x = q;
+ res.y = p;
+ res.z = V;
+ }
+ const unsigned int b = (unsigned int)(::max(0.f, ::min(res.x, 1.f)) * 255.f);
+ const unsigned int g = (unsigned int)(::max(0.f, ::min(res.y, 1.f)) * 255.f);
+ const unsigned int r = (unsigned int)(::max(0.f, ::min(res.z, 1.f)) * 255.f);
+ const unsigned int a = 255U;
+
+ return (a << 24) + (r << 16) + (g << 8) + b;
+ }
+
+ __global__ void drawColorDisp(uchar* disp, size_t disp_step, uchar* out_image, size_t out_step, int width, int height, int ndisp)
+ {
+ const int x = (blockIdx.x * blockDim.x + threadIdx.x) << 2;
+ const int y = blockIdx.y * blockDim.y + threadIdx.y;
+
+ if(x < width && y < height)
+ {
+ uchar4 d4 = *(uchar4*)(disp + y * disp_step + x);
+
+ uint4 res;
+ res.x = cvtPixel(d4.x, ndisp);
+ res.y = cvtPixel(d4.y, ndisp);
+ res.z = cvtPixel(d4.z, ndisp);
+ res.w = cvtPixel(d4.w, ndisp);
+
+ uint4* line = (uint4*)(out_image + y * out_step);
+ line[x >> 2] = res;
+ }
+ }
+
+ __global__ void drawColorDisp(short* disp, size_t disp_step, uchar* out_image, size_t out_step, int width, int height, int ndisp)
+ {
+ const int x = (blockIdx.x * blockDim.x + threadIdx.x) << 1;
+ const int y = blockIdx.y * blockDim.y + threadIdx.y;
+
+ if(x < width && y < height)
+ {
+ short2 d2 = *(short2*)(disp + y * disp_step + x);
+
+ uint2 res;
+ res.x = cvtPixel(d2.x, ndisp);
+ res.y = cvtPixel(d2.y, ndisp);
+
+ uint2* line = (uint2*)(out_image + y * out_step);
+ line[x >> 1] = res;
+ }
+ }
+
+
+ void drawColorDisp_gpu(const PtrStepSzb& src, const PtrStepSzb& dst, int ndisp, const cudaStream_t& stream)
+ {
+ dim3 threads(16, 16, 1);
+ dim3 grid(1, 1, 1);
+ grid.x = divUp(src.cols, threads.x << 2);
+ grid.y = divUp(src.rows, threads.y);
+
+ drawColorDisp<<<grid, threads, 0, stream>>>(src.data, src.step, dst.data, dst.step, src.cols, src.rows, ndisp);
+ cudaSafeCall( cudaGetLastError() );
+
+ if (stream == 0)
+ cudaSafeCall( cudaDeviceSynchronize() );
+ }
+
+ void drawColorDisp_gpu(const PtrStepSz<short>& src, const PtrStepSzb& dst, int ndisp, const cudaStream_t& stream)
+ {
+ dim3 threads(32, 8, 1);
+ dim3 grid(1, 1, 1);
+ grid.x = divUp(src.cols, threads.x << 1);
+ grid.y = divUp(src.rows, threads.y);
+
+ drawColorDisp<<<grid, threads, 0, stream>>>(src.data, src.step / sizeof(short), dst.data, dst.step, src.cols, src.rows, ndisp);
+ cudaSafeCall( cudaGetLastError() );
+
+ if (stream == 0)
+ cudaSafeCall( cudaDeviceSynchronize() );
+ }
}}} // namespace cv { namespace gpu { namespace cudev
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 000000000..3c01a2596
--- /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 000000000..89396fd57
--- /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 <limits>
+
+#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 000000000..eea3d7c00
--- /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 000000000..0fb652180
--- /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 000000000..8c53f4786
--- /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 d602d0a13..a0b1e3094 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 c6dfbb417..d2d0d0f3c 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 <typename T>
- __device__ unsigned int cvtPixel(T d, int ndisp, float S = 1, float V = 1)
- {
- unsigned int H = ((ndisp-d) * 240)/ndisp;
-
- unsigned int hi = (H/60) % 6;
- float f = H/60.f - H/60;
- float p = V * (1 - S);
- float q = V * (1 - f * S);
- float t = V * (1 - (1 - f) * S);
-
- float3 res;
-
- if (hi == 0) //R = V, G = t, B = p
- {
- res.x = p;
- res.y = t;
- res.z = V;
- }
-
- if (hi == 1) // R = q, G = V, B = p
- {
- res.x = p;
- res.y = V;
- res.z = q;
- }
-
- if (hi == 2) // R = p, G = V, B = t
- {
- res.x = t;
- res.y = V;
- res.z = p;
- }
-
- if (hi == 3) // R = p, G = q, B = V
- {
- res.x = V;
- res.y = q;
- res.z = p;
- }
-
- if (hi == 4) // R = t, G = p, B = V
- {
- res.x = V;
- res.y = p;
- res.z = t;
- }
-
- if (hi == 5) // R = V, G = p, B = q
- {
- res.x = q;
- res.y = p;
- res.z = V;
- }
- const unsigned int b = (unsigned int)(::max(0.f, ::min(res.x, 1.f)) * 255.f);
- const unsigned int g = (unsigned int)(::max(0.f, ::min(res.y, 1.f)) * 255.f);
- const unsigned int r = (unsigned int)(::max(0.f, ::min(res.z, 1.f)) * 255.f);
- const unsigned int a = 255U;
-
- return (a << 24) + (r << 16) + (g << 8) + b;
- }
-
- __global__ void drawColorDisp(uchar* disp, size_t disp_step, uchar* out_image, size_t out_step, int width, int height, int ndisp)
- {
- const int x = (blockIdx.x * blockDim.x + threadIdx.x) << 2;
- const int y = blockIdx.y * blockDim.y + threadIdx.y;
-
- if(x < width && y < height)
- {
- uchar4 d4 = *(uchar4*)(disp + y * disp_step + x);
-
- uint4 res;
- res.x = cvtPixel(d4.x, ndisp);
- res.y = cvtPixel(d4.y, ndisp);
- res.z = cvtPixel(d4.z, ndisp);
- res.w = cvtPixel(d4.w, ndisp);
-
- uint4* line = (uint4*)(out_image + y * out_step);
- line[x >> 2] = res;
- }
- }
-
- __global__ void drawColorDisp(short* disp, size_t disp_step, uchar* out_image, size_t out_step, int width, int height, int ndisp)
- {
- const int x = (blockIdx.x * blockDim.x + threadIdx.x) << 1;
- const int y = blockIdx.y * blockDim.y + threadIdx.y;
-
- if(x < width && y < height)
- {
- short2 d2 = *(short2*)(disp + y * disp_step + x);
-
- uint2 res;
- res.x = cvtPixel(d2.x, ndisp);
- res.y = cvtPixel(d2.y, ndisp);
-
- uint2* line = (uint2*)(out_image + y * out_step);
- line[x >> 1] = res;
- }
- }
-
-
- void drawColorDisp_gpu(const PtrStepSzb& src, const PtrStepSzb& dst, int ndisp, const cudaStream_t& stream)
- {
- dim3 threads(16, 16, 1);
- dim3 grid(1, 1, 1);
- grid.x = divUp(src.cols, threads.x << 2);
- grid.y = divUp(src.rows, threads.y);
-
- drawColorDisp<<<grid, threads, 0, stream>>>(src.data, src.step, dst.data, dst.step, src.cols, src.rows, ndisp);
- cudaSafeCall( cudaGetLastError() );
-
- if (stream == 0)
- cudaSafeCall( cudaDeviceSynchronize() );
- }
-
- void drawColorDisp_gpu(const PtrStepSz<short>& src, const PtrStepSzb& dst, int ndisp, const cudaStream_t& stream)
- {
- dim3 threads(32, 8, 1);
- dim3 grid(1, 1, 1);
- grid.x = divUp(src.cols, threads.x << 1);
- grid.y = divUp(src.rows, threads.y);
-
- drawColorDisp<<<grid, threads, 0, stream>>>(src.data, src.step / sizeof(short), dst.data, dst.step, src.cols, src.rows, ndisp);
- cudaSafeCall( cudaGetLastError() );
-
- if (stream == 0)
- cudaSafeCall( cudaDeviceSynchronize() );
- }
-
- /////////////////////////////////// reprojectImageTo3D ///////////////////////////////////////////////
-
- __constant__ float cq[16];
-
- template <typename T, typename D>
- __global__ void reprojectImageTo3D(const PtrStepSz<T> disp, PtrStep<D> xyz)
- {
- const int x = blockIdx.x * blockDim.x + threadIdx.x;
- const int y = blockIdx.y * blockDim.y + threadIdx.y;
-
- if (y >= disp.rows || x >= disp.cols)
- return;
-
- const float qx = x * cq[ 0] + y * cq[ 1] + cq[ 3];
- const float qy = x * cq[ 4] + y * cq[ 5] + cq[ 7];
- const float qz = x * cq[ 8] + y * cq[ 9] + cq[11];
- const float qw = x * cq[12] + y * cq[13] + cq[15];
-
- const T d = disp(y, x);
-
- const float iW = 1.f / (qw + cq[14] * d);
-
- D v = VecTraits<D>::all(1.0f);
- v.x = (qx + cq[2] * d) * iW;
- v.y = (qy + cq[6] * d) * iW;
- v.z = (qz + cq[10] * d) * iW;
-
- xyz(y, x) = v;
- }
-
- template <typename T, typename D>
- void reprojectImageTo3D_gpu(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream)
- {
- dim3 block(32, 8);
- dim3 grid(divUp(disp.cols, block.x), divUp(disp.rows, block.y));
-
- cudaSafeCall( cudaMemcpyToSymbol(cq, q, 16 * sizeof(float)) );
-
- reprojectImageTo3D<T, D><<<grid, block, 0, stream>>>((PtrStepSz<T>)disp, (PtrStepSz<D>)xyz);
- cudaSafeCall( cudaGetLastError() );
-
- if (stream == 0)
- cudaSafeCall( cudaDeviceSynchronize() );
- }
-
- template void reprojectImageTo3D_gpu<uchar, float3>(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream);
- template void reprojectImageTo3D_gpu<uchar, float4>(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream);
- template void reprojectImageTo3D_gpu<short, float3>(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream);
- template void reprojectImageTo3D_gpu<short, float4>(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream);
-
/////////////////////////////////////////// Corner Harris /////////////////////////////////////////////////
texture<float, cudaTextureType2D, cudaReadModeElementType> harrisDxTex(0, cudaFilterModePoint, cudaAddressModeClamp);
diff --git a/modules/gpuimgproc/src/imgproc.cpp b/modules/gpuimgproc/src/imgproc.cpp
index c21a7b837..dabf054b6 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<short>& src, const PtrStepSzb& dst, int ndisp, const cudaStream_t& stream);
- }
-}}}
-
-namespace
-{
- template <typename T>
- void drawColorDisp_caller(const GpuMat& src, GpuMat& dst, int ndisp, const cudaStream_t& stream)
- {
- using namespace ::cv::gpu::cudev::imgproc;
-
- dst.create(src.size(), CV_8UC4);
-
- drawColorDisp_gpu((PtrStepSz<T>)src, dst, ndisp, stream);
- }
-
- typedef void (*drawColorDisp_caller_t)(const GpuMat& src, GpuMat& dst, int ndisp, const cudaStream_t& stream);
-
- const drawColorDisp_caller_t drawColorDisp_callers[] = {drawColorDisp_caller<unsigned char>, 0, 0, drawColorDisp_caller<short>, 0, 0, 0, 0};
-}
-
-void cv::gpu::drawColorDisp(const GpuMat& src, GpuMat& dst, int ndisp, Stream& stream)
-{
- CV_Assert(src.type() == CV_8U || src.type() == CV_16S);
-
- drawColorDisp_callers[src.type()](src, dst, ndisp, StreamAccessor::getStream(stream));
-}
-
-////////////////////////////////////////////////////////////////////////
-// reprojectImageTo3D
-
-namespace cv { namespace gpu { namespace cudev
-{
- namespace imgproc
- {
- template <typename T, typename D>
- void reprojectImageTo3D_gpu(const PtrStepSzb disp, PtrStepSzb xyz, const float* q, cudaStream_t stream);
- }
-}}}
-
-void cv::gpu::reprojectImageTo3D(const GpuMat& disp, GpuMat& xyz, const Mat& Q, int dst_cn, Stream& stream)
-{
- using namespace cv::gpu::cudev::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<uchar, float3>, 0, 0, reprojectImageTo3D_gpu<short, float3>},
- {reprojectImageTo3D_gpu<uchar, float4>, 0, 0, reprojectImageTo3D_gpu<short, float4>}
- };
-
- CV_Assert(disp.type() == CV_8U || disp.type() == CV_16S);
- CV_Assert(Q.type() == CV_32F && Q.rows == 4 && Q.cols == 4 && Q.isContinuous());
- CV_Assert(dst_cn == 3 || dst_cn == 4);
-
- xyz.create(disp.size(), CV_MAKE_TYPE(CV_32F, dst_cn));
-
- funcs[dst_cn == 4][disp.type()](disp, xyz, Q.ptr<float>(), StreamAccessor::getStream(stream));
-}
-
//////////////////////////////////////////////////////////////////////////////
// buildWarpPlaneMaps
diff --git a/samples/cpp/CMakeLists.txt b/samples/cpp/CMakeLists.txt
index 3cf2391e5..be87bebc7 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 9ee7fe8f2..85360f51f 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})