Attempt to allow 192 disparities

components/operators/src/disparity/fixstars_sgm.cpp

@@ -62,7 +62,7 @@ FixstarsSGM::FixstarsSGM(ftl::operators::Graph *g, ftl::Configurable* cfg) :
 	uniqueness_ = cfg->value("uniqueness", 0.95f);
 	P1_ = cfg->value("P1", 10);
 	P2_ = cfg->value("P2", 120);
-	max_disp_ = cfg->value("max_disp", 128);
+	max_disp_ = cfg->value("max_disp", 192);
 
 	if (uniqueness_ < 0.0 || uniqueness_ > 1.0) {
 		uniqueness_ = 1.0;
@@ -79,7 +79,7 @@ FixstarsSGM::FixstarsSGM(ftl::operators::Graph *g, ftl::Configurable* cfg) :
 		LOG(WARNING) << "Invalid value for P2, using value of P1 (" << P1_ << ")";
 	}
 
-	if (!(max_disp_ == 256 || max_disp_ == 128)) {
+	if (!(max_disp_ == 256 || max_disp_ == 128 || max_disp_ == 192)) {
 		max_disp_ = 256;
 		LOG(WARNING) << "Invalid value for max_disp, using default value (256)";
 	}

lib/libsgm/src/horizontal_path_aggregation.cu

@@ -239,6 +239,20 @@ template void enqueue_aggregate_left2right_path<128u>(
 	int min_disp,
 	cudaStream_t stream);
 
+template void enqueue_aggregate_left2right_path<192u>(
+	cost_type *dest,
+	const feature_type *left,
+	const feature_type *right,
+	int width,
+	int height,
+	unsigned int p1,
+	const uint8_t *p2,
+	int p2_pitch,
+	const uint8_t* w,
+	int w_pitch,
+	int min_disp,
+	cudaStream_t stream);
+
 template void enqueue_aggregate_left2right_path<256u>(
 	cost_type *dest,
 	const feature_type *left,
@@ -281,6 +295,20 @@ template void enqueue_aggregate_right2left_path<128u>(
 	int min_disp,
 	cudaStream_t stream);
 
+template void enqueue_aggregate_right2left_path<192u>(
+	cost_type *dest,
+	const feature_type *left,
+	const feature_type *right,
+	int width,
+	int height,
+	unsigned int p1,
+	const uint8_t *p2,
+	int p2_pitch,
+	const uint8_t* w,
+	int w_pitch,
+	int min_disp,
+	cudaStream_t stream);
+
 template void enqueue_aggregate_right2left_path<256u>(
 	cost_type *dest,
 	const feature_type *left,

lib/libsgm/src/oblique_path_aggregation.cu

@@ -243,6 +243,20 @@ template void enqueue_aggregate_upleft2downright_path<128u>(
 	int min_disp,
 	cudaStream_t stream);
 
+template void enqueue_aggregate_upleft2downright_path<192u>(
+	cost_type *dest,
+	const feature_type *left,
+	const feature_type *right,
+	int width,
+	int height,
+	unsigned int p1,
+	const uint8_t *p2,
+	int p2_pitch,
+	const uint8_t* w,
+	int w_pitch,
+	int min_disp,
+	cudaStream_t stream);
+
 template void enqueue_aggregate_upleft2downright_path<256u>(
 	cost_type *dest,
 	const feature_type *left,
@@ -285,6 +299,20 @@ template void enqueue_aggregate_upright2downleft_path<128u>(
 	int min_disp,
 	cudaStream_t stream);
 
+template void enqueue_aggregate_upright2downleft_path<192u>(
+	cost_type *dest,
+	const feature_type *left,
+	const feature_type *right,
+	int width,
+	int height,
+	unsigned int p1,
+	const uint8_t *p2,
+	int p2_pitch,
+	const uint8_t* w,
+	int w_pitch,
+	int min_disp,
+	cudaStream_t stream);
+
 template void enqueue_aggregate_upright2downleft_path<256u>(
 	cost_type *dest,
 	const feature_type *left,
@@ -327,6 +355,20 @@ template void enqueue_aggregate_downright2upleft_path<128u>(
 	int min_disp,
 	cudaStream_t stream);
 
+template void enqueue_aggregate_downright2upleft_path<192u>(
+	cost_type *dest,
+	const feature_type *left,
+	const feature_type *right,
+	int width,
+	int height,
+	unsigned int p1,
+	const uint8_t *p2,
+	int p2_pitch,
+	const uint8_t* w,
+	int w_pitch,
+	int min_disp,
+	cudaStream_t stream);
+
 template void enqueue_aggregate_downright2upleft_path<256u>(
 	cost_type *dest,
 	const feature_type *left,
@@ -369,6 +411,20 @@ template void enqueue_aggregate_downleft2upright_path<128u>(
 	int min_disp,
 	cudaStream_t stream);
 
+template void enqueue_aggregate_downleft2upright_path<192u>(
+	cost_type *dest,
+	const feature_type *left,
+	const feature_type *right,
+	int width,
+	int height,
+	unsigned int p1,
+	const uint8_t *p2,
+	int p2_pitch,
+	const uint8_t* w,
+	int w_pitch,
+	int min_disp,
+	cudaStream_t stream);
+
 template void enqueue_aggregate_downleft2upright_path<256u>(
 	cost_type *dest,
 	const feature_type *left,

lib/libsgm/src/path_aggregation.cu

@@ -101,6 +101,7 @@ void PathAggregation<MAX_DISPARITY>::enqueue(
 
 template class PathAggregation< 64>;
 template class PathAggregation<128>;
+template class PathAggregation<192>;
 template class PathAggregation<256>;
 
 }

lib/libsgm/src/path_aggregation_common.hpp

@@ -31,9 +31,9 @@ struct DynamicProgramming {
 	static_assert(
 		DP_BLOCK_SIZE >= 2,
 		"DP_BLOCK_SIZE must be greater than or equal to 2");
-	static_assert(
+	/*static_assert(
 		(SUBGROUP_SIZE & (SUBGROUP_SIZE - 1)) == 0,
-		"SUBGROUP_SIZE must be a power of 2");
+		"SUBGROUP_SIZE must be a power of 2");*/
 
 	uint32_t last_min;
 	uint32_t dp[DP_BLOCK_SIZE];

lib/libsgm/src/sgm.cu

@@ -158,6 +158,7 @@ void SemiGlobalMatching<T, MAX_DISPARITY>::enqueue(
 
 template class SemiGlobalMatching<uint8_t,   64>;
 template class SemiGlobalMatching<uint8_t,  128>;
+template class SemiGlobalMatching<uint8_t,  192>;
 template class SemiGlobalMatching<uint8_t,  256>;
 template class SemiGlobalMatching<uint16_t,  64>;
 template class SemiGlobalMatching<uint16_t, 128>;

lib/libsgm/src/stereo_sgm.cpp

@@ -63,6 +63,8 @@ namespace sgm {
 				sgm_engine = new SemiGlobalMatchingImpl<uint8_t, 64>();
 			else if (input_depth_bits_ == 8 && disparity_size_ == 128)
 				sgm_engine = new SemiGlobalMatchingImpl<uint8_t, 128>();
+			else if (input_depth_bits_ == 8 && disparity_size_ == 192)
+				sgm_engine = new SemiGlobalMatchingImpl<uint8_t, 192>();
 			else if (input_depth_bits_ == 8 && disparity_size_ == 256)
 				sgm_engine = new SemiGlobalMatchingImpl<uint8_t, 256>();
 			else if (input_depth_bits_ == 16 && disparity_size_ == 64)
@@ -133,9 +135,9 @@ namespace sgm {
 			width_ = height_ = input_depth_bits_ = output_depth_bits_ = disparity_size_ = 0;
 			throw std::logic_error("depth bits must be 8 or 16");
 		}
-		if (disparity_size_ != 64 && disparity_size_ != 128 && disparity_size_ != 256) {
+		if (disparity_size_ != 64 && disparity_size_ != 128 && disparity_size_ != 192 && disparity_size_ != 256) {
 			width_ = height_ = input_depth_bits_ = output_depth_bits_ = disparity_size_ = 0;
-			throw std::logic_error("disparity size must be 64, 128 or 256");
+			throw std::logic_error("disparity size must be 64, 128, 192 or 256");
 		}
 		if (param.subpixel && output_depth_bits != 16) {
 			width_ = height_ = input_depth_bits_ = output_depth_bits_ = disparity_size_ = 0;

lib/libsgm/src/utility.hpp

@@ -189,6 +189,14 @@ __device__ inline void load_uint16_vector<4u>(uint32_t *dest, const uint16_t *pt
 	dest[0] = uint16x4.x; dest[1] = uint16x4.y; dest[2] = uint16x4.z; dest[3] = uint16x4.w;
 }
 
+template <>
+__device__ inline void load_uint16_vector<6u>(uint32_t *dest, const uint16_t *ptr){
+	const auto uint32x3 = load_as<uint3>(ptr);
+	load_uint16_vector<2u>(dest + 0, reinterpret_cast<const uint16_t *>(&uint32x3.x));
+	load_uint16_vector<2u>(dest + 2, reinterpret_cast<const uint16_t *>(&uint32x3.y));
+	load_uint16_vector<2u>(dest + 4, reinterpret_cast<const uint16_t *>(&uint32x3.z));
+}
+
 template <>
 __device__ inline void load_uint16_vector<8u>(uint32_t *dest, const uint16_t *ptr){
 	const auto uint32x4 = load_as<uint4>(ptr);

lib/libsgm/src/vertical_path_aggregation.cu

@@ -206,6 +206,20 @@ template void enqueue_aggregate_up2down_path<256u>(
 	int min_disp,
 	cudaStream_t stream);
 
+template void enqueue_aggregate_up2down_path<192u>(
+	cost_type *dest,
+	const feature_type *left,
+	const feature_type *right,
+	int width,
+	int height,
+	unsigned int p1,
+	const uint8_t *p2,
+	int p2_pitch,
+	const uint8_t* w,
+	int w_pitch,
+	int min_disp,
+	cudaStream_t stream);
+
 template void enqueue_aggregate_down2up_path<64u>(
 	cost_type *dest,
 	const feature_type *left,
@@ -234,6 +248,20 @@ template void enqueue_aggregate_down2up_path<128u>(
 	int min_disp,
 	cudaStream_t stream);
 
+template void enqueue_aggregate_down2up_path<192u>(
+	cost_type *dest,
+	const feature_type *left,
+	const feature_type *right,
+	int width,
+	int height,
+	unsigned int p1,
+	const uint8_t *p2,
+	int p2_pitch,
+	const uint8_t* w,
+	int w_pitch,
+	int min_disp,
+	cudaStream_t stream);
+
 template void enqueue_aggregate_down2up_path<256u>(
 	cost_type *dest,
 	const feature_type *left,

lib/libsgm/src/winner_takes_all.cu

@@ -303,6 +303,7 @@ void WinnerTakesAll<MAX_DISPARITY>::enqueue(
 
 template class WinnerTakesAll< 64>;
 template class WinnerTakesAll<128>;
+template class WinnerTakesAll<192>;
 template class WinnerTakesAll<256>;
 
 }