From dccbdc3b8b5441aba0133a94d19bf063673e91cb Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nicolas.pope@utu.fi>
Date: Mon, 11 Nov 2019 20:51:09 +0200
Subject: [PATCH] Implements #224 to use discontinuity mask
---
applications/reconstruct/src/ilw/ilw.cpp | 8 +-
applications/reconstruct/src/main.cpp | 7 +-
components/operators/CMakeLists.txt | 2 +
.../operators/include/ftl/operators/mask.hpp | 43 ++++++++++
.../include/ftl/operators/operator.hpp | 1 +
components/operators/src/colours.cpp | 4 +-
components/operators/src/filling.cpp | 2 +-
components/operators/src/mask.cpp | 49 ++++++++++++
components/operators/src/mask.cu | 80 +++++++++++++++++++
components/operators/src/mask_cuda.hpp | 59 ++++++++++++++
components/operators/src/normals.cpp | 2 +-
components/operators/src/operator.cpp | 13 ++-
components/operators/src/segmentation.cpp | 10 +--
components/operators/src/smoothing.cpp | 10 +--
14 files changed, 268 insertions(+), 22 deletions(-)
create mode 100644 components/operators/include/ftl/operators/mask.hpp
create mode 100644 components/operators/src/mask.cpp
create mode 100644 components/operators/src/mask.cu
create mode 100644 components/operators/src/mask_cuda.hpp
diff --git a/applications/reconstruct/src/ilw/ilw.cpp b/applications/reconstruct/src/ilw/ilw.cpp
index dfa43267e..55a9539a9 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/applications/reconstruct/src/main.cpp b/applications/reconstruct/src/main.cpp
index 18539c4e9..03c21b142 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 ae282147d..734734256 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 000000000..ef9758e39
--- /dev/null
+++ b/components/operators/include/ftl/operators/mask.hpp
@@ -0,0 +1,43 @@
+#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;
+
+};
+
+/**
+ * Remove depth values marked with the discontinuity mask.
+ */
+class CullDiscontinuity : public ftl::operators::Operator {
+ public:
+ explicit CullDiscontinuity(ftl::Configurable*);
+ ~CullDiscontinuity();
+
+ 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/include/ftl/operators/operator.hpp b/components/operators/include/ftl/operators/operator.hpp
index 729092dd1..bc47afef9 100644
--- a/components/operators/include/ftl/operators/operator.hpp
+++ b/components/operators/include/ftl/operators/operator.hpp
@@ -99,6 +99,7 @@ class Graph : public ftl::Configurable {
private:
std::list<ftl::operators::detail::OperatorNode> operators_;
std::map<std::string, ftl::Configurable*> configs_;
+ cudaStream_t stream_;
ftl::Configurable *_append(ftl::operators::detail::ConstructionHelperBase*);
};
diff --git a/components/operators/src/colours.cpp b/components/operators/src/colours.cpp
index 3174c42e4..474b783c4 100644
--- a/components/operators/src/colours.cpp
+++ b/components/operators/src/colours.cpp
@@ -12,6 +12,8 @@ ColourChannels::~ColourChannels() {
}
bool ColourChannels::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::Source *s, cudaStream_t stream) {
+ auto cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
+
// Convert colour from BGR to BGRA if needed
if (in.get<cv::cuda::GpuMat>(Channel::Colour).type() == CV_8UC3) {
//cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
@@ -19,7 +21,7 @@ bool ColourChannels::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgb
auto &col = in.get<cv::cuda::GpuMat>(Channel::Colour);
temp_.create(col.size(), CV_8UC4);
cv::cuda::swap(col, temp_);
- cv::cuda::cvtColor(temp_,col, cv::COLOR_BGR2BGRA, 0);
+ cv::cuda::cvtColor(temp_,col, cv::COLOR_BGR2BGRA, 0, cvstream);
}
return true;
diff --git a/components/operators/src/filling.cpp b/components/operators/src/filling.cpp
index b1e68b4d7..da7efea6b 100644
--- a/components/operators/src/filling.cpp
+++ b/components/operators/src/filling.cpp
@@ -22,7 +22,7 @@ bool ScanFieldFill::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd
in.createTexture<float>(Channel::Depth),
in.createTexture<float>(Channel::Smoothing),
thresh,
- s->parameters(), 0
+ s->parameters(), stream
);
return true;
diff --git a/components/operators/src/mask.cpp b/components/operators/src/mask.cpp
new file mode 100644
index 000000000..f923f11d0
--- /dev/null
+++ b/components/operators/src/mask.cpp
@@ -0,0 +1,49 @@
+#include <ftl/operators/mask.hpp>
+#include "mask_cuda.hpp"
+
+using ftl::operators::DiscontinuityMask;
+using ftl::operators::CullDiscontinuity;
+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);
+ 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, threshold, stream
+ );
+
+ return true;
+}
+
+
+
+CullDiscontinuity::CullDiscontinuity(ftl::Configurable *cfg) : ftl::operators::Operator(cfg) {
+
+}
+
+CullDiscontinuity::~CullDiscontinuity() {
+
+}
+
+bool CullDiscontinuity::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::Source *s, cudaStream_t stream) {
+ ftl::cuda::cull_discontinuity(
+ in.createTexture<int>(Channel::Mask),
+ out.createTexture<float>(Channel::Depth),
+ stream
+ );
+
+ 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 000000000..e385f41b1
--- /dev/null
+++ b/components/operators/src/mask.cu
@@ -0,0 +1,80 @@
+#include "mask_cuda.hpp"
+
+#define T_PER_BLOCK 8
+
+using ftl::cuda::Mask;
+
+__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, int radius) {
+ 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);
+
+ if (d >= params.minDepth && d <= params.maxDepth) {
+ /* Orts-Escolano S. et al. 2016. Holoportation: Virtual 3D teleportation in real-time. */
+
+ // If colour cross support region terminates within the requested
+ // radius, and the absolute depth difference on the other side is
+ // greater than threshold, then is is a discontinuity.
+ // Repeat for left, right, up and down.
+ const uchar4 sup = support.tex2D((int)x, (int)y);
+ if (sup.x <= radius) {
+ float dS = depth.tex2D((int)x - sup.x - radius, (int)y);
+ if (fabs(dS - d) > threshold) mask.isDiscontinuity(true);
+ }
+ if (sup.y <= radius) {
+ float dS = depth.tex2D((int)x + sup.y + radius, (int)y);
+ if (fabs(dS - d) > threshold) mask.isDiscontinuity(true);
+ }
+ if (sup.z <= radius) {
+ float dS = depth.tex2D((int)x, (int)y - sup.z - radius);
+ if (fabs(dS - d) > threshold) mask.isDiscontinuity(true);
+ }
+ if (sup.w <= radius) {
+ float dS = depth.tex2D((int)x, (int)y + sup.w + radius);
+ if (fabs(dS - d) > threshold) mask.isDiscontinuity(true);
+ }
+ }
+
+ mask_out(x,y) = (int)mask;
+ }
+}
+
+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);
+
+ discontinuity_kernel<<<gridSize, blockSize, 0, stream>>>(mask_out, support, depth, params, thresh, discon);
+ cudaSafeCall( cudaGetLastError() );
+
+#ifdef _DEBUG
+ cudaSafeCall(cudaDeviceSynchronize());
+#endif
+}
+
+
+
+__global__ void cull_discontinuity_kernel(ftl::cuda::TextureObject<int> mask, ftl::cuda::TextureObject<float> depth) {
+ const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
+ const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+ if (x < depth.width() && y < depth.height()) {
+ Mask m(mask.tex2D((int)x,(int)y));
+ if (m.isDiscontinuity()) depth(x,y) = 1000.0f;
+ }
+}
+
+void ftl::cuda::cull_discontinuity(ftl::cuda::TextureObject<int> &mask, ftl::cuda::TextureObject<float> &depth, cudaStream_t stream) {
+ const dim3 gridSize((depth.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
+ const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
+
+ cull_discontinuity_kernel<<<gridSize, blockSize, 0, stream>>>(mask, depth);
+ 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 000000000..6a02aafdb
--- /dev/null
+++ b/components/operators/src/mask_cuda.hpp
@@ -0,0 +1,59 @@
+#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> &mask,
+ ftl::cuda::TextureObject<uchar4> &support,
+ ftl::cuda::TextureObject<float> &depth,
+ const ftl::rgbd::Camera ¶ms,
+ int radius, float threshold,
+ cudaStream_t stream);
+
+void cull_discontinuity(
+ ftl::cuda::TextureObject<int> &mask,
+ ftl::cuda::TextureObject<float> &depth,
+ cudaStream_t stream);
+
+}
+}
+
+#endif
diff --git a/components/operators/src/normals.cpp b/components/operators/src/normals.cpp
index 5f8554c29..57903aa1d 100644
--- a/components/operators/src/normals.cpp
+++ b/components/operators/src/normals.cpp
@@ -28,7 +28,7 @@ bool Normals::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::Sour
ftl::cuda::normals(
out.createTexture<float4>(Channel::Normals, ftl::rgbd::Format<float4>(in.get<cv::cuda::GpuMat>(Channel::Depth).size())),
in.createTexture<float>(Channel::Depth),
- s->parameters(), 0
+ s->parameters(), stream
);
return true;
diff --git a/components/operators/src/operator.cpp b/components/operators/src/operator.cpp
index 91dada28b..e55c181c4 100644
--- a/components/operators/src/operator.cpp
+++ b/components/operators/src/operator.cpp
@@ -31,16 +31,18 @@ bool Operator::apply(FrameSet &in, Frame &out, Source *os, cudaStream_t stream)
Graph::Graph(nlohmann::json &config) : ftl::Configurable(config) {
-
+ cudaSafeCall( cudaStreamCreate(&stream_) );
}
Graph::~Graph() {
-
+ cudaStreamDestroy(stream_);
}
bool Graph::apply(FrameSet &in, FrameSet &out, cudaStream_t stream) {
if (!value("enabled", true)) return false;
+ auto stream_actual = (stream == 0) ? stream_ : stream;
+
if (in.frames.size() != out.frames.size()) return false;
for (auto &i : operators_) {
@@ -53,11 +55,16 @@ bool Graph::apply(FrameSet &in, FrameSet &out, cudaStream_t stream) {
auto *instance = i.instances[j];
if (instance->enabled()) {
- instance->apply(in.frames[j], out.frames[j], in.sources[j], stream);
+ instance->apply(in.frames[j], out.frames[j], in.sources[j], stream_actual);
}
}
}
+ if (stream == 0) {
+ cudaStreamSynchronize(stream_actual);
+ cudaSafeCall( cudaGetLastError() );
+ }
+
return true;
}
diff --git a/components/operators/src/segmentation.cpp b/components/operators/src/segmentation.cpp
index a4bb9b0f7..08b4de467 100644
--- a/components/operators/src/segmentation.cpp
+++ b/components/operators/src/segmentation.cpp
@@ -22,7 +22,7 @@ bool CrossSupport::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd:
out.createTexture<uchar4>(Channel::Support2, ftl::rgbd::Format<uchar4>(in.get<cv::cuda::GpuMat>(Channel::Colour).size())),
config()->value("depth_tau", 0.04f),
config()->value("v_max", 5),
- config()->value("h_max", 5), 0
+ config()->value("h_max", 5), stream
);
} //else {
ftl::cuda::support_region(
@@ -30,7 +30,7 @@ bool CrossSupport::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd:
out.createTexture<uchar4>(Channel::Support1, ftl::rgbd::Format<uchar4>(in.get<cv::cuda::GpuMat>(Channel::Colour).size())),
config()->value("tau", 5.0f),
config()->value("v_max", 5),
- config()->value("h_max", 5), 0
+ config()->value("h_max", 5), stream
);
//}
@@ -62,7 +62,7 @@ bool VisCrossSupport::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rg
in.createTexture<float>(Channel::Depth),
in.createTexture<uchar4>(Channel::Support1),
in.createTexture<uchar4>(Channel::Support2),
- 0
+ stream
);
} else {
ftl::cuda::vis_support_region(
@@ -74,7 +74,7 @@ bool VisCrossSupport::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rg
config()->value("offset_y", 0),
config()->value("spacing_x", 50),
config()->value("spacing_y", 50),
- 0
+ stream
);
if (show_depth) {
@@ -87,7 +87,7 @@ bool VisCrossSupport::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rg
config()->value("offset_y", 0),
config()->value("spacing_x", 50),
config()->value("spacing_y", 50),
- 0
+ stream
);
}
}
diff --git a/components/operators/src/smoothing.cpp b/components/operators/src/smoothing.cpp
index a9041087a..359b17714 100644
--- a/components/operators/src/smoothing.cpp
+++ b/components/operators/src/smoothing.cpp
@@ -33,7 +33,7 @@ bool HFSmoother::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::S
in.createTexture<float>(Channel::Energy, ftl::rgbd::Format<float>(in.get<cv::cuda::GpuMat>(Channel::Depth).size())),
in.createTexture<float>(Channel::Smoothing, ftl::rgbd::Format<float>(in.get<cv::cuda::GpuMat>(Channel::Depth).size())),
var_thresh,
- s->parameters(), 0
+ s->parameters(), stream
);
}
@@ -158,7 +158,7 @@ bool SimpleMLS::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::So
thresh,
radius,
s->parameters(),
- 0
+ stream
);
in.swapChannels(Channel::Depth, Channel::Depth2);
@@ -204,7 +204,7 @@ bool ColourMLS::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::So
col_smooth,
radius,
s->parameters(),
- 0
+ stream
);
} else {
ftl::cuda::colour_mls_smooth_csr(
@@ -218,7 +218,7 @@ bool ColourMLS::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::So
col_smooth,
filling,
s->parameters(),
- 0
+ stream
);
}
@@ -259,7 +259,7 @@ bool AdaptiveMLS::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::
in.createTexture<float>(Channel::Smoothing),
radius,
s->parameters(),
- 0
+ stream
);
in.swapChannels(Channel::Depth, Channel::Depth2);
--
GitLab