diff --git a/components/renderers/cpp/include/ftl/render/splat_render.hpp b/components/renderers/cpp/include/ftl/render/splat_render.hpp index bca5a232b937479234866d42847896f653bd70d9..40bfa459b35e850d2147f74e33eabfaaddb3a261 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 6635d1ab4295d708fef458214e613d0c554f03b4..e66c11b651a6def941ed82ee767d76285ff86e34 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() { @@ -76,6 +81,7 @@ void Splatter::renderChannel( bool is_4chan = scene_->frames[0].get<GpuMat>(channel).type() == CV_32FC4; // Render each camera into virtual view + // TODO: Move out of renderChannel, this is a common step to all channels for (size_t i=0; i < scene_->frames.size(); ++i) { auto &f = scene_->frames[i]; auto *s = scene_->sources[i]; @@ -87,8 +93,9 @@ void Splatter::renderChannel( ftl::cuda::dibr_merge( 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 6f94b7636d4367ae4b465ad58e16163264c3a361..c91673756663d1e2c4aefea5e785ff52d25a487d 100644 --- a/components/renderers/cpp/src/splatter.cu +++ b/components/renderers/cpp/src/splatter.cu @@ -20,13 +20,31 @@ using ftl::render::SplatParams; * 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) { + template <bool CULLING> + __global__ void dibr_merge_kernel(TextureObject<float4> points, + TextureObject<float4> normals, + 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; + // 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); if (camPos.z < params.camera.minDepth) return; @@ -43,11 +61,12 @@ using ftl::render::SplatParams; } } -void ftl::cuda::dibr_merge(TextureObject<float4> &points, 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, 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 7b39e0991686cc288b8046f3cc01fb3cdb41ddad..a90e1445b2bdd45235efff301b16652e5b496f16 100644 --- a/components/renderers/cpp/src/splatter_cuda.hpp +++ b/components/renderers/cpp/src/splatter_cuda.hpp @@ -8,8 +8,10 @@ namespace ftl { namespace cuda { void dibr_merge( ftl::cuda::TextureObject<float4> &points, + ftl::cuda::TextureObject<float4> &normals, ftl::cuda::TextureObject<int> &depth, ftl::render::SplatParams params, + bool culling, cudaStream_t stream); void dibr_attribute(