\r
namespace cv { namespace gpu { namespace csbp \r
{ \r
- void load_constants(int ndisp, float max_data_term, float data_weight, float max_disc_term, float disc_single_jump, \r
- const DevMem2D& left, const DevMem2D& right, const DevMem2D& temp/*, const DevMem2D& temp2*/);\r
+ void load_constants(int ndisp, float max_data_term, float data_weight, float max_disc_term, float disc_single_jump, int min_disp_th,\r
+ const DevMem2D& left, const DevMem2D& right, const DevMem2D& temp);\r
\r
void init_data_cost(int rows, int cols, const DevMem2D& disp_selected_pyr, const DevMem2D& data_cost_selected,\r
size_t msg_step, int msg_type, int h, int w, int level, int nr_plane, int ndisp, int channels, \r
\r
namespace\r
{\r
- const float DEFAULT_MAX_DATA_TERM = 10.0f;\r
- const float DEFAULT_DATA_WEIGHT = 0.07f;\r
- const float DEFAULT_MAX_DISC_TERM = 1.7f;\r
- const float DEFAULT_DISC_SINGLE_JUMP = 1.0f;\r
+ const float DEFAULT_MAX_DATA_TERM = 30.0f;\r
+ const float DEFAULT_DATA_WEIGHT = 1.0f;\r
+ const float DEFAULT_MAX_DISC_TERM = 160.0f;\r
+ const float DEFAULT_DISC_SINGLE_JUMP = 10.0f;\r
}\r
\r
cv::gpu::StereoConstantSpaceBP::StereoConstantSpaceBP(int ndisp_, int iters_, int levels_, int nr_plane_,\r
int msg_type_)\r
: ndisp(ndisp_), iters(iters_), levels(levels_), nr_plane(nr_plane_), \r
max_data_term(DEFAULT_MAX_DATA_TERM), data_weight(DEFAULT_DATA_WEIGHT), \r
- max_disc_term(DEFAULT_MAX_DISC_TERM), disc_single_jump(DEFAULT_DISC_SINGLE_JUMP),\r
+ max_disc_term(DEFAULT_MAX_DISC_TERM), disc_single_jump(DEFAULT_DISC_SINGLE_JUMP), min_disp_th(0),\r
msg_type(msg_type_)\r
{ \r
}\r
\r
cv::gpu::StereoConstantSpaceBP::StereoConstantSpaceBP(int ndisp_, int iters_, int levels_, int nr_plane_,\r
float max_data_term_, float data_weight_, float max_disc_term_, float disc_single_jump_,\r
+ int min_disp_th_,\r
int msg_type_)\r
: ndisp(ndisp_), iters(iters_), levels(levels_), nr_plane(nr_plane_), \r
max_data_term(max_data_term_), data_weight(data_weight_), \r
- max_disc_term(max_disc_term_), disc_single_jump(disc_single_jump_),\r
+ max_disc_term(max_disc_term_), disc_single_jump(disc_single_jump_), min_disp_th(min_disp_th_),\r
msg_type(msg_type_)\r
{ \r
}\r
\r
static void stereo_csbp_gpu_operator(int& ndisp, int& iters, int& levels, int& nr_plane, \r
float& max_data_term, float& data_weight, float& max_disc_term, float& disc_single_jump,\r
+ int& min_disp_th,\r
int& msg_type,\r
GpuMat u[2], GpuMat d[2], GpuMat l[2], GpuMat r[2],\r
GpuMat disp_selected_pyr[2], GpuMat& data_cost, GpuMat& data_cost_selected,\r
////////////////////////////////////////////////////////////////////////////\r
// Compute\r
\r
- csbp::load_constants(ndisp, max_data_term, scale * data_weight, scale * max_disc_term, scale * disc_single_jump, \r
+ csbp::load_constants(ndisp, max_data_term, scale * data_weight, scale * max_disc_term, scale * disc_single_jump, min_disp_th,\r
left, right, temp);\r
\r
l[0] = zero;\r
\r
void cv::gpu::StereoConstantSpaceBP::operator()(const GpuMat& left, const GpuMat& right, GpuMat& disp)\r
{\r
- ::stereo_csbp_gpu_operator(ndisp, iters, levels, nr_plane, max_data_term, data_weight, max_disc_term, disc_single_jump, msg_type,\r
- u, d, l, r, disp_selected_pyr, data_cost, data_cost_selected, temp/*, temp2*/, out, left, right, disp, 0);\r
+ ::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,\r
+ u, d, l, r, disp_selected_pyr, data_cost, data_cost_selected, temp, out, left, right, disp, 0);\r
}\r
\r
void cv::gpu::StereoConstantSpaceBP::operator()(const GpuMat& left, const GpuMat& right, GpuMat& disp, const Stream& stream)\r
{\r
- ::stereo_csbp_gpu_operator(ndisp, iters, levels, nr_plane, max_data_term, data_weight, max_disc_term, disc_single_jump, msg_type,\r
- u, d, l, r, disp_selected_pyr, data_cost, data_cost_selected, temp/*, temp2*/, out, left, right, disp, \r
+ ::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,\r
+ u, d, l, r, disp_selected_pyr, data_cost, data_cost_selected, temp, out, left, right, disp, \r
StreamAccessor::getStream(stream));\r
}\r
\r
namespace csbp_kernels\r
{\r
__constant__ int cndisp;\r
- __constant__ int cth;\r
\r
__constant__ float cmax_data_term;\r
__constant__ float cdata_weight;\r
__constant__ float cmax_disc_term;\r
__constant__ float cdisc_single_jump;\r
+ \r
+ __constant__ int cth;\r
\r
__constant__ size_t cimg_step;\r
__constant__ size_t cmsg_step1;\r
\r
namespace cv { namespace gpu { namespace csbp \r
{\r
- void load_constants(int ndisp, float max_data_term, float data_weight, float max_disc_term, float disc_single_jump, \r
+ void load_constants(int ndisp, float max_data_term, float data_weight, float max_disc_term, float disc_single_jump, int min_disp_th,\r
const DevMem2D& left, const DevMem2D& right, const DevMem2D& temp)\r
{\r
- int th = (int)(ndisp * 0.2); \r
-\r
cudaSafeCall( cudaMemcpyToSymbol(csbp_kernels::cndisp, &ndisp, sizeof(int)) );\r
- cudaSafeCall( cudaMemcpyToSymbol(csbp_kernels::cth, &th, sizeof(int)) );\r
\r
cudaSafeCall( cudaMemcpyToSymbol(csbp_kernels::cmax_data_term, &max_data_term, sizeof(float)) );\r
cudaSafeCall( cudaMemcpyToSymbol(csbp_kernels::cdata_weight, &data_weight, sizeof(float)) );\r
cudaSafeCall( cudaMemcpyToSymbol(csbp_kernels::cmax_disc_term, &max_disc_term, sizeof(float)) );\r
cudaSafeCall( cudaMemcpyToSymbol(csbp_kernels::cdisc_single_jump, &disc_single_jump, sizeof(float)) );\r
\r
+ cudaSafeCall( cudaMemcpyToSymbol(csbp_kernels::cth, &min_disp_th, sizeof(int)) );\r
+ \r
cudaSafeCall( cudaMemcpyToSymbol(csbp_kernels::cimg_step, &left.step, sizeof(size_t)) );\r
\r
cudaSafeCall( cudaMemcpyToSymbol(csbp_kernels::cleft, &left.ptr, sizeof(left.ptr)) );\r