Working support regions

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
 
 

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;
+
+};
+
 }
 }
 

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

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

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> &region,
+        cudaStream_t stream);
+
 }
 }