diff --git a/applications/reconstruct/src/main.cpp b/applications/reconstruct/src/main.cpp index 18539c4e952d4cd77d22ffbfed18ce65eec9c4e6..c99a5380171d25e0da216ab49e2aa50f90fe130e 100644 --- a/applications/reconstruct/src/main.cpp +++ b/applications/reconstruct/src/main.cpp @@ -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 diff --git a/components/operators/CMakeLists.txt b/components/operators/CMakeLists.txt index ae282147dbbd3038d3dbe78981835748de0a81a3..734734256bb930b8f8532e9d308ce39600944bc4 100644 --- a/components/operators/CMakeLists.txt +++ b/components/operators/CMakeLists.txt @@ -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) diff --git a/components/operators/include/ftl/operators/mask.hpp b/components/operators/include/ftl/operators/mask.hpp new file mode 100644 index 0000000000000000000000000000000000000000..f2fc1ad5d9cc494a07c7fe3519b538219439b9d6 --- /dev/null +++ b/components/operators/include/ftl/operators/mask.hpp @@ -0,0 +1,29 @@ +#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_ diff --git a/components/operators/src/mask.cpp b/components/operators/src/mask.cpp new file mode 100644 index 0000000000000000000000000000000000000000..32d3c8c34a91455d668d48df22586f4f7dd63723 --- /dev/null +++ b/components/operators/src/mask.cpp @@ -0,0 +1,26 @@ +#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 diff --git a/components/operators/src/mask.cu b/components/operators/src/mask.cu new file mode 100644 index 0000000000000000000000000000000000000000..d7065b5083b313dad7d89f9dc851634a3581bb9b --- /dev/null +++ b/components/operators/src/mask.cu @@ -0,0 +1,52 @@ +#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 ¶ms, 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 +} diff --git a/components/operators/src/mask_cuda.hpp b/components/operators/src/mask_cuda.hpp new file mode 100644 index 0000000000000000000000000000000000000000..f1eb4972084652f54f092596d86bfa6f4cbdcfb6 --- /dev/null +++ b/components/operators/src/mask_cuda.hpp @@ -0,0 +1,53 @@ +#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 ¶ms, + int radius, + cudaStream_t stream); + +} +} + +#endif