Skip to content
Snippets Groups Projects
Commit f4e3a365 authored by Sebastian Hahta's avatar Sebastian Hahta
Browse files

fix build

parent 650ff94d
No related branches found
No related tags found
1 merge request!105CUDA optical flow smoothing
Pipeline #13575 passed
......@@ -13,12 +13,11 @@ namespace rgbd {
class OFDisparityFilter {
public:
OFDisparityFilter() : n_max_(0), threshold_(0.0), size_(0, 0) {} // TODO: invalid state
OFDisparityFilter() : n_max_(0), threshold_(0.0) {}
OFDisparityFilter(cv::Size size, int n_frames, float threshold);
void filter(ftl::rgbd::Frame &frame, cv::cuda::Stream &stream);
private:
int n_;
int n_max_;
float threshold_;
......
......@@ -21,19 +21,16 @@ __global__ void temporal_median_filter_kernel(
cv::cuda::PtrStepSz<int16_t> optflow,
cv::cuda::PtrStepSz<float> history,
int n_max,
int16_t threshold // fixed point 10.5
uint granularity // 4 for Turing
int16_t threshold, // fixed point 10.5
float granularity // 4 for Turing
)
{
float sorted[max_history]; // TODO: dynamic shared memory
for (STRIDE_Y(y, disp.rows)) {
for (STRIDE_X(x, disp.cols)) {
int flowx = optflow(y, 2 * x);
int flowy = optflow(y, 2 * x + 1);
int flowx = optflow(round(y / granularity), 2 * round(x / granularity));
int flowy = optflow(round(y/ granularity), 2 * round(x / granularity) + 1);
int flowy = optflow(round(y / granularity), 2 * round(x / granularity) + 1);
if ((abs(flowx) + abs(flowy)) > threshold)
{
......@@ -83,8 +80,8 @@ void optflow_filter(cv::cuda::GpuMat &disp, const cv::cuda::GpuMat &optflow,
// 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
4, // TODO: (4 pixels granularity for Turing)
round(threshold * (1 << 5)), // TODO: documentation; 10.5 format
4 // TODO: (4 pixels granularity for Turing)
);
cudaSafeCall(cudaGetLastError());
......
......@@ -15,7 +15,7 @@ using std::vector;
template<typename T> static bool inline isValidDisparity(T d) { return (0.0 < d) && (d < 256.0); } // TODO
OFDisparityFilter::OFDisparityFilter(Size size, int n_frames, float threshold) :
n_(0), n_max_(n_frames + 1), threshold_(threshold)
n_max_(n_frames + 1), threshold_(threshold)
{
CHECK((n_max_ > 1) && (n_max_ >= 32)) << "History length must be between 0 and 31!";
disp_old_ = cv::cuda::GpuMat(cv::Size(size.width * n_max_, size.height), CV_32FC1);
......
0% Loading or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment