Add cull discontinuity operator

applications/reconstruct/src/main.cpp

@@ -312,7 +312,7 @@ static void run(ftl::Configurable *root) {
 	//pipeline1->append<ftl::operators::ScanFieldFill>("filling");  // Generate a smoothing channel
 	pipeline1->append<ftl::operators::CrossSupport>("cross");
 	pipeline1->append<ftl::operators::DiscontinuityMask>("discontinuity");
-	//pipeline1->append<ftl::operators::CullDiscontinuity>("remove_discontinuity");
+	pipeline1->append<ftl::operators::CullDiscontinuity>("remove_discontinuity");
 	pipeline1->append<ftl::operators::ColourMLS>("mls");  // Perform MLS (using smoothing channel)
 	pipeline1->append<ftl::operators::VisCrossSupport>("viscross")->set("enabled", false);
 	// Alignment

components/operators/include/ftl/operators/mask.hpp

@@ -23,6 +23,20 @@ class DiscontinuityMask : public ftl::operators::Operator {
 
 };
 
+/**
+ * Remove depth values marked with the discontinuity mask.
+ */
+class CullDiscontinuity : public ftl::operators::Operator {
+	public:
+    explicit CullDiscontinuity(ftl::Configurable*);
+    ~CullDiscontinuity();
+
+	inline Operator::Type type() const override { return Operator::Type::OneToOne; }
+
+    bool apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::Source *src, cudaStream_t stream) override;
+
+};
+
 }
 }
 

components/operators/src/mask.cpp

@@ -2,6 +2,7 @@
 #include "mask_cuda.hpp"
 
 using ftl::operators::DiscontinuityMask;
+using ftl::operators::CullDiscontinuity;
 using ftl::codecs::Channel;
 using ftl::rgbd::Format;
 
@@ -21,7 +22,27 @@ 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, 0
+		s->parameters(), radius, threshold, stream
+	);
+
+	return true;
+}
+
+
+
+CullDiscontinuity::CullDiscontinuity(ftl::Configurable *cfg) : ftl::operators::Operator(cfg) {
+
+}
+
+CullDiscontinuity::~CullDiscontinuity() {
+
+}
+
+bool CullDiscontinuity::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::Source *s, cudaStream_t stream) {
+	ftl::cuda::cull_discontinuity(
+		in.createTexture<int>(Channel::Mask),
+		out.createTexture<float>(Channel::Depth),
+		stream
 	);
 
 	return true;

components/operators/src/mask.cu

@@ -68,3 +68,27 @@ void ftl::cuda::discontinuity(ftl::cuda::TextureObject<int> &mask_out, ftl::cuda
 	cudaSafeCall(cudaDeviceSynchronize());
 #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;
+
+	if (x < depth.width() && y < depth.height()) {
+		Mask m(mask.tex2D((int)x,(int)y));
+		if (m.isDiscontinuity()) depth(x,y) = 0.0f;
+	}
+}
+
+void ftl::cuda::cull_discontinuity(ftl::cuda::TextureObject<int> &mask, ftl::cuda::TextureObject<float> &depth, cudaStream_t stream) {
+	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);
+	cudaSafeCall( cudaGetLastError() );
+
+#ifdef _DEBUG
+	cudaSafeCall(cudaDeviceSynchronize());
+#endif
+}

components/operators/src/mask_cuda.hpp

@@ -48,6 +48,11 @@ void discontinuity(
 		int radius, float threshold,
 		cudaStream_t stream);
 
+void cull_discontinuity(
+		ftl::cuda::TextureObject<int> &mask,
+		ftl::cuda::TextureObject<float> &depth,
+		cudaStream_t stream);
+
 }
 }