Buggy render of normals

components/renderers/cpp/include/ftl/cuda/normals.hpp

@@ -2,6 +2,8 @@
 #define _FTL_CUDA_NORMALS_HPP_
 
 #include <ftl/cuda_common.hpp>
+#include <ftl/rgbd/camera.hpp>
+#include <ftl/cuda_matrix_util.hpp>
 
 namespace ftl {
 namespace cuda {
@@ -9,6 +11,11 @@ namespace cuda {
 void normals(ftl::cuda::TextureObject<float4> &output,
         ftl::cuda::TextureObject<float4> &input, cudaStream_t stream);
 
+void normal_visualise(ftl::cuda::TextureObject<float4> &norm,
+        ftl::cuda::TextureObject<float> &output,
+        const ftl::rgbd::Camera &camera, const float4x4 &pose,
+        cudaStream_t stream);
+
 }
 }
 

components/renderers/cpp/src/normals.cu

@@ -43,3 +43,38 @@ void ftl::cuda::normals(ftl::cuda::TextureObject<float4> &output,
 	//cutilCheckMsg(__FUNCTION__);
 #endif
 }
+
+//==============================================================================
+
+__global__ void vis_normals_kernel(ftl::cuda::TextureObject<float4> norm,
+        ftl::cuda::TextureObject<float> output,
+        ftl::rgbd::Camera camera, float4x4 pose) {
+    const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
+    const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+    if(x >= norm.width() || y >= norm.height()) return;
+
+    float3 ray = pose * camera.screenToCam(x,y,1.0f);
+    ray = ray / length(ray);
+    float3 n = make_float3(norm.tex2D((int)x,(int)y));
+    n /= length(n);
+
+    output(x,y) = fabs(dot(ray, n))*7.0f;
+}
+
+void ftl::cuda::normal_visualise(ftl::cuda::TextureObject<float4> &norm,
+        ftl::cuda::TextureObject<float> &output,
+        const ftl::rgbd::Camera &camera, const float4x4 &pose,
+        cudaStream_t stream) {
+
+    const dim3 gridSize((norm.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (norm.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
+    const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
+
+    vis_normals_kernel<<<gridSize, blockSize, 0, stream>>>(norm, output, camera, pose);
+
+    cudaSafeCall( cudaGetLastError() );
+#ifdef _DEBUG
+    cudaSafeCall(cudaDeviceSynchronize());
+    //cutilCheckMsg(__FUNCTION__);
+#endif
+}

components/renderers/cpp/src/splat_render.cpp

@@ -268,6 +268,7 @@ bool Splatter::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out, cuda
 		for (auto &f : scene_->frames) {
 			if (!f.hasChannel(Channel::Normals)) {
 				auto &g = f.get<GpuMat>(Channel::Colour);
+				LOG(INFO) << "Make normals channel";
 				ftl::cuda::normals(f.createTexture<float4>(Channel::Normals, Format<float4>(g.cols, g.rows)), f.getTexture<float4>(Channel::Points), stream);
 			}
 		}
@@ -278,7 +279,7 @@ bool Splatter::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out, cuda
 		renderChannel(params, out, Channel::Normals, stream);
 
 		// Convert normal to single float value
-		//ftl::cuda::normal_visualise(out.getTexture<float4>(Channel::Normals), temp_.getTexture<float>(Channel::Contribution), camera);
+		ftl::cuda::normal_visualise(out.getTexture<float4>(Channel::Normals), temp_.getTexture<float>(Channel::Contribution), camera, params.m_viewMatrix, stream);
 
 		// Put in output as single float
 		cv::cuda::swap(temp_.get<GpuMat>(Channel::Contribution), out.create<GpuMat>(Channel::Normals));