Skip to content
Snippets Groups Projects
Commit b683e0ee authored by Nicolas Pope's avatar Nicolas Pope
Browse files

Tweaks to optical flow smoothing

parent 2141e8c0
Branches
Tags
No related merge requests found
......@@ -167,7 +167,16 @@ SourceWindow::SourceWindow(ftl::gui::Screen *screen)
bool SourceWindow::_processFrameset(ftl::rgbd::FrameSet &fs, bool fromstream) {
// Request the channels required by current camera configuration
if (fromstream) interceptor_->select(fs.id, _aggregateChannels(fs.id));
if (fromstream) {
auto cs = _aggregateChannels(fs.id);
auto avail = static_cast<const ftl::stream::Stream*>(interceptor_)->available(fs.id);
if (cs.has(Channel::Depth) && !avail.has(Channel::Depth) && avail.has(Channel::Right)) {
cs -= Channel::Depth;
cs += Channel::Right;
}
interceptor_->select(fs.id, cs);
}
/*if (fs.id > 0) {
LOG(INFO) << "Got frameset: " << fs.id;
......@@ -218,6 +227,7 @@ bool SourceWindow::_processFrameset(ftl::rgbd::FrameSet &fs, bool fromstream) {
void SourceWindow::_checkFrameSets(int id) {
while (framesets_.size() <= id) {
auto *p = ftl::config::create<ftl::operators::Graph>(screen_->root(), "pre_filters");
p->append<ftl::operators::DepthChannel>("depth");
//p->append<ftl::operators::ColourChannels>("colour"); // Convert BGR to BGRA
p->append<ftl::operators::DetectAndTrack>("facedetection")->value("enabled", false);
p->append<ftl::operators::ArUco>("aruco")->value("enabled", false);
......
......@@ -19,7 +19,7 @@ bool ftl::cuda::initialise() {
properties.resize(dev_count);
for (int i=0; i<dev_count; i++) {
cudaSafeCall(cudaGetDeviceProperties(&properties[i], i));
LOG(INFO) << " - " << properties[i].name;
LOG(INFO) << " - " << properties[i].name << " - compute " << properties[i].major << "." << properties[i].minor;
}
return true;
......@@ -29,7 +29,7 @@ bool ftl::cuda::hasCompute(int major, int minor) {
int dev = -1;
cudaSafeCall(cudaGetDevice(&dev));
if (dev > 0) {
if (dev >= 0) {
return properties[dev].major > major ||
(properties[dev].major == major && properties[dev].minor >= minor);
}
......
......@@ -109,6 +109,7 @@ class DepthChannel : public ftl::operators::Operator {
class OpticalFlowTemporalSmoothing : public ftl::operators::Operator {
public:
explicit OpticalFlowTemporalSmoothing(ftl::Configurable*);
OpticalFlowTemporalSmoothing(ftl::Configurable*, const std::tuple<ftl::codecs::Channel> &params);
~OpticalFlowTemporalSmoothing();
inline Operator::Type type() const override { return Operator::Type::OneToOne; }
......@@ -116,9 +117,10 @@ class OpticalFlowTemporalSmoothing : public ftl::operators::Operator {
bool apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, cudaStream_t stream) override;
private:
void _init(ftl::Configurable* cfg);
bool init();
const ftl::codecs::Channel channel_ = ftl::codecs::Channel::Disparity;
ftl::codecs::Channel channel_ = ftl::codecs::Channel::Depth;
cv::cuda::GpuMat history_;
cv::Size size_;
int n_max_;
......
......@@ -8,6 +8,7 @@
#include "ftl/operators/disparity.hpp"
#include "ftl/operators/depth.hpp"
#include "ftl/operators/mask.hpp"
#include "ftl/operators/opticalflow.hpp"
#include "./disparity/opencv/disparity_bilateral_filter.hpp"
......@@ -134,16 +135,21 @@ void DepthChannel::_createPipeline() {
config()->value("height", 720));
pipe_->append<ftl::operators::ColourChannels>("colour"); // Convert BGR to BGRA
pipe_->append<ftl::operators::CrossSupport>("cross");
#ifdef HAVE_LIBSGM
pipe_->append<ftl::operators::FixstarsSGM>("algorithm");
#endif
#ifdef HAVE_OPTFLOW
pipe_->append<ftl::operators::OpticalFlowTemporalSmoothing>("optflow_filter");
pipe_->append<ftl::operators::NVOpticalFlow>("optflow", Channel::Colour, Channel::Flow);
pipe_->append<ftl::operators::OpticalFlowTemporalSmoothing>("optflow_filter", Channel::Disparity);
#endif
pipe_->append<ftl::operators::DisparityBilateralFilter>("bilateral_filter");
//pipe_->append<ftl::operators::OpticalFlowTemporalSmoothing>("optflow_filter", Channel::Disparity);
pipe_->append<ftl::operators::DisparityToDepth>("calculate_depth");
#ifdef HAVE_OPTFLOW
//pipe_->append<ftl::operators::OpticalFlowTemporalSmoothing>("optflow_filter", Channel::Depth); // FIXME: Has a history so not with multiple sources!
#endif
pipe_->append<ftl::operators::Normals>("normals"); // Estimate surface normals
pipe_->append<ftl::operators::CrossSupport>("cross");
pipe_->append<ftl::operators::DiscontinuityMask>("discontinuity_mask");
pipe_->append<ftl::operators::AggreMLS>("mls"); // Perform MLS (using smoothing channel)
}
......@@ -161,7 +167,7 @@ bool DepthChannel::apply(ftl::rgbd::FrameSet &in, ftl::rgbd::FrameSet &out, cuda
cv::cuda::GpuMat& left = f.get<cv::cuda::GpuMat>(Channel::Left);
cv::cuda::GpuMat& right = f.get<cv::cuda::GpuMat>(Channel::Right);
cv::cuda::GpuMat& depth = f.create<cv::cuda::GpuMat>(Channel::Depth);
depth.create(depth_size_, CV_32FC1);
depth.create(left.size(), CV_32FC1);
if (left.empty() || right.empty()) continue;
......
......@@ -11,7 +11,7 @@ void depth_to_disparity(cv::cuda::GpuMat &disparity, const cv::cuda::GpuMat &dep
void optflow_filter(cv::cuda::GpuMat &disp, const cv::cuda::GpuMat &optflow,
cv::cuda::GpuMat &history, int n_max, float threshold,
cv::cuda::GpuMat &history, cv::cuda::GpuMat &support, int n_max, float threshold, bool fill,
cv::cuda::Stream &stream);
}
......
......@@ -18,45 +18,55 @@ using std::vector;
template<typename T> static bool inline isValidDisparity(T d) { return (0.0 < d) && (d < 256.0); } // TODO
OpticalFlowTemporalSmoothing::OpticalFlowTemporalSmoothing(ftl::Configurable* cfg, const std::tuple<ftl::codecs::Channel> &params) :
ftl::operators::Operator(cfg) {
channel_ = std::get<0>(params);
_init(cfg);
}
OpticalFlowTemporalSmoothing::OpticalFlowTemporalSmoothing(ftl::Configurable* cfg) :
ftl::operators::Operator(cfg) {
_init(cfg);
}
void OpticalFlowTemporalSmoothing::_init(ftl::Configurable* cfg) {
size_ = Size(0, 0);
n_max_ = cfg->value("history_size", 7);
n_max_ = cfg->value("history_size", 16);
if (n_max_ < 1) {
LOG(WARNING) << "History length must be larger than 0, using default (0)";
n_max_ = 7;
}
if (n_max_ > 32) {
if (n_max_ > 16) {
// TODO: cuda kernel uses fixed size buffer
LOG(WARNING) << "History length can't be larger than 32 (TODO)";
n_max_ = 32;
LOG(WARNING) << "History length can't be larger than 16 (TODO)";
n_max_ = 16;
}
threshold_ = cfg->value("threshold", 1.0);
threshold_ = cfg->value("threshold", 5.0f);
cfg->on("threshold", [this, &cfg](const ftl::config::Event&) {
float threshold = cfg->value("threshold", 1.0);
cfg->on("threshold", [this](const ftl::config::Event&) {
float threshold = config()->value("threshold", 5.0f);
if (threshold < 0.0) {
LOG(WARNING) << "invalid threshold " << threshold << ", value must be positive";
}
else {
threshold_ = threshold;
init();
//init();
}
});
cfg->on("history_size", [this, &cfg](const ftl::config::Event&) {
int n_max = cfg->value("history_size", 1.0);
int n_max = cfg->value("history_size", 7);
if (n_max < 1) {
LOG(WARNING) << "History length must be larger than 0";
}
else if (n_max_ > 32) {
else if (n_max_ > 16) {
// TODO: cuda kernel uses fixed size buffer
LOG(WARNING) << "History length can't be larger than 32 (TODO)";
LOG(WARNING) << "History length can't be larger than 16 (TODO)";
}
else {
n_max_ = n_max;
......@@ -86,7 +96,7 @@ bool OpticalFlowTemporalSmoothing::apply(Frame &in, Frame &out, cudaStream_t str
if (!init()) { return false; }
}
ftl::cuda::optflow_filter(data, optflow, history_, n_max_, threshold_, cvstream);
ftl::cuda::optflow_filter(data, optflow, history_, in.get<cv::cuda::GpuMat>(Channel::Support1), n_max_, threshold_, config()->value("filling", false), cvstream);
return true;
}
......
......@@ -14,28 +14,55 @@ __device__ void quicksort(float A[], size_t n)
}
template<typename T>
__device__ static bool inline isValidDisparity(T d) { return (0.0 < d) && (d < 256.0); } // TODO
__device__ static bool inline isValidDisparity(T d) { return d > 0.0f; }
static const int max_history = 32; // TODO dynamic shared memory
__device__ inline float supportArea(uchar4 s) {
const float dx = min(s.x,s.y);
const float dy = min(s.z,s.w);
return sqrt(dx*dx + dy*dy);
}
template <int FRAC>
__device__ inline short makeFixed(float v) {
return static_cast<short>(v * (1<<FRAC));
}
static const int MAX_HISTORY = 16; // TODO dynamic shared memory
template <bool FILLING, int HISTORY>
__global__ void temporal_median_filter_kernel(
cv::cuda::PtrStepSz<float> disp,
cv::cuda::PtrStepSz<int16_t> optflow,
cv::cuda::PtrStepSz<short2> optflow,
cv::cuda::PtrStepSz<float> history,
cv::cuda::PtrStepSz<uchar4> support,
int n_max,
int16_t threshold, // fixed point 10.5
float threshold,
float granularity // 4 for Turing
)
{
float sorted[max_history]; // TODO: dynamic shared memory
float sorted[HISTORY]; // TODO: dynamic shared memory
for (STRIDE_Y(y, disp.rows)) {
for (STRIDE_X(x, disp.cols)) {
int flowx = optflow(round(y / granularity), 2 * round(x / granularity));
int flowy = optflow(round(y / granularity), 2 * round(x / granularity) + 1);
float area = supportArea(support(y,x)) / 25.0f;
short2 flow = optflow(round(y / granularity), round(x / granularity));
//int flowy = optflow(round(y / granularity), 2 * round(x / granularity) + 1);
float t = area * threshold + 0.25f; // 0.25 is the 1/4 pixel accuracy NVIDIA claim
if ((abs(flowx) + abs(flowy)) > threshold)
if (max(abs(flow.x),abs(flow.y)) > makeFixed<5>(t))
{
// TODO: Perhaps rather than discard it could follow the optical flow
// This would require the generation of a depth flow also.
// Perhaps do optical flow on the right image and compare motions,
// small differences should indicate a change in depth. Or perhaps
// consider any scale change? But this works less well in more cases
// Most likely the above would have to be a totally separate process
// since the whole history would have to be moved and the idea of
// median breaks a little. Perhaps this operator is for static
// areas and another operator is for motion areas.
// last element in history[x][y][t]
history(y, (x + 1) * n_max - 1) = 0.0;
return;
......@@ -44,13 +71,16 @@ __global__ void temporal_median_filter_kernel(
int count = history(y, (x + 1) * n_max - 1);
int n = count % (n_max - 1);
if (isValidDisparity(disp(y, x)))
const float disparity = disp(y, x);
if (isValidDisparity(disparity))
{
history(y, (x + 1) * n_max - 1) += 1.0;
count++;
history(y, x * n_max + n) = disp(y, x);
history(y, x * n_max + n) = disparity;
}
if (FILLING || isValidDisparity(disparity)) {
int n_end = count;
if (n_end >= n_max) { n_end = n_max - 1; }
......@@ -62,7 +92,9 @@ __global__ void temporal_median_filter_kernel(
}
quicksort(sorted, n_end);
disp(y, x) = sorted[n_end / 2];
const float sd = sorted[n_end / 2];
if (isValidDisparity(sd)) disp(y, x) = sd;
}
}
}}
}
......@@ -71,7 +103,8 @@ namespace ftl {
namespace cuda {
void optflow_filter(cv::cuda::GpuMat &disp, const cv::cuda::GpuMat &optflow,
cv::cuda::GpuMat &history, int n, float threshold,
cv::cuda::GpuMat &history, cv::cuda::GpuMat &support,
int n, float threshold, bool fill,
cv::cuda::Stream &stream)
{
dim3 grid(1, 1, 1);
......@@ -79,12 +112,19 @@ void optflow_filter(cv::cuda::GpuMat &disp, const cv::cuda::GpuMat &optflow,
grid.x = cv::cuda::device::divUp(disp.cols, 128);
grid.y = cv::cuda::device::divUp(disp.rows, 1);
// TODO: dynamic shared memory
temporal_median_filter_kernel<<<grid, threads, 0, cv::cuda::StreamAccessor::getStream(stream)>>>
( disp, optflow, history, n,
round(threshold * (1 << 5)), // TODO: documentation; 10.5 format
if (fill) {
temporal_median_filter_kernel<true, MAX_HISTORY><<<grid, threads, 0, cv::cuda::StreamAccessor::getStream(stream)>>>
( disp, optflow, history, support, n,
threshold,
4 // TODO: (4 pixels granularity for Turing)
);
} else {
temporal_median_filter_kernel<false, MAX_HISTORY><<<grid, threads, 0, cv::cuda::StreamAccessor::getStream(stream)>>>
( disp, optflow, history, support, n,
threshold,
4 // TODO: (4 pixels granularity for Turing)
);
}
cudaSafeCall(cudaGetLastError());
}
......
#include <ftl/operators/opticalflow.hpp>
#include <ftl/exception.hpp>
#include <opencv2/cudaimgproc.hpp>
......@@ -14,6 +15,7 @@ using cv::cuda::GpuMat;
NVOpticalFlow::NVOpticalFlow(ftl::Configurable* cfg) :
ftl::operators::Operator(cfg), channel_in_(ftl::codecs::Channel::Colour), channel_out_(ftl::codecs::Channel::Flow) {
size_ = Size(0, 0);
}
NVOpticalFlow::NVOpticalFlow(ftl::Configurable*cfg, const std::tuple<ftl::codecs::Channel,ftl::codecs::Channel> &channels) : ftl::operators::Operator(cfg) {
......@@ -25,6 +27,10 @@ NVOpticalFlow::~NVOpticalFlow() {
}
bool NVOpticalFlow::init() {
if (!ftl::cuda::hasCompute(7,5)) {
config()->set("enabled", false);
throw FTL_Error("GPU does not support optical flow");
}
nvof_ = cv::cuda::NvidiaOpticalFlow_1_0::create(
size_.width, size_.height,
cv::cuda::NvidiaOpticalFlow_1_0::NV_OF_PERF_LEVEL_SLOW,
......@@ -36,7 +42,8 @@ bool NVOpticalFlow::init() {
}
bool NVOpticalFlow::apply(Frame &in, Frame &out, cudaStream_t stream) {
if (!in.hasChannel(channel_in_)) { return false; }
if (!in.hasChannel(channel_in_)) return false;
if (in.hasChannel(channel_out_)) return true;
if (in.get<GpuMat>(channel_in_).size() != size_) {
size_ = in.get<GpuMat>(channel_in_).size();
......@@ -48,6 +55,8 @@ bool NVOpticalFlow::apply(Frame &in, Frame &out, cudaStream_t stream) {
cv::cuda::cvtColor(in.get<GpuMat>(channel_in_), left_gray_, cv::COLOR_BGRA2GRAY, 0, cvstream);
// TODO: Use optical flow confidence output, perhaps combined with a
// sensitivity adjustment
nvof_->calc(left_gray_, left_gray_prev_, flow, cvstream);
std::swap(left_gray_, left_gray_prev_);
......
......@@ -18,7 +18,7 @@ PixelWeights::~PixelWeights() {
}
bool PixelWeights::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, cudaStream_t stream) {
if (in.hasChannel(Channel::Mask)) return true;
//if (in.hasChannel(Channel::Mask)) return true;
ftl::cuda::PixelWeightingParameters params;
//int radius = config()->value("radius", 2);
......
0% Loading or .
You are about to add 0 people to the discussion. Proceed with caution.
Please register or to comment