diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 573470978b1877cf83d3fbd77d31248d1109ef7f..9fa90878954215b233c22f1a0c087709f5b09ceb 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -441,9 +441,10 @@ namespace cv //! 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 and discontinuity single jump + //! 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, @@ -465,6 +466,8 @@ namespace cv float max_disc_term; float disc_single_jump; + int min_disp_th; + int msg_type; private: GpuMat u[2], d[2], l[2], r[2]; diff --git a/modules/gpu/src/constantspacebp_gpu.cpp b/modules/gpu/src/constantspacebp_gpu.cpp index 2656ed9f68db1001ffe62951689acd30a132a66d..967f94407a59e4db605a6522bc2d0e92f282f52a 100644 --- a/modules/gpu/src/constantspacebp_gpu.cpp +++ b/modules/gpu/src/constantspacebp_gpu.cpp @@ -58,8 +58,8 @@ void cv::gpu::StereoConstantSpaceBP::operator()(const GpuMat&, const GpuMat&, Gp namespace cv { namespace gpu { namespace csbp { - void load_constants(int ndisp, float max_data_term, float data_weight, float max_disc_term, float disc_single_jump, - const DevMem2D& left, const DevMem2D& right, const DevMem2D& temp/*, const DevMem2D& temp2*/); + void load_constants(int ndisp, float max_data_term, float data_weight, float max_disc_term, float disc_single_jump, int min_disp_th, + const DevMem2D& left, const DevMem2D& right, const DevMem2D& temp); void init_data_cost(int rows, int cols, const DevMem2D& disp_selected_pyr, const DevMem2D& data_cost_selected, size_t msg_step, int msg_type, int h, int w, int level, int nr_plane, int ndisp, int channels, @@ -86,33 +86,35 @@ namespace cv { namespace gpu { namespace csbp namespace { - const float DEFAULT_MAX_DATA_TERM = 10.0f; - const float DEFAULT_DATA_WEIGHT = 0.07f; - const float DEFAULT_MAX_DISC_TERM = 1.7f; - const float DEFAULT_DISC_SINGLE_JUMP = 1.0f; + const float DEFAULT_MAX_DATA_TERM = 30.0f; + const float DEFAULT_DATA_WEIGHT = 1.0f; + const float DEFAULT_MAX_DISC_TERM = 160.0f; + const float DEFAULT_DISC_SINGLE_JUMP = 10.0f; } cv::gpu::StereoConstantSpaceBP::StereoConstantSpaceBP(int ndisp_, int iters_, int levels_, int nr_plane_, int msg_type_) : ndisp(ndisp_), iters(iters_), levels(levels_), nr_plane(nr_plane_), max_data_term(DEFAULT_MAX_DATA_TERM), data_weight(DEFAULT_DATA_WEIGHT), - max_disc_term(DEFAULT_MAX_DISC_TERM), disc_single_jump(DEFAULT_DISC_SINGLE_JUMP), + max_disc_term(DEFAULT_MAX_DISC_TERM), disc_single_jump(DEFAULT_DISC_SINGLE_JUMP), min_disp_th(0), msg_type(msg_type_) { } cv::gpu::StereoConstantSpaceBP::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_, int msg_type_) : ndisp(ndisp_), iters(iters_), levels(levels_), nr_plane(nr_plane_), max_data_term(max_data_term_), data_weight(data_weight_), - max_disc_term(max_disc_term_), disc_single_jump(disc_single_jump_), + max_disc_term(max_disc_term_), disc_single_jump(disc_single_jump_), min_disp_th(min_disp_th_), msg_type(msg_type_) { } static void stereo_csbp_gpu_operator(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, GpuMat u[2], GpuMat d[2], GpuMat l[2], GpuMat r[2], GpuMat disp_selected_pyr[2], GpuMat& data_cost, GpuMat& data_cost_selected, @@ -195,7 +197,7 @@ static void stereo_csbp_gpu_operator(int& ndisp, int& iters, int& levels, int& n //////////////////////////////////////////////////////////////////////////// // Compute - csbp::load_constants(ndisp, max_data_term, scale * data_weight, scale * max_disc_term, scale * disc_single_jump, + csbp::load_constants(ndisp, max_data_term, scale * data_weight, scale * max_disc_term, scale * disc_single_jump, min_disp_th, left, right, temp); l[0] = zero; @@ -257,14 +259,14 @@ static void stereo_csbp_gpu_operator(int& ndisp, int& iters, int& levels, int& n void cv::gpu::StereoConstantSpaceBP::operator()(const GpuMat& left, const GpuMat& right, GpuMat& disp) { - ::stereo_csbp_gpu_operator(ndisp, iters, levels, nr_plane, max_data_term, data_weight, max_disc_term, disc_single_jump, msg_type, - u, d, l, r, disp_selected_pyr, data_cost, data_cost_selected, temp/*, temp2*/, out, left, right, disp, 0); + ::stereo_csbp_gpu_operator(ndisp, iters, levels, nr_plane, max_data_term, data_weight, max_disc_term, disc_single_jump, min_disp_th, msg_type, + u, d, l, r, disp_selected_pyr, data_cost, data_cost_selected, temp, out, left, right, disp, 0); } void cv::gpu::StereoConstantSpaceBP::operator()(const GpuMat& left, const GpuMat& right, GpuMat& disp, const Stream& stream) { - ::stereo_csbp_gpu_operator(ndisp, iters, levels, nr_plane, max_data_term, data_weight, max_disc_term, disc_single_jump, msg_type, - u, d, l, r, disp_selected_pyr, data_cost, data_cost_selected, temp/*, temp2*/, out, left, right, disp, + ::stereo_csbp_gpu_operator(ndisp, iters, levels, nr_plane, max_data_term, data_weight, max_disc_term, disc_single_jump, min_disp_th, msg_type, + u, d, l, r, disp_selected_pyr, data_cost, data_cost_selected, temp, out, left, right, disp, StreamAccessor::getStream(stream)); } diff --git a/modules/gpu/src/cuda/constantspacebp.cu b/modules/gpu/src/cuda/constantspacebp.cu index 2bdcb4a0f1ae745340d03ac8a636d82c388507f6..e9bc95dc66e64c027fb649332689ac4a6de8aa52 100644 --- a/modules/gpu/src/cuda/constantspacebp.cu +++ b/modules/gpu/src/cuda/constantspacebp.cu @@ -77,12 +77,13 @@ struct TypeLimits namespace csbp_kernels { __constant__ int cndisp; - __constant__ int cth; __constant__ float cmax_data_term; __constant__ float cdata_weight; __constant__ float cmax_disc_term; __constant__ float cdisc_single_jump; + + __constant__ int cth; __constant__ size_t cimg_step; __constant__ size_t cmsg_step1; @@ -97,19 +98,18 @@ namespace csbp_kernels namespace cv { namespace gpu { namespace csbp { - void load_constants(int ndisp, float max_data_term, float data_weight, float max_disc_term, float disc_single_jump, + void load_constants(int ndisp, float max_data_term, float data_weight, float max_disc_term, float disc_single_jump, int min_disp_th, const DevMem2D& left, const DevMem2D& right, const DevMem2D& temp) { - int th = (int)(ndisp * 0.2); - cudaSafeCall( cudaMemcpyToSymbol(csbp_kernels::cndisp, &ndisp, sizeof(int)) ); - cudaSafeCall( cudaMemcpyToSymbol(csbp_kernels::cth, &th, sizeof(int)) ); cudaSafeCall( cudaMemcpyToSymbol(csbp_kernels::cmax_data_term, &max_data_term, sizeof(float)) ); cudaSafeCall( cudaMemcpyToSymbol(csbp_kernels::cdata_weight, &data_weight, sizeof(float)) ); cudaSafeCall( cudaMemcpyToSymbol(csbp_kernels::cmax_disc_term, &max_disc_term, sizeof(float)) ); cudaSafeCall( cudaMemcpyToSymbol(csbp_kernels::cdisc_single_jump, &disc_single_jump, sizeof(float)) ); + cudaSafeCall( cudaMemcpyToSymbol(csbp_kernels::cth, &min_disp_th, sizeof(int)) ); + cudaSafeCall( cudaMemcpyToSymbol(csbp_kernels::cimg_step, &left.step, sizeof(size_t)) ); cudaSafeCall( cudaMemcpyToSymbol(csbp_kernels::cleft, &left.ptr, sizeof(left.ptr)) );