From 4a68c40a901de07fea5cf3fef53f75b86041aaa2 Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Wed, 2 Oct 2019 10:49:22 +0300
Subject: [PATCH] Buggy render of normals
---
.../cpp/include/ftl/cuda/normals.hpp | 7 ++++
components/renderers/cpp/src/normals.cu | 37 ++++++++++++++++++-
components/renderers/cpp/src/splat_render.cpp | 3 +-
3 files changed, 45 insertions(+), 2 deletions(-)
diff --git a/components/renderers/cpp/include/ftl/cuda/normals.hpp b/components/renderers/cpp/include/ftl/cuda/normals.hpp
index 4cbd760a1..50f2e881a 100644
--- a/components/renderers/cpp/include/ftl/cuda/normals.hpp
+++ b/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);
+
}
}
diff --git a/components/renderers/cpp/src/normals.cu b/components/renderers/cpp/src/normals.cu
index 7ed9bb681..1a5532219 100644
--- a/components/renderers/cpp/src/normals.cu
+++ b/components/renderers/cpp/src/normals.cu
@@ -42,4 +42,39 @@ void ftl::cuda::normals(ftl::cuda::TextureObject<float4> &output,
cudaSafeCall(cudaDeviceSynchronize());
//cutilCheckMsg(__FUNCTION__);
#endif
-}
\ No newline at end of file
+}
+
+//==============================================================================
+
+__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
+}
diff --git a/components/renderers/cpp/src/splat_render.cpp b/components/renderers/cpp/src/splat_render.cpp
index 553f63396..50264124b 100644
--- a/components/renderers/cpp/src/splat_render.cpp
+++ b/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));
--
GitLab