added minimum disparity threshold parameter to StereoConstantSpaceBP
authorVladislav Vinogradov <no@email>
Fri, 13 Aug 2010 11:17:51 +0000 (11:17 +0000)
committerVladislav Vinogradov <no@email>
Fri, 13 Aug 2010 11:17:51 +0000 (11:17 +0000)
modules/gpu/include/opencv2/gpu/gpu.hpp
modules/gpu/src/constantspacebp_gpu.cpp
modules/gpu/src/cuda/constantspacebp.cu

index 5734709..9fa9087 100644 (file)
@@ -441,9 +441,10 @@ namespace cv
 \r
             //! the full constructor taking the number of disparities, number of BP iterations on each level,\r
             //! number of levels, number of active disparity on the first level, truncation of data cost, data weight, \r
-            //! truncation of discontinuity cost and discontinuity single jump\r
+            //! truncation of discontinuity cost, discontinuity single jump and minimum disparity threshold\r
             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 = 0,\r
                                   int msg_type = CV_32F);\r
 \r
             //! the stereo correspondence operator. Finds the disparity for the specified rectified stereo pair,\r
@@ -465,6 +466,8 @@ namespace cv
             float max_disc_term; \r
             float disc_single_jump;\r
 \r
+            int min_disp_th;\r
+\r
             int msg_type;\r
         private:\r
             GpuMat u[2], d[2], l[2], r[2];\r
index 2656ed9..967f944 100644 (file)
@@ -58,8 +58,8 @@ void cv::gpu::StereoConstantSpaceBP::operator()(const GpuMat&, const GpuMat&, Gp
 \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
@@ -86,33 +86,35 @@ namespace cv { namespace gpu { namespace csbp
 \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
@@ -195,7 +197,7 @@ static void stereo_csbp_gpu_operator(int& ndisp, int& iters, int& levels, int& n
     ////////////////////////////////////////////////////////////////////////////\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
@@ -257,14 +259,14 @@ static void stereo_csbp_gpu_operator(int& ndisp, int& iters, int& levels, int& n
 \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
index 2bdcb4a..e9bc95d 100644 (file)
@@ -77,12 +77,13 @@ struct TypeLimits<float>
 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
@@ -97,19 +98,18 @@ namespace csbp_kernels
 \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