diff --git a/applications/reconstruct/src/ilw/ilw.cpp b/applications/reconstruct/src/ilw/ilw.cpp index 737a67d3834944a26ec0c2736f030e4370f637ab..087aa100c03157b7d3f17601fd919f9cbf7fb8f8 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 e60778d6a9bad74cbfa29cc377ee5ce28176adb1..bb300d327ee62aa7964d6ca0642fea8d5c95fb5b 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 06a8f0d2332d847e86d0033684b59698bcf980bd..5e10cc7cf6f2ff76143369848d7d6ee7618b194d 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 a4ba19713dbb492bcfc25232c61c176537397823..b979ff7c08334552aed9c70eac5edd3a512b9a5d 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 );