fixstars / libSGM

Stereo Semi Global Matching by cuda
Apache License 2.0
597 stars 187 forks source link

Changing NUM_PATH to another value #35

Closed P-Light closed 5 years ago

P-Light commented 5 years ago

I am trying to launch SGM-algorithm with default NUM_PATH and it works fine. But I have noticed that buffer_size in path_aggregation.cu (53 string) depends on NUM_PATH, MAX_DISPARITY. So I want to change NUM_PATH value from 8 to 4 for reduction occupied device memory. After changing this variable program is crashing: segmentation fault. There are ways for changing NUM_PATHS and after SGM will works fine? Thank you very much!

Best Regards,

sotsuka-fixstars commented 5 years ago

@P-Light Thank you for using libSGM!

So I want to change NUM_PATH value from 8 to 4 for reduction occupied device memory. After changing this variable program is crashing: segmentation fault.

I think that you forgot changing NUM_PATHS in winner_takes_all.cu. A following patch I checked works fine. Please try it.

diff --git a/src/path_aggregation.cu b/src/path_aggregation.cu
index de713de..f7643f5 100644
--- a/src/path_aggregation.cu
+++ b/src/path_aggregation.cu
@@ -17,7 +17,6 @@ limitations under the License.
 #include "path_aggregation.hpp"
 #include "vertical_path_aggregation.hpp"
 #include "horizontal_path_aggregation.hpp"
-#include "oblique_path_aggregation.hpp"

 namespace sgm {

@@ -68,18 +67,6 @@ void PathAggregation<MAX_DISPARITY>::enqueue(
    path_aggregation::enqueue_aggregate_right2left_path<MAX_DISPARITY>(
        m_cost_buffer.data() + 3 * buffer_step,
        left, right, width, height, p1, p2, m_streams[3]);
-   path_aggregation::enqueue_aggregate_upleft2downright_path<MAX_DISPARITY>(
-       m_cost_buffer.data() + 4 * buffer_step,
-       left, right, width, height, p1, p2, m_streams[4]);
-   path_aggregation::enqueue_aggregate_upright2downleft_path<MAX_DISPARITY>(
-       m_cost_buffer.data() + 5 * buffer_step,
-       left, right, width, height, p1, p2, m_streams[5]);
-   path_aggregation::enqueue_aggregate_downright2upleft_path<MAX_DISPARITY>(
-       m_cost_buffer.data() + 6 * buffer_step,
-       left, right, width, height, p1, p2, m_streams[6]);
-   path_aggregation::enqueue_aggregate_downleft2upright_path<MAX_DISPARITY>(
-       m_cost_buffer.data() + 7 * buffer_step,
-       left, right, width, height, p1, p2, m_streams[7]);
    for(unsigned int i = 0; i < NUM_PATHS; ++i){
        cudaEventRecord(m_events[i], m_streams[i]);
        cudaStreamWaitEvent(stream, m_events[i], 0);
diff --git a/src/path_aggregation.hpp b/src/path_aggregation.hpp
index 3d2b35d..db119b6 100644
--- a/src/path_aggregation.hpp
+++ b/src/path_aggregation.hpp
@@ -26,7 +26,7 @@ template <size_t MAX_DISPARITY>
 class PathAggregation {

 private:
-   static const unsigned int NUM_PATHS = 8;
+   static const unsigned int NUM_PATHS = 4;

    DeviceBuffer<cost_type> m_cost_buffer;
    cudaStream_t m_streams[NUM_PATHS];
diff --git a/src/winner_takes_all.cu b/src/winner_takes_all.cu
index 6269b85..b6dbb67 100644
--- a/src/winner_takes_all.cu
+++ b/src/winner_takes_all.cu
@@ -23,7 +23,7 @@ namespace sgm {

 namespace {

-static constexpr unsigned int NUM_PATHS = 8u;
+static constexpr unsigned int NUM_PATHS = 4u;

 static constexpr unsigned int WARPS_PER_BLOCK = 8u;
 static constexpr unsigned int BLOCK_SIZE = WARPS_PER_BLOCK * WARP_SIZE;

Sincerely,

P-Light commented 5 years ago

@sotsuka-fixstars thank you! It works)

Thank you for using libSGM!

libSGM is awesome implementation of Semi-Global Block Matching. I have the pleasure to use libSGM. Best Regards