From 937e6024dcb7c1133341adb8d82e08ce43c5696a Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nicolas.pope@utu.fi>
Date: Sat, 29 Feb 2020 22:51:59 +0200
Subject: [PATCH] Implements #315 border masking
---
applications/gui/src/camera.cpp | 5 +++
applications/gui/src/src_window.cpp | 1 +
.../operators/include/ftl/operators/mask.hpp | 14 +++++++
.../include/ftl/operators/mask_cuda.hpp | 5 +++
components/operators/src/mask.cpp | 27 ++++++++++++-
components/operators/src/mask.cu | 38 ++++++++++++++++++-
6 files changed, 88 insertions(+), 2 deletions(-)
diff --git a/applications/gui/src/camera.cpp b/applications/gui/src/camera.cpp
index 4076c1ee6..3afbe4264 100644
--- a/applications/gui/src/camera.cpp
+++ b/applications/gui/src/camera.cpp
@@ -150,6 +150,11 @@ void ftl::gui::Camera::draw(std::vector<ftl::rgbd::FrameSet*> &fss) {
ftl::cuda::flip<uchar4>(dst1, 0);
texture1_.unmap(0);
+ depth1_.make(buf.width(), buf.height());
+ dst1 = depth1_.map(0);
+ dst1.setTo(cv::Scalar(0.5f));
+ depth1_.unmap(0);
+
width_ = texture1_.width();
height_ = texture1_.height();
return;
diff --git a/applications/gui/src/src_window.cpp b/applications/gui/src/src_window.cpp
index 15e9e96f7..084d7992c 100644
--- a/applications/gui/src/src_window.cpp
+++ b/applications/gui/src/src_window.cpp
@@ -241,6 +241,7 @@ void SourceWindow::_checkFrameSets(int id) {
p->append<ftl::operators::CullWeight>("remove_weights")->value("enabled", false);
p->append<ftl::operators::DegradeWeight>("degrade");
p->append<ftl::operators::VisCrossSupport>("viscross")->set("enabled", false);
+ p->append<ftl::operators::BorderMask>("border_mask");
p->append<ftl::operators::CullDiscontinuity>("remove_discontinuity");
p->append<ftl::operators::MultiViewMLS>("mvmls")->value("enabled", false);
diff --git a/components/operators/include/ftl/operators/mask.hpp b/components/operators/include/ftl/operators/mask.hpp
index 5c542207d..294caf54f 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 {
};
+/**
+ * Generate a depth border mask.
+ */
+class BorderMask : public ftl::operators::Operator {
+ public:
+ explicit BorderMask(ftl::Configurable*);
+ ~BorderMask();
+
+ inline Operator::Type type() const override { return Operator::Type::OneToOne; }
+
+ bool apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, cudaStream_t stream) override;
+
+};
+
/**
* Remove depth values marked with the discontinuity mask.
*/
diff --git a/components/operators/include/ftl/operators/mask_cuda.hpp b/components/operators/include/ftl/operators/mask_cuda.hpp
index f780d5178..aadab5e5a 100644
--- a/components/operators/include/ftl/operators/mask_cuda.hpp
+++ b/components/operators/include/ftl/operators/mask_cuda.hpp
@@ -69,6 +69,11 @@ void discontinuity(
float area_max,
cudaStream_t stream);
+void border_mask(
+ ftl::cuda::TextureObject<ftl::cuda::Mask::type> &mask,
+ int left, int right, int top, int bottom,
+ cudaStream_t stream);
+
void cull_mask(
ftl::cuda::TextureObject<ftl::cuda::Mask::type> &mask,
ftl::cuda::TextureObject<float> &depth,
diff --git a/components/operators/src/mask.cpp b/components/operators/src/mask.cpp
index beec2bd28..31038b7bb 100644
--- a/components/operators/src/mask.cpp
+++ b/components/operators/src/mask.cpp
@@ -2,6 +2,7 @@
#include <ftl/operators/mask_cuda.hpp>
using ftl::operators::DiscontinuityMask;
+using ftl::operators::BorderMask;
using ftl::operators::CullDiscontinuity;
using ftl::codecs::Channel;
using ftl::rgbd::Format;
@@ -49,6 +50,30 @@ bool DiscontinuityMask::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, cudaS
+BorderMask::BorderMask(ftl::Configurable *cfg) : ftl::operators::Operator(cfg) {
+
+}
+
+BorderMask::~BorderMask() {
+
+}
+
+bool BorderMask::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, cudaStream_t stream) {
+ int leftm = config()->value("left", 100);
+ int rightm = config()->value("right", 5);
+ int topm = config()->value("top",5);
+ int bottomm = config()->value("bottom",5);
+
+ ftl::cuda::border_mask(
+ out.createTexture<uint8_t>(Channel::Mask, ftl::rgbd::Format<uint8_t>(in.get<cv::cuda::GpuMat>(Channel::Depth).size())),
+ leftm, rightm, topm, bottomm, stream
+ );
+
+ return true;
+}
+
+
+
CullDiscontinuity::CullDiscontinuity(ftl::Configurable *cfg) : ftl::operators::Operator(cfg) {
}
@@ -60,7 +85,7 @@ CullDiscontinuity::~CullDiscontinuity() {
bool CullDiscontinuity::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, cudaStream_t stream) {
if (!in.hasChannel(Channel::Depth) || !in.hasChannel(Channel::Mask)) return false;
- uint8_t maskID = config()->value("mask_id", (unsigned int)ftl::cuda::Mask::kMask_Discontinuity);
+ uint8_t maskID = config()->value("mask_id", (unsigned int)(ftl::cuda::Mask::kMask_Discontinuity | ftl::cuda::Mask::kMask_Bad));
unsigned int radius = config()->value("radius", 2);
bool inverted = config()->value("invert", false);
diff --git a/components/operators/src/mask.cu b/components/operators/src/mask.cu
index 7c7ce32ad..85fde5190 100644
--- a/components/operators/src/mask.cu
+++ b/components/operators/src/mask.cu
@@ -68,12 +68,48 @@ void ftl::cuda::discontinuity( ftl::cuda::TextureObject<uint8_t> &mask_out, ftl:
// =============================================================================
+__global__ void border_mask_kernel(uint8_t* __restrict__ mask_out,
+ int pitch, int width, int height,
+ int left, int right, int top, int bottom) {
+
+ const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
+ const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+ if (x < width && y < height) {
+ Mask mask(mask_out[x+y*pitch]);
+ if (x < left || x >= width-right || y < top || y >= height-bottom) {
+ mask.isBad(true);
+ mask_out[x+y*pitch] = (int)mask;
+ }
+ }
+}
+
+void ftl::cuda::border_mask(ftl::cuda::TextureObject<uint8_t> &mask_out,
+ int left, int right, int top, int bottom, cudaStream_t stream) {
+
+ static constexpr int THREADS_X = 128;
+ static constexpr int THREADS_Y = 4;
+
+ const dim3 gridSize((mask_out.width() + THREADS_X - 1)/THREADS_X, (mask_out.height() + THREADS_Y - 1)/THREADS_Y);
+ const dim3 blockSize(THREADS_X, THREADS_Y);
+
+ border_mask_kernel<<<gridSize, blockSize, 0, stream>>>(mask_out.devicePtr(), mask_out.pixelPitch(),
+ mask_out.width(), mask_out.height(), left, right, top, bottom);
+ cudaSafeCall( cudaGetLastError() );
+
+ #ifdef _DEBUG
+ cudaSafeCall(cudaDeviceSynchronize());
+ #endif
+}
+
+// =============================================================================
+
template <int RADIUS, bool INVERT>
__global__ void cull_mask_kernel(ftl::cuda::TextureObject<uint8_t> mask, ftl::cuda::TextureObject<float> depth, uint8_t id) {
const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
- if (x < depth.width()-RADIUS && y < depth.height()-RADIUS) {
+ if (x < depth.width() && y < depth.height()) {
bool isdiscon = false;
#pragma unroll
--
GitLab