diff --git a/components/renderers/cpp/src/splatter.cu b/components/renderers/cpp/src/splatter.cu index 6346daa883019453bef1d2f80be36ba5afa0aaf2..d939c3eb09729f60a0723fa9c322303d4a989649 100644 --- a/components/renderers/cpp/src/splatter.cu +++ b/components/renderers/cpp/src/splatter.cu @@ -441,9 +441,11 @@ __global__ void dibr_normalise_kernel( const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y; if (x < in.width() && y < in.height()) { + const float contrib = contribs.tex2D((int)x,(int)y); const A a = in.tex2D((int)x,(int)y); //const float4 normal = normals.tex2D((int)x,(int)y); - const float contrib = contribs.tex2D((int)x,(int)y); + + //out(x,y) = (contrib == 0.0f) ? make<B>(a) : make<B>(a / contrib); if (contrib > 0.0f) { out(x,y) = make<B>(a / contrib); @@ -464,3 +466,39 @@ void ftl::cuda::dibr_normalise(TextureObject<A> &in, TextureObject<B> &out, Text template void ftl::cuda::dibr_normalise<float4,uchar4>(TextureObject<float4> &in, TextureObject<uchar4> &out, TextureObject<float> &contribs, cudaStream_t stream); template void ftl::cuda::dibr_normalise<float,float>(TextureObject<float> &in, TextureObject<float> &out, TextureObject<float> &contribs, cudaStream_t stream); template void ftl::cuda::dibr_normalise<float4,float4>(TextureObject<float4> &in, TextureObject<float4> &out, TextureObject<float> &contribs, cudaStream_t stream); + + +// ===== Show bad colour normalise ============================================= + +__global__ void show_missing_colour_kernel( + TextureObject<float> depth, + TextureObject<uchar4> out, + TextureObject<float> contribs, + uchar4 bad_colour, + ftl::rgbd::Camera cam) { + const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x; + const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y; + + if (x < out.width() && y < out.height()) { + const float contrib = contribs.tex2D((int)x,(int)y); + const float d = depth.tex2D(x,y); + + if (contrib < 0.0000001f && d > cam.minDepth && d < cam.maxDepth) { + out(x,y) = bad_colour; + } + } +} + +void ftl::cuda::show_missing_colour( + TextureObject<float> &depth, + TextureObject<uchar4> &out, + TextureObject<float> &contribs, + uchar4 bad_colour, + const ftl::rgbd::Camera &cam, + cudaStream_t stream) { + const dim3 gridSize((out.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (out.height() + T_PER_BLOCK - 1)/T_PER_BLOCK); + const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK); + + show_missing_colour_kernel<<<gridSize, blockSize, 0, stream>>>(depth, out, contribs, bad_colour, cam); + cudaSafeCall( cudaGetLastError() ); +} diff --git a/components/renderers/cpp/src/splatter_cuda.hpp b/components/renderers/cpp/src/splatter_cuda.hpp index 010a637188523030ea2a40ac3c6cdd720c3a0148..778aea4e6ab68634a807c52b668abc15a31285b4 100644 --- a/components/renderers/cpp/src/splatter_cuda.hpp +++ b/components/renderers/cpp/src/splatter_cuda.hpp @@ -97,6 +97,14 @@ namespace cuda { ftl::cuda::TextureObject<float> &contribs, cudaStream_t stream); + void show_missing_colour( + ftl::cuda::TextureObject<float> &depth, + ftl::cuda::TextureObject<uchar4> &out, + ftl::cuda::TextureObject<float> &contribs, + uchar4 bad_colour, + const ftl::rgbd::Camera &cam, + cudaStream_t stream); + void show_mask( ftl::cuda::TextureObject<uchar4> &colour, ftl::cuda::TextureObject<int> &mask, diff --git a/components/renderers/cpp/src/tri_render.cpp b/components/renderers/cpp/src/tri_render.cpp index b3fba82a7de676f2aa963e80467a45dbf3f0bd13..fdc167f39dd3e6696335856a8d9426cb5f4f83cb 100644 --- a/components/renderers/cpp/src/tri_render.cpp +++ b/components/renderers/cpp/src/tri_render.cpp @@ -603,6 +603,17 @@ bool Triangular::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out) { // Reprojection of colours onto surface _renderChannel(out, Channel::Colour, Channel::Colour, stream_); + + if (value("show_bad_colour", false)) { + ftl::cuda::show_missing_colour( + out.getTexture<float>(Channel::Depth), + out.getTexture<uchar4>(Channel::Colour), + temp_.getTexture<float>(Channel::Contribution), + make_uchar4(255,0,0,0), + camera, + stream_ + ); + } if (chan == Channel::Depth) {