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

Discon mask operator

parent 4025ac5d
No related branches found
No related tags found
1 merge request!165Implements #224 to use discontinuity mask
......@@ -35,6 +35,7 @@
#include <ftl/operators/normals.hpp>
#include <ftl/operators/filling.hpp>
#include <ftl/operators/segmentation.hpp>
#include <ftl/operators/mask.hpp>
#include <ftl/cuda/normals.hpp>
#include <ftl/registration.hpp>
......@@ -305,13 +306,15 @@ static void run(ftl::Configurable *root) {
// Create the source depth map pipeline
auto *pipeline1 = ftl::config::create<ftl::operators::Graph>(root, "pre_filters");
pipeline1->append<ftl::operators::ColourChannels>("colour"); // Convert BGR to BGRA
pipeline1->append<ftl::operators::HFSmoother>("hfnoise"); // Remove high-frequency noise
//pipeline1->append<ftl::operators::HFSmoother>("hfnoise"); // Remove high-frequency noise
pipeline1->append<ftl::operators::Normals>("normals"); // Estimate surface normals
//pipeline1->append<ftl::operators::SmoothChannel>("smoothing"); // Generate a smoothing channel
//pipeline1->append<ftl::operators::ScanFieldFill>("filling"); // Generate a smoothing channel
pipeline1->append<ftl::operators::CrossSupport>("cross");
pipeline1->append<ftl::operators::DiscontinuityMask>("discontinuity");
//pipeline1->append<ftl::operators::CullDiscontinuity>("remove_discontinuity");
pipeline1->append<ftl::operators::ColourMLS>("mls"); // Perform MLS (using smoothing channel)
pipeline1->append<ftl::operators::VisCrossSupport>("viscross");
pipeline1->append<ftl::operators::VisCrossSupport>("viscross")->set("enabled", false);
// Alignment
......
......@@ -14,6 +14,8 @@ set(OPERSRC
src/disparity/bilateral_filter.cpp
src/segmentation.cu
src/segmentation.cpp
src/mask.cu
src/mask.cpp
)
if (HAVE_OPTFLOW)
......
#ifndef _FTL_OPERATORS_MASK_HPP_
#define _FTL_OPERATORS_MASK_HPP_
#include <ftl/operators/operator.hpp>
#include <ftl/cuda_common.hpp>
namespace ftl {
namespace operators {
/**
* Generate a masking channel that indicates depth discontinuities within a
* specified radius from a pixel. This is useful for culling potentially bad
* depth and colour values when merging and smoothing.
*/
class DiscontinuityMask : public ftl::operators::Operator {
public:
explicit DiscontinuityMask(ftl::Configurable*);
~DiscontinuityMask();
inline Operator::Type type() const override { return Operator::Type::OneToOne; }
bool apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::Source *src, cudaStream_t stream) override;
};
}
}
#endif // _FTL_OPERATORS_MASK_HPP_
#include <ftl/operators/mask.hpp>
#include "mask_cuda.hpp"
using ftl::operators::DiscontinuityMask;
using ftl::codecs::Channel;
using ftl::rgbd::Format;
DiscontinuityMask::DiscontinuityMask(ftl::Configurable *cfg) : ftl::operators::Operator(cfg) {
}
DiscontinuityMask::~DiscontinuityMask() {
}
bool DiscontinuityMask::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::Source *s, cudaStream_t stream) {
int radius = config()->value("radius", 2);
ftl::cuda::discontinuity(
out.createTexture<int>(Channel::Mask, ftl::rgbd::Format<int>(in.get<cv::cuda::GpuMat>(Channel::Depth).size())),
in.createTexture<float>(Channel::Depth),
s->parameters(), radius, 0
);
return true;
}
\ No newline at end of file
#include "mask_cuda.hpp"
#define T_PER_BLOCK 8
using ftl::cuda::Mask;
template <int RADIUS>
__global__ void discontinuity_kernel(ftl::cuda::TextureObject<int> mask_out, ftl::cuda::TextureObject<float> depth, ftl::rgbd::Camera params) {
const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
if (x < params.width && y < params.height) {
Mask mask(0);
const float d = depth.tex2D((int)x, (int)y);
// Calculate depth between 0.0 and 1.0
//float p = (d - params.minDepth) / (params.maxDepth - params.minDepth);
if (d >= params.minDepth && d <= params.maxDepth) {
/* Orts-Escolano S. et al. 2016. Holoportation: Virtual 3D teleportation in real-time. */
// Is there a discontinuity nearby?
for (int u=-RADIUS; u<=RADIUS; ++u) {
for (int v=-RADIUS; v<=RADIUS; ++v) {
// If yes, the flag using w = -1
if (fabs(depth.tex2D((int)x+u, (int)y+v) - d) > 0.1f) mask.isDiscontinuity(true);
}
}
}
mask_out(x,y) = (int)mask;
}
}
void ftl::cuda::discontinuity(ftl::cuda::TextureObject<int> &mask_out, ftl::cuda::TextureObject<float> &depth, const ftl::rgbd::Camera &params, int discon, cudaStream_t stream) {
const dim3 gridSize((params.width + T_PER_BLOCK - 1)/T_PER_BLOCK, (params.height + T_PER_BLOCK - 1)/T_PER_BLOCK);
const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
switch (discon) {
case 5 : discontinuity_kernel<5><<<gridSize, blockSize, 0, stream>>>(mask_out, depth, params); break;
case 4 : discontinuity_kernel<4><<<gridSize, blockSize, 0, stream>>>(mask_out, depth, params); break;
case 3 : discontinuity_kernel<3><<<gridSize, blockSize, 0, stream>>>(mask_out, depth, params); break;
case 2 : discontinuity_kernel<2><<<gridSize, blockSize, 0, stream>>>(mask_out, depth, params); break;
case 1 : discontinuity_kernel<1><<<gridSize, blockSize, 0, stream>>>(mask_out, depth, params); break;
default: break;
}
cudaSafeCall( cudaGetLastError() );
#ifdef _DEBUG
cudaSafeCall(cudaDeviceSynchronize());
#endif
}
#ifndef _FTL_CUDA_MASK_HPP_
#define _FTL_CUDA_MASK_HPP_
#include <ftl/cuda_common.hpp>
#include <ftl/rgbd/camera.hpp>
namespace ftl {
namespace cuda {
/**
* Wrap an int mask value used to flag individual depth pixels.
*/
class Mask {
public:
__device__ inline Mask() : v_(0) {}
__device__ explicit inline Mask(int v) : v_(v) {}
#ifdef __CUDACC__
__device__ inline Mask(const ftl::cuda::TextureObject<int> &m, int x, int y) : v_(m.tex2D(x,y)) {}
#endif
__device__ inline operator int() const { return v_; }
__device__ inline bool is(int m) const { return v_ & m; }
__device__ inline bool isFilled() const { return v_ & kMask_Filled; }
__device__ inline bool isDiscontinuity() const { return v_ & kMask_Discontinuity; }
__device__ inline bool hasCorrespondence() const { return v_ & kMask_Correspondence; }
__device__ inline bool isBad() const { return v_ & kMask_Bad; }
__device__ inline void isFilled(bool v) { v_ = (v) ? v_ | kMask_Filled : v_ & (~kMask_Filled); }
__device__ inline void isDiscontinuity(bool v) { v_ = (v) ? v_ | kMask_Discontinuity : v_ & (~kMask_Discontinuity); }
__device__ inline void hasCorrespondence(bool v) { v_ = (v) ? v_ | kMask_Correspondence : v_ & (~kMask_Correspondence); }
__device__ inline void isBad(bool v) { v_ = (v) ? v_ | kMask_Bad : v_ & (~kMask_Bad); }
static constexpr int kMask_Filled = 0x0001;
static constexpr int kMask_Discontinuity = 0x0002;
static constexpr int kMask_Correspondence = 0x0004;
static constexpr int kMask_Bad = 0x0008;
private:
int v_;
};
void discontinuity(
ftl::cuda::TextureObject<int> &mast,
ftl::cuda::TextureObject<float> &depth,
const ftl::rgbd::Camera &params,
int radius,
cudaStream_t stream);
}
}
#endif
0% Loading or .
You are about to add 0 people to the discussion. Proceed with caution.
Please register or to comment