diff --git a/components/renderers/cpp/src/screen.cu b/components/renderers/cpp/src/screen.cu index 91f24c840e031e4aef0487274f2ed05f8315cd51..8610359d2bbec8eaa6f6972116a5f58cf6433b7a 100644 --- a/components/renderers/cpp/src/screen.cu +++ b/components/renderers/cpp/src/screen.cu @@ -13,6 +13,7 @@ using ftl::render::SplatParams; * Convert source screen position to output screen coordinates. */ __global__ void screen_coord_kernel(TextureObject<float> depth, + TextureObject<float> depth_out, TextureObject<short2> screen_out, SplatParams params, float4x4 pose, Camera camera) { const int x = blockIdx.x*blockDim.x + threadIdx.x; const int y = blockIdx.y*blockDim.y + threadIdx.y; @@ -31,12 +32,13 @@ using ftl::render::SplatParams; if (camPos.z < params.camera.minDepth || camPos.z > params.camera.maxDepth || screenPos.x >= params.camera.width || screenPos.y >= params.camera.height) screenPos = make_uint2(30000,30000); screen_out(x,y) = make_short2(screenPos.x, screenPos.y); + depth_out(x,y) = camPos.z; } -void ftl::cuda::screen_coord(TextureObject<float> &depth, TextureObject<short2> &screen_out, const SplatParams ¶ms, const float4x4 &pose, const Camera &camera, cudaStream_t stream) { +void ftl::cuda::screen_coord(TextureObject<float> &depth, TextureObject<float> &depth_out, TextureObject<short2> &screen_out, const SplatParams ¶ms, const float4x4 &pose, const Camera &camera, 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); - screen_coord_kernel<<<gridSize, blockSize, 0, stream>>>(depth, screen_out, params, pose, camera); + screen_coord_kernel<<<gridSize, blockSize, 0, stream>>>(depth, depth_out, screen_out, params, pose, camera); cudaSafeCall( cudaGetLastError() ); } diff --git a/components/renderers/cpp/src/splat_render.cpp b/components/renderers/cpp/src/splat_render.cpp index 1de452f9304943cb8e7f7d418a519aa35a3ba99f..64d987c1026dd59255a94b7d5e17d09704829946 100644 --- a/components/renderers/cpp/src/splat_render.cpp +++ b/components/renderers/cpp/src/splat_render.cpp @@ -278,12 +278,13 @@ void Splatter::_dibr(cudaStream_t stream) { ftl::cuda::screen_coord( f.createTexture<float>(Channel::Depth), + f.createTexture<float>(Channel::Depth2, Format<float>(f.get<GpuMat>(Channel::Depth).size())), f.createTexture<short2>(Channel::Screen, Format<short2>(f.get<GpuMat>(Channel::Depth).size())), params_, pose, s->parameters(), stream ); ftl::cuda::triangle_render1( - f.getTexture<float>(Channel::Depth), + f.getTexture<float>(Channel::Depth2), temp_.createTexture<int>(Channel::Depth2), f.getTexture<short2>(Channel::Screen), params_, stream diff --git a/components/renderers/cpp/src/splatter_cuda.hpp b/components/renderers/cpp/src/splatter_cuda.hpp index 2f036714a7a4c54f6932317d8d2bfb13408c8589..5e70225296866111164f7d1197b3650b726495fe 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 screen_coord( ftl::cuda::TextureObject<float> &depth, + ftl::cuda::TextureObject<float> &depth_out, ftl::cuda::TextureObject<short2> &screen_out, const ftl::render::SplatParams ¶ms, const float4x4 &pose,