Skip to content
Snippets Groups Projects
Commit 9899dc7b authored by Nicolas Pope's avatar Nicolas Pope
Browse files

Add other ilw params

parent aa2d2089
No related branches found
No related tags found
2 merge requests!116Implements #133 point alignment,!114Ongoing #133 improvements
...@@ -19,6 +19,9 @@ ILW::ILW(nlohmann::json &config) : ftl::Configurable(config) { ...@@ -19,6 +19,9 @@ ILW::ILW(nlohmann::json &config) : ftl::Configurable(config) {
motion_rate_ = value("motion_rate", 0.4f); motion_rate_ = value("motion_rate", 0.4f);
motion_window_ = value("motion_window", 3); motion_window_ = value("motion_window", 3);
use_lab_ = value("use_Lab", false); 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) { on("ilw_align", [this](const ftl::config::Event &e) {
enabled_ = value("ilw_align", true); enabled_ = value("ilw_align", true);
...@@ -40,18 +43,30 @@ ILW::ILW(nlohmann::json &config) : ftl::Configurable(config) { ...@@ -40,18 +43,30 @@ ILW::ILW(nlohmann::json &config) : ftl::Configurable(config) {
use_lab_ = value("use_Lab", false); use_lab_ = value("use_Lab", false);
}); });
flags_ = 0; on("colour_smooth", [this](const ftl::config::Event &e) {
if (value("ignore_bad", false)) flags_ |= ftl::cuda::kILWFlag_IgnoreBad; params_.colour_smooth = value("colour_smooth", 50.0f);
if (value("restrict_z", true)) flags_ |= ftl::cuda::kILWFlag_RestrictZ; });
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) { on("ignore_bad", [this](const ftl::config::Event &e) {
if (value("ignore_bad", false)) flags_ |= ftl::cuda::kILWFlag_IgnoreBad; if (value("ignore_bad", false)) params_.flags |= ftl::cuda::kILWFlag_IgnoreBad;
else flags_ &= ~ftl::cuda::kILWFlag_IgnoreBad; else params_.flags &= ~ftl::cuda::kILWFlag_IgnoreBad;
}); });
on("restrict_z", [this](const ftl::config::Event &e) { on("restrict_z", [this](const ftl::config::Event &e) {
if (value("restrict_z", false)) flags_ |= ftl::cuda::kILWFlag_RestrictZ; if (value("restrict_z", false)) params_.flags |= ftl::cuda::kILWFlag_RestrictZ;
else 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) { ...@@ -163,7 +178,7 @@ bool ILW::_phase1(ftl::rgbd::FrameSet &fs, int win, cudaStream_t stream) {
pose1, pose1,
pose2, pose2,
s2->parameters(), s2->parameters(),
flags_, params_,
win, win,
stream stream
); );
......
...@@ -37,7 +37,7 @@ __global__ void correspondence_energy_vector_kernel( ...@@ -37,7 +37,7 @@ __global__ void correspondence_energy_vector_kernel(
TextureObject<float> eout, TextureObject<float> eout,
float4x4 pose1, // Inverse float4x4 pose1, // Inverse
float4x4 pose2, // Inverse float4x4 pose2, // Inverse
Camera cam2, uint flags) { Camera cam2, ftl::cuda::ILWParams params) {
// Each warp picks point in p1 // Each warp picks point in p1
const int tid = (threadIdx.x + threadIdx.y * blockDim.x); const int tid = (threadIdx.x + threadIdx.y * blockDim.x);
...@@ -63,16 +63,16 @@ __global__ void correspondence_energy_vector_kernel( ...@@ -63,16 +63,16 @@ __global__ void correspondence_energy_vector_kernel(
const float v = (i / COR_WIN_RADIUS) - (COR_WIN_RADIUS / 2); const float v = (i / COR_WIN_RADIUS) - (COR_WIN_RADIUS / 2);
const float3 world2 = make_float3(p2.tex2D(screen2.x+u, screen2.y+v)); 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); const uchar4 colour2 = c2.tex2D(screen2.x+u, screen2.y+v);
// Determine degree of correspondence // 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 // Point is too far away to even count
if (world2.x != MINF && cost == 1.0f) continue; if (world2.x != MINF && cost == 1.0f) continue;
// Mix ratio of colour and distance costs // 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; //cost /= 2.0f;
++count; ++count;
...@@ -92,7 +92,7 @@ __global__ void correspondence_energy_vector_kernel( ...@@ -92,7 +92,7 @@ __global__ void correspondence_energy_vector_kernel(
if (best && mincost < 1.0f) { if (best && mincost < 1.0f) {
float3 tvecA = pose1 * bestpoint; float3 tvecA = pose1 * bestpoint;
float3 tvecB = pose1 * world1; float3 tvecB = pose1 * world1;
if (flags & ftl::cuda::kILWFlag_RestrictZ) { if (params.flags & ftl::cuda::kILWFlag_RestrictZ) {
tvecA.x = tvecB.x; tvecA.x = tvecB.x;
tvecA.y = tvecB.y; tvecA.y = tvecB.y;
} }
...@@ -123,7 +123,7 @@ void ftl::cuda::correspondence_energy_vector( ...@@ -123,7 +123,7 @@ void ftl::cuda::correspondence_energy_vector(
TextureObject<float> &eout, TextureObject<float> &eout,
float4x4 &pose1, float4x4 &pose1,
float4x4 &pose2, float4x4 &pose2,
const Camera &cam2, uint flags, int win, const Camera &cam2, const ILWParams &params, int win,
cudaStream_t stream) { cudaStream_t stream) {
const dim3 gridSize((p1.width() + 2 - 1)/2, (p1.height() + T_PER_BLOCK - 1)/T_PER_BLOCK); 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( ...@@ -132,9 +132,9 @@ void ftl::cuda::correspondence_energy_vector(
//printf("COR SIZE %d,%d\n", p1.width(), p1.height()); //printf("COR SIZE %d,%d\n", p1.width(), p1.height());
switch (win) { 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 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, flags); 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, flags); 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() ); cudaSafeCall( cudaGetLastError() );
} }
......
...@@ -6,6 +6,8 @@ ...@@ -6,6 +6,8 @@
#include <ftl/configurable.hpp> #include <ftl/configurable.hpp>
#include <vector> #include <vector>
#include "ilw_cuda.hpp"
namespace ftl { namespace ftl {
namespace detail { namespace detail {
...@@ -62,7 +64,7 @@ class ILW : public ftl::Configurable { ...@@ -62,7 +64,7 @@ class ILW : public ftl::Configurable {
std::vector<detail::ILWData> data_; std::vector<detail::ILWData> data_;
bool enabled_; bool enabled_;
unsigned int flags_; ftl::cuda::ILWParams params_;
int iterations_; int iterations_;
float motion_rate_; float motion_rate_;
int motion_window_; int motion_window_;
......
...@@ -8,6 +8,14 @@ ...@@ -8,6 +8,14 @@
namespace ftl { namespace ftl {
namespace cuda { 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_IgnoreBad = 0x0001;
static const uint kILWFlag_RestrictZ = 0x0002; static const uint kILWFlag_RestrictZ = 0x0002;
...@@ -21,7 +29,7 @@ void correspondence_energy_vector( ...@@ -21,7 +29,7 @@ void correspondence_energy_vector(
float4x4 &pose1, float4x4 &pose1,
float4x4 &pose2, float4x4 &pose2,
const ftl::rgbd::Camera &cam2, const ftl::rgbd::Camera &cam2,
uint flags, int win, const ILWParams &params, int win,
cudaStream_t stream cudaStream_t stream
); );
......
0% Loading or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment