diff --git a/applications/reconstruct/src/ilw/ilw.cpp b/applications/reconstruct/src/ilw/ilw.cpp index 65e1a19dfe35e89019f5b3aa9dfbe85310b03f33..778a3d36e9178e90e5a5ba90c029be244c302565 100644 --- a/applications/reconstruct/src/ilw/ilw.cpp +++ b/applications/reconstruct/src/ilw/ilw.cpp @@ -116,7 +116,7 @@ bool ILW::_phase0(ftl::rgbd::FrameSet &fs, cudaStream_t stream) { auto &t = f.createTexture<float4>(Channel::Points, Format<float4>(f.get<GpuMat>(Channel::Colour).size())); auto pose = MatrixConversion::toCUDA(s->getPose().cast<float>()); //.inverse()); - ftl::cuda::point_cloud(t, f.createTexture<float>(Channel::Depth), s->parameters(), pose, stream); + ftl::cuda::point_cloud(t, f.createTexture<float>(Channel::Depth), s->parameters(), pose, 2, stream); // TODO: Create energy vector texture and clear it // Create energy and clear it diff --git a/components/renderers/cpp/include/ftl/cuda/points.hpp b/components/renderers/cpp/include/ftl/cuda/points.hpp index cb996c4ec68a794b2b9ede7a03cf2b2c5d92047d..a705e39e2c9f4e94f3c216f6f9168586ede4811c 100644 --- a/components/renderers/cpp/include/ftl/cuda/points.hpp +++ b/components/renderers/cpp/include/ftl/cuda/points.hpp @@ -16,7 +16,7 @@ struct ClipSpace { void point_cloud(ftl::cuda::TextureObject<float4> &output, ftl::cuda::TextureObject<float> &depth, const ftl::rgbd::Camera ¶ms, - const float4x4 &pose, cudaStream_t stream); + const float4x4 &pose, uint discon, cudaStream_t stream); void clipping(ftl::cuda::TextureObject<float4> &points, const ClipSpace &clip, cudaStream_t stream); diff --git a/components/renderers/cpp/src/points.cu b/components/renderers/cpp/src/points.cu index abcb49f2e7e044ed811e3cdf45946640730a86ab..b2e83577d4fe1d9e7c3d82ecfb615724904d980f 100644 --- a/components/renderers/cpp/src/points.cu +++ b/components/renderers/cpp/src/points.cu @@ -11,26 +11,51 @@ __global__ void point_cloud_kernel(ftl::cuda::TextureObject<float4> output, ftl: if (x < params.width && y < params.height) { output(x,y) = make_float4(MINF, MINF, MINF, MINF); - float d = depth.tex2D((int)x, (int)y); + const float d = depth.tex2D((int)x, (int)y); + float p = d; if (d >= params.minDepth && d <= params.maxDepth) { + /* Orts-Escolano S. et al. 2016. Holoportation: Virtual 3D teleportation in real-time. */ // Is there a discontinuity nearby? for (int u=-RADIUS; u<=RADIUS; ++u) { for (int v=-RADIUS; v<=RADIUS; ++v) { - if (fabs(depth.tex2D((int)x+u, (int)y+v) - d) > 0.1f) return; + // If yes, the flag using w = -1 + if (fabs(depth.tex2D((int)x+u, (int)y+v) - d) > 0.1f) p = -1.0f; } } + output(x,y) = make_float4(pose * params.screenToCam(x, y, d), p); + } + } +} + +template <> +__global__ void point_cloud_kernel<0>(ftl::cuda::TextureObject<float4> output, ftl::cuda::TextureObject<float> depth, ftl::rgbd::Camera params, float4x4 pose) +{ + const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x; + const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y; + + if (x < params.width && y < params.height) { + output(x,y) = make_float4(MINF, MINF, MINF, MINF); + + float d = depth.tex2D((int)x, (int)y); + + if (d >= params.minDepth && d <= params.maxDepth) { output(x,y) = make_float4(pose * params.screenToCam(x, y, d), d); } } } -void ftl::cuda::point_cloud(ftl::cuda::TextureObject<float4> &output, ftl::cuda::TextureObject<float> &depth, const ftl::rgbd::Camera ¶ms, const float4x4 &pose, cudaStream_t stream) { +void ftl::cuda::point_cloud(ftl::cuda::TextureObject<float4> &output, ftl::cuda::TextureObject<float> &depth, const ftl::rgbd::Camera ¶ms, const float4x4 &pose, uint discon, cudaStream_t stream) { const dim3 gridSize((params.width + T_PER_BLOCK - 1)/T_PER_BLOCK, (params.height + T_PER_BLOCK - 1)/T_PER_BLOCK); const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK); - point_cloud_kernel<3><<<gridSize, blockSize, 0, stream>>>(output, depth, params, pose); + switch (discon) { + case 3 : point_cloud_kernel<3><<<gridSize, blockSize, 0, stream>>>(output, depth, params, pose); break; + case 2 : point_cloud_kernel<2><<<gridSize, blockSize, 0, stream>>>(output, depth, params, pose); break; + case 1 : point_cloud_kernel<1><<<gridSize, blockSize, 0, stream>>>(output, depth, params, pose); break; + default: point_cloud_kernel<0><<<gridSize, blockSize, 0, stream>>>(output, depth, params, pose); + } cudaSafeCall( cudaGetLastError() ); #ifdef _DEBUG diff --git a/components/renderers/cpp/src/splat_render.cpp b/components/renderers/cpp/src/splat_render.cpp index f5b5a46256495ada872ccd1317abb57bd9769e0e..68d60599f9e2b6868faf78abbdfff669210e41cb 100644 --- a/components/renderers/cpp/src/splat_render.cpp +++ b/components/renderers/cpp/src/splat_render.cpp @@ -91,7 +91,7 @@ void Splatter::renderChannel( auto &t = f.createTexture<float4>(Channel::Points, Format<float4>(f.get<GpuMat>(Channel::Colour).size())); auto pose = MatrixConversion::toCUDA(s->getPose().cast<float>()); //.inverse()); - ftl::cuda::point_cloud(t, f.createTexture<float>(Channel::Depth), s->parameters(), pose, stream); + ftl::cuda::point_cloud(t, f.createTexture<float>(Channel::Depth), s->parameters(), pose, 0, stream); //LOG(INFO) << "POINTS Added"; } diff --git a/components/renderers/cpp/src/splatter.cu b/components/renderers/cpp/src/splatter.cu index ddea639bd3c28df2e99043a16cb1c03caa6aed23..41777d49d4a35888f297bc14bfc9c99d034ee67c 100644 --- a/components/renderers/cpp/src/splatter.cu +++ b/components/renderers/cpp/src/splatter.cu @@ -24,11 +24,11 @@ using ftl::render::SplatParams; const int x = blockIdx.x*blockDim.x + threadIdx.x; const int y = blockIdx.y*blockDim.y + threadIdx.y; - const float3 worldPos = make_float3(points.tex2D(x, y)); - if (worldPos.x == MINF) return; + const float4 worldPos = points.tex2D(x, y); + if (worldPos.x == MINF || worldPos.w < 0.0f) return; // Find the virtual screen position of current point - const float3 camPos = params.m_viewMatrix * worldPos; + const float3 camPos = params.m_viewMatrix * make_float3(worldPos); if (camPos.z < params.camera.minDepth) return; if (camPos.z > params.camera.maxDepth) return;