Merge branch '49-32bit_input_support' into 'master'

入力画像の24ビット対応 See merge request tech/adaskit/libSGM!64
fixstars · Aug 1, 2023 · 339b6bb · 339b6bb
2 parents e23afca + 80cb3a4
commit 339b6bb
Show file tree

Hide file tree

Showing 8 changed files with 138 additions and 13 deletions.
diff --git a/include/libsgm.h b/include/libsgm.h
@@ -117,7 +117,7 @@ class StereoSGM
 	* @param width Processed image's width.
 	* @param height Processed image's height.
 	* @param disparity_size It must be 64, 128 or 256.
-	* @param input_depth_bits Processed image's bits per pixel. It must be 8 or 16.
+	* @param input_depth_bits Processed image's bits per pixel. It must be 8, 16 or 32.
 	* @param output_depth_bits Disparity image's bits per pixel. It must be 8 or 16.
 	* @param inout_type Specify input/output pointer type. See sgm::EXECUTE_TYPE.
 	* @attention
@@ -130,7 +130,7 @@ class StereoSGM
 	* @param width Processed image's width.
 	* @param height Processed image's height.
 	* @param disparity_size It must be 64, 128 or 256.
-	* @param input_depth_bits Processed image's bits per pixel. It must be 8 or 16.
+	* @param input_depth_bits Processed image's bits per pixel. It must be 8, 16 or 32.
 	* @param output_depth_bits Disparity image's bits per pixel. It must be 8 or 16.
 	* @param src_pitch Source image's pitch (pixels).
 	* @param dst_pitch Destination image's pitch (pixels).

diff --git a/include/libsgm_wrapper.h b/include/libsgm_wrapper.h
@@ -48,7 +48,7 @@ class LibSGMWrapper
 
 	/**
 	* Execute stereo semi global matching via wrapper class.
-	* @param I1        Input left image.  Image's type is must be CV_8U or CV_16U
+	* @param I1        Input left image.  Image's type is must be CV_8U, CV_16U or CV_32S
 	* @param I2        Input right image.  Image's size and type must be same with I1.
 	* @param disparity Output image.  Its memory will be allocated automatically dependent on input image size.
 	* @attention
@@ -59,7 +59,7 @@ class LibSGMWrapper
 
 	/**
 	* Execute stereo semi global matching via wrapper class.
-	* @param I1        Input left image.  Image's type is must be CV_8U or CV_16U.
+	* @param I1        Input left image.  Image's type is must be CV_8U, CV_16U or CV_32S.
 	* @param I2        Input right image.  Image's size and type must be same with I1.
 	* @param disparity Output image.  Its memory will be allocated automatically dependent on input image size.
 	* @attention

diff --git a/src/census_transform.cu b/src/census_transform.cu
@@ -191,14 +191,18 @@ void census_transform(const DeviceImage& src, DeviceImage& dst, CensusType type)
 	if (type == CensusType::CENSUS_9x7) {
 		if (src.type == SGM_8U)
 			census_transform_kernel<<<gdim, bdim>>>(dst.ptr<uint64_t>(), src.ptr<uint8_t>(), w, h, src.step);
-		else
+		else if (src.type == SGM_16U)
 			census_transform_kernel<<<gdim, bdim>>>(dst.ptr<uint64_t>(), src.ptr<uint16_t>(), w, h, src.step);
+		else
+			census_transform_kernel<<<gdim, bdim>>>(dst.ptr<uint64_t>(), src.ptr<uint32_t>(), w, h, src.step);
 	}
 	else if (type == CensusType::SYMMETRIC_CENSUS_9x7) {
 		if (src.type == SGM_8U)
 			symmetric_census_kernel<<<gdim, bdim>>>(dst.ptr<uint32_t>(), src.ptr<uint8_t>(), w, h, src.step);
-		else
+		else if (src.type == SGM_16U)
 			symmetric_census_kernel<<<gdim, bdim>>>(dst.ptr<uint32_t>(), src.ptr<uint16_t>(), w, h, src.step);
+		else
+			symmetric_census_kernel<<<gdim, bdim>>>(dst.ptr<uint32_t>(), src.ptr<uint32_t>(), w, h, src.step);
 	}
 
 	CUDA_CHECK(cudaGetLastError());

diff --git a/src/check_consistency.cu b/src/check_consistency.cu
@@ -69,11 +69,16 @@ void check_consistency(DeviceImage& dispL, const DeviceImage& dispR, const Devic
 		check_consistency_kernel<SRC_T><<<grid, block>>>(dispL.ptr<uint16_t>(), dispR.ptr<uint16_t>(),
 			srcL.ptr<SRC_T>(), w, h, srcL.step, dispL.step, subpixel, LR_max_diff);
 	}
-	else {
+	else if (srcL.type == SGM_16U) {
 		using SRC_T = uint16_t;
 		check_consistency_kernel<SRC_T><<<grid, block>>>(dispL.ptr<uint16_t>(), dispR.ptr<uint16_t>(),
 			srcL.ptr<SRC_T>(), w, h, srcL.step, dispL.step, subpixel, LR_max_diff);
 	}
+	else {
+		using SRC_T = uint32_t;
+		check_consistency_kernel<SRC_T><<<grid, block>>>(dispL.ptr<uint16_t>(), dispR.ptr<uint16_t>(),
+			srcL.ptr<SRC_T>(), w, h, srcL.step, dispL.step, subpixel, LR_max_diff);
+	}
 
 	CUDA_CHECK(cudaGetLastError());
 }

diff --git a/src/libsgm.cpp b/src/libsgm.cpp
@@ -65,13 +65,13 @@ class StereoSGM::Impl
 		param_(param)
 	{
 		// check values
-		SGM_ASSERT(src_depth == 8 || src_depth == 16, "src depth bits must be 8 or 16");
+		SGM_ASSERT(src_depth == 8 || src_depth == 16 || src_depth == 32, "src depth bits must be 8, 16 or 32");
 		SGM_ASSERT(dst_depth == 8 || dst_depth == 16, "dst depth bits must be 8 or 16");
 		SGM_ASSERT(disparity_size == 64 || disparity_size == 128 || disparity_size == 256, "disparity size must be 64 or 128 or 256");
 		SGM_ASSERT(has_enough_depth(dst_depth, disparity_size, param_.min_disp, param_.subpixel),
 			"output depth bits must be sufficient for representing output value");
 
-		src_type_ = src_depth == 8 ? SGM_8U : SGM_16U;
+		src_type_ = src_depth == 8 ? SGM_8U : src_depth == 16 ? SGM_16U : SGM_32U;
 		dst_type_ = dst_depth == 8 ? SGM_8U : SGM_16U;
 
 		is_src_devptr_ = (inout_type & 0x01) > 0;

diff --git a/src/libsgm_wrapper.cpp b/src/libsgm_wrapper.cpp
@@ -74,7 +74,7 @@ struct LibSGMWrapper::Creator
 	Creator(const cv::cuda::GpuMat& src, const cv::cuda::GpuMat& dst)
 	{
 		const int depth = src.depth();
-		CV_Assert(depth == CV_8U || depth == CV_16U);
+		CV_Assert(depth == CV_8U || depth == CV_16U || depth == CV_32S);
 		width = src.cols;
 		height = src.rows;
 		src_pitch = static_cast<int>(src.step1());
@@ -87,7 +87,7 @@ struct LibSGMWrapper::Creator
 	Creator(const cv::Mat& src, const cv::Mat& dst)
 	{
 		const int depth = src.depth();
-		CV_Assert(depth == CV_8U || depth == CV_16U);
+		CV_Assert(depth == CV_8U || depth == CV_16U || depth == CV_32S);
 		width = src.cols;
 		height = src.rows;
 		src_pitch = static_cast<int>(src.step1());
@@ -108,7 +108,7 @@ void LibSGMWrapper::execute(const cv::cuda::GpuMat& I1, const cv::cuda::GpuMat&
 	CV_Assert(size == I2.size());
 	CV_Assert(I1.type() == I2.type());
 	const int depth = I1.depth();
-	CV_Assert(depth == CV_8U || depth == CV_16U);
+	CV_Assert(depth == CV_8U || depth == CV_16U || depth == CV_32S);
 	if (disparity.size() != size || disparity.depth() != CV_16S) {
 		disparity.create(size, CV_16S);
 	}
@@ -127,7 +127,7 @@ void LibSGMWrapper::execute(const cv::Mat& I1, const cv::Mat& I2, cv::Mat& dispa
 	CV_Assert(size == I2.size());
 	CV_Assert(I1.type() == I2.type());
 	const int depth = I1.depth();
-	CV_Assert(depth == CV_8U || depth == CV_16U);
+	CV_Assert(depth == CV_8U || depth == CV_16U || depth == CV_32S);
 	if (disparity.size() != size || disparity.depth() != CV_16S) {
 		disparity.create(size, CV_16S);
 	}

diff --git a/test/census_transform_test.cpp b/test/census_transform_test.cpp
@@ -64,13 +64,17 @@ void census_transform(const HostImage& src, HostImage& dst, CensusType type)
 			census_transform_9x7_<uint8_t>(src, dst);
 		if (src.type == SGM_16U)
 			census_transform_9x7_<uint16_t>(src, dst);
+		if (src.type == SGM_32U)
+			census_transform_9x7_<uint32_t>(src, dst);
 	}
 	if (type == CensusType::SYMMETRIC_CENSUS_9x7) {
 		dst.create(src.rows, src.cols, SGM_32U);
 		if (src.type == SGM_8U)
 			symmetric_census_9x7_<uint8_t>(src, dst);
 		if (src.type == SGM_16U)
 			symmetric_census_9x7_<uint16_t>(src, dst);
+		if (src.type == SGM_32U)
+			symmetric_census_9x7_<uint32_t>(src, dst);
 	}
 }
 
@@ -126,6 +130,31 @@ TEST(CensusTransformTest, RandomU16)
 	EXPECT_TRUE(equals(h_dst, d_dst));
 }
 
+TEST(CensusTransformTest, RandomU32)
+{
+	using namespace sgm;
+	using namespace details;
+
+	const int w = 631;
+	const int h = 479;
+	const int pitch = 640;
+	const ImageType stype = SGM_32U;
+	const ImageType dtype = SGM_64U;
+	const CensusType censusType = CensusType::CENSUS_9x7;
+
+	HostImage h_src(h, w, stype, pitch), h_dst(h, w, dtype);
+	DeviceImage d_src(h, w, stype, pitch), d_dst(h, w, dtype);
+
+	random_fill(h_src);
+	d_src.upload(h_src.data);
+	d_dst.fill_zero();
+
+	census_transform(h_src, h_dst, censusType);
+	census_transform(d_src, d_dst, censusType);
+
+	EXPECT_TRUE(equals(h_dst, d_dst));
+}
+
 TEST(SymmetricCensusTest, RandomU8)
 {
 	using namespace sgm;
@@ -175,3 +204,28 @@ TEST(SymmetricCensusTest, Random16U)
 
 	EXPECT_TRUE(equals(h_dst, d_dst));
 }
+
+TEST(SymmetricCensusTest, Random32U)
+{
+	using namespace sgm;
+	using namespace details;
+
+	const int w = 631;
+	const int h = 479;
+	const int pitch = 640;
+	const ImageType stype = SGM_32U;
+	const ImageType dtype = SGM_32U;
+	const CensusType censusType = CensusType::SYMMETRIC_CENSUS_9x7;
+
+	HostImage h_src(h, w, stype, pitch), h_dst(h, w, dtype);
+	DeviceImage d_src(h, w, stype, pitch), d_dst(h, w, dtype);
+
+	random_fill(h_src);
+	d_src.upload(h_src.data);
+	d_dst.fill_zero();
+
+	census_transform(h_src, h_dst, censusType);
+	census_transform(d_src, d_dst, censusType);
+
+	EXPECT_TRUE(equals(h_dst, d_dst));
+}
diff --git a/test/check_consistency_test.cpp b/test/check_consistency_test.cpp
@@ -49,6 +49,8 @@ void check_consistency(HostImage& dispL, const HostImage& dispR, const HostImage
 		check_consistency_<uint8_t>(dispL, dispR, srcL, subpixel, LR_max_diff);
 	if (srcL.type == SGM_16U)
 		check_consistency_<uint16_t>(dispL, dispR, srcL, subpixel, LR_max_diff);
+	if (srcL.type == SGM_32U)
+		check_consistency_<uint32_t>(dispL, dispR, srcL, subpixel, LR_max_diff);
 }
 
 } // namespace sgm
@@ -113,6 +115,36 @@ TEST(CheckConsistencyTest, RandomU16)
 	EXPECT_TRUE(equals(h_dispL, d_dispL));
 }
 
+TEST(CheckConsistencyTest, RandomU32)
+{
+	using namespace sgm;
+	using namespace details;
+
+	const int w = 631;
+	const int h = 479;
+	const int pitch = 640;
+	const ImageType stype = SGM_32U;
+	const ImageType dtype = SGM_16U;
+	const int LR_max_diff = 5;
+	const bool subpixel = false;
+
+	HostImage h_srcL(h, w, stype, pitch), h_dispL(h, w, dtype, pitch), h_dispR(h, w, dtype, pitch);
+	DeviceImage d_srcL(h, w, stype, pitch), d_dispL(h, w, dtype, pitch), d_dispR(h, w, dtype, pitch);
+
+	random_fill(h_srcL);
+	random_fill(h_dispL);
+	random_fill(h_dispR);
+
+	d_srcL.upload(h_srcL.data);
+	d_dispL.upload(h_dispL.data);
+	d_dispR.upload(h_dispR.data);
+
+	check_consistency(h_dispL, h_dispR, h_srcL, subpixel, LR_max_diff);
+	check_consistency(d_dispL, d_dispR, d_srcL, subpixel, LR_max_diff);
+
+	EXPECT_TRUE(equals(h_dispL, d_dispL));
+}
+
 TEST(CheckConsistencyTest, RandomU8_Subpixel)
 {
 	using namespace sgm;
@@ -172,3 +204,33 @@ TEST(CheckConsistencyTest, RandomU16_Subpixel)
 
 	EXPECT_TRUE(equals(h_dispL, d_dispL));
 }
+
+TEST(CheckConsistencyTest, RandomU32_Subpixel)
+{
+	using namespace sgm;
+	using namespace details;
+
+	const int w = 631;
+	const int h = 479;
+	const int pitch = 640;
+	const ImageType stype = SGM_32U;
+	const ImageType dtype = SGM_16U;
+	const int LR_max_diff = 5;
+	const bool subpixel = true;
+
+	HostImage h_srcL(h, w, stype, pitch), h_dispL(h, w, dtype, pitch), h_dispR(h, w, dtype, pitch);
+	DeviceImage d_srcL(h, w, stype, pitch), d_dispL(h, w, dtype, pitch), d_dispR(h, w, dtype, pitch);
+
+	random_fill(h_srcL);
+	random_fill(h_dispL);
+	random_fill(h_dispR);
+
+	d_srcL.upload(h_srcL.data);
+	d_dispL.upload(h_dispL.data);
+	d_dispR.upload(h_dispR.data);
+
+	check_consistency(h_dispL, h_dispR, h_srcL, subpixel, LR_max_diff);
+	check_consistency(d_dispL, d_dispR, d_srcL, subpixel, LR_max_diff);
+
+	EXPECT_TRUE(equals(h_dispL, d_dispL));
+}