From 758099dafe180a55cafc3cdbbb8f27c0fe8dd3e3 Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Wed, 30 Oct 2019 16:06:48 +0200
Subject: [PATCH] Iterate smoothing kernel
---
.../reconstruct/src/filters/smoothing.cu | 38 ++++++++++---------
.../reconstruct/src/filters/smoothing.hpp | 2 +-
applications/reconstruct/src/main.cpp | 15 +++++---
3 files changed, 30 insertions(+), 25 deletions(-)
diff --git a/applications/reconstruct/src/filters/smoothing.cu b/applications/reconstruct/src/filters/smoothing.cu
index b5557ed0b..9ea08594c 100644
--- a/applications/reconstruct/src/filters/smoothing.cu
+++ b/applications/reconstruct/src/filters/smoothing.cu
@@ -52,30 +52,32 @@ void ftl::cuda::depth_smooth(
ftl::cuda::TextureObject<uchar4> &colour_in,
ftl::cuda::TextureObject<float> &depth_out,
const ftl::rgbd::Camera &camera,
- int radius, float factor, float thresh, cudaStream_t stream) {
+ int radius, float factor, float thresh, int iters, 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);
- switch (radius) {
- case 5 : depth_smooth_kernel<5><<<gridSize, blockSize, 0, stream>>>(depth_in, colour_in, depth_out, camera, factor, thresh); break;
- case 4 : depth_smooth_kernel<4><<<gridSize, blockSize, 0, stream>>>(depth_in, colour_in, depth_out, camera, factor, thresh); break;
- case 3 : depth_smooth_kernel<3><<<gridSize, blockSize, 0, stream>>>(depth_in, colour_in, depth_out, camera, factor, thresh); break;
- case 2 : depth_smooth_kernel<2><<<gridSize, blockSize, 0, stream>>>(depth_in, colour_in, depth_out, camera, factor, thresh); break;
- case 1 : depth_smooth_kernel<1><<<gridSize, blockSize, 0, stream>>>(depth_in, colour_in, depth_out, camera, factor, thresh); break;
- default: break;
- }
- cudaSafeCall( cudaGetLastError() );
+ for (int n=0; n<iters; ++n) {
+ switch (radius) {
+ case 5 : depth_smooth_kernel<5><<<gridSize, blockSize, 0, stream>>>(depth_in, colour_in, depth_out, camera, factor, thresh); break;
+ case 4 : depth_smooth_kernel<4><<<gridSize, blockSize, 0, stream>>>(depth_in, colour_in, depth_out, camera, factor, thresh); break;
+ case 3 : depth_smooth_kernel<3><<<gridSize, blockSize, 0, stream>>>(depth_in, colour_in, depth_out, camera, factor, thresh); break;
+ case 2 : depth_smooth_kernel<2><<<gridSize, blockSize, 0, stream>>>(depth_in, colour_in, depth_out, camera, factor, thresh); break;
+ case 1 : depth_smooth_kernel<1><<<gridSize, blockSize, 0, stream>>>(depth_in, colour_in, depth_out, camera, factor, thresh); break;
+ default: break;
+ }
+ cudaSafeCall( cudaGetLastError() );
- switch (radius) {
- case 5 : depth_smooth_kernel<5><<<gridSize, blockSize, 0, stream>>>(depth_out, colour_in, depth_in, camera, factor, thresh); break;
- case 4 : depth_smooth_kernel<4><<<gridSize, blockSize, 0, stream>>>(depth_out, colour_in, depth_in, camera, factor, thresh); break;
- case 3 : depth_smooth_kernel<3><<<gridSize, blockSize, 0, stream>>>(depth_out, colour_in, depth_in, camera, factor, thresh); break;
- case 2 : depth_smooth_kernel<2><<<gridSize, blockSize, 0, stream>>>(depth_out, colour_in, depth_in, camera, factor, thresh); break;
- case 1 : depth_smooth_kernel<1><<<gridSize, blockSize, 0, stream>>>(depth_out, colour_in, depth_in, camera, factor, thresh); break;
- default: break;
+ switch (radius) {
+ case 5 : depth_smooth_kernel<5><<<gridSize, blockSize, 0, stream>>>(depth_out, colour_in, depth_in, camera, factor, thresh); break;
+ case 4 : depth_smooth_kernel<4><<<gridSize, blockSize, 0, stream>>>(depth_out, colour_in, depth_in, camera, factor, thresh); break;
+ case 3 : depth_smooth_kernel<3><<<gridSize, blockSize, 0, stream>>>(depth_out, colour_in, depth_in, camera, factor, thresh); break;
+ case 2 : depth_smooth_kernel<2><<<gridSize, blockSize, 0, stream>>>(depth_out, colour_in, depth_in, camera, factor, thresh); break;
+ case 1 : depth_smooth_kernel<1><<<gridSize, blockSize, 0, stream>>>(depth_out, colour_in, depth_in, camera, factor, thresh); break;
+ default: break;
+ }
+ cudaSafeCall( cudaGetLastError() );
}
- cudaSafeCall( cudaGetLastError() );
#ifdef _DEBUG
cudaSafeCall(cudaDeviceSynchronize());
diff --git a/applications/reconstruct/src/filters/smoothing.hpp b/applications/reconstruct/src/filters/smoothing.hpp
index 61f85e272..d5d6a4705 100644
--- a/applications/reconstruct/src/filters/smoothing.hpp
+++ b/applications/reconstruct/src/filters/smoothing.hpp
@@ -12,7 +12,7 @@ void depth_smooth(
ftl::cuda::TextureObject<uchar4> &colour_in,
ftl::cuda::TextureObject<float> &depth_out,
const ftl::rgbd::Camera &camera,
- int radius, float factor, float thresh,
+ int radius, float factor, float thresh, int iters,
cudaStream_t stream);
}
diff --git a/applications/reconstruct/src/main.cpp b/applications/reconstruct/src/main.cpp
index 4da112c8a..f3272e0d9 100644
--- a/applications/reconstruct/src/main.cpp
+++ b/applications/reconstruct/src/main.cpp
@@ -245,9 +245,11 @@ static void run(ftl::Configurable *root) {
bool busy = false;
+ auto *filter = ftl::config::create<ftl::Configurable>(root, "filters");
+
group->setLatency(4);
group->setName("ReconGroup");
- group->sync([splat,virt,&busy,&slave,&scene_A,&scene_B,&align,controls](ftl::rgbd::FrameSet &fs) -> bool {
+ group->sync([splat,virt,&busy,&slave,&scene_A,&scene_B,&align,controls,filter](ftl::rgbd::FrameSet &fs) -> bool {
//cudaSetDevice(scene->getCUDADevice());
//if (slave.isPaused()) return true;
@@ -262,7 +264,7 @@ static void run(ftl::Configurable *root) {
// Swap the entire frameset to allow rapid return
fs.swapTo(scene_A);
- ftl::pool.push([&scene_B,&scene_A,&busy,&slave,&align](int id) {
+ ftl::pool.push([&scene_B,&scene_A,&busy,&slave,&align, filter](int id) {
//cudaSetDevice(scene->getCUDADevice());
// TODO: Release frameset here...
//cudaSafeCall(cudaStreamSynchronize(scene->getIntegrationStream()));
@@ -270,9 +272,10 @@ static void run(ftl::Configurable *root) {
UNIQUE_LOCK(scene_A.mtx, lk);
cv::cuda::GpuMat tmp;
- float factor = align->value("smooth_factor", 0.04f);
- float colour_limit = align->value("colour_limit", 50.0f);
- bool do_smooth = align->value("pre_smooth", false);
+ float factor = filter->value("smooth_factor", 0.04f);
+ float colour_limit = filter->value("colour_limit", 50.0f);
+ bool do_smooth = filter->value("pre_smooth", false);
+ int iters = filter->value("iterations", 1);
if (do_smooth) {
// Presmooth...
@@ -295,7 +298,7 @@ static void run(ftl::Configurable *root) {
f.createTexture<uchar4>(Channel::Colour),
f.createTexture<float>(Channel::Depth2, ftl::rgbd::Format<float>(f.get<cv::cuda::GpuMat>(Channel::Depth).size())),
s->parameters(),
- 1, factor, colour_limit, 0
+ 1, factor, colour_limit, iters, 0
);
}
}
--
GitLab