From 8e8ca4f2ff00849127c7c0271d72bf24df91dd3d Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Tue, 29 Oct 2019 11:03:24 +0200
Subject: [PATCH] Fix for use of wrong depth in render

---
 components/renderers/cpp/src/screen.cu         | 6 ++++--
 components/renderers/cpp/src/splat_render.cpp  | 3 ++-
 components/renderers/cpp/src/splatter_cuda.hpp | 1 +
 3 files changed, 7 insertions(+), 3 deletions(-)

diff --git a/components/renderers/cpp/src/screen.cu b/components/renderers/cpp/src/screen.cu
index 91f24c840..8610359d2 100644
--- a/components/renderers/cpp/src/screen.cu
+++ b/components/renderers/cpp/src/screen.cu
@@ -13,6 +13,7 @@ using ftl::render::SplatParams;
  * Convert source screen position to output screen coordinates.
  */
  __global__ void screen_coord_kernel(TextureObject<float> depth,
+        TextureObject<float> depth_out,
 		TextureObject<short2> screen_out, SplatParams params, float4x4 pose, Camera camera) {
 	const int x = blockIdx.x*blockDim.x + threadIdx.x;
 	const int y = blockIdx.y*blockDim.y + threadIdx.y;
@@ -31,12 +32,13 @@ using ftl::render::SplatParams;
     if (camPos.z < params.camera.minDepth || camPos.z > params.camera.maxDepth || screenPos.x >= params.camera.width || screenPos.y >= params.camera.height)
         screenPos = make_uint2(30000,30000);
     screen_out(x,y) = make_short2(screenPos.x, screenPos.y);
+    depth_out(x,y) = camPos.z;
 }
 
-void ftl::cuda::screen_coord(TextureObject<float> &depth, TextureObject<short2> &screen_out, const SplatParams &params, const float4x4 &pose, const Camera &camera, cudaStream_t stream) {
+void ftl::cuda::screen_coord(TextureObject<float> &depth, TextureObject<float> &depth_out, TextureObject<short2> &screen_out, const SplatParams &params, const float4x4 &pose, const Camera &camera, 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);
 
-	screen_coord_kernel<<<gridSize, blockSize, 0, stream>>>(depth, screen_out, params, pose, camera);
+	screen_coord_kernel<<<gridSize, blockSize, 0, stream>>>(depth, depth_out, screen_out, params, pose, camera);
     cudaSafeCall( cudaGetLastError() );
 }
diff --git a/components/renderers/cpp/src/splat_render.cpp b/components/renderers/cpp/src/splat_render.cpp
index 1de452f93..64d987c10 100644
--- a/components/renderers/cpp/src/splat_render.cpp
+++ b/components/renderers/cpp/src/splat_render.cpp
@@ -278,12 +278,13 @@ void Splatter::_dibr(cudaStream_t stream) {
 
 		ftl::cuda::screen_coord(
 			f.createTexture<float>(Channel::Depth),
+			f.createTexture<float>(Channel::Depth2, Format<float>(f.get<GpuMat>(Channel::Depth).size())),
 			f.createTexture<short2>(Channel::Screen, Format<short2>(f.get<GpuMat>(Channel::Depth).size())),
 			params_, pose, s->parameters(), stream
 		);
 
 		ftl::cuda::triangle_render1(
-			f.getTexture<float>(Channel::Depth),
+			f.getTexture<float>(Channel::Depth2),
 			temp_.createTexture<int>(Channel::Depth2),
 			f.getTexture<short2>(Channel::Screen),
 			params_, stream
diff --git a/components/renderers/cpp/src/splatter_cuda.hpp b/components/renderers/cpp/src/splatter_cuda.hpp
index 2f036714a..5e7022529 100644
--- a/components/renderers/cpp/src/splatter_cuda.hpp
+++ b/components/renderers/cpp/src/splatter_cuda.hpp
@@ -8,6 +8,7 @@ namespace ftl {
 namespace cuda {
 	void screen_coord(
 		ftl::cuda::TextureObject<float> &depth,
+		ftl::cuda::TextureObject<float> &depth_out,
 		ftl::cuda::TextureObject<short2> &screen_out,
 		const ftl::render::SplatParams &params,
 		const float4x4 &pose,
-- 
GitLab