diff --git a/components/renderers/cpp/src/splat_render.cpp b/components/renderers/cpp/src/splat_render.cpp index 6635d1ab4295d708fef458214e613d0c554f03b4..3d6e95e2c7f79a40867d660bbe87fcc550e591b8 100644 --- a/components/renderers/cpp/src/splat_render.cpp +++ b/components/renderers/cpp/src/splat_render.cpp @@ -76,6 +76,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,6 +88,7 @@ void Splatter::renderChannel( ftl::cuda::dibr_merge( f.createTexture<float4>(Channel::Points), + f.createTexture<float4>(Channel::Normals), temp_.getTexture<int>(Channel::Depth), params, stream ); diff --git a/components/renderers/cpp/src/splatter.cu b/components/renderers/cpp/src/splatter.cu index 6f94b7636d4367ae4b465ad58e16163264c3a361..a7392f15a5bc4aa344f70277960744eb3f34bf41 100644 --- a/components/renderers/cpp/src/splatter.cu +++ b/components/renderers/cpp/src/splatter.cu @@ -20,13 +20,27 @@ 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) { + __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; + 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 +57,11 @@ 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, 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); + dibr_merge_kernel<<<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..b47d5559394918400ececbfce262c686295ed861 100644 --- a/components/renderers/cpp/src/splatter_cuda.hpp +++ b/components/renderers/cpp/src/splatter_cuda.hpp @@ -8,6 +8,7 @@ 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, cudaStream_t stream);