From 3b6c8137feea6a4030b901747d11885206258311 Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Wed, 2 Oct 2019 21:01:15 +0300
Subject: [PATCH] Allow config of back point cull
---
.../cpp/include/ftl/render/splat_render.hpp | 1 +
components/renderers/cpp/src/splat_render.cpp | 7 ++++-
components/renderers/cpp/src/splatter.cu | 31 +++++++++++--------
.../renderers/cpp/src/splatter_cuda.hpp | 1 +
4 files changed, 26 insertions(+), 14 deletions(-)
diff --git a/components/renderers/cpp/include/ftl/render/splat_render.hpp b/components/renderers/cpp/include/ftl/render/splat_render.hpp
index bca5a232b..40bfa459b 100644
--- a/components/renderers/cpp/include/ftl/render/splat_render.hpp
+++ b/components/renderers/cpp/include/ftl/render/splat_render.hpp
@@ -44,6 +44,7 @@ class Splatter : public ftl::render::Renderer {
ftl::cuda::ClipSpace clip_;
bool clipping_;
float norm_filter_;
+ bool backcull_;
};
}
diff --git a/components/renderers/cpp/src/splat_render.cpp b/components/renderers/cpp/src/splat_render.cpp
index 3d6e95e2c..e66c11b65 100644
--- a/components/renderers/cpp/src/splat_render.cpp
+++ b/components/renderers/cpp/src/splat_render.cpp
@@ -54,6 +54,11 @@ Splatter::Splatter(nlohmann::json &config, ftl::rgbd::FrameSet *fs) : ftl::rende
on("normal_filter", [this](const ftl::config::Event &e) {
norm_filter_ = value("normal_filter", -1.0f);
});
+
+ backcull_ = value("back_cull", true);
+ on("back_cull", [this](const ftl::config::Event &e) {
+ backcull_ = value("back_cull", true);
+ });
}
Splatter::~Splatter() {
@@ -90,7 +95,7 @@ void Splatter::renderChannel(
f.createTexture<float4>(Channel::Points),
f.createTexture<float4>(Channel::Normals),
temp_.getTexture<int>(Channel::Depth),
- params, stream
+ params, backcull_, stream
);
//LOG(INFO) << "DIBR DONE";
diff --git a/components/renderers/cpp/src/splatter.cu b/components/renderers/cpp/src/splatter.cu
index a7392f15a..c91673756 100644
--- a/components/renderers/cpp/src/splatter.cu
+++ b/components/renderers/cpp/src/splatter.cu
@@ -20,6 +20,7 @@ using ftl::render::SplatParams;
* Pass 1: Directly render each camera into virtual view but with no upsampling
* for sparse points.
*/
+ template <bool CULLING>
__global__ void dibr_merge_kernel(TextureObject<float4> points,
TextureObject<float4> normals,
TextureObject<int> depth, SplatParams params) {
@@ -29,17 +30,20 @@ using ftl::render::SplatParams;
const float4 worldPos = points.tex2D(x, y);
if (worldPos.x == MINF || (!(params.m_flags & ftl::render::kShowDisconMask) && worldPos.w < 0.0f)) return;
- float3 ray = params.m_viewMatrixInverse.getFloat3x3() * params.camera.screenToCam(x,y,1.0f);
- ray = ray / length(ray);
- float3 n = make_float3(normals.tex2D((int)x,(int)y));
- float l = length(n);
- if (l == 0) {
- return;
- }
- n /= l;
-
- const float facing = dot(ray, n);
- if (facing <= 0.0f) return;
+ // Compile time enable/disable of culling back facing points
+ if (CULLING) {
+ float3 ray = params.m_viewMatrixInverse.getFloat3x3() * params.camera.screenToCam(x,y,1.0f);
+ ray = ray / length(ray);
+ float3 n = make_float3(normals.tex2D((int)x,(int)y));
+ float l = length(n);
+ if (l == 0) {
+ return;
+ }
+ n /= l;
+
+ const float facing = dot(ray, n);
+ if (facing <= 0.0f) return;
+ }
// Find the virtual screen position of current point
const float3 camPos = params.m_viewMatrix * make_float3(worldPos);
@@ -57,11 +61,12 @@ using ftl::render::SplatParams;
}
}
-void ftl::cuda::dibr_merge(TextureObject<float4> &points, TextureObject<float4> &normals, TextureObject<int> &depth, SplatParams params, cudaStream_t stream) {
+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);
- dibr_merge_kernel<<<gridSize, blockSize, 0, stream>>>(points, normals, depth, params);
+ if (culling) dibr_merge_kernel<true><<<gridSize, blockSize, 0, stream>>>(points, normals, depth, params);
+ else dibr_merge_kernel<false><<<gridSize, blockSize, 0, stream>>>(points, normals, depth, params);
cudaSafeCall( cudaGetLastError() );
}
diff --git a/components/renderers/cpp/src/splatter_cuda.hpp b/components/renderers/cpp/src/splatter_cuda.hpp
index b47d55593..a90e1445b 100644
--- a/components/renderers/cpp/src/splatter_cuda.hpp
+++ b/components/renderers/cpp/src/splatter_cuda.hpp
@@ -11,6 +11,7 @@ namespace cuda {
ftl::cuda::TextureObject<float4> &normals,
ftl::cuda::TextureObject<int> &depth,
ftl::render::SplatParams params,
+ bool culling,
cudaStream_t stream);
void dibr_attribute(
--
GitLab