diff --git a/components/renderers/cpp/src/splatter.cu b/components/renderers/cpp/src/splatter.cu index e79b3d32038985743d9e657d6df6eb0aceb7231a..6346daa883019453bef1d2f80be36ba5afa0aaf2 100644 --- a/components/renderers/cpp/src/splatter.cu +++ b/components/renderers/cpp/src/splatter.cu @@ -75,6 +75,34 @@ using ftl::cuda::warpSum; } } +/* + * 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) { + 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; + + // Find the virtual screen position of current point + const float3 camPos = params.m_viewMatrix * make_float3(worldPos); + if (camPos.z < params.camera.minDepth) return; + if (camPos.z > params.camera.maxDepth) return; + + const float d = camPos.z; + + const uint2 screenPos = params.camera.camToScreen<uint2>(camPos); + const unsigned int cx = screenPos.x; + const unsigned int cy = screenPos.y; + if (d > params.camera.minDepth && d < params.camera.maxDepth && cx < depth.width() && cy < depth.height()) { + // Transform estimated point to virtual cam space and output z + atomicMin(&depth(cx,cy), d * 1000.0f); + } +} + 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); @@ -84,6 +112,14 @@ void ftl::cuda::dibr_merge(TextureObject<float4> &points, TextureObject<float4> cudaSafeCall( cudaGetLastError() ); } +void ftl::cuda::dibr_merge(TextureObject<float4> &points, 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); + cudaSafeCall( cudaGetLastError() ); +} + //============================================================================== diff --git a/components/renderers/cpp/src/splatter_cuda.hpp b/components/renderers/cpp/src/splatter_cuda.hpp index 14fb2dc64fc37351fd2fbd8d8ac36110505026d6..6e3d22cbf5ef5e1583079da56fe45fb6f4d34bf3 100644 --- a/components/renderers/cpp/src/splatter_cuda.hpp +++ b/components/renderers/cpp/src/splatter_cuda.hpp @@ -30,6 +30,12 @@ namespace cuda { bool culling, cudaStream_t stream); + void dibr_merge( + ftl::cuda::TextureObject<float4> &points, + ftl::cuda::TextureObject<int> &depth, + ftl::render::SplatParams params, + cudaStream_t stream); + template <typename T> void splat( ftl::cuda::TextureObject<float4> &normals, diff --git a/components/renderers/cpp/src/tri_render.cpp b/components/renderers/cpp/src/tri_render.cpp index 58e21c3672343ad3a4a0a73a84ef8833f85da497..7f89ec2311d3c3beb900b6b4deca5344a301ca63 100644 --- a/components/renderers/cpp/src/tri_render.cpp +++ b/components/renderers/cpp/src/tri_render.cpp @@ -270,9 +270,8 @@ void Triangular::_dibr(cudaStream_t stream) { ftl::cuda::dibr_merge( f.createTexture<float4>(Channel::Points), - f.createTexture<float4>(Channel::Normals), temp_.createTexture<int>(Channel::Depth2), - params_, backcull_, stream + params_, stream ); } }