diff --git a/lib/libsgm/src/horizontal_path_aggregation.cu b/lib/libsgm/src/horizontal_path_aggregation.cu index acd53e1147746dc64b2d7fec7cf6da5e9f2c3883..11a94bc07081bbea4d29967b5390d9037c7319e7 100644 --- a/lib/libsgm/src/horizontal_path_aggregation.cu +++ b/lib/libsgm/src/horizontal_path_aggregation.cu @@ -88,7 +88,7 @@ __global__ void aggregate_horizontal_path_kernel( }else{ for(unsigned int i = 0; i < DP_BLOCKS_PER_THREAD; ++i){ for(unsigned int j = 0; j < DP_BLOCK_SIZE; ++j){ - const int x = static_cast<int>(width - min_disp - (j + dp_offset)); + const int x = static_cast<int>(width - (j + dp_offset)) - min_disp; if(0 <= x && x < static_cast<int>(width)){ right_buffer[i][j] = __ldg(&right[i * feature_step + x]); }else{ @@ -121,9 +121,9 @@ __global__ void aggregate_horizontal_path_kernel( #else right_buffer[j][0] = __shfl_up(t, 1, SUBGROUP_SIZE); #endif - if(lane_id == 0){ + if(lane_id == 0 && x >= min_disp + dp_offset){ right_buffer[j][0] = - __ldg(&right[j * feature_step + x - dp_offset]); + __ldg(&right[j * feature_step + x - min_disp - dp_offset]); } }else{ const feature_type t = right_buffer[j][0]; @@ -137,9 +137,9 @@ __global__ void aggregate_horizontal_path_kernel( right_buffer[j][DP_BLOCK_SIZE - 1] = __shfl_down(t, 1, SUBGROUP_SIZE); #endif if(lane_id + 1 == SUBGROUP_SIZE){ - if(x >= dp_offset + DP_BLOCK_SIZE - 1){ + if(x >= min_disp + dp_offset + DP_BLOCK_SIZE - 1){ right_buffer[j][DP_BLOCK_SIZE - 1] = - __ldg(&right[j * feature_step + x - (dp_offset + DP_BLOCK_SIZE - 1)]); + __ldg(&right[j * feature_step + x - min_disp - (dp_offset + DP_BLOCK_SIZE - 1)]); }else{ right_buffer[j][DP_BLOCK_SIZE - 1] = 0; }