diff --git a/applications/reconstruct/src/ilw/ilw_cuda.hpp b/applications/reconstruct/src/ilw/ilw_cuda.hpp index e4a099031e7c2da558384c694034ccd516380baa..60ea20539eee55d4f8abab4d68d644fa6a941618 100644 --- a/applications/reconstruct/src/ilw/ilw_cuda.hpp +++ b/applications/reconstruct/src/ilw/ilw_cuda.hpp @@ -22,6 +22,36 @@ static const uint kILWFlag_RestrictZ = 0x0002; static const uint kILWFlag_SkipBadColour = 0x0004; static const uint kILWFlag_ColourConfidenceOnly = 0x0008; +/** + * Wrap an int mask value used to flag individual depth pixels. + */ +class ILWMask { + public: + __device__ explicit inline ILWMask(int v) : v_(v) {} + #ifdef __CUDACC__ + __device__ inline ILWMask(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 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); } + + private: + int v_; + + static const int kMask_Filled = 0x0001; + static const int kMask_Discontinuity = 0x0002; + static const int kMask_Correspondence = 0x0004; + static const int kMask_Bad = 0x0008; +}; + void correspondence_energy_vector( ftl::cuda::TextureObject<float> &d1, ftl::cuda::TextureObject<float> &d2, diff --git a/components/rgbd-sources/include/ftl/rgbd/channels.hpp b/components/rgbd-sources/include/ftl/rgbd/channels.hpp index e89a42ce3c808e254e8ae6085b4517a1f74bd79d..551507730ffbee1c38784e8db8882d1c5ceaac8f 100644 --- a/components/rgbd-sources/include/ftl/rgbd/channels.hpp +++ b/components/rgbd-sources/include/ftl/rgbd/channels.hpp @@ -24,6 +24,7 @@ enum struct Channel : int { EnergyVector, // 32FC4 Flow, // 32F Energy, // 32F + Mask, // 32U LeftGray, RightGray, Overlay1