Add motion window setting

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
         );
     }

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() );
 }

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_;
 };
 
 }

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
 );