From f35c50485fb1d7d2dcbd3a2a97bab050053af4e5 Mon Sep 17 00:00:00 2001
From: Sebastian Hahta <joseha@utu.fi>
Date: Fri, 29 Nov 2019 13:10:08 +0200
Subject: [PATCH] do not use width&height from params

---
 .../reconstruct/src/ilw/discontinuity.cu      | 30 ++++++-----
 applications/reconstruct/src/ilw/ilw_cuda.hpp | 54 ++++++++++---------
 components/operators/src/mask.cpp             |  4 +-
 components/operators/src/mask.cu              | 31 ++++++-----
 components/operators/src/mask_cuda.hpp        |  8 +--
 .../src/sources/stereovideo/stereovideo.cpp   | 27 +++++-----
 6 files changed, 87 insertions(+), 67 deletions(-)

diff --git a/applications/reconstruct/src/ilw/discontinuity.cu b/applications/reconstruct/src/ilw/discontinuity.cu
index fe78d4715..fcadde03e 100644
--- a/applications/reconstruct/src/ilw/discontinuity.cu
+++ b/applications/reconstruct/src/ilw/discontinuity.cu
@@ -5,11 +5,12 @@
 using ftl::cuda::Mask;
 
 template <int RADIUS>
-__global__ void discontinuity_kernel(ftl::cuda::TextureObject<int> mask_out, ftl::cuda::TextureObject<float> depth, ftl::rgbd::Camera params) {
+__global__ void discontinuity_kernel(ftl::cuda::TextureObject<int> mask_out, ftl::cuda::TextureObject<float> depth,
+										const cv::Size size, const double minDepth, const double maxDepth) {
 	const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
 	const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
 
-	if (x < params.width && y < params.height) {
+	if (x < size.width && y < size.height) {
 		Mask mask(0);
 
 		const float d = depth.tex2D((int)x, (int)y);
@@ -17,7 +18,7 @@ __global__ void discontinuity_kernel(ftl::cuda::TextureObject<int> mask_out, ftl
 		// Calculate depth between 0.0 and 1.0
 		//float p = (d - params.minDepth) / (params.maxDepth - params.minDepth);
 
-		if (d >= params.minDepth && d <= params.maxDepth) {
+		if (d >= minDepth && d <= maxDepth) {
 			/* Orts-Escolano S. et al. 2016. Holoportation: Virtual 3D teleportation in real-time. */
 			// Is there a discontinuity nearby?
 			for (int u=-RADIUS; u<=RADIUS; ++u) {
@@ -26,22 +27,25 @@ __global__ void discontinuity_kernel(ftl::cuda::TextureObject<int> mask_out, ftl
 					if (fabs(depth.tex2D((int)x+u, (int)y+v) - d) > 0.1f) mask.isDiscontinuity(true);
 				}
 			}
-        }
-        
-        mask_out(x,y) = (int)mask;
+		}
+		
+		mask_out(x,y) = (int)mask;
 	}
 }
 
-void ftl::cuda::discontinuity(ftl::cuda::TextureObject<int> &mask_out, ftl::cuda::TextureObject<float> &depth, const ftl::rgbd::Camera &params, uint discon, cudaStream_t stream) {
-	const dim3 gridSize((params.width + T_PER_BLOCK - 1)/T_PER_BLOCK, (params.height + T_PER_BLOCK - 1)/T_PER_BLOCK);
+void ftl::cuda::discontinuity(ftl::cuda::TextureObject<int> &mask_out, ftl::cuda::TextureObject<float> &depth,
+								const cv::Size size, const double minDepth, const double maxDepth,
+								uint discon, cudaStream_t stream) {
+			
+	const dim3 gridSize((size.width + T_PER_BLOCK - 1)/T_PER_BLOCK, (size.height + T_PER_BLOCK - 1)/T_PER_BLOCK);
 	const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
 
 	switch (discon) {
-    case 5 :	discontinuity_kernel<5><<<gridSize, blockSize, 0, stream>>>(mask_out, depth, params); break;
-	case 4 :	discontinuity_kernel<4><<<gridSize, blockSize, 0, stream>>>(mask_out, depth, params); break;
-	case 3 :	discontinuity_kernel<3><<<gridSize, blockSize, 0, stream>>>(mask_out, depth, params); break;
-	case 2 :	discontinuity_kernel<2><<<gridSize, blockSize, 0, stream>>>(mask_out, depth, params); break;
-	case 1 :	discontinuity_kernel<1><<<gridSize, blockSize, 0, stream>>>(mask_out, depth, params); break;
+	case 5 :	discontinuity_kernel<5><<<gridSize, blockSize, 0, stream>>>(mask_out, depth, size, minDepth, maxDepth); break;
+	case 4 :	discontinuity_kernel<4><<<gridSize, blockSize, 0, stream>>>(mask_out, depth, size, minDepth, maxDepth); break;
+	case 3 :	discontinuity_kernel<3><<<gridSize, blockSize, 0, stream>>>(mask_out, depth, size, minDepth, maxDepth); break;
+	case 2 :	discontinuity_kernel<2><<<gridSize, blockSize, 0, stream>>>(mask_out, depth, size, minDepth, maxDepth); break;
+	case 1 :	discontinuity_kernel<1><<<gridSize, blockSize, 0, stream>>>(mask_out, depth, size, minDepth, maxDepth); break;
 	default:	break;
 	}
 	cudaSafeCall( cudaGetLastError() );
diff --git a/applications/reconstruct/src/ilw/ilw_cuda.hpp b/applications/reconstruct/src/ilw/ilw_cuda.hpp
index fad97afbd..94e522347 100644
--- a/applications/reconstruct/src/ilw/ilw_cuda.hpp
+++ b/applications/reconstruct/src/ilw/ilw_cuda.hpp
@@ -10,15 +10,15 @@ namespace ftl {
 namespace cuda {
 
 struct ILWParams {
-    float spatial_smooth;
-    float colour_smooth;
+	float spatial_smooth;
+	float colour_smooth;
 	float fill_match;
 	float fill_threshold;
 	float match_threshold;
-    float cost_ratio;
-    float cost_threshold;
+	float cost_ratio;
+	float cost_threshold;
 	float range;
-    uint flags;
+	uint flags;
 };
 
 static const uint kILWFlag_IgnoreBad = 0x0001;
@@ -29,7 +29,9 @@ static const uint kILWFlag_ColourConfidenceOnly = 0x0008;
 void discontinuity(
 	ftl::cuda::TextureObject<int> &mask_out,
 	ftl::cuda::TextureObject<float> &depth,
-	const ftl::rgbd::Camera &params,
+	const cv::Size size,
+	const double minDepth,
+	const double maxDepth,
 	uint discon, cudaStream_t stream
 );
 
@@ -49,32 +51,32 @@ void preprocess_depth(
 );
 
 void correspondence(
-    ftl::cuda::TextureObject<float> &d1,
-    ftl::cuda::TextureObject<float> &d2,
-    ftl::cuda::TextureObject<uchar4> &c1,
-    ftl::cuda::TextureObject<uchar4> &c2,
-    ftl::cuda::TextureObject<float> &dout,
-    ftl::cuda::TextureObject<float> &conf,
+	ftl::cuda::TextureObject<float> &d1,
+	ftl::cuda::TextureObject<float> &d2,
+	ftl::cuda::TextureObject<uchar4> &c1,
+	ftl::cuda::TextureObject<uchar4> &c2,
+	ftl::cuda::TextureObject<float> &dout,
+	ftl::cuda::TextureObject<float> &conf,
 	ftl::cuda::TextureObject<int> &mask,
-    float4x4 &pose1,
-    float4x4 &pose1_inv,
-    float4x4 &pose2,
-    const ftl::rgbd::Camera &cam1,
-    const ftl::rgbd::Camera &cam2,
-    const ILWParams &params, int win,
-    cudaStream_t stream
+	float4x4 &pose1,
+	float4x4 &pose1_inv,
+	float4x4 &pose2,
+	const ftl::rgbd::Camera &cam1,
+	const ftl::rgbd::Camera &cam2,
+	const ILWParams &params, int win,
+	cudaStream_t stream
 );
 
 void move_points(
-    ftl::cuda::TextureObject<float> &d_old,
-    ftl::cuda::TextureObject<float> &d_new,
+	ftl::cuda::TextureObject<float> &d_old,
+	ftl::cuda::TextureObject<float> &d_new,
 	ftl::cuda::TextureObject<float> &conf,
-    const ftl::rgbd::Camera &camera,
-    const float4x4 &pose,
+	const ftl::rgbd::Camera &camera,
+	const float4x4 &pose,
 	const ILWParams &params,
-    float rate,
-    int radius,
-    cudaStream_t stream
+	float rate,
+	int radius,
+	cudaStream_t stream
 );
 
 }
diff --git a/components/operators/src/mask.cpp b/components/operators/src/mask.cpp
index f923f11d0..c7dcbb2ac 100644
--- a/components/operators/src/mask.cpp
+++ b/components/operators/src/mask.cpp
@@ -22,7 +22,9 @@ bool DiscontinuityMask::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::
 		out.createTexture<int>(Channel::Mask, ftl::rgbd::Format<int>(in.get<cv::cuda::GpuMat>(Channel::Depth).size())),
 		in.createTexture<uchar4>(Channel::Support1),
 		in.createTexture<float>(Channel::Depth),
-		s->parameters(), radius, threshold, stream
+		in.get<cv::cuda::GpuMat>(Channel::Depth).size(),
+		s->parameters().minDepth, s->parameters().maxDepth,
+		radius, threshold, stream
 	);
 
 	return true;
diff --git a/components/operators/src/mask.cu b/components/operators/src/mask.cu
index e385f41b1..91ddf19dd 100644
--- a/components/operators/src/mask.cu
+++ b/components/operators/src/mask.cu
@@ -4,16 +4,21 @@
 
 using ftl::cuda::Mask;
 
-__global__ void discontinuity_kernel(ftl::cuda::TextureObject<int> mask_out, ftl::cuda::TextureObject<uchar4> support, ftl::cuda::TextureObject<float> depth, ftl::rgbd::Camera params, float threshold, int radius) {
+__global__ void discontinuity_kernel(ftl::cuda::TextureObject<int> mask_out,
+										ftl::cuda::TextureObject<uchar4> support,
+										ftl::cuda::TextureObject<float> depth, 
+										const cv::Size size, const double minDepth, const double maxDepth,
+										float threshold, int radius) {
+	
 	const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
 	const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
 
-	if (x < params.width && y < params.height) {
+	if (x < size.width && y < size.height) {
 		Mask mask(0);
 
 		const float d = depth.tex2D((int)x, (int)y);
 
-		if (d >= params.minDepth && d <= params.maxDepth) {
+		if (d >= minDepth && d <= maxDepth) {
 			/* Orts-Escolano S. et al. 2016. Holoportation: Virtual 3D teleportation in real-time. */
 
 			// If colour cross support region terminates within the requested
@@ -37,17 +42,21 @@ __global__ void discontinuity_kernel(ftl::cuda::TextureObject<int> mask_out, ftl
 				float dS = depth.tex2D((int)x, (int)y + sup.w + radius);
 				if (fabs(dS - d) > threshold) mask.isDiscontinuity(true);
 			}
-        }
-        
-        mask_out(x,y) = (int)mask;
+		}
+		
+		mask_out(x,y) = (int)mask;
 	}
 }
 
-void ftl::cuda::discontinuity(ftl::cuda::TextureObject<int> &mask_out, ftl::cuda::TextureObject<uchar4> &support, ftl::cuda::TextureObject<float> &depth, const ftl::rgbd::Camera &params, int discon, float thresh, cudaStream_t stream) {
-	const dim3 gridSize((params.width + T_PER_BLOCK - 1)/T_PER_BLOCK, (params.height + T_PER_BLOCK - 1)/T_PER_BLOCK);
+void ftl::cuda::discontinuity(	ftl::cuda::TextureObject<int> &mask_out, ftl::cuda::TextureObject<uchar4> &support,
+								ftl::cuda::TextureObject<float> &depth,
+								const cv::Size size, const double minDepth, const double maxDepth,
+								int discon, float thresh, cudaStream_t stream) {
+	
+	const dim3 gridSize((size.width + T_PER_BLOCK - 1)/T_PER_BLOCK, (size.height + T_PER_BLOCK - 1)/T_PER_BLOCK);
 	const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
 
-    discontinuity_kernel<<<gridSize, blockSize, 0, stream>>>(mask_out, support, depth, params, thresh, discon);
+	discontinuity_kernel<<<gridSize, blockSize, 0, stream>>>(mask_out, support, depth, size, minDepth, maxDepth, thresh, discon);
 	cudaSafeCall( cudaGetLastError() );
 
 #ifdef _DEBUG
@@ -55,8 +64,6 @@ void ftl::cuda::discontinuity(ftl::cuda::TextureObject<int> &mask_out, ftl::cuda
 #endif
 }
 
-
-
 __global__ void cull_discontinuity_kernel(ftl::cuda::TextureObject<int> mask, ftl::cuda::TextureObject<float> depth) {
 	const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
 	const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
@@ -71,7 +78,7 @@ void ftl::cuda::cull_discontinuity(ftl::cuda::TextureObject<int> &mask, ftl::cud
 	const dim3 gridSize((depth.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
 	const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
 
-    cull_discontinuity_kernel<<<gridSize, blockSize, 0, stream>>>(mask, depth);
+	cull_discontinuity_kernel<<<gridSize, blockSize, 0, stream>>>(mask, depth);
 	cudaSafeCall( cudaGetLastError() );
 
 #ifdef _DEBUG
diff --git a/components/operators/src/mask_cuda.hpp b/components/operators/src/mask_cuda.hpp
index 6a02aafdb..20c266290 100644
--- a/components/operators/src/mask_cuda.hpp
+++ b/components/operators/src/mask_cuda.hpp
@@ -19,7 +19,7 @@ class Mask {
 	#endif
 	__device__ inline operator int() const { return v_; }
 
-    __device__ inline bool is(int m) const { return v_ & m; }
+	__device__ inline bool is(int m) const { return v_ & m; }
 
 	__device__ inline bool isFilled() const { return v_ & kMask_Filled; }
 	__device__ inline bool isDiscontinuity() const { return v_ & kMask_Discontinuity; }
@@ -31,7 +31,7 @@ class Mask {
 	__device__ inline void hasCorrespondence(bool v) { v_ = (v) ? v_ | kMask_Correspondence : v_ & (~kMask_Correspondence); }
 	__device__ inline void isBad(bool v) { v_ = (v) ? v_ | kMask_Bad : v_ & (~kMask_Bad); }
 
-    static constexpr int kMask_Filled = 0x0001;
+	static constexpr int kMask_Filled = 0x0001;
 	static constexpr int kMask_Discontinuity = 0x0002;
 	static constexpr int kMask_Correspondence = 0x0004;
 	static constexpr int kMask_Bad = 0x0008;
@@ -44,7 +44,9 @@ void discontinuity(
 		ftl::cuda::TextureObject<int> &mask,
 		ftl::cuda::TextureObject<uchar4> &support,
 		ftl::cuda::TextureObject<float> &depth,
-		const ftl::rgbd::Camera &params,
+		const cv::Size size,
+		const double minDepth,
+		const double maxDepth,
 		int radius, float threshold,
 		cudaStream_t stream);
 
diff --git a/components/rgbd-sources/src/sources/stereovideo/stereovideo.cpp b/components/rgbd-sources/src/sources/stereovideo/stereovideo.cpp
index b467099f5..b4f365101 100644
--- a/components/rgbd-sources/src/sources/stereovideo/stereovideo.cpp
+++ b/components/rgbd-sources/src/sources/stereovideo/stereovideo.cpp
@@ -130,12 +130,12 @@ void StereoVideoSource::init(const string &file) {
 	pipeline_input_->append<ftl::operators::NVOpticalFlow>("optflow");
 	#endif
 
-	depth_size_ = cv::Size(	host_->value("width", color_size_.width),
-							host_->value("height", color_size_.height));
-
 	pipeline_depth_ = ftl::config::create<ftl::operators::Graph>(host_, "disparity");
-	pipeline_depth_->append<ftl::operators::FixstarsSGM>("algorithm");
 
+	depth_size_ = cv::Size(	pipeline_depth_->value("width", color_size_.width),
+							pipeline_depth_->value("height", color_size_.height));
+
+	pipeline_depth_->append<ftl::operators::FixstarsSGM>("algorithm");
 	#ifdef HAVE_OPTFLOW
 	pipeline_depth_->append<ftl::operators::OpticalFlowTemporalSmoothing>("optflow_filter");
 	#endif
@@ -207,24 +207,27 @@ bool StereoVideoSource::compute(int n, int b) {
 	}
 
 	if (chan == Channel::Depth) {
-		// stereo algorithms assume input same size as output, resize if
-		// necessary
+		// stereo algorithms assume input same size as output
 		bool resize = (depth_size_ != color_size_);
+
+		cv::cuda::GpuMat& left = frame.get<cv::cuda::GpuMat>(Channel::Left);
+		cv::cuda::GpuMat& right = frame.get<cv::cuda::GpuMat>(Channel::Right);
+
+		if (left.empty() || right.empty()) {
+			return false;
+		}
+
 		if (resize) {
-			cv::cuda::GpuMat &left = frame.get<cv::cuda::GpuMat>(Channel::Left);
-			cv::cuda::GpuMat &right = frame.get<cv::cuda::GpuMat>(Channel::Right);
 			std::swap(fullres_left_, left);
 			std::swap(fullres_right_, right);
-			cv::cuda::resize(fullres_left_, left, depth_size_, 0.0, 0.0, cv::INTER_CUBIC, stream_);
-			cv::cuda::resize(fullres_right_, right, depth_size_, 0.0, 0.0, cv::INTER_CUBIC, stream_);
+			cv::cuda::resize(fullres_left_, left, depth_size_, 0, 0, cv::INTER_CUBIC, stream_);
+			cv::cuda::resize(fullres_right_, right, depth_size_, 0, 0, cv::INTER_CUBIC, stream_);
 		}
 
 		pipeline_depth_->apply(frame, frame, host_, cv::cuda::StreamAccessor::getStream(stream_));
 		stream_.waitForCompletion();
 		
 		if (resize) {
-			cv::cuda::GpuMat &left = frame.get<cv::cuda::GpuMat>(Channel::Left);
-			cv::cuda::GpuMat &right = frame.get<cv::cuda::GpuMat>(Channel::Right);
 			std::swap(fullres_left_, left);
 			std::swap(fullres_right_, right);
 		}
-- 
GitLab