From 297f0c06706c8c4d268d1404a28becdd0164e1cd Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Tue, 29 Oct 2019 20:16:07 +0200
Subject: [PATCH] Allow for point cloud mode
---
components/renderers/cpp/src/reprojection.cu | 101 ++++++++++++++++++
.../renderers/cpp/src/splatter_cuda.hpp | 11 ++
components/renderers/cpp/src/tri_render.cpp | 36 +++++--
3 files changed, 137 insertions(+), 11 deletions(-)
diff --git a/components/renderers/cpp/src/reprojection.cu b/components/renderers/cpp/src/reprojection.cu
index c157f6fe6..c35edd4c9 100644
--- a/components/renderers/cpp/src/reprojection.cu
+++ b/components/renderers/cpp/src/reprojection.cu
@@ -171,3 +171,104 @@ template void ftl::cuda::reproject(
const ftl::render::SplatParams ¶ms,
const ftl::rgbd::Camera &camera,
const float4x4 &poseInv, cudaStream_t stream);
+
+//==============================================================================
+// Without normals
+//==============================================================================
+
+/*
+ * Pass 2: Accumulate attribute contributions if the points pass a visibility test.
+ */
+ template <typename A, typename B>
+__global__ void reprojection_kernel(
+ TextureObject<A> in, // Attribute input
+ TextureObject<float> depth_src,
+ TextureObject<int> depth_in, // Virtual depth map
+ TextureObject<B> out, // Accumulated output
+ TextureObject<float> contrib,
+ SplatParams params,
+ Camera camera, float4x4 poseInv) {
+
+ const int x = (blockIdx.x*blockDim.x + threadIdx.x);
+ const int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+ const float d = (float)depth_in.tex2D((int)x, (int)y) / 1000.0f;
+ if (d < params.camera.minDepth || d > params.camera.maxDepth) return;
+
+ const float3 worldPos = params.m_viewMatrixInverse * params.camera.screenToCam(x, y, d);
+ //if (worldPos.x == MINF || (!(params.m_flags & ftl::render::kShowDisconMask) && worldPos.w < 0.0f)) return;
+
+ const float3 camPos = poseInv * worldPos;
+ if (camPos.z < camera.minDepth) return;
+ if (camPos.z > camera.maxDepth) return;
+ const uint2 screenPos = camera.camToScreen<uint2>(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);
+ float weight = ftl::cuda::weighting(fabs(camPos.z - d2), 0.02f);
+ const B weighted = make<B>(input) * weight;
+
+ if (weight > 0.0f) {
+ accumulateOutput(out, contrib, make_uint2(x,y), weighted, weight);
+ //out(screenPos.x, screenPos.y) = input;
+ }
+}
+
+
+template <typename A, typename B>
+void ftl::cuda::reproject(
+ TextureObject<A> &in,
+ TextureObject<float> &depth_src, // Original 3D points
+ TextureObject<int> &depth_in, // Virtual depth map
+ TextureObject<B> &out, // Accumulated output
+ TextureObject<float> &contrib,
+ const SplatParams ¶ms,
+ const Camera &camera, const float4x4 &poseInv, cudaStream_t stream) {
+ const dim3 gridSize((out.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (out.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
+ const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
+
+ reprojection_kernel<<<gridSize, blockSize, 0, stream>>>(
+ in,
+ depth_src,
+ depth_in,
+ out,
+ contrib,
+ params,
+ camera,
+ poseInv
+ );
+ cudaSafeCall( cudaGetLastError() );
+}
+
+template void ftl::cuda::reproject(
+ ftl::cuda::TextureObject<uchar4> &in, // Original colour image
+ ftl::cuda::TextureObject<float> &depth_src, // Original 3D points
+ ftl::cuda::TextureObject<int> &depth_in, // Virtual depth map
+ ftl::cuda::TextureObject<float4> &out, // Accumulated output
+ ftl::cuda::TextureObject<float> &contrib,
+ const ftl::render::SplatParams ¶ms,
+ const ftl::rgbd::Camera &camera,
+ const float4x4 &poseInv, cudaStream_t stream);
+
+template void ftl::cuda::reproject(
+ ftl::cuda::TextureObject<float> &in, // Original colour image
+ ftl::cuda::TextureObject<float> &depth_src, // Original 3D points
+ ftl::cuda::TextureObject<int> &depth_in, // Virtual depth map
+ ftl::cuda::TextureObject<float> &out, // Accumulated output
+ ftl::cuda::TextureObject<float> &contrib,
+ const ftl::render::SplatParams ¶ms,
+ const ftl::rgbd::Camera &camera,
+ const float4x4 &poseInv, cudaStream_t stream);
+
+template void ftl::cuda::reproject(
+ ftl::cuda::TextureObject<float4> &in, // Original colour image
+ ftl::cuda::TextureObject<float> &depth_src, // Original 3D points
+ ftl::cuda::TextureObject<int> &depth_in, // Virtual depth map
+ ftl::cuda::TextureObject<float4> &out, // Accumulated output
+ ftl::cuda::TextureObject<float> &contrib,
+ const ftl::render::SplatParams ¶ms,
+ const ftl::rgbd::Camera &camera,
+ const float4x4 &poseInv, cudaStream_t stream);
diff --git a/components/renderers/cpp/src/splatter_cuda.hpp b/components/renderers/cpp/src/splatter_cuda.hpp
index 6e3d22cbf..46ef5012d 100644
--- a/components/renderers/cpp/src/splatter_cuda.hpp
+++ b/components/renderers/cpp/src/splatter_cuda.hpp
@@ -67,6 +67,17 @@ namespace cuda {
const ftl::rgbd::Camera &camera,
const float4x4 &poseInv, cudaStream_t stream);
+ template <typename A, typename B>
+ void reproject(
+ ftl::cuda::TextureObject<A> &in, // Original colour image
+ ftl::cuda::TextureObject<float> &depth_src, // Original 3D points
+ ftl::cuda::TextureObject<int> &depth_in, // Virtual depth map
+ ftl::cuda::TextureObject<B> &out, // Accumulated output
+ ftl::cuda::TextureObject<float> &contrib,
+ const ftl::render::SplatParams ¶ms,
+ const ftl::rgbd::Camera &camera,
+ const float4x4 &poseInv, cudaStream_t stream);
+
template <typename A, typename B>
void dibr_normalise(
ftl::cuda::TextureObject<A> &in,
diff --git a/components/renderers/cpp/src/tri_render.cpp b/components/renderers/cpp/src/tri_render.cpp
index 7f89ec231..972e97fc9 100644
--- a/components/renderers/cpp/src/tri_render.cpp
+++ b/components/renderers/cpp/src/tri_render.cpp
@@ -212,17 +212,31 @@ void Triangular::__reprojectChannel(ftl::rgbd::Frame &output, ftl::codecs::Chann
auto poseInv = MatrixConversion::toCUDA(s->getPose().cast<float>().inverse());
- ftl::cuda::reproject(
- f.createTexture<T>(in),
- f.createTexture<float>(Channel::Depth), // TODO: Use depth?
- temp_.getTexture<int>(Channel::Depth2),
- accum_.getTexture<float4>(Channel::Normals),
- temp_.createTexture<typename AccumSelector<T>::type>(AccumSelector<T>::channel),
- temp_.getTexture<float>(Channel::Contribution),
- params_,
- s->parameters(),
- poseInv, stream
- );
+ if (splat_) {
+ ftl::cuda::reproject(
+ f.createTexture<T>(in),
+ f.createTexture<float>(Channel::Depth), // TODO: Use depth?
+ temp_.getTexture<int>(Channel::Depth2),
+ accum_.getTexture<float4>(Channel::Normals),
+ temp_.createTexture<typename AccumSelector<T>::type>(AccumSelector<T>::channel),
+ temp_.getTexture<float>(Channel::Contribution),
+ params_,
+ s->parameters(),
+ poseInv, stream
+ );
+ } else {
+ // Can't use normals with point cloud version
+ ftl::cuda::reproject(
+ f.createTexture<T>(in),
+ f.createTexture<float>(Channel::Depth), // TODO: Use depth?
+ temp_.getTexture<int>(Channel::Depth2),
+ temp_.createTexture<typename AccumSelector<T>::type>(AccumSelector<T>::channel),
+ temp_.getTexture<float>(Channel::Contribution),
+ params_,
+ s->parameters(),
+ poseInv, stream
+ );
+ }
}
ftl::cuda::dibr_normalise(
--
GitLab