Allow any radius

components/operators/src/mask.cu

@@ -4,8 +4,7 @@
 
 using ftl::cuda::Mask;
 
-template <int RADIUS>
-__global__ void discontinuity_kernel(ftl::cuda::TextureObject<int> mask_out, ftl::cuda::TextureObject<uchar4> support, ftl::cuda::TextureObject<float> depth, ftl::rgbd::Camera params, float threshold) {
+__global__ void discontinuity_kernel(ftl::cuda::TextureObject<int> mask_out, ftl::cuda::TextureObject<uchar4> support, ftl::cuda::TextureObject<float> depth, ftl::rgbd::Camera params, float threshold, int radius) {
 	const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
 	const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
 
@@ -28,20 +27,20 @@ __global__ void discontinuity_kernel(ftl::cuda::TextureObject<int> mask_out, ftl
 			//}
 
 			const uchar4 sup = support.tex2D((int)x, (int)y);
-			if (sup.x <= RADIUS) {
-				float dS = depth.tex2D((int)x - sup.x - RADIUS, (int)y);
+			if (sup.x <= radius) {
+				float dS = depth.tex2D((int)x - sup.x - radius, (int)y);
 				if (fabs(dS - d) > threshold) mask.isDiscontinuity(true);
 			}
-			if (sup.y <= RADIUS) {
-				float dS = depth.tex2D((int)x + sup.y + RADIUS, (int)y);
+			if (sup.y <= radius) {
+				float dS = depth.tex2D((int)x + sup.y + radius, (int)y);
 				if (fabs(dS - d) > threshold) mask.isDiscontinuity(true);
 			}
-			if (sup.z <= RADIUS) {
-				float dS = depth.tex2D((int)x, (int)y - sup.z - RADIUS);
+			if (sup.z <= radius) {
+				float dS = depth.tex2D((int)x, (int)y - sup.z - radius);
 				if (fabs(dS - d) > threshold) mask.isDiscontinuity(true);
 			}
-			if (sup.w <= RADIUS) {
-				float dS = depth.tex2D((int)x, (int)y + sup.w + RADIUS);
+			if (sup.w <= radius) {
+				float dS = depth.tex2D((int)x, (int)y + sup.w + radius);
 				if (fabs(dS - d) > threshold) mask.isDiscontinuity(true);
 			}
         }
@@ -54,14 +53,7 @@ void ftl::cuda::discontinuity(ftl::cuda::TextureObject<int> &mask_out, ftl::cuda
 	const dim3 gridSize((params.width + T_PER_BLOCK - 1)/T_PER_BLOCK, (params.height + T_PER_BLOCK - 1)/T_PER_BLOCK);
 	const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
 
-	switch (discon) {
-    case 5 :	discontinuity_kernel<5><<<gridSize, blockSize, 0, stream>>>(mask_out, support, depth, params, thresh); break;
-	case 4 :	discontinuity_kernel<4><<<gridSize, blockSize, 0, stream>>>(mask_out, support, depth, params, thresh); break;
-	case 3 :	discontinuity_kernel<3><<<gridSize, blockSize, 0, stream>>>(mask_out, support, depth, params, thresh); break;
-	case 2 :	discontinuity_kernel<2><<<gridSize, blockSize, 0, stream>>>(mask_out, support, depth, params, thresh); break;
-	case 1 :	discontinuity_kernel<1><<<gridSize, blockSize, 0, stream>>>(mask_out, support, depth, params, thresh); break;
-	default:	break;
-	}
+    discontinuity_kernel<<<gridSize, blockSize, 0, stream>>>(mask_out, support, depth, params, thresh, discon);
 	cudaSafeCall( cudaGetLastError() );
 
 #ifdef _DEBUG
@@ -77,7 +69,7 @@ __global__ void cull_discontinuity_kernel(ftl::cuda::TextureObject<int> mask, ft
 
 	if (x < depth.width() && y < depth.height()) {
 		Mask m(mask.tex2D((int)x,(int)y));
-		if (m.isDiscontinuity()) depth(x,y) = 0.0f;
+		if (m.isDiscontinuity()) depth(x,y) = 1000.0f;
 	}
 }