Be more selective on discon mask

applications/reconstruct/src/ilw/ilw.cpp

@@ -215,11 +215,11 @@ bool ILW::_phase0(ftl::rgbd::FrameSet &fs, cudaStream_t stream) {
 
         f.createTexture<float>(Channel::Depth2, Format<float>(f.get<GpuMat>(Channel::Colour).size()));
         f.createTexture<float>(Channel::Confidence, Format<float>(f.get<GpuMat>(Channel::Colour).size()));
-		f.createTexture<int>(Channel::Mask, Format<int>(f.get<GpuMat>(Channel::Colour).size()));
+		//f.createTexture<int>(Channel::Mask, Format<int>(f.get<GpuMat>(Channel::Colour).size()));
         f.createTexture<uchar4>(Channel::Colour);
 		f.createTexture<float>(Channel::Depth);
 
-		f.get<GpuMat>(Channel::Mask).setTo(cv::Scalar(0));
+		//f.get<GpuMat>(Channel::Mask).setTo(cv::Scalar(0));
     }
 
 	//cudaSafeCall(cudaStreamSynchronize(stream_));
@@ -232,7 +232,7 @@ bool ILW::_phase1(ftl::rgbd::FrameSet &fs, int win, cudaStream_t stream) {
     cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
 
     // Find discontinuity mask
-    for (size_t i=0; i<fs.frames.size(); ++i) {
+    /*for (size_t i=0; i<fs.frames.size(); ++i) {
         auto &f = fs.frames[i];
         auto s = fs.sources[i];
 
@@ -243,7 +243,7 @@ bool ILW::_phase1(ftl::rgbd::FrameSet &fs, int win, cudaStream_t stream) {
             discon_mask_,
             stream
         );
-    }
+    }*/
 
 	// First do any preprocessing
 	if (fill_depth_) {

components/operators/src/mask.cpp

@@ -15,11 +15,13 @@ DiscontinuityMask::~DiscontinuityMask() {
 
 bool DiscontinuityMask::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::Source *s, cudaStream_t stream) {
 	int radius = config()->value("radius", 2);
+	float threshold = config()->value("threshold", 0.1f);
 
 	ftl::cuda::discontinuity(
 		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, 0
+		s->parameters(), radius, threshold, 0
 	);
 
 	return true;

components/operators/src/mask.cu

@@ -5,7 +5,7 @@
 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<uchar4> support, ftl::cuda::TextureObject<float> depth, ftl::rgbd::Camera params, float threshold) {
 	const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
 	const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
 
@@ -20,11 +20,29 @@ __global__ void discontinuity_kernel(ftl::cuda::TextureObject<int> mask_out, ftl
 		if (d >= params.minDepth && d <= params.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) {
-				for (int v=-RADIUS; v<=RADIUS; ++v) {
+			//for (int u=-RADIUS; u<=RADIUS; ++u) {
+			//	for (int v=-RADIUS; v<=RADIUS; ++v) {
 					// If yes, the flag using w = -1
-					if (fabs(depth.tex2D((int)x+u, (int)y+v) - d) > 0.1f) mask.isDiscontinuity(true);
+			//		if (fabs(depth.tex2D((int)x+u, (int)y+v) - d) > 0.1f) mask.isDiscontinuity(true);
+			//	}
+			//}
+
+			const uchar4 sup = support.tex2D((int)x, (int)y);
+			if (sup.x <= RADIUS) {
+				float dS = depth.tex2D((int)x - sup.x - 1, (int)y);
+				if (fabs(dS - d) > threshold) mask.isDiscontinuity(true);
+			}
+			if (sup.y <= RADIUS) {
+				float dS = depth.tex2D((int)x + sup.y + 1, (int)y);
+				if (fabs(dS - d) > threshold) mask.isDiscontinuity(true);
+			}
+			if (sup.z <= RADIUS) {
+				float dS = depth.tex2D((int)x, (int)y - sup.z - 1);
+				if (fabs(dS - d) > threshold) mask.isDiscontinuity(true);
 			}
+			if (sup.w <= RADIUS) {
+				float dS = depth.tex2D((int)x, (int)y + sup.w + 1);
+				if (fabs(dS - d) > threshold) mask.isDiscontinuity(true);
 			}
         }
         
@@ -32,16 +50,16 @@ __global__ void discontinuity_kernel(ftl::cuda::TextureObject<int> mask_out, ftl
 	}
 }
 
-void ftl::cuda::discontinuity(ftl::cuda::TextureObject<int> &mask_out, ftl::cuda::TextureObject<float> &depth, const ftl::rgbd::Camera &params, int discon, cudaStream_t stream) {
+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);
 	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, support, depth, params, thresh); break;
+	case 4 :	discontinuity_kernel<4><<<gridSize, blockSize, 0, stream>>>(mask_out, support, depth, params, thresh); break;
+	case 3 :	discontinuity_kernel<3><<<gridSize, blockSize, 0, stream>>>(mask_out, support, depth, params, thresh); break;
+	case 2 :	discontinuity_kernel<2><<<gridSize, blockSize, 0, stream>>>(mask_out, support, depth, params, thresh); break;
+	case 1 :	discontinuity_kernel<1><<<gridSize, blockSize, 0, stream>>>(mask_out, support, depth, params, thresh); break;
 	default:	break;
 	}
 	cudaSafeCall( cudaGetLastError() );

components/operators/src/mask_cuda.hpp

@@ -41,10 +41,11 @@ class Mask {
 };
 
 void discontinuity(
-		ftl::cuda::TextureObject<int> &mast,
+		ftl::cuda::TextureObject<int> &mask,
+		ftl::cuda::TextureObject<uchar4> &support,
 		ftl::cuda::TextureObject<float> &depth,
 		const ftl::rgbd::Camera &params,
-		int radius,
+		int radius, float threshold,
 		cudaStream_t stream);
 
 }