diff --git a/components/renderers/cpp/include/ftl/render/splat_render.hpp b/components/renderers/cpp/include/ftl/render/splat_render.hpp
index bd83d8bd1bda862d71e2ef12f5239a578e2d86e6..3b36e8ec98dd37aaf40b0e3a7e12cad6b3b5a14e 100644
--- a/components/renderers/cpp/include/ftl/render/splat_render.hpp
+++ b/components/renderers/cpp/include/ftl/render/splat_render.hpp
@@ -53,6 +53,7 @@ class Splatter : public ftl::render::Renderer {
uchar4 light_ambient_;
ftl::render::SplatParams params_;
cudaStream_t stream_;
+ float3 light_pos_;
template <typename T>
void __blendChannel(ftl::rgbd::Frame &, ftl::codecs::Channel in, ftl::codecs::Channel out, cudaStream_t);
diff --git a/components/renderers/cpp/src/normals.cu b/components/renderers/cpp/src/normals.cu
index 626015ce7ad9f9471079917b0678369bbb23f5cb..7b385a8c5c8713ed8bc0469ead391061f874cce0 100644
--- a/components/renderers/cpp/src/normals.cu
+++ b/components/renderers/cpp/src/normals.cu
@@ -149,6 +149,31 @@ __global__ void smooth_normals_kernel(ftl::cuda::TextureObject<float4> norms,
output(x,y) = (contrib > 0.0f) ? make_float4(pose*nsum, 1.0f) : make_float4(0.0f);
}
+template <>
+__global__ void smooth_normals_kernel<0>(ftl::cuda::TextureObject<float4> norms,
+ ftl::cuda::TextureObject<float4> output,
+ ftl::cuda::TextureObject<int> depth,
+ ftl::rgbd::Camera camera, float3x3 pose, float smoothing) {
+ const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
+ const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+ if(x >= depth.width() || y >= depth.height()) return;
+
+ const float3 p0 = camera.screenToCam(x,y, (float)depth.tex2D((int)x,(int)y) / 1000.0f);
+
+ if (p0.z < camera.minDepth || p0.z > camera.maxDepth) return;
+
+ // 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);
+
+ const float4 n = norms.tex2D((int)x,(int)y);
+ output(x,y) = n;
+}
+
void ftl::cuda::normals(ftl::cuda::TextureObject<float4> &output,
ftl::cuda::TextureObject<float4> &temp,
ftl::cuda::TextureObject<float4> &input,
@@ -166,7 +191,10 @@ void ftl::cuda::normals(ftl::cuda::TextureObject<float4> &output,
case 9: smooth_normals_kernel<9><<<gridSize, blockSize, 0, stream>>>(temp, output, input, camera, pose, smoothing); break;
case 7: smooth_normals_kernel<7><<<gridSize, blockSize, 0, stream>>>(temp, output, input, camera, pose, smoothing); break;
case 5: smooth_normals_kernel<5><<<gridSize, blockSize, 0, stream>>>(temp, output, input, camera, pose, smoothing); break;
- case 3: smooth_normals_kernel<3><<<gridSize, blockSize, 0, stream>>>(temp, output, input, camera, pose, smoothing); break;
+ case 3: smooth_normals_kernel<3><<<gridSize, blockSize, 0, stream>>>(temp, output, input, camera, pose, smoothing); break;
+ case 2: smooth_normals_kernel<2><<<gridSize, blockSize, 0, stream>>>(temp, output, input, camera, pose, smoothing); break;
+ case 1: smooth_normals_kernel<1><<<gridSize, blockSize, 0, stream>>>(temp, output, input, camera, pose, smoothing); break;
+ case 0: smooth_normals_kernel<0><<<gridSize, blockSize, 0, stream>>>(temp, output, input, camera, pose, smoothing); break;
}
cudaSafeCall( cudaGetLastError() );
diff --git a/components/renderers/cpp/src/splat_render.cpp b/components/renderers/cpp/src/splat_render.cpp
index 80a6f2a1ba1b1131173451a1f6dc24de86b080ee..4b94e29c355fdd675566c44c52876fdff7866760 100644
--- a/components/renderers/cpp/src/splat_render.cpp
+++ b/components/renderers/cpp/src/splat_render.cpp
@@ -123,6 +123,11 @@ Splatter::Splatter(nlohmann::json &config, ftl::rgbd::FrameSet *fs) : ftl::rende
light_ambient_ = parseCUDAColour(value("ambient", std::string("#0e0e0e")));
});
+ light_pos_ = make_float3(value("light_x", 0.3f), value("light_y", 0.2f), value("light_z", 1.0f));
+ on("light_x", [this](const ftl::config::Event &e) { light_pos_.x = value("light_x", 0.3f); });
+ on("light_y", [this](const ftl::config::Event &e) { light_pos_.y = value("light_y", 0.3f); });
+ on("light_z", [this](const ftl::config::Event &e) { light_pos_.z = value("light_z", 0.3f); });
+
cudaSafeCall(cudaStreamCreate(&stream_));
}
@@ -386,7 +391,7 @@ bool Splatter::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out) {
ftl::cuda::normals(f.createTexture<float4>(Channel::Normals, Format<float4>(g.cols, g.rows)),
temp_.getTexture<float4>(Channel::Normals),
f.getTexture<float4>(Channel::Points),
- 3, 0.04f,
+ 1, 0.02f,
s->parameters(), pose.getFloat3x3(), stream_);
if (norm_filter_ > -0.1f) {
@@ -395,18 +400,33 @@ bool Splatter::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out) {
}
}
+ Channel chan = src->getChannel();
+
int aligned_source = value("aligned_source",-1);
if (aligned_source >= 0 && aligned_source < scene_->frames.size()) {
// FIXME: Output may not be same resolution as source!
cudaSafeCall(cudaStreamSynchronize(stream_));
scene_->frames[aligned_source].copyTo(Channel::Depth + Channel::Colour, out);
+
+ if (chan == Channel::Normals) {
+ // Convert normal to single float value
+ temp_.create<GpuMat>(Channel::Colour, Format<uchar4>(camera.width, camera.height)).setTo(cv::Scalar(0,0,0,0), cvstream);
+ ftl::cuda::normal_visualise(scene_->frames[aligned_source].getTexture<float4>(Channel::Normals), temp_.createTexture<uchar4>(Channel::Colour),
+ light_pos_,
+ light_diffuse_,
+ light_ambient_, stream_);
+
+ // Put in output as single float
+ cv::cuda::swap(temp_.get<GpuMat>(Channel::Colour), out.create<GpuMat>(Channel::Normals));
+ out.resetTexture(Channel::Normals);
+ }
+
return true;
}
_dibr(stream_);
_renderChannel(out, Channel::Colour, Channel::Colour, stream_);
- Channel chan = src->getChannel();
if (chan == Channel::Depth)
{
//temp_.get<GpuMat>(Channel::Depth).convertTo(out.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 1000.0f, cvstream);
@@ -417,9 +437,9 @@ bool Splatter::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out) {
_renderChannel(out, Channel::Normals, Channel::Normals, stream_);
// Convert normal to single float value
- temp_.create<GpuMat>(Channel::Colour, Format<uchar4>(camera.width, camera.height));
+ temp_.create<GpuMat>(Channel::Colour, Format<uchar4>(camera.width, camera.height)).setTo(cv::Scalar(0,0,0,0), cvstream);
ftl::cuda::normal_visualise(out.getTexture<float4>(Channel::Normals), temp_.createTexture<uchar4>(Channel::Colour),
- make_float3(0.3f, 0.2f, 1.0f),
+ light_pos_,
light_diffuse_,
light_ambient_, stream_);
diff --git a/components/renderers/cpp/src/splatter.cu b/components/renderers/cpp/src/splatter.cu
index cd69098ea6f80c269f66deb935e652ffc33398c8..3a9270e542ee58e836410e0de598cbd4ab9e269c 100644
--- a/components/renderers/cpp/src/splatter.cu
+++ b/components/renderers/cpp/src/splatter.cu
@@ -44,8 +44,10 @@ using ftl::cuda::warpSum;
// Compile time enable/disable of culling back facing points
if (CULLING) {
float3 ray = params.m_viewMatrixInverse.getFloat3x3() * params.camera.screenToCam(x,y,1.0f);
- ray = ray / length(ray);
- float3 n = make_float3(normals.tex2D((int)x,(int)y));
+ ray = ray / length(ray);
+ float4 n4 = normals.tex2D((int)x,(int)y);
+ if (n4.w == 0.0f) return;
+ float3 n = make_float3(n4);
float l = length(n);
if (l == 0) {
return;