Merge branch 'feature/255/disconcross' into 'master'

Implements #255 discon support window

Closes #255

See merge request nicolas.pope/ftl!183

applications/reconstruct/src/main.cpp

@@ -331,9 +331,10 @@ static void run(ftl::Configurable *root) {
 	//pipeline1->append<ftl::operators::ScanFieldFill>("filling");  // Generate a smoothing channel
 	pipeline1->append<ftl::operators::CrossSupport>("cross");
 	pipeline1->append<ftl::operators::DiscontinuityMask>("discontinuity");
+	pipeline1->append<ftl::operators::CrossSupport>("cross2")->set("discon_support", true);
 	pipeline1->append<ftl::operators::CullDiscontinuity>("remove_discontinuity");
 	//pipeline1->append<ftl::operators::AggreMLS>("mls");  // Perform MLS (using smoothing channel)
-	pipeline1->append<ftl::operators::VisCrossSupport>("viscross")->set("enabled", false);
+	pipeline1->append<ftl::operators::VisCrossSupport>("viscross"); //->set("enabled", false);
 	pipeline1->append<ftl::operators::MultiViewMLS>("mvmls");
 	// Alignment
 

components/operators/src/segmentation.cpp

@@ -14,18 +14,17 @@ CrossSupport::~CrossSupport() {
 }
 
 bool CrossSupport::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::Source *s, cudaStream_t stream) {
-	bool use_depth = config()->value("depth_region", false);
+	bool use_mask = config()->value("discon_support", false);
 
-	if (use_depth) {
+	if (use_mask) {
 		ftl::cuda::support_region(
-			in.createTexture<float>(Channel::Depth),
+			in.createTexture<int>(Channel::Mask),
 			out.createTexture<uchar4>(Channel::Support2, ftl::rgbd::Format<uchar4>(in.get<cv::cuda::GpuMat>(Channel::Colour).size())),
-			config()->value("depth_tau", 0.04f),
 			config()->value("v_max", 5),
 			config()->value("h_max", 5),
-			config()->value("symmetric", true), stream
+			config()->value("symmetric", false), stream
 		);
-	} //else {
+	} else {
 		ftl::cuda::support_region(
 			in.createTexture<uchar4>(Channel::Colour),
 			out.createTexture<uchar4>(Channel::Support1, ftl::rgbd::Format<uchar4>(in.get<cv::cuda::GpuMat>(Channel::Colour).size())),
@@ -34,7 +33,7 @@ bool CrossSupport::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd:
 			config()->value("h_max", 5),
 			config()->value("symmetric", true), stream
 		);
-	//}
+	}
 
 	return true;
 }

components/operators/src/segmentation.cu

 #include "segmentation_cuda.hpp"
+#include "mask_cuda.hpp"
 
 #define T_PER_BLOCK 8
 
 using ftl::cuda::TextureObject;
+using ftl::cuda::Mask;
 
 template <typename T>
 __device__ inline float cross(T p1, T p2);
@@ -89,6 +91,64 @@ __device__ uchar4 calculate_support_region(const TextureObject<T> &img, int x, i
     return result;
 }
 
+__device__ uchar4 calculate_support_region(const TextureObject<int> &img, int x, int y, 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(0, 0, 0, 0);
+
+	Mask m1(img.tex2D(x,y));
+
+	int u;
+    for (u=x-1; u >= x_min; --u) {
+		Mask m2(img.tex2D(u,y));
+        if (m2.isDiscontinuity()) {
+            result.x = x - u - 1;
+            break;
+		}
+	}
+	if (u < x_min) result.x = x - x_min;
+	
+    for (u=x+1; u <= x_max; ++u) {
+		Mask m2(img.tex2D(u,y));
+        if (m2.isDiscontinuity()) {
+            result.y = u - x - 1;
+            break;
+		}
+	}
+	if (u > x_max) result.y = x_max - x;
+
+	int v;
+    for (v=y-1; v >= y_min; --v) {
+		Mask m2(img.tex2D(x,v));
+        if (m2.isDiscontinuity()) {
+            result.z = y - v - 1;
+            break;
+		}
+	}
+	if (v < y_min) result.z = y - y_min;
+
+    for (v=y+1; v <= y_max; ++v) {
+		Mask m2(img.tex2D(x,v));
+        if (m2.isDiscontinuity()) {
+            result.w = v - y - 1;
+            break;
+		}
+	}
+	if (v > y_max) result.w = y_max - y;
+
+	// Make symetric left/right and up/down
+	if (false) {
+		result.x = min(result.x, result.y);
+		result.y = result.x;
+		result.z = min(result.z, result.w);
+		result.w = result.z;
+	}
+    return result;
+}
+
 template <typename T, bool SYM>
 __global__ void support_region_kernel(TextureObject<T> img, TextureObject<uchar4> region, float tau, int v_max, int h_max) {
     const int x = blockIdx.x*blockDim.x + threadIdx.x;
@@ -99,6 +159,15 @@ __global__ void support_region_kernel(TextureObject<T> img, TextureObject<uchar4
     region(x,y) = calculate_support_region<T,SYM>(img, x, y, tau, v_max, h_max);
 }
 
+__global__ void support_region_kernel(TextureObject<int> img, TextureObject<uchar4> region, 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 >= img.width() || y >= img.height()) return;
+
+    region(x,y) = calculate_support_region(img, x, y, v_max, h_max);
+}
+
 void ftl::cuda::support_region(
         ftl::cuda::TextureObject<uchar4> &colour,
         ftl::cuda::TextureObject<uchar4> &region,
@@ -142,6 +211,26 @@ void ftl::cuda::support_region(
 	#endif
 }
 
+void ftl::cuda::support_region(
+		ftl::cuda::TextureObject<int> &mask,
+		ftl::cuda::TextureObject<uchar4> &region,
+		int v_max,
+		int h_max,
+		bool sym,
+		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>>>(mask, region, v_max, h_max);
+	cudaSafeCall( cudaGetLastError() );
+
+
+	#ifdef _DEBUG
+	cudaSafeCall(cudaDeviceSynchronize());
+	#endif
+}
+
 __global__ void vis_support_region_kernel(TextureObject<uchar4> colour, TextureObject<uchar4> region, uchar4 bcolour, uchar4 acolour,
 		int ox, int oy, int dx, int dy) {
     const int x = blockIdx.x*blockDim.x + threadIdx.x;

components/operators/src/segmentation_cuda.hpp

@@ -18,6 +18,12 @@ void support_region(
 		float tau, int v_max, int h_max, bool sym,
 		cudaStream_t stream);
 
+void support_region(
+		ftl::cuda::TextureObject<int> &mask,
+		ftl::cuda::TextureObject<uchar4> &region,
+		int v_max, int h_max, bool sym,
+		cudaStream_t stream);
+
 void vis_support_region(
         ftl::cuda::TextureObject<uchar4> &colour,
         ftl::cuda::TextureObject<uchar4> &region,