From e19d15417208a6faa346b6d572d806bbcca4ce76 Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Mon, 30 Sep 2019 13:10:06 +0300
Subject: [PATCH] Add motion window setting
---
applications/reconstruct/src/ilw/ilw.cpp | 6 ++++++
applications/reconstruct/src/ilw/ilw.cu | 12 ++++++++++--
applications/reconstruct/src/ilw/ilw.hpp | 1 +
applications/reconstruct/src/ilw/ilw_cuda.hpp | 1 +
4 files changed, 18 insertions(+), 2 deletions(-)
diff --git a/applications/reconstruct/src/ilw/ilw.cpp b/applications/reconstruct/src/ilw/ilw.cpp
index 1148762a3..23ba44f8e 100644
--- a/applications/reconstruct/src/ilw/ilw.cpp
+++ b/applications/reconstruct/src/ilw/ilw.cpp
@@ -17,6 +17,7 @@ ILW::ILW(nlohmann::json &config) : ftl::Configurable(config) {
enabled_ = value("ilw_align", true);
iterations_ = value("iterations", 1);
motion_rate_ = value("motion_rate", 0.4f);
+ motion_window_ = value("motion_window", 3);
on("ilw_align", [this](const ftl::config::Event &e) {
enabled_ = value("ilw_align", true);
@@ -30,6 +31,10 @@ ILW::ILW(nlohmann::json &config) : ftl::Configurable(config) {
motion_rate_ = value("motion_rate", 0.4f);
});
+ on("motion_window", [this](const ftl::config::Event &e) {
+ motion_window_ = value("motion_window", 3);
+ });
+
flags_ = 0;
if (value("ignore_bad", false)) flags_ |= ftl::cuda::kILWFlag_IgnoreBad;
if (value("restrict_z", true)) flags_ |= ftl::cuda::kILWFlag_RestrictZ;
@@ -181,6 +186,7 @@ bool ILW::_phase2(ftl::rgbd::FrameSet &fs, float rate, cudaStream_t stream) {
f.getTexture<float4>(Channel::EnergyVector),
fs.sources[i]->parameters(),
rate,
+ motion_window_,
stream
);
}
diff --git a/applications/reconstruct/src/ilw/ilw.cu b/applications/reconstruct/src/ilw/ilw.cu
index 96fd0727d..9ac1666d4 100644
--- a/applications/reconstruct/src/ilw/ilw.cu
+++ b/applications/reconstruct/src/ilw/ilw.cu
@@ -141,8 +141,9 @@ void ftl::cuda::correspondence_energy_vector(
//==============================================================================
-#define MOTION_RADIUS 3
+//#define MOTION_RADIUS 9
+template <int MOTION_RADIUS>
__global__ void move_points_kernel(
ftl::cuda::TextureObject<float4> p,
ftl::cuda::TextureObject<float4> ev,
@@ -184,12 +185,19 @@ void ftl::cuda::move_points(
ftl::cuda::TextureObject<float4> &v,
const ftl::rgbd::Camera &camera,
float rate,
+ int radius,
cudaStream_t stream) {
const dim3 gridSize((p.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (p.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
- move_points_kernel<<<gridSize, blockSize, 0, stream>>>(p,v,camera,rate);
+ switch (radius) {
+ case 9 : move_points_kernel<9><<<gridSize, blockSize, 0, stream>>>(p,v,camera,rate); break;
+ case 5 : move_points_kernel<5><<<gridSize, blockSize, 0, stream>>>(p,v,camera,rate); break;
+ case 3 : move_points_kernel<3><<<gridSize, blockSize, 0, stream>>>(p,v,camera,rate); break;
+ case 1 : move_points_kernel<1><<<gridSize, blockSize, 0, stream>>>(p,v,camera,rate); break;
+ case 0 : move_points_kernel<0><<<gridSize, blockSize, 0, stream>>>(p,v,camera,rate); break;
+ }
cudaSafeCall( cudaGetLastError() );
}
diff --git a/applications/reconstruct/src/ilw/ilw.hpp b/applications/reconstruct/src/ilw/ilw.hpp
index ce3b09795..068e18cc1 100644
--- a/applications/reconstruct/src/ilw/ilw.hpp
+++ b/applications/reconstruct/src/ilw/ilw.hpp
@@ -63,6 +63,7 @@ class ILW : public ftl::Configurable {
unsigned int flags_;
int iterations_;
float motion_rate_;
+ int motion_window_;
};
}
diff --git a/applications/reconstruct/src/ilw/ilw_cuda.hpp b/applications/reconstruct/src/ilw/ilw_cuda.hpp
index b9b9f8988..35d5977b1 100644
--- a/applications/reconstruct/src/ilw/ilw_cuda.hpp
+++ b/applications/reconstruct/src/ilw/ilw_cuda.hpp
@@ -30,6 +30,7 @@ void move_points(
ftl::cuda::TextureObject<float4> &v,
const ftl::rgbd::Camera &camera,
float rate,
+ int radius,
cudaStream_t stream
);
--
GitLab