diff --git a/applications/reconstruct/src/ilw/ilw.cpp b/applications/reconstruct/src/ilw/ilw.cpp index 8d8f938678e58d1a01f40b64c1fbd96e86b08014..d175090cf70cb011b293e696c479509eb37629e0 100644 --- a/applications/reconstruct/src/ilw/ilw.cpp +++ b/applications/reconstruct/src/ilw/ilw.cpp @@ -64,6 +64,7 @@ ILW::ILW(nlohmann::json &config) : ftl::Configurable(config) { if (value("ignore_bad", false)) params_.flags |= ftl::cuda::kILWFlag_IgnoreBad; if (value("ignore_bad_colour", false)) params_.flags |= ftl::cuda::kILWFlag_SkipBadColour; if (value("restrict_z", true)) params_.flags |= ftl::cuda::kILWFlag_RestrictZ; + if (value("colour_confidence_only", false)) params_.flags |= ftl::cuda::kILWFlag_ColourConfidenceOnly; on("ignore_bad", [this](const ftl::config::Event &e) { if (value("ignore_bad", false)) params_.flags |= ftl::cuda::kILWFlag_IgnoreBad; @@ -79,6 +80,11 @@ ILW::ILW(nlohmann::json &config) : ftl::Configurable(config) { if (value("restrict_z", false)) params_.flags |= ftl::cuda::kILWFlag_RestrictZ; else params_.flags &= ~ftl::cuda::kILWFlag_RestrictZ; }); + + on("colour_confidence_only", [this](const ftl::config::Event &e) { + if (value("colour_confidence_only", false)) params_.flags |= ftl::cuda::kILWFlag_ColourConfidenceOnly; + else params_.flags &= ~ftl::cuda::kILWFlag_ColourConfidenceOnly; + }); } ILW::~ILW() { diff --git a/applications/reconstruct/src/ilw/ilw.cu b/applications/reconstruct/src/ilw/ilw.cu index 298075fda70abfccc27c2fdca24734e7237a97f3..b33a3e7feba861fa5923e5ddb245812ebe330c70 100644 --- a/applications/reconstruct/src/ilw/ilw.cu +++ b/applications/reconstruct/src/ilw/ilw.cu @@ -27,7 +27,7 @@ __device__ inline float warpSum(float e) { //#define COR_WIN_RADIUS 17 //#define COR_WIN_SIZE (COR_WIN_RADIUS * COR_WIN_RADIUS) -#define WINDOW_RADIUS 2 +#define WINDOW_RADIUS 1 template<int COR_STEPS> __global__ void correspondence_energy_vector_kernel( @@ -97,7 +97,7 @@ __global__ void correspondence_energy_vector_kernel( //cost /= 2.0f; ++count; - avgcost += cost; + avgcost += (params.flags & ftl::cuda::kILWFlag_ColourConfidenceOnly) ? ccost : cost; if (world2.x != MINF && cost < bestcost) { bestdepth = depth_adjust; bestcost = cost; @@ -111,7 +111,7 @@ __global__ void correspondence_energy_vector_kernel( const float mincost = warpMin(bestcost); bool best = mincost == bestcost; avgcost = warpSum(avgcost) / count; - const float confidence = (avgcost - mincost); + const float confidence = (params.flags & ftl::cuda::kILWFlag_ColourConfidenceOnly) ? avgcost : (avgcost - mincost); // FIXME: Multiple threads in warp could match this. if (best && mincost < 1.0f) { @@ -130,12 +130,14 @@ __global__ void correspondence_energy_vector_kernel( 0.0f, // * (1.0f - mincost) * confidence, bestdepth, // * (1.0f - mincost) * confidence, (1.0f - mincost) * confidence); + + eout(x,y) = max(eout(x, y), (1.0f - mincost) * confidence * 12.0f); } //eout(x,y) = max(eout(x,y), (length(bestpoint-world1) / 0.04f) * 7.0f); //eout(x,y) = max(eout(x,y), (1.0f - mincost) * 7.0f); //eout(x,y) = max(eout(x, y), (1.0f - mincost) * confidence * (length(bestpoint-world1) / 0.04f) * 12.0f); - eout(x,y) = max(eout(x, y), (1.0f - mincost) * confidence * 12.0f); + //eout(x,y) = max(eout(x, y), confidence * 12.0f); } else if (mincost >= 1.0f && lane == 0) { //vout(x,y) = make_float4(0.0f); @@ -162,7 +164,7 @@ void ftl::cuda::correspondence_energy_vector( //printf("COR SIZE %d,%d\n", p1.width(), p1.height()); - correspondence_energy_vector_kernel<64><<<gridSize, blockSize, 0, stream>>>(p1, p2, c1, c2, vout, eout, pose1, pose1_inv, pose2, cam1, cam2, params); + correspondence_energy_vector_kernel<32><<<gridSize, blockSize, 0, stream>>>(p1, p2, c1, c2, vout, eout, pose1, pose1_inv, pose2, cam1, cam2, params); //switch (win) { //case 17 : correspondence_energy_vector_kernel<17><<<gridSize, blockSize, 0, stream>>>(p1, p2, c1, c2, vout, eout, pose1, pose1_inv, pose2, cam1, cam2, params); break; diff --git a/applications/reconstruct/src/ilw/ilw_cuda.hpp b/applications/reconstruct/src/ilw/ilw_cuda.hpp index 03bde33d4db58b19597da4f30417c49a97461e60..4a7ed8ac635a80d20af07255fc47474810babc6f 100644 --- a/applications/reconstruct/src/ilw/ilw_cuda.hpp +++ b/applications/reconstruct/src/ilw/ilw_cuda.hpp @@ -19,6 +19,7 @@ struct ILWParams { static const uint kILWFlag_IgnoreBad = 0x0001; static const uint kILWFlag_RestrictZ = 0x0002; static const uint kILWFlag_SkipBadColour = 0x0004; +static const uint kILWFlag_ColourConfidenceOnly = 0x0008; void correspondence_energy_vector( ftl::cuda::TextureObject<float4> &p1,