From f98afd111e39d0bfea5416b38a68ef629c0a3a97 Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Mon, 30 Sep 2019 09:14:41 +0300
Subject: [PATCH] Add flags and ignore bad points
---
applications/reconstruct/src/ilw/ilw.cpp | 9 +++++++++
applications/reconstruct/src/ilw/ilw.cu | 12 ++++++------
applications/reconstruct/src/ilw/ilw.hpp | 1 +
applications/reconstruct/src/ilw/ilw_cuda.hpp | 3 +++
4 files changed, 19 insertions(+), 6 deletions(-)
diff --git a/applications/reconstruct/src/ilw/ilw.cpp b/applications/reconstruct/src/ilw/ilw.cpp
index 737a67d38..087aa100c 100644
--- a/applications/reconstruct/src/ilw/ilw.cpp
+++ b/applications/reconstruct/src/ilw/ilw.cpp
@@ -19,6 +19,14 @@ ILW::ILW(nlohmann::json &config) : ftl::Configurable(config) {
on("ilw_align", [this](const ftl::config::Event &e) {
enabled_ = value("ilw_align", true);
});
+
+ flags_ = 0;
+ if (value("ignore_bad", false)) flags_ |= ftl::cuda::kILWFlag_IgnoreBad;
+
+ 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;
+ });
}
ILW::~ILW() {
@@ -120,6 +128,7 @@ bool ILW::_phase1(ftl::rgbd::FrameSet &fs, cudaStream_t stream) {
f1.getTexture<float>(Channel::Energy),
pose,
s2->parameters(),
+ flags_,
stream
);
} catch (ftl::exception &e) {
diff --git a/applications/reconstruct/src/ilw/ilw.cu b/applications/reconstruct/src/ilw/ilw.cu
index e60778d6a..bb300d327 100644
--- a/applications/reconstruct/src/ilw/ilw.cu
+++ b/applications/reconstruct/src/ilw/ilw.cu
@@ -35,7 +35,7 @@ __global__ void correspondence_energy_vector_kernel(
TextureObject<float4> vout,
TextureObject<float> eout,
float4x4 pose2, // Inverse
- Camera cam2) {
+ Camera cam2, uint flags) {
// Each warp picks point in p1
const int tid = (threadIdx.x + threadIdx.y * blockDim.x);
@@ -61,13 +61,13 @@ __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;
const uchar4 colour2 = c2.tex2D(screen2.x+u, screen2.y+v);
- if (world2.x == MINF) continue;
// Determine degree of correspondence
float cost = 1.0f - ftl::cuda::spatialWeighting(world1, world2, 0.04f);
// Point is too far away to even count
- if (cost == 1.0f) continue;
+ 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;
@@ -75,7 +75,7 @@ __global__ void correspondence_energy_vector_kernel(
++count;
avgcost += cost;
- if (cost < bestcost) {
+ if (world2.x != MINF && cost < bestcost) {
bestpoint = world2;
bestcost = cost;
}
@@ -113,7 +113,7 @@ void ftl::cuda::correspondence_energy_vector(
TextureObject<float4> &vout,
TextureObject<float> &eout,
float4x4 &pose2,
- const Camera &cam2,
+ const Camera &cam2, uint flags,
cudaStream_t stream) {
const dim3 gridSize((p1.width() + 2 - 1)/2, (p1.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
@@ -122,7 +122,7 @@ void ftl::cuda::correspondence_energy_vector(
//printf("COR SIZE %d,%d\n", p1.width(), p1.height());
correspondence_energy_vector_kernel<<<gridSize, blockSize, 0, stream>>>(
- p1, p2, c1, c2, vout, eout, pose2, cam2
+ p1, p2, c1, c2, vout, eout, pose2, cam2, flags
);
cudaSafeCall( cudaGetLastError() );
}
diff --git a/applications/reconstruct/src/ilw/ilw.hpp b/applications/reconstruct/src/ilw/ilw.hpp
index 06a8f0d23..5e10cc7cf 100644
--- a/applications/reconstruct/src/ilw/ilw.hpp
+++ b/applications/reconstruct/src/ilw/ilw.hpp
@@ -60,6 +60,7 @@ class ILW : public ftl::Configurable {
std::vector<detail::ILWData> data_;
bool enabled_;
+ unsigned int flags_;
};
}
diff --git a/applications/reconstruct/src/ilw/ilw_cuda.hpp b/applications/reconstruct/src/ilw/ilw_cuda.hpp
index a4ba19713..b979ff7c0 100644
--- a/applications/reconstruct/src/ilw/ilw_cuda.hpp
+++ b/applications/reconstruct/src/ilw/ilw_cuda.hpp
@@ -8,6 +8,8 @@
namespace ftl {
namespace cuda {
+static const uint kILWFlag_IgnoreBad = 0x0001;
+
void correspondence_energy_vector(
ftl::cuda::TextureObject<float4> &p1,
ftl::cuda::TextureObject<float4> &p2,
@@ -17,6 +19,7 @@ void correspondence_energy_vector(
ftl::cuda::TextureObject<float> &eout,
float4x4 &pose2,
const ftl::rgbd::Camera &cam2,
+ uint flags,
cudaStream_t stream
);
--
GitLab