From 7f0af72337bb05a43f7fe8a1cf50c43f6639dae0 Mon Sep 17 00:00:00 2001
From: Sebastian Hahta <joseha@utu.fi>
Date: Sun, 22 Mar 2020 18:46:30 +0200
Subject: [PATCH] ground truth analysis
---
applications/gui/src/camera.cpp | 28 +--
applications/gui/src/src_window.cpp | 8 +-
.../include/ftl/operators/cuda/gt.hpp | 23 ++-
components/operators/src/gt_analysis.cpp | 85 +++++---
components/operators/src/gt_analysis.cu | 188 ++++++++++++------
5 files changed, 225 insertions(+), 107 deletions(-)
diff --git a/applications/gui/src/camera.cpp b/applications/gui/src/camera.cpp
index a2b9ef5b8..24fb0aedb 100644
--- a/applications/gui/src/camera.cpp
+++ b/applications/gui/src/camera.cpp
@@ -48,11 +48,11 @@ static Eigen::Affine3d create_rotation_matrix(float ax, float ay, float az) {
ftl::gui::Camera::Camera(ftl::gui::Screen *screen, int fsmask, int fid, ftl::codecs::Channel c)
: screen_(screen), fsmask_(fsmask), fid_(fid), texture1_(GLTexture::Type::BGRA), texture2_(GLTexture::Type::BGRA), depth1_(GLTexture::Type::Float), channel_(c),channels_(0u) {
-
+
eye_ = Eigen::Vector3d::Zero();
neye_ = Eigen::Vector4d::Zero();
rotmat_.setIdentity();
-
+
//up_ = Eigen::Vector3f(0,1.0f,0);
lerpSpeed_ = 0.999f;
sdepth_ = false;
@@ -98,7 +98,7 @@ ftl::gui::Camera::Camera(ftl::gui::Screen *screen, int fsmask, int fid, ftl::cod
});
intrinsics_ = ftl::create<ftl::Configurable>(renderer_, "intrinsics");
-
+
state_.getLeft() = ftl::rgbd::Camera::from(intrinsics_);
state_.getRight() = state_.getLeft();
@@ -389,7 +389,7 @@ void ftl::gui::Camera::_draw(std::vector<ftl::rgbd::FrameSet*> &fss) {
if (!post_pipe_) {
post_pipe_ = ftl::config::create<ftl::operators::Graph>(screen_->root(), "post_filters");
post_pipe_->append<ftl::operators::FXAA>("fxaa");
- post_pipe_->append<ftl::operators::GTAnalysis>("gtanal");
+ post_pipe_->append<ftl::operators::GTAnalysis>("gtanalyse");
}
post_pipe_->apply(frame_, frame_, 0);
@@ -587,14 +587,14 @@ bool ftl::gui::Camera::setVR(bool on) {
state_.getLeft().height = size_y;
state_.getRight().width = size_x;
state_.getRight().height = size_y;
-
+
intrinsic = getCameraMatrix(screen_->getVR(), vr::Eye_Left);
CHECK(intrinsic(0, 2) < 0 && intrinsic(1, 2) < 0);
state_.getLeft().fx = intrinsic(0,0);
state_.getLeft().fy = intrinsic(0,0);
state_.getLeft().cx = intrinsic(0,2);
state_.getLeft().cy = intrinsic(1,2);
-
+
intrinsic = getCameraMatrix(screen_->getVR(), vr::Eye_Right);
CHECK(intrinsic(0, 2) < 0 && intrinsic(1, 2) < 0);
state_.getRight().fx = intrinsic(0,0);
@@ -671,7 +671,7 @@ const void ftl::gui::Camera::captureFrame() {
if (screen_->isVR()) {
#ifdef HAVE_OPENVR
-
+
vr::VRCompositor()->SetTrackingSpace(vr::TrackingUniverseStanding);
vr::VRCompositor()->WaitGetPoses(rTrackedDevicePose_, vr::k_unMaxTrackedDeviceCount, NULL, 0 );
@@ -679,12 +679,12 @@ const void ftl::gui::Camera::captureFrame() {
{
Eigen::Matrix4d eye_l = ConvertSteamVRMatrixToMatrix4(
vr::VRSystem()->GetEyeToHeadTransform(vr::Eye_Left));
-
+
//Eigen::Matrix4d eye_r = ConvertSteamVRMatrixToMatrix4(
// vr::VRSystem()->GetEyeToHeadTransform(vr::Eye_Left));
float baseline_in = 2.0 * eye_l(0, 3);
-
+
if (baseline_in != baseline_) {
baseline_ = baseline_in;
//src_->set("baseline", baseline_);
@@ -698,17 +698,17 @@ const void ftl::gui::Camera::captureFrame() {
vreye[0] = pose(0, 3);
vreye[1] = -pose(1, 3);
vreye[2] = -pose(2, 3);
-
+
// NOTE: If modified, should be verified with VR headset!
Eigen::Matrix3d R;
R = Eigen::AngleAxisd(ea[0], Eigen::Vector3d::UnitX()) *
Eigen::AngleAxisd(-ea[1], Eigen::Vector3d::UnitY()) *
- Eigen::AngleAxisd(-ea[2], Eigen::Vector3d::UnitZ());
-
+ Eigen::AngleAxisd(-ea[2], Eigen::Vector3d::UnitZ());
+
//double rd = 180.0 / 3.141592;
//LOG(INFO) << "rotation x: " << ea[0] *rd << ", y: " << ea[1] * rd << ", z: " << ea[2] * rd;
// pose.block<3, 3>(0, 0) = R;
-
+
rotmat_.block(0, 0, 3, 3) = R;
// TODO: Apply a rotation to orient also
@@ -720,7 +720,7 @@ const void ftl::gui::Camera::captureFrame() {
Eigen::Translation3d trans(eye_ + vreye);
Eigen::Affine3d t(trans);
viewPose = t.matrix() * rotmat_;
-
+
} else {
//LOG(ERROR) << "No VR Pose";
}
diff --git a/applications/gui/src/src_window.cpp b/applications/gui/src/src_window.cpp
index b3b71fabf..d397285fe 100644
--- a/applications/gui/src/src_window.cpp
+++ b/applications/gui/src/src_window.cpp
@@ -55,7 +55,7 @@ using std::vector;
using ftl::config::json_t;
static ftl::rgbd::Generator *createSourceGenerator(ftl::Configurable *root, const std::vector<ftl::rgbd::Source*> &srcs) {
-
+
auto *grp = new ftl::rgbd::Group();
/*auto pipeline = ftl::config::create<ftl::operators::Graph>(root, "pipeline");
pipeline->append<ftl::operators::DetectAndTrack>("facedetection")->value("enabled", false);
@@ -75,7 +75,7 @@ SourceWindow::SourceWindow(ftl::gui::Screen *screen)
setLayout(new nanogui::BoxLayout(nanogui::Orientation::Vertical, nanogui::Alignment::Fill, 20, 5));
using namespace nanogui;
-
+
new Label(this, "Select Camera","sans-bold",20);
// FIXME: Reallocating the vector may currently causes thread issues since
@@ -209,7 +209,7 @@ bool SourceWindow::_processFrameset(ftl::rgbd::FrameSet &fs, bool fromstream) {
}
if (!paused_) {
- if (!fs.test(ftl::data::FSFlag::PARTIAL) || !screen_->root()->value("drop_partial_framesets", false)) {
+ if (!fs.test(ftl::data::FSFlag::PARTIAL) || !screen_->root()->value("drop_partial_framesets", false)) {
// Enforce interpolated colour and GPU upload
for (size_t i=0; i<fs.frames.size(); ++i) {
if (!fs.hasFrame(i)) continue;
@@ -274,7 +274,7 @@ void SourceWindow::_checkFrameSets(size_t id) {
p->append<ftl::operators::CullDiscontinuity>("remove_discontinuity");
p->append<ftl::operators::MultiViewMLS>("mvmls")->value("enabled", false);
p->append<ftl::operators::Poser>("poser")->value("enabled", true);
- p->append<ftl::operators::GTAnalysis>("gtanal");
+ p->append<ftl::operators::GTAnalysis>("gtanalyse");
pre_pipelines_.push_back(p);
framesets_.push_back(new ftl::rgbd::FrameSet);
diff --git a/components/operators/include/ftl/operators/cuda/gt.hpp b/components/operators/include/ftl/operators/cuda/gt.hpp
index cd317a615..17ef34532 100644
--- a/components/operators/include/ftl/operators/cuda/gt.hpp
+++ b/components/operators/include/ftl/operators/cuda/gt.hpp
@@ -8,10 +8,13 @@ namespace ftl {
namespace cuda {
struct GTAnalysisData {
- int invalid; // Count of invalid (missing depth)
- int bad; // Count bad (above x disparity error)
- float totalerror; // Summed disparity error (of valid values)
- int masked; // Count of pixels masked.
+ 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 bad; // number of bad pixels (error outside min/max threshold)
+ float err; // sum of absolute error
+ float err_sq; // sum of squared error
};
void gt_analysis(
@@ -20,8 +23,10 @@ void gt_analysis(
ftl::cuda::TextureObject<float> >,
ftl::cuda::GTAnalysisData *out,
const ftl::rgbd::Camera &cam,
- float threshold,
- float outmax,
+ float t_min,
+ float t_max,
+ uchar4 colour_value,
+ bool use_disparity,
cudaStream_t stream
);
@@ -30,11 +35,13 @@ void gt_analysis(
ftl::cuda::TextureObject<float> >,
ftl::cuda::GTAnalysisData *out,
const ftl::rgbd::Camera &cam,
- float threshold,
+ float t_min,
+ float t_max,
+ bool use_disparity,
cudaStream_t stream
);
}
}
-#endif
\ No newline at end of file
+#endif
\ No newline at end of file
diff --git a/components/operators/src/gt_analysis.cpp b/components/operators/src/gt_analysis.cpp
index d4a29a589..7010d1ead 100644
--- a/components/operators/src/gt_analysis.cpp
+++ b/components/operators/src/gt_analysis.cpp
@@ -21,8 +21,53 @@ std::string to_string_with_precision(const T a_value, const int n = 6) {
return out.str();
}
+struct Options {
+ float t_min;
+ float t_max;
+ uchar4 colour;
+};
+
+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
+};
+
+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
+};
+
+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) + "%");
+
+ msgs.push_back( "RMS: "
+ + to_string_with_precision(sqrt(data.err_sq/data.bad) * scale, 2)
+ + (unit.empty() ? "" : " " + unit));
+}
+
bool GTAnalysis::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, cudaStream_t stream) {
- if (in.hasChannel(Channel::Depth) && in.hasChannel(Channel::GroundTruth)) {
+
+ if (!in.hasChannel(Channel::Depth) || !in.hasChannel(Channel::GroundTruth)) {
+ return true;
+ }
+
+ std::vector<std::string> msgs;
+ if (in.hasChannel(Channel::Messages)) { in.get(Channel::Messages, msgs); }
+
+ bool use_disp = config()->value("use_disparity", true);
+ auto &dmat = in.get<cv::cuda::GpuMat>(Channel::Depth);
+ const float npixels = dmat.rows * dmat.cols;
+ ftl::cuda::GTAnalysisData err;
+
+ for (const auto &o : (use_disp ? options_disparity : options_depth)) {
if (config()->value("show_colour", false)) {
ftl::cuda::gt_analysis(
in.createTexture<uchar4>(Channel::Colour),
@@ -30,41 +75,35 @@ bool GTAnalysis::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, cudaStream_t
in.createTexture<float>(Channel::GroundTruth),
output_,
in.getLeft(),
- config()->value("bad_threshold", 2.0f),
- config()->value("viz_threshold", 5.0f),
+ o.t_min,
+ o.t_max,
+ o.colour,
+ use_disp,
stream
);
- } else {
+ }
+ else {
ftl::cuda::gt_analysis(
in.createTexture<float>(Channel::Depth),
in.createTexture<float>(Channel::GroundTruth),
output_,
in.getLeft(),
- config()->value("bad_threshold", 2.0f),
+ o.t_min,
+ o.t_max,
+ use_disp,
stream
);
}
- ftl::cuda::GTAnalysisData anal;
- cudaMemcpy(&anal, output_, sizeof(anal), cudaMemcpyDeviceToHost);
-
- auto &dmat = in.get<cv::cuda::GpuMat>(Channel::Depth);
- int totalvalid = dmat.cols*dmat.rows - anal.invalid - anal.masked;
- //int totaltested = dmat.cols*dmat.rows - anal.masked;
-
- float pbad = float(anal.bad) / float(totalvalid) * 100.0f;
- float pinvalid = float(anal.invalid) / float(dmat.cols*dmat.rows - anal.masked) * 100.0f;
- float avgerr = anal.totalerror / float(totalvalid) * 100.0f;
-
- std::vector<std::string> msgs;
- if (in.hasChannel(Channel::Messages)) in.get(Channel::Messages, msgs);
+ cudaMemcpy(&err, output_, sizeof(err), cudaMemcpyDeviceToHost);
+ msgs.push_back(" ");
+ if (use_disp) { report(msgs, err, o, npixels, "px", 1.0); }
+ else { report(msgs, err, o, npixels, "mm", 1000.0); }
+ }
- msgs.push_back(string("Bad %: ") + to_string_with_precision(pbad, 1));
- msgs.push_back(string("Invalid %: ") + to_string_with_precision(pinvalid,1));
- msgs.push_back(string("Avg Error: ") + to_string_with_precision(avgerr, 2));
+ //float pmasked = float(err.masked) / npixels;
- in.create(Channel::Messages, msgs);
- }
+ in.create(Channel::Messages, msgs);
return true;
}
diff --git a/components/operators/src/gt_analysis.cu b/components/operators/src/gt_analysis.cu
index 53b76cd63..4cc83ad2b 100644
--- a/components/operators/src/gt_analysis.cu
+++ b/components/operators/src/gt_analysis.cu
@@ -6,8 +6,8 @@
#define FULL_MASK 0xffffffff
-template <bool COLOUR>
-__global__ void gt_anal_kernel(
+template <bool DISPARITY, bool VISUALISE>
+__global__ void gt_analysis_kernel(
uchar4* __restrict__ colour,
int cpitch,
int width,
@@ -18,29 +18,36 @@ __global__ void gt_anal_kernel(
int gpitch,
ftl::cuda::GTAnalysisData *out,
ftl::rgbd::Camera cam,
- float threshold,
- float outmax
+ float t_min,
+ float t_max,
+ uchar4 colour_value
) {
__shared__ int sinvalid;
- __shared__ int sbad;
+ __shared__ int sinvalid_masked;
__shared__ int smasked;
+ __shared__ int sbad;
__shared__ float serr;
+ __shared__ float serr_sq;
if (threadIdx.x == 0 && threadIdx.y == 0) {
sinvalid = 0;
- sbad = 0;
+ sinvalid_masked = 0;
smasked = 0;
+ sbad = 0;
serr = 0.0f;
+ serr_sq = 0.0f;
}
__syncthreads();
const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
int invalid = 0;
- int bad = 0;
+ int invalid_masked = 0;
int masked = 0;
+ int bad = 0;
float err = 0.0f;
+ float err_sq = 0.0f;
const float numer = cam.baseline*cam.fx;
@@ -54,57 +61,77 @@ __global__ void gt_anal_kernel(
float dval = d_ptr[y*dpitch];
const int tmasked = (gtval > cam.minDepth && gtval < cam.maxDepth) ? 0 : 1;
- const int tinvalid = (tmasked == 0 && (dval <= cam.minDepth || dval >= cam.maxDepth)) ? 1 : 0;
-
- uchar4 c = make_uchar4((tinvalid==1)?255:0,0,0,255);
+ const int tinvalid = (dval <= cam.minDepth || dval >= cam.maxDepth) ? 1 : 0;
- // Convert both to disparity...
if (tinvalid == 0 && tmasked == 0) {
- dval = (numer / dval);
- gtval = (numer / gtval);
+ if (DISPARITY) {
+ dval = (numer / dval);
+ gtval = (numer / gtval);
+ }
const float e = fabsf(dval-gtval);
- bad += (e >= threshold) ? 1 : 0;
- err += e;
- if (COLOUR) {
- float nerr = min(1.0f, e / outmax);
- c.z = min(255.0f, 255.0f * nerr);
+ if ((t_min < e) && (e <= t_max)) {
+ bad += 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;
}
}
- invalid += tinvalid;
masked += tmasked;
-
- if (COLOUR) colour[x+y*cpitch] = c;
}
}
// Warp aggregate
#pragma unroll
for (int i = WARP_SIZE/2; i > 0; i /= 2) {
- bad += __shfl_xor_sync(FULL_MASK, bad, i, WARP_SIZE);
invalid += __shfl_xor_sync(FULL_MASK, invalid, i, WARP_SIZE);
+ invalid_masked += __shfl_xor_sync(FULL_MASK, invalid_masked, i, WARP_SIZE);
masked += __shfl_xor_sync(FULL_MASK, masked, i, WARP_SIZE);
+ bad += __shfl_xor_sync(FULL_MASK, bad, 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(&serr, err);
- atomicAdd(&sbad, bad);
atomicAdd(&sinvalid, invalid);
+ atomicAdd(&sinvalid_masked, invalid_masked);
atomicAdd(&smasked, masked);
+ atomicAdd(&sbad, bad);
+ atomicAdd(&serr, err);
+ atomicAdd(&serr_sq, err_sq);
}
__syncthreads();
// Global aggregate
if (threadIdx.x == 0 && threadIdx.y == 0) {
- atomicAdd(&out->totalerror, serr);
- atomicAdd(&out->bad, sbad);
+ //atomicAdd(&out->totalerror, serr);
atomicAdd(&out->invalid, sinvalid);
+ atomicAdd(&out->invalid_masked, sinvalid_masked);
atomicAdd(&out->masked, smasked);
+ atomicAdd(&out->bad, sbad);
+ atomicAdd(&out->err, serr);
+ atomicAdd(&out->err_sq, serr_sq);
}
}
@@ -114,8 +141,10 @@ void ftl::cuda::gt_analysis(
ftl::cuda::TextureObject<float> >,
ftl::cuda::GTAnalysisData *out,
const ftl::rgbd::Camera &cam,
- float threshold,
- float outmax,
+ float t_min,
+ float t_max,
+ uchar4 colour_value,
+ bool use_disparity,
cudaStream_t stream
) {
static constexpr int THREADS_X = 128;
@@ -126,21 +155,41 @@ void ftl::cuda::gt_analysis(
cudaMemsetAsync(out, 0, sizeof(ftl::cuda::GTAnalysisData), stream);
- gt_anal_kernel<true><<<gridSize, blockSize, 0, stream>>>(
- colour.devicePtr(),
- colour.pixelPitch(),
- colour.width(),
- colour.height(),
- depth.devicePtr(),
- depth.pixelPitch(),
- gt.devicePtr(),
- gt.pixelPitch(),
- out,
- cam,
- threshold,
- outmax
- );
- cudaSafeCall( cudaGetLastError() );
+ if (use_disparity) {
+ gt_analysis_kernel<true, true><<<gridSize, blockSize, 0, stream>>>(
+ colour.devicePtr(),
+ colour.pixelPitch(),
+ colour.width(),
+ colour.height(),
+ depth.devicePtr(),
+ depth.pixelPitch(),
+ gt.devicePtr(),
+ gt.pixelPitch(),
+ out,
+ cam,
+ t_min,
+ t_max,
+ colour_value
+ );
+ }
+ else {
+ gt_analysis_kernel<false, true><<<gridSize, blockSize, 0, stream>>>(
+ colour.devicePtr(),
+ colour.pixelPitch(),
+ colour.width(),
+ colour.height(),
+ depth.devicePtr(),
+ depth.pixelPitch(),
+ gt.devicePtr(),
+ gt.pixelPitch(),
+ out,
+ cam,
+ t_min,
+ t_max,
+ colour_value
+ );
+ }
+ cudaSafeCall(cudaGetLastError());
#ifdef _DEBUG
cudaSafeCall(cudaDeviceSynchronize());
@@ -152,7 +201,9 @@ void ftl::cuda::gt_analysis(
ftl::cuda::TextureObject<float> >,
ftl::cuda::GTAnalysisData *out,
const ftl::rgbd::Camera &cam,
- float threshold,
+ float t_min,
+ float t_max,
+ bool use_disparity,
cudaStream_t stream
) {
static constexpr int THREADS_X = 128;
@@ -163,20 +214,41 @@ void ftl::cuda::gt_analysis(
cudaMemsetAsync(out, 0, sizeof(ftl::cuda::GTAnalysisData), stream);
- gt_anal_kernel<false><<<gridSize, blockSize, 0, stream>>>(
- nullptr,
- 0,
- depth.width(),
- depth.height(),
- depth.devicePtr(),
- depth.pixelPitch(),
- gt.devicePtr(),
- gt.pixelPitch(),
- out,
- cam,
- threshold,
- 1.0f
- );
+ if (use_disparity) {
+ gt_analysis_kernel<true, false><<<gridSize, blockSize, 0, stream>>>(
+ nullptr,
+ 0,
+ depth.width(),
+ depth.height(),
+ depth.devicePtr(),
+ depth.pixelPitch(),
+ gt.devicePtr(),
+ gt.pixelPitch(),
+ out,
+ cam,
+ t_min,
+ t_max,
+ {0,0,0,0}
+ );
+ }
+ else {
+ gt_analysis_kernel<false, false><<<gridSize, blockSize, 0, stream>>>(
+ nullptr,
+ 0,
+ depth.width(),
+ depth.height(),
+ depth.devicePtr(),
+ depth.pixelPitch(),
+ gt.devicePtr(),
+ gt.pixelPitch(),
+ out,
+ cam,
+ t_min,
+ t_max,
+ {0,0,0,0}
+ );
+ }
+
cudaSafeCall( cudaGetLastError() );
#ifdef _DEBUG
--
GitLab