diff --git a/components/renderers/cpp/src/reprojection.cu b/components/renderers/cpp/src/reprojection.cu index 8a075b65382f7b574c6c4496557364a7a3aef1ee..dc636f35980e49ade7f0da4832554770d5704e65 100644 --- a/components/renderers/cpp/src/reprojection.cu +++ b/components/renderers/cpp/src/reprojection.cu @@ -203,13 +203,13 @@ __global__ void reprojection_kernel( const float3 camPos = poseInv * worldPos; if (camPos.z < camera.minDepth) return; if (camPos.z > camera.maxDepth) return; - const uint2 screenPos = camera.camToScreen<uint2>(camPos); + const float2 screenPos = camera.camToScreen<float2>(camPos); // Not on screen so stop now... if (screenPos.x >= depth_src.width() || screenPos.y >= depth_src.height()) return; - const float d2 = depth_src.tex2D((int)screenPos.x, (int)screenPos.y); - const A input = in.tex2D((int)screenPos.x, (int)screenPos.y); //generateInput(in.tex2D((int)screenPos.x, (int)screenPos.y), params, worldPos); + const float d2 = depth_src.tex2D((int)(screenPos.x+0.5f), (int)(screenPos.y+0.5f)); + const auto input = in.tex2D(screenPos.x, screenPos.y); //generateInput(in.tex2D((int)screenPos.x, (int)screenPos.y), params, worldPos); float weight = ftl::cuda::weighting(fabs(camPos.z - d2), 0.02f); const B weighted = make<B>(input) * weight; diff --git a/components/renderers/cpp/src/splatter.cu b/components/renderers/cpp/src/splatter.cu index d939c3eb09729f60a0723fa9c322303d4a989649..2986234bb3bbc24f762ff5ba0103ba173f4a0093 100644 --- a/components/renderers/cpp/src/splatter.cu +++ b/components/renderers/cpp/src/splatter.cu @@ -103,6 +103,40 @@ 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<float> depth, + TextureObject<int> depth_out, + float4x4 transform, + ftl::rgbd::Camera cam, + SplatParams params) { + const int x = blockIdx.x*blockDim.x + threadIdx.x; + const int y = blockIdx.y*blockDim.y + threadIdx.y; + + const float d0 = depth.tex2D(x, y); + if (d0 <= cam.minDepth || d0 >= cam.maxDepth) return; + + const float3 camPos = transform * cam.screenToCam(x,y,d0); + //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_out(cx,cy), d * 100000.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); @@ -120,6 +154,14 @@ void ftl::cuda::dibr_merge(TextureObject<float4> &points, TextureObject<int> &de cudaSafeCall( cudaGetLastError() ); } +void ftl::cuda::dibr_merge(TextureObject<float> &depth, TextureObject<int> &depth_out, const float4x4 &transform, const ftl::rgbd::Camera &cam, 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>>>(depth, depth_out, transform, cam, params); + cudaSafeCall( cudaGetLastError() ); +} + //============================================================================== diff --git a/components/renderers/cpp/src/splatter_cuda.hpp b/components/renderers/cpp/src/splatter_cuda.hpp index 778aea4e6ab68634a807c52b668abc15a31285b4..3918a0eaba7e1e88abb8c0ec8a3f5dd4d687137e 100644 --- a/components/renderers/cpp/src/splatter_cuda.hpp +++ b/components/renderers/cpp/src/splatter_cuda.hpp @@ -43,6 +43,14 @@ namespace cuda { ftl::render::SplatParams params, cudaStream_t stream); + void dibr_merge( + ftl::cuda::TextureObject<float> &depth, + ftl::cuda::TextureObject<int> &depth_out, + const float4x4 &transform, + const ftl::rgbd::Camera &cam, + 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 fdc167f39dd3e6696335856a8d9426cb5f4f83cb..d42d5a0e2e1a3862dfdf856efc8c2d72ebca542f 100644 --- a/components/renderers/cpp/src/tri_render.cpp +++ b/components/renderers/cpp/src/tri_render.cpp @@ -300,9 +300,13 @@ void Triangular::_dibr(ftl::rgbd::Frame &out, cudaStream_t stream) { continue; } + auto transform = params_.m_viewMatrix * MatrixConversion::toCUDA(s->getPose().cast<float>()); + ftl::cuda::dibr_merge( - f.createTexture<float4>(Channel::Points), + f.createTexture<float>(Channel::Depth), temp_.createTexture<int>(Channel::Depth2), + transform, + s->parameters(), params_, stream ); }