Resolves #251 rendering of point clouds

components/renderers/cpp/src/reprojection.cu

@@ -203,13 +203,13 @@ __global__ void reprojection_kernel(
 	const float3 camPos = poseInv * worldPos;
 	if (camPos.z < camera.minDepth) return;
 	if (camPos.z > camera.maxDepth) return;
-	const uint2 screenPos = camera.camToScreen<uint2>(camPos);
+	const float2 screenPos = camera.camToScreen<float2>(camPos);
 
 	// Not on screen so stop now...
 	if (screenPos.x >= depth_src.width() || screenPos.y >= depth_src.height()) return;
     
-	const float d2 = depth_src.tex2D((int)screenPos.x, (int)screenPos.y);
-	const A input = in.tex2D((int)screenPos.x, (int)screenPos.y); //generateInput(in.tex2D((int)screenPos.x, (int)screenPos.y), params, worldPos);
+	const float d2 = depth_src.tex2D((int)(screenPos.x+0.5f), (int)(screenPos.y+0.5f));
+	const auto input = in.tex2D(screenPos.x, screenPos.y); //generateInput(in.tex2D((int)screenPos.x, (int)screenPos.y), params, worldPos);
 	float weight = ftl::cuda::weighting(fabs(camPos.z - d2), 0.02f);
 	const B weighted = make<B>(input) * weight;
 

components/renderers/cpp/src/splatter.cu

@@ -103,6 +103,40 @@ using ftl::cuda::warpSum;
 	}
 }
 
+/*
+ * Pass 1: Directly render each camera into virtual view but with no upsampling
+ * for sparse points.
+ */
+ __global__ void dibr_merge_kernel(TextureObject<float> depth,
+		TextureObject<int> depth_out,
+		float4x4 transform,
+		ftl::rgbd::Camera cam,
+		SplatParams params) {
+	const int x = blockIdx.x*blockDim.x + threadIdx.x;
+	const int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+	const float d0 = depth.tex2D(x, y);
+	if (d0 <= cam.minDepth || d0 >= cam.maxDepth) return;
+
+	const float3 camPos = transform * cam.screenToCam(x,y,d0);
+	//if (worldPos.x == MINF || (!(params.m_flags & ftl::render::kShowDisconMask) && worldPos.w < 0.0f)) return;
+
+	// Find the virtual screen position of current point
+	//const float3 camPos = params.m_viewMatrix * make_float3(worldPos);
+	//if (camPos.z < params.camera.minDepth) return;
+	//if (camPos.z > params.camera.maxDepth) return;
+
+	const float d = camPos.z;
+
+	const uint2 screenPos = params.camera.camToScreen<uint2>(camPos);
+	const unsigned int cx = screenPos.x;
+	const unsigned int cy = screenPos.y;
+	if (d > params.camera.minDepth && d < params.camera.maxDepth && cx < depth.width() && cy < depth.height()) {
+		// Transform estimated point to virtual cam space and output z
+		atomicMin(&depth_out(cx,cy), d * 100000.0f);
+	}
+}
+
 void ftl::cuda::dibr_merge(TextureObject<float4> &points, TextureObject<float4> &normals, TextureObject<int> &depth, SplatParams params, bool culling, cudaStream_t stream) {
     const dim3 gridSize((depth.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
     const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
@@ -120,6 +154,14 @@ void ftl::cuda::dibr_merge(TextureObject<float4> &points, TextureObject<int> &de
     cudaSafeCall( cudaGetLastError() );
 }
 
+void ftl::cuda::dibr_merge(TextureObject<float> &depth, TextureObject<int> &depth_out, const float4x4 &transform, const ftl::rgbd::Camera &cam, SplatParams params, cudaStream_t stream) {
+    const dim3 gridSize((depth.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
+    const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
+
+	dibr_merge_kernel<<<gridSize, blockSize, 0, stream>>>(depth, depth_out, transform, cam, params);
+    cudaSafeCall( cudaGetLastError() );
+}
+
 //==============================================================================
 
 

components/renderers/cpp/src/splatter_cuda.hpp

@@ -43,6 +43,14 @@ namespace cuda {
 		ftl::render::SplatParams params,
 		cudaStream_t stream);
 
+	void dibr_merge(
+		ftl::cuda::TextureObject<float> &depth,
+		ftl::cuda::TextureObject<int> &depth_out,
+		const float4x4 &transform,
+		const ftl::rgbd::Camera &cam,
+		ftl::render::SplatParams params,
+		cudaStream_t stream);
+
 	template <typename T>
 	void splat(
         ftl::cuda::TextureObject<float4> &normals,

components/renderers/cpp/src/tri_render.cpp

@@ -300,9 +300,13 @@ void Triangular::_dibr(ftl::rgbd::Frame &out, cudaStream_t stream) {
 			continue;
 		}
 
+		auto transform = params_.m_viewMatrix * MatrixConversion::toCUDA(s->getPose().cast<float>());
+
 		ftl::cuda::dibr_merge(
-			f.createTexture<float4>(Channel::Points),
+			f.createTexture<float>(Channel::Depth),
 			temp_.createTexture<int>(Channel::Depth2),
+			transform,
+			s->parameters(),
 			params_, stream
 		);
 	}