From 5509f29fa0183c37cfb8cf90bac471d64a8d5a8b Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Tue, 29 Oct 2019 18:28:00 +0200
Subject: [PATCH] WIP Point cloud view again
---
components/renderers/cpp/src/splatter.cu | 36 +++++++++++++++++++
.../renderers/cpp/src/splatter_cuda.hpp | 6 ++++
components/renderers/cpp/src/tri_render.cpp | 3 +-
3 files changed, 43 insertions(+), 2 deletions(-)
diff --git a/components/renderers/cpp/src/splatter.cu b/components/renderers/cpp/src/splatter.cu
index e79b3d320..6346daa88 100644
--- a/components/renderers/cpp/src/splatter.cu
+++ b/components/renderers/cpp/src/splatter.cu
@@ -75,6 +75,34 @@ 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<float4> points,
+ TextureObject<int> depth, SplatParams params) {
+ const int x = blockIdx.x*blockDim.x + threadIdx.x;
+ const int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+ const float4 worldPos = points.tex2D(x, y);
+ 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(cx,cy), d * 1000.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);
@@ -84,6 +112,14 @@ void ftl::cuda::dibr_merge(TextureObject<float4> &points, TextureObject<float4>
cudaSafeCall( cudaGetLastError() );
}
+void ftl::cuda::dibr_merge(TextureObject<float4> &points, TextureObject<int> &depth, 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>>>(points, depth, params);
+ cudaSafeCall( cudaGetLastError() );
+}
+
//==============================================================================
diff --git a/components/renderers/cpp/src/splatter_cuda.hpp b/components/renderers/cpp/src/splatter_cuda.hpp
index 14fb2dc64..6e3d22cbf 100644
--- a/components/renderers/cpp/src/splatter_cuda.hpp
+++ b/components/renderers/cpp/src/splatter_cuda.hpp
@@ -30,6 +30,12 @@ namespace cuda {
bool culling,
cudaStream_t stream);
+ void dibr_merge(
+ ftl::cuda::TextureObject<float4> &points,
+ ftl::cuda::TextureObject<int> &depth,
+ ftl::render::SplatParams params,
+ cudaStream_t stream);
+
template <typename T>
void splat(
ftl::cuda::TextureObject<float4> &normals,
diff --git a/components/renderers/cpp/src/tri_render.cpp b/components/renderers/cpp/src/tri_render.cpp
index 58e21c367..7f89ec231 100644
--- a/components/renderers/cpp/src/tri_render.cpp
+++ b/components/renderers/cpp/src/tri_render.cpp
@@ -270,9 +270,8 @@ void Triangular::_dibr(cudaStream_t stream) {
ftl::cuda::dibr_merge(
f.createTexture<float4>(Channel::Points),
- f.createTexture<float4>(Channel::Normals),
temp_.createTexture<int>(Channel::Depth2),
- params_, backcull_, stream
+ params_, stream
);
}
}
--
GitLab