diff --git a/applications/reconstruct/src/ilw/ilw.cpp b/applications/reconstruct/src/ilw/ilw.cpp index 1148762a3b37522da9056d444b684aa7aa35b6e4..23ba44f8e9d45a1b8d705fe8e4bd37289df0823b 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 96fd0727d992ec31ef3e79f7ccc6da827dbaf32b..9ac1666d49281658fef27562d6f8ee577cada9c9 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 ce3b0979543c6f4204df9aad3b3f2d29aad14dfa..068e18cc1e15dc4ae79cfe30172e985e2fbb92a0 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 b9b9f89889825b61ff153d9f0c27f32c69f14e1c..35d5977b10707b13b5c34e96395060b5267d6295 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 );