diff --git a/applications/reconstruct/src/ilw/ilw.cpp b/applications/reconstruct/src/ilw/ilw.cpp
index dfa43267ed5d3ac09309a466b150c0498fb72967..55a9539a9cc9766bf49849164099ee5f620c4c8d 100644
--- a/applications/reconstruct/src/ilw/ilw.cpp
+++ b/applications/reconstruct/src/ilw/ilw.cpp
@@ -215,11 +215,11 @@ bool ILW::_phase0(ftl::rgbd::FrameSet &fs, cudaStream_t stream) {
f.createTexture<float>(Channel::Depth2, Format<float>(f.get<GpuMat>(Channel::Colour).size()));
f.createTexture<float>(Channel::Confidence, Format<float>(f.get<GpuMat>(Channel::Colour).size()));
- f.createTexture<int>(Channel::Mask, Format<int>(f.get<GpuMat>(Channel::Colour).size()));
+ //f.createTexture<int>(Channel::Mask, Format<int>(f.get<GpuMat>(Channel::Colour).size()));
f.createTexture<uchar4>(Channel::Colour);
f.createTexture<float>(Channel::Depth);
- f.get<GpuMat>(Channel::Mask).setTo(cv::Scalar(0));
+ //f.get<GpuMat>(Channel::Mask).setTo(cv::Scalar(0));
}
//cudaSafeCall(cudaStreamSynchronize(stream_));
@@ -232,7 +232,7 @@ bool ILW::_phase1(ftl::rgbd::FrameSet &fs, int win, cudaStream_t stream) {
cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
// Find discontinuity mask
- for (size_t i=0; i<fs.frames.size(); ++i) {
+ /*for (size_t i=0; i<fs.frames.size(); ++i) {
auto &f = fs.frames[i];
auto s = fs.sources[i];
@@ -243,7 +243,7 @@ bool ILW::_phase1(ftl::rgbd::FrameSet &fs, int win, cudaStream_t stream) {
discon_mask_,
stream
);
- }
+ }*/
// First do any preprocessing
if (fill_depth_) {
diff --git a/components/operators/src/mask.cpp b/components/operators/src/mask.cpp
index 32d3c8c34a91455d668d48df22586f4f7dd63723..1ac8ba92ca0b4e8d5f37ea4b993d75effb8cf825 100644
--- a/components/operators/src/mask.cpp
+++ b/components/operators/src/mask.cpp
@@ -15,11 +15,13 @@ 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);
+ float threshold = config()->value("threshold", 0.1f);
ftl::cuda::discontinuity(
out.createTexture<int>(Channel::Mask, ftl::rgbd::Format<int>(in.get<cv::cuda::GpuMat>(Channel::Depth).size())),
+ in.createTexture<uchar4>(Channel::Support1),
in.createTexture<float>(Channel::Depth),
- s->parameters(), radius, 0
+ s->parameters(), radius, threshold, 0
);
return true;
diff --git a/components/operators/src/mask.cu b/components/operators/src/mask.cu
index d7065b5083b313dad7d89f9dc851634a3581bb9b..ac6467a704e9697c63fd3528e61dd36cd3be35a2 100644
--- a/components/operators/src/mask.cu
+++ b/components/operators/src/mask.cu
@@ -5,7 +5,7 @@
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) {
+__global__ void discontinuity_kernel(ftl::cuda::TextureObject<int> mask_out, ftl::cuda::TextureObject<uchar4> support, ftl::cuda::TextureObject<float> depth, ftl::rgbd::Camera params, float threshold) {
const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
@@ -20,11 +20,29 @@ __global__ void discontinuity_kernel(ftl::cuda::TextureObject<int> mask_out, ftl
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) {
+ //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);
- }
+ // if (fabs(depth.tex2D((int)x+u, (int)y+v) - d) > 0.1f) mask.isDiscontinuity(true);
+ // }
+ //}
+
+ const uchar4 sup = support.tex2D((int)x, (int)y);
+ if (sup.x <= RADIUS) {
+ float dS = depth.tex2D((int)x - sup.x - 1, (int)y);
+ if (fabs(dS - d) > threshold) mask.isDiscontinuity(true);
+ }
+ if (sup.y <= RADIUS) {
+ float dS = depth.tex2D((int)x + sup.y + 1, (int)y);
+ if (fabs(dS - d) > threshold) mask.isDiscontinuity(true);
+ }
+ if (sup.z <= RADIUS) {
+ float dS = depth.tex2D((int)x, (int)y - sup.z - 1);
+ if (fabs(dS - d) > threshold) mask.isDiscontinuity(true);
+ }
+ if (sup.w <= RADIUS) {
+ float dS = depth.tex2D((int)x, (int)y + sup.w + 1);
+ if (fabs(dS - d) > threshold) mask.isDiscontinuity(true);
}
}
@@ -32,16 +50,16 @@ __global__ void discontinuity_kernel(ftl::cuda::TextureObject<int> mask_out, ftl
}
}
-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) {
+void ftl::cuda::discontinuity(ftl::cuda::TextureObject<int> &mask_out, ftl::cuda::TextureObject<uchar4> &support, ftl::cuda::TextureObject<float> &depth, const ftl::rgbd::Camera ¶ms, int discon, float thresh, 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;
+ case 5 : discontinuity_kernel<5><<<gridSize, blockSize, 0, stream>>>(mask_out, support, depth, params, thresh); break;
+ case 4 : discontinuity_kernel<4><<<gridSize, blockSize, 0, stream>>>(mask_out, support, depth, params, thresh); break;
+ case 3 : discontinuity_kernel<3><<<gridSize, blockSize, 0, stream>>>(mask_out, support, depth, params, thresh); break;
+ case 2 : discontinuity_kernel<2><<<gridSize, blockSize, 0, stream>>>(mask_out, support, depth, params, thresh); break;
+ case 1 : discontinuity_kernel<1><<<gridSize, blockSize, 0, stream>>>(mask_out, support, depth, params, thresh); break;
default: break;
}
cudaSafeCall( cudaGetLastError() );
diff --git a/components/operators/src/mask_cuda.hpp b/components/operators/src/mask_cuda.hpp
index f1eb4972084652f54f092596d86bfa6f4cbdcfb6..64e4121139f3050198142ae898f696fa04930fe4 100644
--- a/components/operators/src/mask_cuda.hpp
+++ b/components/operators/src/mask_cuda.hpp
@@ -41,10 +41,11 @@ class Mask {
};
void discontinuity(
- ftl::cuda::TextureObject<int> &mast,
+ ftl::cuda::TextureObject<int> &mask,
+ ftl::cuda::TextureObject<uchar4> &support,
ftl::cuda::TextureObject<float> &depth,
const ftl::rgbd::Camera ¶ms,
- int radius,
+ int radius, float threshold,
cudaStream_t stream);
}