From cabdd52fd8a52e12fc8c4030a50e0b2ed99f7f92 Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Mon, 11 Nov 2019 20:16:18 +0200
Subject: [PATCH] Add cull discontinuity operator
---
applications/reconstruct/src/main.cpp | 2 +-
.../operators/include/ftl/operators/mask.hpp | 14 +++++++++++
components/operators/src/mask.cpp | 23 +++++++++++++++++-
components/operators/src/mask.cu | 24 +++++++++++++++++++
components/operators/src/mask_cuda.hpp | 5 ++++
5 files changed, 66 insertions(+), 2 deletions(-)
diff --git a/applications/reconstruct/src/main.cpp b/applications/reconstruct/src/main.cpp
index c99a53801..03c21b142 100644
--- a/applications/reconstruct/src/main.cpp
+++ b/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
diff --git a/components/operators/include/ftl/operators/mask.hpp b/components/operators/include/ftl/operators/mask.hpp
index f2fc1ad5d..ef9758e39 100644
--- a/components/operators/include/ftl/operators/mask.hpp
+++ b/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;
+
+};
+
}
}
diff --git a/components/operators/src/mask.cpp b/components/operators/src/mask.cpp
index 1ac8ba92c..f923f11d0 100644
--- a/components/operators/src/mask.cpp
+++ b/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;
diff --git a/components/operators/src/mask.cu b/components/operators/src/mask.cu
index aac736b1c..397935c5d 100644
--- a/components/operators/src/mask.cu
+++ b/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
+}
diff --git a/components/operators/src/mask_cuda.hpp b/components/operators/src/mask_cuda.hpp
index 64e412113..6a02aafdb 100644
--- a/components/operators/src/mask_cuda.hpp
+++ b/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);
+
}
}
--
GitLab