diff --git a/components/operators/include/ftl/operators/cuda/gt.hpp b/components/operators/include/ftl/operators/cuda/gt.hpp
index 17ef345321b1f1c114bd2378adce665eb0789780..e98600e6d7f7d4bb00d47cf9ecaeadd1bbca89f8 100644
--- a/components/operators/include/ftl/operators/cuda/gt.hpp
+++ b/components/operators/include/ftl/operators/cuda/gt.hpp
@@ -8,11 +8,12 @@ namespace ftl {
namespace cuda {
struct GTAnalysisData {
- int invalid; // number of invalid (missing depth and not masked)
- int invalid_masked; // number of masked pixels with value
- int masked; // number of pixels masked
+ int valid; // number of pixels with valid pixels
+ int missing; // number of missing non-masked pixels
+ int missing_masked; // number of missing masked pixels
+ int masked; // number of pixels masked (in gt)
- int bad; // number of bad pixels (error outside min/max threshold)
+ int good; // number of good pixels (valid value and error within min/max threshold)
float err; // sum of absolute error
float err_sq; // sum of squared error
};
@@ -21,6 +22,7 @@ void gt_analysis(
ftl::cuda::TextureObject<uchar4> &colour,
ftl::cuda::TextureObject<float> &depth,
ftl::cuda::TextureObject<float> >,
+ ftl::cuda::TextureObject<uchar> &mask,
ftl::cuda::GTAnalysisData *out,
const ftl::rgbd::Camera &cam,
float t_min,
@@ -33,6 +35,7 @@ void gt_analysis(
void gt_analysis(
ftl::cuda::TextureObject<float> &depth,
ftl::cuda::TextureObject<float> >,
+ ftl::cuda::TextureObject<uchar> &mask,
ftl::cuda::GTAnalysisData *out,
const ftl::rgbd::Camera &cam,
float t_min,
diff --git a/components/operators/src/gt_analysis.cpp b/components/operators/src/gt_analysis.cpp
index 7010d1ead3b8c4f018984c2d9e8abc341f7b3890..066f04d94c1ec96c2a01d8094e99cf86a66136bf 100644
--- a/components/operators/src/gt_analysis.cpp
+++ b/components/operators/src/gt_analysis.cpp
@@ -15,10 +15,10 @@ GTAnalysis::~GTAnalysis() {
template <typename T>
std::string to_string_with_precision(const T a_value, const int n = 6) {
- std::ostringstream out;
- out.precision(n);
- out << std::fixed << a_value;
- return out.str();
+ std::ostringstream out;
+ out.precision(n);
+ out << std::fixed << a_value;
+ return out.str();
}
struct Options {
@@ -28,28 +28,28 @@ struct Options {
};
static const std::vector<Options> options_disparity = {
- {-INFINITY, INFINITY, {0,0,224,255}}, // bad: over 2px error
- {-INFINITY, 2.0, {66,174,255,255}}, // 1 to 2px error
- {-INFINITY, 1.0, {16,192,16,255}}, // good: 0.1 to 1px error
- {-INFINITY, 0.25, {64,255,64,255}}, // less than 0.25px error
+ {-INFINITY, INFINITY, {0,0,224,255}},
+ {-INFINITY, 2.0, {66,174,255,255}},
+ {-INFINITY, 1.0, {16,192,16,255}},
+ {-INFINITY, 0.25, {64,255,64,255}},
};
static const std::vector<Options> options_depth = {
- {-INFINITY, INFINITY, {0,0,224,255}}, // over 100mm
- {-INFINITY, 0.1, {66,174,255,255}}, // 50 to 100mm
- {-INFINITY, 0.05, {16,192,16,255}}, // 10 to 50mm
- {-INFINITY, 0.01, {64,255,64,255}}, // less than 10 mm
+ {-INFINITY, INFINITY, {0,0,224,255}},
+ {-INFINITY, 0.1, {66,174,255,255}},
+ {-INFINITY, 0.025, {16,192,16,255}},
+ {-INFINITY, 0.01, {64,255,64,255}},
};
static void report(std::vector<std::string> &msgs, const ftl::cuda::GTAnalysisData &data,
const Options &o, float npixels, const std::string &unit="", float scale=1.0f) {
msgs.push_back( "(" + to_string_with_precision(o.t_min, 2)
- + "," + to_string_with_precision(o.t_max, 2) + "] "
- + to_string_with_precision(100.0f*data.bad/npixels, 1) + "%");
-
+ + "," + to_string_with_precision(o.t_max, 2) + "] ");
+ msgs.push_back("valid: " + to_string_with_precision(100.0f*data.good/data.valid, 1) + "%, "
+ + "all: " + to_string_with_precision(100.0f*data.good/npixels, 1) + "%");
msgs.push_back( "RMS: "
- + to_string_with_precision(sqrt(data.err_sq/data.bad) * scale, 2)
+ + to_string_with_precision(sqrt(data.err_sq/data.good) * scale, 2)
+ (unit.empty() ? "" : " " + unit));
}
@@ -73,6 +73,7 @@ bool GTAnalysis::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, cudaStream_t
in.createTexture<uchar4>(Channel::Colour),
in.createTexture<float>(Channel::Depth),
in.createTexture<float>(Channel::GroundTruth),
+ in.createTexture<uchar>(Channel::Mask),
output_,
in.getLeft(),
o.t_min,
@@ -86,6 +87,7 @@ bool GTAnalysis::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, cudaStream_t
ftl::cuda::gt_analysis(
in.createTexture<float>(Channel::Depth),
in.createTexture<float>(Channel::GroundTruth),
+ in.createTexture<uchar>(Channel::Mask),
output_,
in.getLeft(),
o.t_min,
@@ -101,8 +103,6 @@ bool GTAnalysis::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, cudaStream_t
else { report(msgs, err, o, npixels, "mm", 1000.0); }
}
- //float pmasked = float(err.masked) / npixels;
-
in.create(Channel::Messages, msgs);
return true;
diff --git a/components/operators/src/gt_analysis.cu b/components/operators/src/gt_analysis.cu
index 4cc83ad2b76b3a7488304c3b4a620ec28a7dc436..e3e524057868947c4220f7cd157ac2c939c4784b 100644
--- a/components/operators/src/gt_analysis.cu
+++ b/components/operators/src/gt_analysis.cu
@@ -16,25 +16,28 @@ __global__ void gt_analysis_kernel(
int dpitch,
const float* __restrict__ gt,
int gpitch,
+ const uchar* __restrict__ mask,
+ int mpitch,
ftl::cuda::GTAnalysisData *out,
ftl::rgbd::Camera cam,
float t_min,
float t_max,
uchar4 colour_value
) {
-
- __shared__ int sinvalid;
- __shared__ int sinvalid_masked;
+ __shared__ int svalid;
+ __shared__ int smissing;
+ __shared__ int smissing_masked;
__shared__ int smasked;
- __shared__ int sbad;
+ __shared__ int sgood;
__shared__ float serr;
__shared__ float serr_sq;
if (threadIdx.x == 0 && threadIdx.y == 0) {
- sinvalid = 0;
- sinvalid_masked = 0;
+ svalid = 0;
+ smissing = 0;
+ smissing_masked = 0;
smasked = 0;
- sbad = 0;
+ sgood = 0;
serr = 0.0f;
serr_sq = 0.0f;
}
@@ -42,10 +45,11 @@ __global__ void gt_analysis_kernel(
const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
- int invalid = 0;
- int invalid_masked = 0;
+ int valid = 0;
+ int missing = 0;
+ int missing_masked = 0;
int masked = 0;
- int bad = 0;
+ int good = 0;
float err = 0.0f;
float err_sq = 0.0f;
@@ -54,16 +58,21 @@ __global__ void gt_analysis_kernel(
if (x < width) {
const float* __restrict__ gt_ptr = gt+x;
const float* __restrict__ d_ptr = depth+x;
+ const uchar* __restrict__ m_ptr = mask+x;
for (STRIDE_Y(y, height)) {
// TODO: Verify gt and depth pitch are same
float gtval = gt_ptr[y*dpitch];
float dval = d_ptr[y*dpitch];
- const int tmasked = (gtval > cam.minDepth && gtval < cam.maxDepth) ? 0 : 1;
+ const int tmasked = (m_ptr[y*mpitch] == 0) ? 0 : 1;
const int tinvalid = (dval <= cam.minDepth || dval >= cam.maxDepth) ? 1 : 0;
+ const int tgtinvalid = (gtval > cam.minDepth && gtval < cam.maxDepth) ? 0 : 1;
+
+ if (tinvalid == 0 && tgtinvalid == 0) {
+ // if there is valid value in both (gt and depth)
+ valid += 1;
- if (tinvalid == 0 && tmasked == 0) {
if (DISPARITY) {
dval = (numer / dval);
gtval = (numer / gtval);
@@ -72,51 +81,51 @@ __global__ void gt_analysis_kernel(
const float e = fabsf(dval-gtval);
if ((t_min < e) && (e <= t_max)) {
- bad += 1;
+ good += 1;
err += e;
err_sq += e*e;
if (VISUALISE) { colour[x+y*cpitch] = colour_value; }
}
}
- if (VISUALISE) {
- if (tinvalid == 1 && tmasked == 1) {
- // correctly missing
- colour[x+y*cpitch] = {0, 0, 0, 255};
- }
- else if (tinvalid == 0 && tmasked == 1) {
- // masked and not missing
- colour[x+y*cpitch] = {128, 0, 128, 255};
- invalid_masked += 1;
- }
- else if (tinvalid == 1 && tmasked == 0) {
- // missing non-masked
- colour[x+y*cpitch] = {192, 0, 0, 255};
- invalid += 1;
- }
+ else if (tinvalid == 0 && tmasked == 1 && tgtinvalid == 1) {
+ // masked and not missing (but no gt value)
+ if (VISUALISE) { colour[x+y*cpitch] = {192, 0, 192, 255}; } // magenta
+ }
+ else if (tinvalid == 1 && (tmasked == 1 || tgtinvalid == 1)) {
+ // missing and (masked or missing gt)
+ if (VISUALISE) { colour[x+y*cpitch] = {0, 0, 0, 255}; } // black
+ missing_masked += 1;
+ }
+ else if (tinvalid == 1) {
+ // missing value (not masked)
+ if (VISUALISE) { colour[x+y*cpitch] = {224, 32, 32, 255}; } // blue
+ missing += 1;
}
- masked += tmasked;
+ masked += (tmasked == 1 || tgtinvalid == 1) ? 1 : 0;
}
}
// Warp aggregate
#pragma unroll
for (int i = WARP_SIZE/2; i > 0; i /= 2) {
- invalid += __shfl_xor_sync(FULL_MASK, invalid, i, WARP_SIZE);
- invalid_masked += __shfl_xor_sync(FULL_MASK, invalid_masked, i, WARP_SIZE);
+ valid += __shfl_xor_sync(FULL_MASK, valid, i, WARP_SIZE);
+ missing += __shfl_xor_sync(FULL_MASK, missing, i, WARP_SIZE);
+ missing_masked += __shfl_xor_sync(FULL_MASK, missing_masked, i, WARP_SIZE);
masked += __shfl_xor_sync(FULL_MASK, masked, i, WARP_SIZE);
- bad += __shfl_xor_sync(FULL_MASK, bad, i, WARP_SIZE);
+ good += __shfl_xor_sync(FULL_MASK, good, i, WARP_SIZE);
err += __shfl_xor_sync(FULL_MASK, err, i, WARP_SIZE);
err_sq += __shfl_xor_sync(FULL_MASK, err_sq, i, WARP_SIZE);
}
// Block aggregate
if (threadIdx.x % WARP_SIZE == 0) {
- atomicAdd(&sinvalid, invalid);
- atomicAdd(&sinvalid_masked, invalid_masked);
+ atomicAdd(&svalid, valid);
+ atomicAdd(&smissing, missing);
+ atomicAdd(&smissing_masked, missing_masked);
atomicAdd(&smasked, masked);
- atomicAdd(&sbad, bad);
+ atomicAdd(&sgood, good);
atomicAdd(&serr, err);
atomicAdd(&serr_sq, err_sq);
}
@@ -125,11 +134,11 @@ __global__ void gt_analysis_kernel(
// Global aggregate
if (threadIdx.x == 0 && threadIdx.y == 0) {
- //atomicAdd(&out->totalerror, serr);
- atomicAdd(&out->invalid, sinvalid);
- atomicAdd(&out->invalid_masked, sinvalid_masked);
+ atomicAdd(&out->valid, svalid);
+ atomicAdd(&out->missing, smissing);
+ atomicAdd(&out->missing_masked, smissing_masked);
atomicAdd(&out->masked, smasked);
- atomicAdd(&out->bad, sbad);
+ atomicAdd(&out->good, sgood);
atomicAdd(&out->err, serr);
atomicAdd(&out->err_sq, serr_sq);
}
@@ -139,6 +148,7 @@ void ftl::cuda::gt_analysis(
ftl::cuda::TextureObject<uchar4> &colour,
ftl::cuda::TextureObject<float> &depth,
ftl::cuda::TextureObject<float> >,
+ ftl::cuda::TextureObject<uchar> &mask,
ftl::cuda::GTAnalysisData *out,
const ftl::rgbd::Camera &cam,
float t_min,
@@ -165,6 +175,8 @@ void ftl::cuda::gt_analysis(
depth.pixelPitch(),
gt.devicePtr(),
gt.pixelPitch(),
+ mask.devicePtr(),
+ mask.pixelPitch(),
out,
cam,
t_min,
@@ -182,6 +194,8 @@ void ftl::cuda::gt_analysis(
depth.pixelPitch(),
gt.devicePtr(),
gt.pixelPitch(),
+ mask.devicePtr(),
+ mask.pixelPitch(),
out,
cam,
t_min,
@@ -199,6 +213,7 @@ void ftl::cuda::gt_analysis(
void ftl::cuda::gt_analysis(
ftl::cuda::TextureObject<float> &depth,
ftl::cuda::TextureObject<float> >,
+ ftl::cuda::TextureObject<uchar> &mask,
ftl::cuda::GTAnalysisData *out,
const ftl::rgbd::Camera &cam,
float t_min,
@@ -224,6 +239,8 @@ void ftl::cuda::gt_analysis(
depth.pixelPitch(),
gt.devicePtr(),
gt.pixelPitch(),
+ mask.devicePtr(),
+ mask.pixelPitch(),
out,
cam,
t_min,
@@ -241,6 +258,8 @@ void ftl::cuda::gt_analysis(
depth.pixelPitch(),
gt.devicePtr(),
gt.pixelPitch(),
+ mask.devicePtr(),
+ mask.pixelPitch(),
out,
cam,
t_min,