From d0572bb2ce9aa0eb945980e3f1c4f96d1bf4182f Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Sat, 2 Nov 2019 14:14:46 +0200
Subject: [PATCH] Experiment with post render filters
---
components/renderers/cpp/CMakeLists.txt | 2 +-
.../cpp/include/ftl/cuda/normals.hpp | 8 ++
.../cpp/include/ftl/render/tri_render.hpp | 7 +-
components/renderers/cpp/src/normals.cu | 73 +++++++++++++++++++
components/renderers/cpp/src/reprojection.cu | 24 +++---
.../renderers/cpp/src/splatter_cuda.hpp | 4 +-
components/renderers/cpp/src/tri_render.cpp | 57 +++++++++------
7 files changed, 137 insertions(+), 38 deletions(-)
diff --git a/components/renderers/cpp/CMakeLists.txt b/components/renderers/cpp/CMakeLists.txt
index 14bacda71..89cf408a3 100644
--- a/components/renderers/cpp/CMakeLists.txt
+++ b/components/renderers/cpp/CMakeLists.txt
@@ -18,6 +18,6 @@ target_include_directories(ftlrender PUBLIC
$<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}/include>
$<INSTALL_INTERFACE:include>
PRIVATE src)
-target_link_libraries(ftlrender ftlrgbd ftlcommon Eigen3::Eigen Threads::Threads ${OpenCV_LIBS})
+target_link_libraries(ftlrender ftlrgbd ftlcommon ftlfilter Eigen3::Eigen Threads::Threads ${OpenCV_LIBS})
#ADD_SUBDIRECTORY(test)
diff --git a/components/renderers/cpp/include/ftl/cuda/normals.hpp b/components/renderers/cpp/include/ftl/cuda/normals.hpp
index 227ab5362..dc3d0265c 100644
--- a/components/renderers/cpp/include/ftl/cuda/normals.hpp
+++ b/components/renderers/cpp/include/ftl/cuda/normals.hpp
@@ -24,6 +24,14 @@ void normals(ftl::cuda::TextureObject<float4> &output,
const ftl::rgbd::Camera &camera,
const float3x3 &pose_inv, const float3x3 &pose, cudaStream_t stream);
+void normals(ftl::cuda::TextureObject<float4> &output,
+ ftl::cuda::TextureObject<float4> &temp,
+ ftl::cuda::TextureObject<float> &input,
+ int radius,
+ float smoothing,
+ const ftl::rgbd::Camera &camera,
+ const float3x3 &pose_inv, const float3x3 &pose, cudaStream_t stream);
+
void normals(ftl::cuda::TextureObject<float4> &output,
ftl::cuda::TextureObject<float> &input, // Integer depth values
const ftl::rgbd::Camera &camera,
diff --git a/components/renderers/cpp/include/ftl/render/tri_render.hpp b/components/renderers/cpp/include/ftl/render/tri_render.hpp
index 57dc68a7b..ede35069a 100644
--- a/components/renderers/cpp/include/ftl/render/tri_render.hpp
+++ b/components/renderers/cpp/include/ftl/render/tri_render.hpp
@@ -5,6 +5,7 @@
#include <ftl/rgbd/frameset.hpp>
#include <ftl/render/splat_params.hpp>
#include <ftl/cuda/points.hpp>
+#include <ftl/filters/filter.hpp>
namespace ftl {
namespace render {
@@ -42,11 +43,13 @@ class Triangular : public ftl::render::Renderer {
cudaStream_t stream_;
float3 light_pos_;
+ ftl::Filters *filters_;
+
template <typename T>
void __reprojectChannel(ftl::rgbd::Frame &, ftl::codecs::Channel in, ftl::codecs::Channel out, cudaStream_t);
void _reprojectChannel(ftl::rgbd::Frame &, ftl::codecs::Channel in, ftl::codecs::Channel out, cudaStream_t);
- void _dibr(cudaStream_t);
- void _mesh(cudaStream_t);
+ void _dibr(ftl::rgbd::Frame &, cudaStream_t);
+ void _mesh(ftl::rgbd::Frame &, ftl::rgbd::Source *, cudaStream_t);
};
}
diff --git a/components/renderers/cpp/src/normals.cu b/components/renderers/cpp/src/normals.cu
index 7dcdf2f5f..97452555d 100644
--- a/components/renderers/cpp/src/normals.cu
+++ b/components/renderers/cpp/src/normals.cu
@@ -208,6 +208,51 @@ __global__ void smooth_normals_kernel<0>(ftl::cuda::TextureObject<float4> norms,
output(x,y) = n;
}
+template <int RADIUS>
+__global__ void smooth_normals_kernel(ftl::cuda::TextureObject<float4> norms,
+ ftl::cuda::TextureObject<float4> output,
+ ftl::cuda::TextureObject<float> 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, depth.tex2D((int)x,(int)y));
+ float3 nsum = make_float3(0.0f);
+ float contrib = 0.0f;
+
+ output(x,y) = make_float4(0.0f,0.0f,0.0f,0.0f);
+
+ if (p0.z < camera.minDepth || p0.z > camera.maxDepth) return;
+
+ for (int v=-RADIUS; v<=RADIUS; ++v) {
+ for (int u=-RADIUS; u<=RADIUS; ++u) {
+ const float3 p = camera.screenToCam(x+u,y+v, depth.tex2D((int)x+u,(int)y+v));
+ if (p.z < camera.minDepth || p.z > camera.maxDepth) continue;
+ const float s = ftl::cuda::spatialWeighting(p0, p, smoothing);
+ //const float s = 1.0f;
+
+ //if (s > 0.0f) {
+ const float4 n = norms.tex2D((int)x+u,(int)y+v);
+ if (n.w > 0.0f) {
+ nsum += make_float3(n) * s;
+ contrib += s;
+ }
+ //}
+ }
+ }
+
+ // 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(pose*nsum, 1.0f) : make_float4(0.0f);
+}
+
void ftl::cuda::normals(ftl::cuda::TextureObject<float4> &output,
ftl::cuda::TextureObject<float4> &temp,
ftl::cuda::TextureObject<float4> &input,
@@ -266,6 +311,34 @@ void ftl::cuda::normals(ftl::cuda::TextureObject<float4> &output,
#endif
}
+void ftl::cuda::normals(ftl::cuda::TextureObject<float4> &output,
+ ftl::cuda::TextureObject<float4> &temp,
+ ftl::cuda::TextureObject<float> &input,
+ int radius,
+ float smoothing,
+ const ftl::rgbd::Camera &camera,
+ const float3x3 &pose_inv, const float3x3 &pose,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, camera);
+ cudaSafeCall( cudaGetLastError() );
+
+ switch (radius) {
+ case 7: smooth_normals_kernel<7><<<gridSize, blockSize, 0, stream>>>(temp, output, input, camera, pose, smoothing);
+ case 5: smooth_normals_kernel<5><<<gridSize, blockSize, 0, stream>>>(temp, output, input, camera, pose, smoothing);
+ case 3: smooth_normals_kernel<3><<<gridSize, blockSize, 0, stream>>>(temp, output, input, camera, pose, smoothing);
+ case 2: smooth_normals_kernel<2><<<gridSize, blockSize, 0, stream>>>(temp, output, input, camera, pose, smoothing);
+ case 1: smooth_normals_kernel<1><<<gridSize, blockSize, 0, stream>>>(temp, output, input, camera, pose, smoothing);
+ }
+ cudaSafeCall( cudaGetLastError() );
+
+ #ifdef _DEBUG
+ cudaSafeCall(cudaDeviceSynchronize());
+ //cutilCheckMsg(__FUNCTION__);
+ #endif
+}
+
void ftl::cuda::normals(ftl::cuda::TextureObject<float4> &output,
ftl::cuda::TextureObject<float> &input,
const ftl::rgbd::Camera &camera,
diff --git a/components/renderers/cpp/src/reprojection.cu b/components/renderers/cpp/src/reprojection.cu
index c11d23dd1..e74b4bb60 100644
--- a/components/renderers/cpp/src/reprojection.cu
+++ b/components/renderers/cpp/src/reprojection.cu
@@ -61,7 +61,7 @@ __device__ inline void accumulateOutput(TextureObject<float4> &out, TextureObjec
__global__ void reprojection_kernel(
TextureObject<A> in, // Attribute input
TextureObject<float> depth_src,
- TextureObject<int> depth_in, // Virtual depth map
+ TextureObject<float> depth_in, // Virtual depth map
TextureObject<float4> normals,
TextureObject<B> out, // Accumulated output
TextureObject<float> contrib,
@@ -71,7 +71,7 @@ __global__ void reprojection_kernel(
const int x = (blockIdx.x*blockDim.x + threadIdx.x);
const int y = blockIdx.y*blockDim.y + threadIdx.y;
- const float d = (float)depth_in.tex2D((int)x, (int)y) / 100000.0f;
+ const float d = depth_in.tex2D((int)x, (int)y);
if (d < params.camera.minDepth || d > params.camera.maxDepth) return;
const float3 worldPos = params.m_viewMatrixInverse * params.camera.screenToCam(x, y, d);
@@ -118,7 +118,7 @@ template <typename A, typename B>
void ftl::cuda::reproject(
TextureObject<A> &in,
TextureObject<float> &depth_src, // Original 3D points
- TextureObject<int> &depth_in, // Virtual depth map
+ TextureObject<float> &depth_in, // Virtual depth map
TextureObject<float4> &normals,
TextureObject<B> &out, // Accumulated output
TextureObject<float> &contrib,
@@ -144,7 +144,7 @@ void ftl::cuda::reproject(
template void ftl::cuda::reproject(
ftl::cuda::TextureObject<uchar4> &in, // Original colour image
ftl::cuda::TextureObject<float> &depth_src, // Original 3D points
- ftl::cuda::TextureObject<int> &depth_in, // Virtual depth map
+ ftl::cuda::TextureObject<float> &depth_in, // Virtual depth map
ftl::cuda::TextureObject<float4> &normals,
ftl::cuda::TextureObject<float4> &out, // Accumulated output
ftl::cuda::TextureObject<float> &contrib,
@@ -155,7 +155,7 @@ template void ftl::cuda::reproject(
template void ftl::cuda::reproject(
ftl::cuda::TextureObject<float> &in, // Original colour image
ftl::cuda::TextureObject<float> &depth_src, // Original 3D points
- ftl::cuda::TextureObject<int> &depth_in, // Virtual depth map
+ ftl::cuda::TextureObject<float> &depth_in, // Virtual depth map
ftl::cuda::TextureObject<float4> &normals,
ftl::cuda::TextureObject<float> &out, // Accumulated output
ftl::cuda::TextureObject<float> &contrib,
@@ -166,7 +166,7 @@ template void ftl::cuda::reproject(
template void ftl::cuda::reproject(
ftl::cuda::TextureObject<float4> &in, // Original colour image
ftl::cuda::TextureObject<float> &depth_src, // Original 3D points
- ftl::cuda::TextureObject<int> &depth_in, // Virtual depth map
+ ftl::cuda::TextureObject<float> &depth_in, // Virtual depth map
ftl::cuda::TextureObject<float4> &normals,
ftl::cuda::TextureObject<float4> &out, // Accumulated output
ftl::cuda::TextureObject<float> &contrib,
@@ -185,7 +185,7 @@ template void ftl::cuda::reproject(
__global__ void reprojection_kernel(
TextureObject<A> in, // Attribute input
TextureObject<float> depth_src,
- TextureObject<int> depth_in, // Virtual depth map
+ TextureObject<float> depth_in, // Virtual depth map
TextureObject<B> out, // Accumulated output
TextureObject<float> contrib,
SplatParams params,
@@ -194,7 +194,7 @@ __global__ void reprojection_kernel(
const int x = (blockIdx.x*blockDim.x + threadIdx.x);
const int y = blockIdx.y*blockDim.y + threadIdx.y;
- const float d = (float)depth_in.tex2D((int)x, (int)y) / 100000.0f;
+ const float d = depth_in.tex2D((int)x, (int)y);
if (d < params.camera.minDepth || d > params.camera.maxDepth) return;
const float3 worldPos = params.m_viewMatrixInverse * params.camera.screenToCam(x, y, d);
@@ -224,7 +224,7 @@ template <typename A, typename B>
void ftl::cuda::reproject(
TextureObject<A> &in,
TextureObject<float> &depth_src, // Original 3D points
- TextureObject<int> &depth_in, // Virtual depth map
+ TextureObject<float> &depth_in, // Virtual depth map
TextureObject<B> &out, // Accumulated output
TextureObject<float> &contrib,
const SplatParams ¶ms,
@@ -248,7 +248,7 @@ void ftl::cuda::reproject(
template void ftl::cuda::reproject(
ftl::cuda::TextureObject<uchar4> &in, // Original colour image
ftl::cuda::TextureObject<float> &depth_src, // Original 3D points
- ftl::cuda::TextureObject<int> &depth_in, // Virtual depth map
+ ftl::cuda::TextureObject<float> &depth_in, // Virtual depth map
ftl::cuda::TextureObject<float4> &out, // Accumulated output
ftl::cuda::TextureObject<float> &contrib,
const ftl::render::SplatParams ¶ms,
@@ -258,7 +258,7 @@ template void ftl::cuda::reproject(
template void ftl::cuda::reproject(
ftl::cuda::TextureObject<float> &in, // Original colour image
ftl::cuda::TextureObject<float> &depth_src, // Original 3D points
- ftl::cuda::TextureObject<int> &depth_in, // Virtual depth map
+ ftl::cuda::TextureObject<float> &depth_in, // Virtual depth map
ftl::cuda::TextureObject<float> &out, // Accumulated output
ftl::cuda::TextureObject<float> &contrib,
const ftl::render::SplatParams ¶ms,
@@ -268,7 +268,7 @@ template void ftl::cuda::reproject(
template void ftl::cuda::reproject(
ftl::cuda::TextureObject<float4> &in, // Original colour image
ftl::cuda::TextureObject<float> &depth_src, // Original 3D points
- ftl::cuda::TextureObject<int> &depth_in, // Virtual depth map
+ ftl::cuda::TextureObject<float> &depth_in, // Virtual depth map
ftl::cuda::TextureObject<float4> &out, // Accumulated output
ftl::cuda::TextureObject<float> &contrib,
const ftl::render::SplatParams ¶ms,
diff --git a/components/renderers/cpp/src/splatter_cuda.hpp b/components/renderers/cpp/src/splatter_cuda.hpp
index 038f5a3e5..838eda409 100644
--- a/components/renderers/cpp/src/splatter_cuda.hpp
+++ b/components/renderers/cpp/src/splatter_cuda.hpp
@@ -66,7 +66,7 @@ namespace cuda {
void reproject(
ftl::cuda::TextureObject<A> &in, // Original colour image
ftl::cuda::TextureObject<float> &depth_src, // Original 3D points
- ftl::cuda::TextureObject<int> &depth_in, // Virtual depth map
+ ftl::cuda::TextureObject<float> &depth_in, // Virtual depth map
ftl::cuda::TextureObject<float4> &normals,
ftl::cuda::TextureObject<B> &out, // Accumulated output
ftl::cuda::TextureObject<float> &contrib,
@@ -78,7 +78,7 @@ namespace cuda {
void reproject(
ftl::cuda::TextureObject<A> &in, // Original colour image
ftl::cuda::TextureObject<float> &depth_src, // Original 3D points
- ftl::cuda::TextureObject<int> &depth_in, // Virtual depth map
+ ftl::cuda::TextureObject<float> &depth_in, // Virtual depth map
ftl::cuda::TextureObject<B> &out, // Accumulated output
ftl::cuda::TextureObject<float> &contrib,
const ftl::render::SplatParams ¶ms,
diff --git a/components/renderers/cpp/src/tri_render.cpp b/components/renderers/cpp/src/tri_render.cpp
index 7f93a8f68..5804da220 100644
--- a/components/renderers/cpp/src/tri_render.cpp
+++ b/components/renderers/cpp/src/tri_render.cpp
@@ -7,6 +7,8 @@
#include <opencv2/core/cuda_stream_accessor.hpp>
+#include <ftl/filters/smoothing.hpp>
+
#include <string>
using ftl::render::Triangular;
@@ -129,6 +131,9 @@ Triangular::Triangular(nlohmann::json &config, ftl::rgbd::FrameSet *fs) : ftl::r
on("light_z", [this](const ftl::config::Event &e) { light_pos_.z = value("light_z", 0.3f); });
cudaSafeCall(cudaStreamCreate(&stream_));
+
+ filters_ = ftl::create<ftl::Filters>(this, "filters");
+ filters_->create<ftl::filters::DepthSmoother>("hfnoise");
}
Triangular::~Triangular() {
@@ -215,9 +220,9 @@ void Triangular::__reprojectChannel(ftl::rgbd::Frame &output, ftl::codecs::Chann
if (mesh_) {
ftl::cuda::reproject(
f.createTexture<T>(in),
- f.createTexture<float>(Channel::Depth), // TODO: Use depth?
- temp_.getTexture<int>(Channel::Depth2),
- accum_.getTexture<float4>(Channel::Normals),
+ f.createTexture<float>(Channel::Depth),
+ output.getTexture<float>(Channel::Depth),
+ output.getTexture<float4>(Channel::Normals),
temp_.createTexture<typename AccumSelector<T>::type>(AccumSelector<T>::channel),
temp_.getTexture<float>(Channel::Contribution),
params_,
@@ -228,8 +233,8 @@ void Triangular::__reprojectChannel(ftl::rgbd::Frame &output, ftl::codecs::Chann
// Can't use normals with point cloud version
ftl::cuda::reproject(
f.createTexture<T>(in),
- f.createTexture<float>(Channel::Depth), // TODO: Use depth?
- temp_.getTexture<int>(Channel::Depth2),
+ f.createTexture<float>(Channel::Depth),
+ output.getTexture<float>(Channel::Depth),
temp_.createTexture<typename AccumSelector<T>::type>(AccumSelector<T>::channel),
temp_.getTexture<float>(Channel::Contribution),
params_,
@@ -269,7 +274,7 @@ void Triangular::_reprojectChannel(ftl::rgbd::Frame &output, ftl::codecs::Channe
}
}
-void Triangular::_dibr(cudaStream_t stream) {
+void Triangular::_dibr(ftl::rgbd::Frame &out, cudaStream_t stream) {
cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
temp_.get<GpuMat>(Channel::Depth2).setTo(cv::Scalar(0x7FFFFFFF), cvstream);
@@ -288,9 +293,12 @@ void Triangular::_dibr(cudaStream_t stream) {
params_, stream
);
}
+
+ // Convert from int depth to float depth
+ temp_.get<GpuMat>(Channel::Depth2).convertTo(out.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 100000.0f, cvstream);
}
-void Triangular::_mesh(cudaStream_t stream) {
+void Triangular::_mesh(ftl::rgbd::Frame &out, ftl::rgbd::Source *src, cudaStream_t stream) {
cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
bool do_blend = value("mesh_blend", true);
@@ -346,6 +354,18 @@ void Triangular::_mesh(cudaStream_t stream) {
);
}
}
+
+ // Convert from int depth to float depth
+ temp_.get<GpuMat>(Channel::Depth2).convertTo(out.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 100000.0f, cvstream);
+
+ //filters_->filter(out, src, stream);
+
+ // Generate normals for final virtual image
+ ftl::cuda::normals(out.createTexture<float4>(Channel::Normals, Format<float4>(params_.camera.width, params_.camera.height)),
+ temp_.createTexture<float4>(Channel::Normals),
+ out.createTexture<float>(Channel::Depth),
+ value("normal_radius", 1), value("normal_smoothing", 0.02f),
+ params_.camera, params_.m_viewMatrix.getFloat3x3(), params_.m_viewMatrixInverse.getFloat3x3(), stream_);
}
void Triangular::_renderChannel(
@@ -500,17 +520,10 @@ bool Triangular::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out) {
// Create and render triangles for depth
if (mesh_) {
- _mesh(stream_);
+ _mesh(out, src, stream_);
} else {
- _dibr(stream_);
+ _dibr(out, stream_);
}
-
- // Generate normals for final virtual image
- ftl::cuda::normals(accum_.createTexture<float4>(Channel::Normals, Format<float4>(camera.width, camera.height)),
- temp_.createTexture<float4>(Channel::Normals),
- temp_.getTexture<int>(Channel::Depth2),
- value("normal_radius", 1), value("normal_smoothing", 0.02f),
- params_.camera, params_.m_viewMatrix.getFloat3x3(), params_.m_viewMatrixInverse.getFloat3x3(), stream_);
// Reprojection of colours onto surface
_renderChannel(out, Channel::Colour, Channel::Colour, stream_);
@@ -518,14 +531,16 @@ bool Triangular::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out) {
if (chan == Channel::Depth)
{
// Just convert int depth to float depth
- temp_.get<GpuMat>(Channel::Depth2).convertTo(out.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 100000.0f, cvstream);
+ //temp_.get<GpuMat>(Channel::Depth2).convertTo(out.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 100000.0f, cvstream);
} else if (chan == Channel::Normals) {
// Visualise normals to RGBA
- out.create<GpuMat>(Channel::Normals, Format<uchar4>(camera.width, camera.height)).setTo(cv::Scalar(0,0,0,0), cvstream);
- ftl::cuda::normal_visualise(accum_.getTexture<float4>(Channel::Normals), out.createTexture<uchar4>(Channel::Normals),
+ accum_.create<GpuMat>(Channel::Normals, Format<uchar4>(camera.width, camera.height)).setTo(cv::Scalar(0,0,0,0), cvstream);
+ ftl::cuda::normal_visualise(out.getTexture<float4>(Channel::Normals), accum_.createTexture<uchar4>(Channel::Normals),
light_pos_,
light_diffuse_,
light_ambient_, stream_);
+
+ accum_.swapTo(Channels(Channel::Normals), out);
}
//else if (chan == Channel::Contribution)
//{
@@ -558,9 +573,9 @@ bool Triangular::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out) {
// Need to re-dibr due to pose change
if (mesh_) {
- _mesh(stream_);
+ _mesh(out, src, stream_);
} else {
- _dibr(stream_);
+ _dibr(out, stream_);
}
_renderChannel(out, Channel::Left, Channel::Right, stream_);
--
GitLab