From ba827213fb49fcfce4ec3ea8950643824b441e8f Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Sat, 2 Nov 2019 12:20:33 +0200
Subject: [PATCH] All mls radius option
---
components/filters/src/smoothing.cpp | 2 ++
components/filters/src/smoothing.cu | 9 ++++++++-
components/filters/src/smoothing_cuda.hpp | 1 +
3 files changed, 11 insertions(+), 1 deletion(-)
diff --git a/components/filters/src/smoothing.cpp b/components/filters/src/smoothing.cpp
index 19360840e..b288d3d5b 100644
--- a/components/filters/src/smoothing.cpp
+++ b/components/filters/src/smoothing.cpp
@@ -71,6 +71,7 @@ void MLSSmoother::filter(ftl::rgbd::Frame &f, ftl::rgbd::Source *s, cudaStream_t
//bool do_smooth = value("mls_smooth", false);
float thresh = value("mls_threshold", 0.04f);
int iters = value("mls_iterations", 1);
+ int radius = value("mls_radius",2);
//if (!do_smooth) return;
@@ -89,6 +90,7 @@ void MLSSmoother::filter(ftl::rgbd::Frame &f, ftl::rgbd::Source *s, cudaStream_t
f.createTexture<float>(Channel::Depth),
f.createTexture<float>(Channel::Depth2, ftl::rgbd::Format<float>(f.get<cv::cuda::GpuMat>(Channel::Depth).size())),
thresh,
+ radius,
s->parameters(),
0
);
diff --git a/components/filters/src/smoothing.cu b/components/filters/src/smoothing.cu
index b334d7c0e..903f90a19 100644
--- a/components/filters/src/smoothing.cu
+++ b/components/filters/src/smoothing.cu
@@ -77,13 +77,20 @@ void ftl::cuda::mls_smooth(
ftl::cuda::TextureObject<float> &depth_in,
ftl::cuda::TextureObject<float> &depth_out,
float smoothing,
+ int radius,
const ftl::rgbd::Camera &camera,
cudaStream_t stream) {
const dim3 gridSize((depth_out.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth_out.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
- mls_smooth_kernel<2><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, smoothing, camera);
+ switch (radius) {
+ case 5: mls_smooth_kernel<5><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, smoothing, camera); break;
+ case 4: mls_smooth_kernel<4><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, smoothing, camera); break;
+ case 3: mls_smooth_kernel<3><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, smoothing, camera); break;
+ case 2: mls_smooth_kernel<2><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, smoothing, camera); break;
+ case 1: mls_smooth_kernel<1><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, smoothing, camera); break;
+ }
cudaSafeCall( cudaGetLastError() );
diff --git a/components/filters/src/smoothing_cuda.hpp b/components/filters/src/smoothing_cuda.hpp
index d1ad1ce63..6ebe397e7 100644
--- a/components/filters/src/smoothing_cuda.hpp
+++ b/components/filters/src/smoothing_cuda.hpp
@@ -13,6 +13,7 @@ void mls_smooth(
ftl::cuda::TextureObject<float> &depth_in,
ftl::cuda::TextureObject<float> &depth_out,
float smoothing,
+ int radius,
const ftl::rgbd::Camera &camera,
cudaStream_t stream);
--
GitLab