At dot product with camera to normals

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,

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

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>();