From 90bfabd4fe3db7d7873dfa648288c460cdecc37e Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Thu, 7 Nov 2019 16:31:59 +0200
Subject: [PATCH] Working support regions
---
applications/reconstruct/src/main.cpp | 7 ++-
.../include/ftl/operators/segmentation.hpp | 14 ++++++
components/operators/src/segmentation.cpp | 22 +++++++++
components/operators/src/segmentation.cu | 48 +++++++++++++++++--
.../operators/src/segmentation_cuda.hpp | 5 ++
5 files changed, 89 insertions(+), 7 deletions(-)
diff --git a/applications/reconstruct/src/main.cpp b/applications/reconstruct/src/main.cpp
index 73445da8c..d5f103387 100644
--- a/applications/reconstruct/src/main.cpp
+++ b/applications/reconstruct/src/main.cpp
@@ -34,6 +34,7 @@
#include <ftl/operators/colours.hpp>
#include <ftl/operators/normals.hpp>
#include <ftl/operators/filling.hpp>
+#include <ftl/operators/segmentation.hpp>
#include <ftl/cuda/normals.hpp>
#include <ftl/registration.hpp>
@@ -165,7 +166,7 @@ static void run(ftl::Configurable *root) {
}
auto configproxy = ConfigProxy(net);
- configproxy.add(root, "source/disparity", "disparity");
+ //configproxy.add(root, "source/disparity", "disparity");
// Create scene transform, intended for axis aligning the walls and floor
Eigen::Matrix4d transform;
@@ -300,9 +301,11 @@ static void run(ftl::Configurable *root) {
pipeline1->append<ftl::operators::ColourChannels>("colour"); // Convert BGR to BGRA
pipeline1->append<ftl::operators::HFSmoother>("hfnoise"); // Remove high-frequency noise
pipeline1->append<ftl::operators::Normals>("normals"); // Estimate surface normals
- pipeline1->append<ftl::operators::SmoothChannel>("smoothing"); // Generate a smoothing channel
+ //pipeline1->append<ftl::operators::SmoothChannel>("smoothing"); // Generate a smoothing channel
//pipeline1->append<ftl::operators::ScanFieldFill>("filling"); // Generate a smoothing channel
pipeline1->append<ftl::operators::ColourMLS>("mls"); // Perform MLS (using smoothing channel)
+ pipeline1->append<ftl::operators::CrossSupport>("cross");
+ pipeline1->append<ftl::operators::VisCrossSupport>("viscross");
// Alignment
diff --git a/components/operators/include/ftl/operators/segmentation.hpp b/components/operators/include/ftl/operators/segmentation.hpp
index dbd601e07..f905ce969 100644
--- a/components/operators/include/ftl/operators/segmentation.hpp
+++ b/components/operators/include/ftl/operators/segmentation.hpp
@@ -20,6 +20,20 @@ class CrossSupport : public ftl::operators::Operator {
};
+/**
+ * Visualise the cross support regions channel.
+ */
+class VisCrossSupport : public ftl::operators::Operator {
+ public:
+ explicit VisCrossSupport(ftl::Configurable*);
+ ~VisCrossSupport();
+
+ 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/segmentation.cpp b/components/operators/src/segmentation.cpp
index 29cab58bf..1ca3d6907 100644
--- a/components/operators/src/segmentation.cpp
+++ b/components/operators/src/segmentation.cpp
@@ -2,6 +2,7 @@
#include "segmentation_cuda.hpp"
using ftl::operators::CrossSupport;
+using ftl::operators::VisCrossSupport;
using ftl::codecs::Channel;
CrossSupport::CrossSupport(ftl::Configurable *cfg) : ftl::operators::Operator(cfg) {
@@ -21,5 +22,26 @@ bool CrossSupport::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd:
config()->value("h_max", 10), 0
);
+ return true;
+}
+
+
+
+
+VisCrossSupport::VisCrossSupport(ftl::Configurable *cfg) : ftl::operators::Operator(cfg) {
+
+}
+
+VisCrossSupport::~VisCrossSupport() {
+
+}
+
+bool VisCrossSupport::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::Source *s, cudaStream_t stream) {
+ ftl::cuda::vis_support_region(
+ in.createTexture<uchar4>(Channel::Colour),
+ in.createTexture<uchar4>(Channel::Colour2),
+ 0
+ );
+
return true;
}
\ No newline at end of file
diff --git a/components/operators/src/segmentation.cu b/components/operators/src/segmentation.cu
index 3dd05eb13..6228180a0 100644
--- a/components/operators/src/segmentation.cu
+++ b/components/operators/src/segmentation.cu
@@ -11,11 +11,11 @@ __device__ inline int cross(uchar4 p1, uchar4 p2) {
__device__ uchar4 calculate_support_region(const TextureObject<uchar4> &img, int x, int y, int tau, int v_max, int h_max) {
int x_min = max(0, x - h_max);
- int x_max = max(img.width()-1, x + h_max);
+ int x_max = min(img.width()-1, x + h_max);
int y_min = max(0, y - v_max);
- int y_max = max(img.height()-1, y + v_max);
+ int y_max = min(img.height()-1, y + v_max);
- uchar4 result = make_uchar4(x - x_min, x_max - x, y - y_min, y_max - y);
+ uchar4 result = make_uchar4(x - x_min, x_max - x, y - y_min, y_max - y);
uchar4 colour = img.tex2D(x,y);
@@ -28,7 +28,7 @@ __device__ uchar4 calculate_support_region(const TextureObject<uchar4> &img, int
for (int u=x+1; u <= x_max; ++u) {
if (cross(colour, img.tex2D(u,y)) > tau) {
- result.y = x - u;
+ result.y = u - x;
break;
}
}
@@ -42,7 +42,7 @@ __device__ uchar4 calculate_support_region(const TextureObject<uchar4> &img, int
for (int v=y+1; v <= y_max; ++v) {
if (cross(colour, img.tex2D(x,v)) > tau) {
- result.w = y - v;
+ result.w = v - y;
break;
}
}
@@ -74,6 +74,44 @@ void ftl::cuda::support_region(
cudaSafeCall( cudaGetLastError() );
+ #ifdef _DEBUG
+ cudaSafeCall(cudaDeviceSynchronize());
+ #endif
+}
+
+__global__ void vis_support_region_kernel(TextureObject<uchar4> colour, TextureObject<uchar4> region) {
+ const int x = blockIdx.x*blockDim.x + threadIdx.x;
+ const int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+ if (x < 0 || y < 0 || x >= colour.width() || y >= colour.height()) return;
+
+ // Grid pattern
+ if (x % 50 != 0 || y % 50 != 0) return;
+
+ uchar4 base = region.tex2D(x,y);
+
+ for (int v=-base.z; v<=base.w; ++v) {
+ uchar4 baseY = region.tex2D(x,y+v);
+
+ for (int u=-baseY.x; u<=baseY.y; ++u) {
+ if (x+u < 0 || y+v < 0 || x+u >= colour.width() || y+v >= colour.height()) continue;
+ colour(x+u, y+v) = make_uchar4(255,0,0,0);
+ }
+ }
+}
+
+void ftl::cuda::vis_support_region(
+ ftl::cuda::TextureObject<uchar4> &colour,
+ ftl::cuda::TextureObject<uchar4> ®ion,
+ cudaStream_t stream) {
+
+ const dim3 gridSize((region.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (region.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
+ const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
+
+ vis_support_region_kernel<<<gridSize, blockSize, 0, stream>>>(colour, region);
+ cudaSafeCall( cudaGetLastError() );
+
+
#ifdef _DEBUG
cudaSafeCall(cudaDeviceSynchronize());
#endif
diff --git a/components/operators/src/segmentation_cuda.hpp b/components/operators/src/segmentation_cuda.hpp
index ef8a2458c..566b9d033 100644
--- a/components/operators/src/segmentation_cuda.hpp
+++ b/components/operators/src/segmentation_cuda.hpp
@@ -12,6 +12,11 @@ void support_region(
int tau, int v_max, int h_max,
cudaStream_t stream);
+void vis_support_region(
+ ftl::cuda::TextureObject<uchar4> &colour,
+ ftl::cuda::TextureObject<uchar4> ®ion,
+ cudaStream_t stream);
+
}
}
--
GitLab