Cull disconts at render instead

applications/reconstruct/src/ilw/ilw.cpp

@@ -116,7 +116,7 @@ bool ILW::_phase0(ftl::rgbd::FrameSet &fs, cudaStream_t stream) {
 			
         auto &t = f.createTexture<float4>(Channel::Points, Format<float4>(f.get<GpuMat>(Channel::Colour).size()));
         auto pose = MatrixConversion::toCUDA(s->getPose().cast<float>()); //.inverse());
-        ftl::cuda::point_cloud(t, f.createTexture<float>(Channel::Depth), s->parameters(), pose, stream);
+        ftl::cuda::point_cloud(t, f.createTexture<float>(Channel::Depth), s->parameters(), pose, 2, stream);
 
         // TODO: Create energy vector texture and clear it
         // Create energy and clear it

components/renderers/cpp/include/ftl/cuda/points.hpp

@@ -16,7 +16,7 @@ struct ClipSpace {
 void point_cloud(ftl::cuda::TextureObject<float4> &output,
 		ftl::cuda::TextureObject<float> &depth,
 		const ftl::rgbd::Camera &params,
-		const float4x4 &pose, cudaStream_t stream);
+		const float4x4 &pose, uint discon, cudaStream_t stream);
 
 void clipping(ftl::cuda::TextureObject<float4> &points,
 		const ClipSpace &clip, cudaStream_t stream);

components/renderers/cpp/src/points.cu

@@ -11,26 +11,51 @@ __global__ void point_cloud_kernel(ftl::cuda::TextureObject<float4> output, ftl:
 	if (x < params.width && y < params.height) {
 		output(x,y) = make_float4(MINF, MINF, MINF, MINF);
 
-		float d = depth.tex2D((int)x, (int)y);
+		const float d = depth.tex2D((int)x, (int)y);
+		float p = d;
 
 		if (d >= params.minDepth && d <= params.maxDepth) {
+			/* Orts-Escolano S. et al. 2016. Holoportation: Virtual 3D teleportation in real-time. */
 			// Is there a discontinuity nearby?
 			for (int u=-RADIUS; u<=RADIUS; ++u) {
 				for (int v=-RADIUS; v<=RADIUS; ++v) {
-					if (fabs(depth.tex2D((int)x+u, (int)y+v) - d) > 0.1f) return;
+					// If yes, the flag using w = -1
+					if (fabs(depth.tex2D((int)x+u, (int)y+v) - d) > 0.1f) p = -1.0f;
+				}
+			}
+
+			output(x,y) = make_float4(pose * params.screenToCam(x, y, d), p);
 		}
 	}
+}
+
+template <>
+__global__ void point_cloud_kernel<0>(ftl::cuda::TextureObject<float4> output, ftl::cuda::TextureObject<float> depth, ftl::rgbd::Camera params, float4x4 pose)
+{
+	const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
+	const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+	if (x < params.width && y < params.height) {
+		output(x,y) = make_float4(MINF, MINF, MINF, MINF);
 
+		float d = depth.tex2D((int)x, (int)y);
+
+		if (d >= params.minDepth && d <= params.maxDepth) {
 			output(x,y) = make_float4(pose * params.screenToCam(x, y, d), d);
 		}
 	}
 }
 
-void ftl::cuda::point_cloud(ftl::cuda::TextureObject<float4> &output, ftl::cuda::TextureObject<float> &depth, const ftl::rgbd::Camera &params, const float4x4 &pose, cudaStream_t stream) {
+void ftl::cuda::point_cloud(ftl::cuda::TextureObject<float4> &output, ftl::cuda::TextureObject<float> &depth, const ftl::rgbd::Camera &params, const float4x4 &pose, uint discon, cudaStream_t stream) {
 	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);
 
-	point_cloud_kernel<3><<<gridSize, blockSize, 0, stream>>>(output, depth, params, pose);
+	switch (discon) {
+	case 3 :	point_cloud_kernel<3><<<gridSize, blockSize, 0, stream>>>(output, depth, params, pose); break;
+	case 2 :	point_cloud_kernel<2><<<gridSize, blockSize, 0, stream>>>(output, depth, params, pose); break;
+	case 1 :	point_cloud_kernel<1><<<gridSize, blockSize, 0, stream>>>(output, depth, params, pose); break;
+	default:	point_cloud_kernel<0><<<gridSize, blockSize, 0, stream>>>(output, depth, params, pose);
+	}
 	cudaSafeCall( cudaGetLastError() );
 
 #ifdef _DEBUG

components/renderers/cpp/src/splat_render.cpp

@@ -91,7 +91,7 @@ void Splatter::renderChannel(
 			
 			auto &t = f.createTexture<float4>(Channel::Points, Format<float4>(f.get<GpuMat>(Channel::Colour).size()));
 			auto pose = MatrixConversion::toCUDA(s->getPose().cast<float>()); //.inverse());
-			ftl::cuda::point_cloud(t, f.createTexture<float>(Channel::Depth), s->parameters(), pose, stream);
+			ftl::cuda::point_cloud(t, f.createTexture<float>(Channel::Depth), s->parameters(), pose, 0, stream);
 
 			//LOG(INFO) << "POINTS Added";
 		}

components/renderers/cpp/src/splatter.cu

@@ -24,11 +24,11 @@ using ftl::render::SplatParams;
 	const int x = blockIdx.x*blockDim.x + threadIdx.x;
 	const int y = blockIdx.y*blockDim.y + threadIdx.y;
 
-	const float3 worldPos = make_float3(points.tex2D(x, y));
-	if (worldPos.x == MINF) return;
+	const float4 worldPos = points.tex2D(x, y);
+	if (worldPos.x == MINF || worldPos.w < 0.0f) return;
 
     // Find the virtual screen position of current point
-	const float3 camPos = params.m_viewMatrix * worldPos;
+	const float3 camPos = params.m_viewMatrix * make_float3(worldPos);
 	if (camPos.z < params.camera.minDepth) return;
 	if (camPos.z > params.camera.maxDepth) return;