From 5d30c1142f9b79121cfd423686f027f79fa1fab6 Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Wed, 2 Oct 2019 18:16:09 +0300
Subject: [PATCH] Cull disconts at render instead
---
applications/reconstruct/src/ilw/ilw.cpp | 2 +-
.../renderers/cpp/include/ftl/cuda/points.hpp | 2 +-
components/renderers/cpp/src/points.cu | 33 ++++++++++++++++---
components/renderers/cpp/src/splat_render.cpp | 2 +-
components/renderers/cpp/src/splatter.cu | 6 ++--
5 files changed, 35 insertions(+), 10 deletions(-)
diff --git a/applications/reconstruct/src/ilw/ilw.cpp b/applications/reconstruct/src/ilw/ilw.cpp
index 65e1a19df..778a3d36e 100644
--- a/applications/reconstruct/src/ilw/ilw.cpp
+++ b/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
diff --git a/components/renderers/cpp/include/ftl/cuda/points.hpp b/components/renderers/cpp/include/ftl/cuda/points.hpp
index cb996c4ec..a705e39e2 100644
--- a/components/renderers/cpp/include/ftl/cuda/points.hpp
+++ b/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 ¶ms,
- 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);
diff --git a/components/renderers/cpp/src/points.cu b/components/renderers/cpp/src/points.cu
index abcb49f2e..b2e83577d 100644
--- a/components/renderers/cpp/src/points.cu
+++ b/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 ¶ms, 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 ¶ms, 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
diff --git a/components/renderers/cpp/src/splat_render.cpp b/components/renderers/cpp/src/splat_render.cpp
index f5b5a4625..68d60599f 100644
--- a/components/renderers/cpp/src/splat_render.cpp
+++ b/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";
}
diff --git a/components/renderers/cpp/src/splatter.cu b/components/renderers/cpp/src/splatter.cu
index ddea639bd..41777d49d 100644
--- a/components/renderers/cpp/src/splatter.cu
+++ b/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;
--
GitLab