Cross support region code

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

components/operators/include/ftl/operators/segmentation.hpp

+#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_

components/operators/src/segmentation.cpp

+#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

components/operators/src/segmentation.cu

+#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> &region,
+        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
+}

components/operators/src/segmentation_cuda.hpp

+#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> &region,
+		int tau, int v_max, int h_max,
+		cudaStream_t stream);
+
+}
+}
+
+#endif