From 0151f8d8aaa80eb674b10d3ab47dcf946dec44ee Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Thu, 7 Nov 2019 15:46:55 +0200
Subject: [PATCH] Cross support region code
---
components/operators/CMakeLists.txt | 2 +
.../include/ftl/operators/segmentation.hpp | 26 ++++++
components/operators/src/segmentation.cpp | 25 ++++++
components/operators/src/segmentation.cu | 80 +++++++++++++++++++
.../operators/src/segmentation_cuda.hpp | 18 +++++
5 files changed, 151 insertions(+)
create mode 100644 components/operators/include/ftl/operators/segmentation.hpp
create mode 100644 components/operators/src/segmentation.cpp
create mode 100644 components/operators/src/segmentation.cu
create mode 100644 components/operators/src/segmentation_cuda.hpp
diff --git a/components/operators/CMakeLists.txt b/components/operators/CMakeLists.txt
index d946fae15..a1f4fc1eb 100644
--- a/components/operators/CMakeLists.txt
+++ b/components/operators/CMakeLists.txt
@@ -8,6 +8,8 @@ add_library(ftloperators
src/normals.cpp
src/filling.cpp
src/filling.cu
+ src/segmentation.cu
+ src/segmentation.cpp
)
# These cause errors in CI build and are being removed from PCL in newer versions
diff --git a/components/operators/include/ftl/operators/segmentation.hpp b/components/operators/include/ftl/operators/segmentation.hpp
new file mode 100644
index 000000000..dbd601e07
--- /dev/null
+++ b/components/operators/include/ftl/operators/segmentation.hpp
@@ -0,0 +1,26 @@
+#ifndef _FTL_OPERATORS_SEGMENTATION_HPP_
+#define _FTL_OPERATORS_SEGMENTATION_HPP_
+
+#include <ftl/operators/operator.hpp>
+
+namespace ftl {
+namespace operators {
+
+/**
+ * Generate the cross support regions channel.
+ */
+class CrossSupport : public ftl::operators::Operator {
+ public:
+ explicit CrossSupport(ftl::Configurable*);
+ ~CrossSupport();
+
+ 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;
+
+};
+
+}
+}
+
+#endif // _FTL_OPERATORS_SEGMENTATION_HPP_
diff --git a/components/operators/src/segmentation.cpp b/components/operators/src/segmentation.cpp
new file mode 100644
index 000000000..29cab58bf
--- /dev/null
+++ b/components/operators/src/segmentation.cpp
@@ -0,0 +1,25 @@
+#include <ftl/operators/segmentation.hpp>
+#include "segmentation_cuda.hpp"
+
+using ftl::operators::CrossSupport;
+using ftl::codecs::Channel;
+
+CrossSupport::CrossSupport(ftl::Configurable *cfg) : ftl::operators::Operator(cfg) {
+
+}
+
+CrossSupport::~CrossSupport() {
+
+}
+
+bool CrossSupport::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::Source *s, cudaStream_t stream) {
+ ftl::cuda::support_region(
+ in.createTexture<uchar4>(Channel::Colour),
+ out.createTexture<uchar4>(Channel::Colour2, ftl::rgbd::Format<uchar4>(in.get<cv::cuda::GpuMat>(Channel::Colour).size())),
+ config()->value("tau", 5),
+ config()->value("v_max", 10),
+ config()->value("h_max", 10), 0
+ );
+
+ return true;
+}
\ No newline at end of file
diff --git a/components/operators/src/segmentation.cu b/components/operators/src/segmentation.cu
new file mode 100644
index 000000000..3dd05eb13
--- /dev/null
+++ b/components/operators/src/segmentation.cu
@@ -0,0 +1,80 @@
+#include "segmentation_cuda.hpp"
+
+#define T_PER_BLOCK 8
+
+using ftl::cuda::TextureObject;
+
+
+__device__ inline int cross(uchar4 p1, uchar4 p2) {
+ return max(max(__sad(p1.x,p2.x,0),__sad(p1.y,p2.y,0)), __sad(p1.z,p2.z,0));
+}
+
+__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 y_min = max(0, y - v_max);
+ int y_max = max(img.height()-1, y + v_max);
+
+ uchar4 result = make_uchar4(x - x_min, x_max - x, y - y_min, y_max - y);
+
+ uchar4 colour = img.tex2D(x,y);
+
+ for (int u=x-1; u >= x_min; --u) {
+ if (cross(colour, img.tex2D(u,y)) > tau) {
+ result.x = x - u;
+ break;
+ }
+ }
+
+ for (int u=x+1; u <= x_max; ++u) {
+ if (cross(colour, img.tex2D(u,y)) > tau) {
+ result.y = x - u;
+ break;
+ }
+ }
+
+ for (int v=y-1; v >= y_min; --v) {
+ if (cross(colour, img.tex2D(x,v)) > tau) {
+ result.z = y - v;
+ break;
+ }
+ }
+
+ for (int v=y+1; v <= y_max; ++v) {
+ if (cross(colour, img.tex2D(x,v)) > tau) {
+ result.w = y - v;
+ break;
+ }
+ }
+
+ return result;
+}
+
+__global__ void support_region_kernel(TextureObject<uchar4> colour, TextureObject<uchar4> region, int tau, int v_max, int h_max) {
+ 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;
+
+ region(x,y) = calculate_support_region(colour, x, y, tau, v_max, h_max);
+}
+
+void ftl::cuda::support_region(
+ ftl::cuda::TextureObject<uchar4> &colour,
+ ftl::cuda::TextureObject<uchar4> ®ion,
+ int tau,
+ int v_max,
+ int h_max,
+ 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);
+
+ support_region_kernel<<<gridSize, blockSize, 0, stream>>>(colour, region, tau, v_max, h_max);
+ cudaSafeCall( cudaGetLastError() );
+
+
+ #ifdef _DEBUG
+ cudaSafeCall(cudaDeviceSynchronize());
+ #endif
+}
diff --git a/components/operators/src/segmentation_cuda.hpp b/components/operators/src/segmentation_cuda.hpp
new file mode 100644
index 000000000..ef8a2458c
--- /dev/null
+++ b/components/operators/src/segmentation_cuda.hpp
@@ -0,0 +1,18 @@
+#ifndef _FTL_CUDA_SEGMENTATION_HPP_
+#define _FTL_CUDA_SEGMENTATION_HPP_
+
+#include <ftl/cuda_common.hpp>
+
+namespace ftl {
+namespace cuda {
+
+void support_region(
+ ftl::cuda::TextureObject<uchar4> &colour,
+ ftl::cuda::TextureObject<uchar4> ®ion,
+ int tau, int v_max, int h_max,
+ cudaStream_t stream);
+
+}
+}
+
+#endif
--
GitLab