Merge branch 'master' into feature/vision-operator

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>
@@ -306,9 +307,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::CrossSupport>("cross")->set("enabled", false);
 	pipeline1->append<ftl::operators::ColourMLS>("mls");  // Perform MLS (using smoothing channel)
+	pipeline1->append<ftl::operators::VisCrossSupport>("viscross")->set("enabled", false);
 	// Alignment
 
 

components/operators/CMakeLists.txt

-add_library(ftloperators
+set(OPERSRC
 	src/smoothing.cpp
 	src/smoothing.cu
 	src/mls.cu
@@ -15,8 +15,16 @@ add_library(ftloperators
 	src/disparity/disparity_to_depth.cpp
 	src/disparity/fixstars_sgm.cpp
 	src/disparity/bilateral_filter.cpp
+	src/segmentation.cu
+	src/segmentation.cpp
 )
 
+if (HAVE_OPTFLOW)
+	list(APPEND OPERSRC src/nvopticalflow.cpp)
+endif()
+
+add_library(ftloperators ${OPERSRC})
+
 # These cause errors in CI build and are being removed from PCL in newer versions
 # target_compile_options(ftlrender PUBLIC ${PCL_DEFINITIONS})
 

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;
+
+};
+
+/**
+ * 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;
+
+};
+
+}
+}
+
+#endif  // _FTL_OPERATORS_SEGMENTATION_HPP_

components/operators/src/mls.cu

@@ -212,6 +212,117 @@ void ftl::cuda::colour_mls_smooth(
 }
 
 
+// ===== Colour MLS using cross support region =================================
+
+/*
+ * Smooth depth map using Moving Least Squares. This version uses colour
+ * similarity weights to adjust the spatial smoothing factor. It is naive in
+ * that similar colours may exist on both sides of an edge and colours at the
+ * level of single pixels can be subject to noise. Colour noise should first
+ * be removed from the image.
+ */
+ __global__ void colour_mls_smooth_csr_kernel(
+	 	TextureObject<uchar4> region,
+		TextureObject<float4> normals_in,
+		TextureObject<float4> normals_out,
+        TextureObject<float> depth_in,        // Virtual depth map
+		TextureObject<float> depth_out,   // Accumulated output
+		TextureObject<uchar4> colour_in,
+		float smoothing,
+		float colour_smoothing,
+        ftl::rgbd::Camera camera) {
+        
+    const int x = blockIdx.x*blockDim.x + threadIdx.x;
+    const int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+    if (x < 0 || y < 0 || x >= depth_in.width() || y >= depth_in.height()) return;
+
+	float3 aX = make_float3(0.0f,0.0f,0.0f);
+	float3 nX = make_float3(0.0f,0.0f,0.0f);
+    float contrib = 0.0f;
+
+	float d0 = depth_in.tex2D(x, y);
+	depth_out(x,y) = d0;
+	normals_out(x,y) = normals_in(x,y);
+	if (d0 < camera.minDepth || d0 > camera.maxDepth) return;
+	float3 X = camera.screenToCam((int)(x),(int)(y),d0);
+
+	uchar4 c0 = colour_in.tex2D(x, y);
+
+    // Neighbourhood
+	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) {
+			const float d = depth_in.tex2D(x+u, y+v);
+			if (d < camera.minDepth || d > camera.maxDepth) continue;
+
+			// Point and normal of neighbour
+			const float3 Xi = camera.screenToCam((int)(x)+u,(int)(y)+v,d);
+			const float3 Ni = make_float3(normals_in.tex2D((int)(x)+u, (int)(y)+v));
+
+			if (Ni.x+Ni.y+Ni.z == 0.0f) continue;
+
+			const uchar4 c = colour_in.tex2D(x+u, y+v);
+			const float cw = ftl::cuda::colourWeighting(c0,c,colour_smoothing);
+
+			// Gauss approx weighting function using point distance
+			const float w = ftl::cuda::spatialWeighting(X,Xi,smoothing*cw);
+
+			aX += Xi*w;
+			nX += Ni*w;
+			contrib += w;
+		}
+	}
+
+	if (contrib == 0.0f) return;
+	
+	nX /= contrib;  // Weighted average normal
+	aX /= contrib;  // Weighted average point (centroid)
+
+	// Signed-Distance Field function
+	float fX = nX.x * (X.x - aX.x) + nX.y * (X.y - aX.y) + nX.z * (X.z - aX.z);
+
+	// Calculate new point using SDF function to adjust depth (and position)
+	X = X - nX * fX;
+	
+	//uint2 screen = camera.camToScreen<uint2>(X);
+
+    //if (screen.x < depth_out.width() && screen.y < depth_out.height()) {
+    //    depth_out(screen.x,screen.y) = X.z;
+	//}
+	depth_out(x,y) = X.z;
+	normals_out(x,y) = make_float4(nX / length(nX), 0.0f);
+}
+
+void ftl::cuda::colour_mls_smooth_csr(
+		ftl::cuda::TextureObject<uchar4> &region,
+		ftl::cuda::TextureObject<float4> &normals_in,
+		ftl::cuda::TextureObject<float4> &normals_out,
+		ftl::cuda::TextureObject<float> &depth_in,
+		ftl::cuda::TextureObject<float> &depth_out,
+		ftl::cuda::TextureObject<uchar4> &colour_in,
+		float smoothing,
+		float colour_smoothing,
+		const ftl::rgbd::Camera &camera,
+		cudaStream_t stream) {
+
+	const dim3 gridSize((depth_out.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth_out.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
+	const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
+
+	colour_mls_smooth_csr_kernel<<<gridSize, blockSize, 0, stream>>>(region, normals_in, normals_out, depth_in, depth_out, colour_in, smoothing, colour_smoothing, camera);
+		
+	cudaSafeCall( cudaGetLastError() );
+
+
+	#ifdef _DEBUG
+	cudaSafeCall(cudaDeviceSynchronize());
+	#endif
+}
+
+
 // ===== Adaptive MLS Smooth =====================================================
 
 /*

components/operators/src/segmentation.cpp

+#include <ftl/operators/segmentation.hpp>
+#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) {
+
+}
+
+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", 5),
+        config()->value("h_max", 5), 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

+#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 = min(img.width()-1, x + h_max);
+    int y_min = max(0, 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 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 = u - x;
+            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 = v - y;
+            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
+}
+
+__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;
+			uchar4 col = colour.tex2D(x+u, y+v);
+			colour(x+u, y+v) = make_uchar4(255,col.y,col.z,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

+#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);
+
+void vis_support_region(
+        ftl::cuda::TextureObject<uchar4> &colour,
+        ftl::cuda::TextureObject<uchar4> &region,
+        cudaStream_t stream);
+
+}
+}
+
+#endif 

components/operators/src/smoothing.cpp

@@ -183,6 +183,7 @@ bool ColourMLS::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::So
 	float col_smooth = config()->value("mls_colour_smoothing", 30.0f);
 	int iters = config()->value("mls_iterations", 10);
 	int radius = config()->value("mls_radius",3);
+	bool crosssup = config()->value("cross_support", false);
 
 	if (!in.hasChannel(Channel::Normals)) {
 		LOG(ERROR) << "Required normals channel missing for MLS";
@@ -191,18 +192,33 @@ bool ColourMLS::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::So
 
 	// FIXME: Assume in and out are the same frame.
 	for (int i=0; i<iters; ++i) {
-		ftl::cuda::colour_mls_smooth(
-			in.createTexture<float4>(Channel::Normals),
-			in.createTexture<float4>(Channel::Points, ftl::rgbd::Format<float4>(in.get<cv::cuda::GpuMat>(Channel::Depth).size())),
-			in.createTexture<float>(Channel::Depth),
-			in.createTexture<float>(Channel::Depth2, ftl::rgbd::Format<float>(in.get<cv::cuda::GpuMat>(Channel::Depth).size())),
-			in.createTexture<uchar4>(Channel::Colour),
-			thresh,
-			col_smooth,
-			radius,
-			s->parameters(),
-			0
-		);
+		if (!crosssup) {
+			ftl::cuda::colour_mls_smooth(
+				in.createTexture<float4>(Channel::Normals),
+				in.createTexture<float4>(Channel::Points, ftl::rgbd::Format<float4>(in.get<cv::cuda::GpuMat>(Channel::Depth).size())),
+				in.createTexture<float>(Channel::Depth),
+				in.createTexture<float>(Channel::Depth2, ftl::rgbd::Format<float>(in.get<cv::cuda::GpuMat>(Channel::Depth).size())),
+				in.createTexture<uchar4>(Channel::Colour),
+				thresh,
+				col_smooth,
+				radius,
+				s->parameters(),
+				0
+			);
+		} else {
+			ftl::cuda::colour_mls_smooth_csr(
+				in.createTexture<uchar4>(Channel::Colour2),
+				in.createTexture<float4>(Channel::Normals),
+				in.createTexture<float4>(Channel::Points, ftl::rgbd::Format<float4>(in.get<cv::cuda::GpuMat>(Channel::Depth).size())),
+				in.createTexture<float>(Channel::Depth),
+				in.createTexture<float>(Channel::Depth2, ftl::rgbd::Format<float>(in.get<cv::cuda::GpuMat>(Channel::Depth).size())),
+				in.createTexture<uchar4>(Channel::Colour),
+				thresh,
+				col_smooth,
+				s->parameters(),
+				0
+			);
+		}
 
 		in.swapChannels(Channel::Depth, Channel::Depth2);
 		in.swapChannels(Channel::Normals, Channel::Points);

components/operators/src/smoothing_cuda.hpp

@@ -29,6 +29,18 @@ void colour_mls_smooth(
 		const ftl::rgbd::Camera &camera,
 		cudaStream_t stream);
 
+void colour_mls_smooth_csr(
+		ftl::cuda::TextureObject<uchar4> &region,
+		ftl::cuda::TextureObject<float4> &normals_in,
+		ftl::cuda::TextureObject<float4> &normals_out,
+		ftl::cuda::TextureObject<float> &depth_in,
+		ftl::cuda::TextureObject<float> &depth_out,
+		ftl::cuda::TextureObject<uchar4> &colour_in,
+		float smoothing,
+		float colour_smoothing,
+		const ftl::rgbd::Camera &camera,
+		cudaStream_t stream);
+
 void adaptive_mls_smooth(
 		ftl::cuda::TextureObject<float4> &normals_in,
 		ftl::cuda::TextureObject<float4> &normals_out,