diff --git a/components/renderers/cpp/include/ftl/cuda/normals.hpp b/components/renderers/cpp/include/ftl/cuda/normals.hpp index b4d2ced197c383a6d2dc185bf1c8883d6dde356a..1b4dbc86c46a3fddb3d42b3cc0f8023ab3c62483 100644 --- a/components/renderers/cpp/include/ftl/cuda/normals.hpp +++ b/components/renderers/cpp/include/ftl/cuda/normals.hpp @@ -10,7 +10,8 @@ namespace cuda { void normals(ftl::cuda::TextureObject<float4> &output, ftl::cuda::TextureObject<float4> &temp, - ftl::cuda::TextureObject<float4> &input, cudaStream_t stream); + ftl::cuda::TextureObject<float4> &input, + const ftl::rgbd::Camera &camera, cudaStream_t stream); void normal_visualise(ftl::cuda::TextureObject<float4> &norm, ftl::cuda::TextureObject<uchar4> &output, diff --git a/components/renderers/cpp/src/normals.cu b/components/renderers/cpp/src/normals.cu index cbbd0e96797ee11fffd43de0df399a03efcf8db8..de31404a62b8cf0f9d248f7b6f6a6a53cec8b01b 100644 --- a/components/renderers/cpp/src/normals.cu +++ b/components/renderers/cpp/src/normals.cu @@ -34,7 +34,8 @@ __global__ void computeNormals_kernel(ftl::cuda::TextureObject<float4> output, template <int RADIUS> __global__ void smooth_normals_kernel(ftl::cuda::TextureObject<float4> norms, ftl::cuda::TextureObject<float4> output, - ftl::cuda::TextureObject<float4> points, float smoothing) { + ftl::cuda::TextureObject<float4> points, + ftl::rgbd::Camera camera, float smoothing) { const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x; const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y; @@ -63,21 +64,27 @@ __global__ void smooth_normals_kernel(ftl::cuda::TextureObject<float4> norms, } } - // FIXME: USE A DIFFERENT OUTPUT BUFFER - //__syncthreads(); - output(x,y) = (contrib > 0.0f) ? make_float4(nsum / contrib, 1.0f) : make_float4(0.0f); + // Compute dot product of normal with camera to obtain measure of how + // well this point faces the source camera, a measure of confidence + float3 ray = camera.screenToCam(x, y, 1.0f); + ray = ray / length(ray); + nsum /= contrib; + nsum /= length(nsum); + + output(x,y) = (contrib > 0.0f) ? make_float4(nsum, dot(nsum, ray)) : make_float4(0.0f); } void ftl::cuda::normals(ftl::cuda::TextureObject<float4> &output, ftl::cuda::TextureObject<float4> &temp, - ftl::cuda::TextureObject<float4> &input, cudaStream_t stream) { + ftl::cuda::TextureObject<float4> &input, + const ftl::rgbd::Camera &camera, cudaStream_t stream) { const dim3 gridSize((input.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (input.height() + T_PER_BLOCK - 1)/T_PER_BLOCK); const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK); computeNormals_kernel<<<gridSize, blockSize, 0, stream>>>(temp, input); cudaSafeCall( cudaGetLastError() ); - smooth_normals_kernel<3><<<gridSize, blockSize, 0, stream>>>(temp, output, input, 0.04f); + smooth_normals_kernel<3><<<gridSize, blockSize, 0, stream>>>(temp, output, input, camera, 0.04f); cudaSafeCall( cudaGetLastError() ); #ifdef _DEBUG diff --git a/components/renderers/cpp/src/splat_render.cpp b/components/renderers/cpp/src/splat_render.cpp index d993e145856840e7af545de7365e1e77af7f0866..c801c27ce573e988daf7a0087b305e882c267230 100644 --- a/components/renderers/cpp/src/splat_render.cpp +++ b/components/renderers/cpp/src/splat_render.cpp @@ -331,7 +331,7 @@ bool Splatter::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out, cuda auto &g = f.get<GpuMat>(Channel::Colour); ftl::cuda::normals(f.createTexture<float4>(Channel::Normals, Format<float4>(g.cols, g.rows)), temp_.getTexture<float4>(Channel::Normals), // FIXME: Uses assumption of vcam res same as input res - f.getTexture<float4>(Channel::Points), stream); + f.getTexture<float4>(Channel::Points), s->parameters(), stream); if (norm_filter_ > -0.1f) { Eigen::Matrix4f matrix = s->getPose().cast<float>();