diff --git a/applications/reconstruct/src/ilw/ilw.cpp b/applications/reconstruct/src/ilw/ilw.cpp
index a1828a472c3978e55d0b485bde8b65535c443ea9..a7fbdd4004bc09d241890dbc4f1a1cfabb0c11cb 100644
--- a/applications/reconstruct/src/ilw/ilw.cpp
+++ b/applications/reconstruct/src/ilw/ilw.cpp
@@ -19,6 +19,9 @@ ILW::ILW(nlohmann::json &config) : ftl::Configurable(config) {
motion_rate_ = value("motion_rate", 0.4f);
motion_window_ = value("motion_window", 3);
use_lab_ = value("use_Lab", false);
+ params_.colour_smooth = value("colour_smooth", 50.0f);
+ params_.spatial_smooth = value("spatial_smooth", 0.04f);
+ params_.cost_ratio = value("cost_ratio", 75.0f);
on("ilw_align", [this](const ftl::config::Event &e) {
enabled_ = value("ilw_align", true);
@@ -40,18 +43,30 @@ ILW::ILW(nlohmann::json &config) : ftl::Configurable(config) {
use_lab_ = value("use_Lab", false);
});
- flags_ = 0;
- if (value("ignore_bad", false)) flags_ |= ftl::cuda::kILWFlag_IgnoreBad;
- if (value("restrict_z", true)) flags_ |= ftl::cuda::kILWFlag_RestrictZ;
+ on("colour_smooth", [this](const ftl::config::Event &e) {
+ params_.colour_smooth = value("colour_smooth", 50.0f);
+ });
+
+ on("spatial_smooth", [this](const ftl::config::Event &e) {
+ params_.spatial_smooth = value("spatial_smooth", 0.04f);
+ });
+
+ on("cost_ratio", [this](const ftl::config::Event &e) {
+ params_.cost_ratio = value("cost_ratio", 75.0f);
+ });
+
+ params_.flags = 0;
+ if (value("ignore_bad", false)) params_.flags |= ftl::cuda::kILWFlag_IgnoreBad;
+ if (value("restrict_z", true)) params_.flags |= ftl::cuda::kILWFlag_RestrictZ;
on("ignore_bad", [this](const ftl::config::Event &e) {
- if (value("ignore_bad", false)) flags_ |= ftl::cuda::kILWFlag_IgnoreBad;
- else flags_ &= ~ftl::cuda::kILWFlag_IgnoreBad;
+ if (value("ignore_bad", false)) params_.flags |= ftl::cuda::kILWFlag_IgnoreBad;
+ else params_.flags &= ~ftl::cuda::kILWFlag_IgnoreBad;
});
on("restrict_z", [this](const ftl::config::Event &e) {
- if (value("restrict_z", false)) flags_ |= ftl::cuda::kILWFlag_RestrictZ;
- else flags_ &= ~ftl::cuda::kILWFlag_RestrictZ;
+ if (value("restrict_z", false)) params_.flags |= ftl::cuda::kILWFlag_RestrictZ;
+ else params_.flags &= ~ftl::cuda::kILWFlag_RestrictZ;
});
}
@@ -163,7 +178,7 @@ bool ILW::_phase1(ftl::rgbd::FrameSet &fs, int win, cudaStream_t stream) {
pose1,
pose2,
s2->parameters(),
- flags_,
+ params_,
win,
stream
);
diff --git a/applications/reconstruct/src/ilw/ilw.cu b/applications/reconstruct/src/ilw/ilw.cu
index 9ac1666d49281658fef27562d6f8ee577cada9c9..9da781d91994fa6dd0214ddd4df1ccd1d1cdaecb 100644
--- a/applications/reconstruct/src/ilw/ilw.cu
+++ b/applications/reconstruct/src/ilw/ilw.cu
@@ -37,7 +37,7 @@ __global__ void correspondence_energy_vector_kernel(
TextureObject<float> eout,
float4x4 pose1, // Inverse
float4x4 pose2, // Inverse
- Camera cam2, uint flags) {
+ Camera cam2, ftl::cuda::ILWParams params) {
// Each warp picks point in p1
const int tid = (threadIdx.x + threadIdx.y * blockDim.x);
@@ -63,16 +63,16 @@ __global__ void correspondence_energy_vector_kernel(
const float v = (i / COR_WIN_RADIUS) - (COR_WIN_RADIUS / 2);
const float3 world2 = make_float3(p2.tex2D(screen2.x+u, screen2.y+v));
- if (flags & ftl::cuda::kILWFlag_IgnoreBad && world2.x == MINF) continue;
+ if (params.flags & ftl::cuda::kILWFlag_IgnoreBad && world2.x == MINF) continue;
const uchar4 colour2 = c2.tex2D(screen2.x+u, screen2.y+v);
// Determine degree of correspondence
- float cost = 1.0f - ftl::cuda::spatialWeighting(world1, world2, 0.04f);
+ float cost = 1.0f - ftl::cuda::spatialWeighting(world1, world2, params.spatial_smooth);
// Point is too far away to even count
if (world2.x != MINF && cost == 1.0f) continue;
// Mix ratio of colour and distance costs
- cost = 0.75f * (1.0f - ftl::cuda::colourWeighting(colour1, colour2, 50.0f)) + 0.25 * cost;
+ cost = params.cost_ratio * (1.0f - ftl::cuda::colourWeighting(colour1, colour2, params.colour_smooth)) + (1.0f - params.cost_ratio) * cost;
//cost /= 2.0f;
++count;
@@ -92,7 +92,7 @@ __global__ void correspondence_energy_vector_kernel(
if (best && mincost < 1.0f) {
float3 tvecA = pose1 * bestpoint;
float3 tvecB = pose1 * world1;
- if (flags & ftl::cuda::kILWFlag_RestrictZ) {
+ if (params.flags & ftl::cuda::kILWFlag_RestrictZ) {
tvecA.x = tvecB.x;
tvecA.y = tvecB.y;
}
@@ -123,7 +123,7 @@ void ftl::cuda::correspondence_energy_vector(
TextureObject<float> &eout,
float4x4 &pose1,
float4x4 &pose2,
- const Camera &cam2, uint flags, int win,
+ const Camera &cam2, const ILWParams ¶ms, int win,
cudaStream_t stream) {
const dim3 gridSize((p1.width() + 2 - 1)/2, (p1.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
@@ -132,9 +132,9 @@ void ftl::cuda::correspondence_energy_vector(
//printf("COR SIZE %d,%d\n", p1.width(), p1.height());
switch (win) {
- case 17 : correspondence_energy_vector_kernel<17><<<gridSize, blockSize, 0, stream>>>(p1, p2, c1, c2, vout, eout, pose1, pose2, cam2, flags); break;
- case 9 : correspondence_energy_vector_kernel<9><<<gridSize, blockSize, 0, stream>>>(p1, p2, c1, c2, vout, eout, pose1, pose2, cam2, flags); break;
- case 5 : correspondence_energy_vector_kernel<5><<<gridSize, blockSize, 0, stream>>>(p1, p2, c1, c2, vout, eout, pose1, pose2, cam2, flags); break;
+ case 17 : correspondence_energy_vector_kernel<17><<<gridSize, blockSize, 0, stream>>>(p1, p2, c1, c2, vout, eout, pose1, pose2, cam2, params); break;
+ case 9 : correspondence_energy_vector_kernel<9><<<gridSize, blockSize, 0, stream>>>(p1, p2, c1, c2, vout, eout, pose1, pose2, cam2, params); break;
+ case 5 : correspondence_energy_vector_kernel<5><<<gridSize, blockSize, 0, stream>>>(p1, p2, c1, c2, vout, eout, pose1, pose2, cam2, params); break;
}
cudaSafeCall( cudaGetLastError() );
}
diff --git a/applications/reconstruct/src/ilw/ilw.hpp b/applications/reconstruct/src/ilw/ilw.hpp
index c6602586b4c10c05ef8574bcd8c564d108843aa7..66de927a285716b2a1d9efd1b2fdb8577a1f6149 100644
--- a/applications/reconstruct/src/ilw/ilw.hpp
+++ b/applications/reconstruct/src/ilw/ilw.hpp
@@ -6,6 +6,8 @@
#include <ftl/configurable.hpp>
#include <vector>
+#include "ilw_cuda.hpp"
+
namespace ftl {
namespace detail {
@@ -62,7 +64,7 @@ class ILW : public ftl::Configurable {
std::vector<detail::ILWData> data_;
bool enabled_;
- unsigned int flags_;
+ ftl::cuda::ILWParams params_;
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 35d5977b10707b13b5c34e96395060b5267d6295..63c5a57af17bdd9e9088326fd9f5bcd421777ae6 100644
--- a/applications/reconstruct/src/ilw/ilw_cuda.hpp
+++ b/applications/reconstruct/src/ilw/ilw_cuda.hpp
@@ -8,6 +8,14 @@
namespace ftl {
namespace cuda {
+struct ILWParams {
+ float spatial_smooth;
+ float colour_smooth;
+ float cost_ratio;
+ float threshold;
+ uint flags;
+};
+
static const uint kILWFlag_IgnoreBad = 0x0001;
static const uint kILWFlag_RestrictZ = 0x0002;
@@ -21,7 +29,7 @@ void correspondence_energy_vector(
float4x4 &pose1,
float4x4 &pose2,
const ftl::rgbd::Camera &cam2,
- uint flags, int win,
+ const ILWParams ¶ms, int win,
cudaStream_t stream
);