diff --git a/components/renderers/cpp/src/reprojection.cu b/components/renderers/cpp/src/reprojection.cu index 72b7cd07275f3c9b41c207009dd4b7eef6ad7c9b..b5bac1ed14f23b3009de17ccaaba1e115417f5e2 100644 --- a/components/renderers/cpp/src/reprojection.cu +++ b/components/renderers/cpp/src/reprojection.cu @@ -333,3 +333,57 @@ void ftl::cuda::equirectangular_reproject( equirectangular_kernel<<<gridSize, blockSize, 0, stream>>>(image_in, image_out, camera, pose); cudaSafeCall( cudaGetLastError() ); } + +// ==== Correct for bad colours ================================================ + +__device__ inline uchar4 make_uchar4(const float4 v) { + return make_uchar4(v.x,v.y,v.z,v.w); +} + +template <int RADIUS> +__global__ void fix_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) { + float4 sumcol = make_float4(0.0f); + float count = 0.0f; + + for (int v=-RADIUS; v<=RADIUS; ++v) { + for (int u=-RADIUS; u<=RADIUS; ++u) { + const float contrib = contribs.tex2D((int)x+u,(int)y+v); + const float4 c = make_float4(out(int(x)+u,int(y)+v)); + if (contrib > 0.0000001f) { + sumcol += c; + count += 1.0f; + } + } + } + + out(x,y) = (count > 0.0f) ? make_uchar4(sumcol / count) : bad_colour; + } + } +} + +void ftl::cuda::fix_bad_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); + + fix_colour_kernel<1><<<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 53b9f675e9f26144e9f2d31fff9c889b3ef58720..5d027b2c6ff0c12efa28ca277d0404c11c5eb49c 100644 --- a/components/renderers/cpp/src/splatter_cuda.hpp +++ b/components/renderers/cpp/src/splatter_cuda.hpp @@ -122,6 +122,14 @@ namespace cuda { ftl::cuda::TextureObject<int> &d1, ftl::cuda::TextureObject<float> &d2, float factor, cudaStream_t stream); + + void fix_bad_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); } } diff --git a/components/renderers/cpp/src/tri_render.cpp b/components/renderers/cpp/src/tri_render.cpp index 659bcc8a99c929840506604d1a2dddb4b5dbc58f..23934f5f2bd8a06291d6ba9a3834d518ec15b1ce 100644 --- a/components/renderers/cpp/src/tri_render.cpp +++ b/components/renderers/cpp/src/tri_render.cpp @@ -542,7 +542,7 @@ void Triangular::_postprocessColours(ftl::rgbd::Frame &out) { ); } - if (value("show_bad_colour", true)) { + if (value("show_bad_colour", false)) { ftl::cuda::show_missing_colour( out.getTexture<float>(Channel::Depth), out.getTexture<uchar4>(Channel::Colour), @@ -551,6 +551,15 @@ void Triangular::_postprocessColours(ftl::rgbd::Frame &out) { params_.camera, stream_ ); + } else { + ftl::cuda::fix_bad_colour( + out.getTexture<float>(Channel::Depth), + out.getTexture<uchar4>(Channel::Colour), + temp_.getTexture<float>(Channel::Contribution), + make_uchar4(255,0,0,0), + params_.camera, + stream_ + ); } }