diff --git a/applications/reconstruct/src/main.cpp b/applications/reconstruct/src/main.cpp
index 4deb6256e622a1cd9283aeefe2e3520fc72023e3..18539c4e952d4cd77d22ffbfed18ce65eec9c4e6 100644
--- a/applications/reconstruct/src/main.cpp
+++ b/applications/reconstruct/src/main.cpp
@@ -309,9 +309,9 @@ static void run(ftl::Configurable *root) {
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")->set("enabled", false);
+ pipeline1->append<ftl::operators::CrossSupport>("cross");
pipeline1->append<ftl::operators::ColourMLS>("mls"); // Perform MLS (using smoothing channel)
- pipeline1->append<ftl::operators::VisCrossSupport>("viscross")->set("enabled", false);
+ pipeline1->append<ftl::operators::VisCrossSupport>("viscross");
// Alignment
diff --git a/components/codecs/include/ftl/codecs/channels.hpp b/components/codecs/include/ftl/codecs/channels.hpp
index 85cd9fffc45ddcda1b6cc13a160c74c30797b804..3a486cb32cb1c220737818affefbe9d0328ac20e 100644
--- a/components/codecs/include/ftl/codecs/channels.hpp
+++ b/components/codecs/include/ftl/codecs/channels.hpp
@@ -19,7 +19,7 @@ enum struct Channel : int {
Deviation = 4,
Screen = 4,
Normals = 5, // 32FC4
- Points = 6, // 32FC4
+ Points = 6, // 32FC4 (should be deprecated)
Confidence = 7, // 32F
Contribution = 7, // 32F
EnergyVector = 8, // 32FC4
@@ -28,9 +28,9 @@ enum struct Channel : int {
Energy = 10, // 32F
Mask = 11, // 32U
Density = 12, // 32F
- LeftGray = 13, // Deprecated
- RightGray = 14, // Deprecated
- Overlay1 = 15, // Unused
+ Support1 = 13, // 8UC4 (currently)
+ Support2 = 14, // 8UC4 (currently)
+ Segmentation = 15, // 32S?
AudioLeft = 32,
AudioRight = 33,
diff --git a/components/codecs/src/channels.cpp b/components/codecs/src/channels.cpp
index 25b06b533b5c95d16bd60365ccbcc4373e09bd45..effaf9eb261dfe460f2d6f5bf6f08849e4f93538 100644
--- a/components/codecs/src/channels.cpp
+++ b/components/codecs/src/channels.cpp
@@ -21,9 +21,9 @@ static ChannelInfo info[] = {
"Energy", CV_32F,
"Mask", CV_32S,
"Density", CV_32F,
- "LeftGray", CV_8U,
- "RightGray", CV_8U,
- "Overlay1", CV_8UC3,
+ "Support1", CV_8UC4,
+ "Support2", CV_8UC4,
+ "Segmentation", CV_32S,
"NoName", 0,
"NoName", 0,
diff --git a/components/common/cpp/include/ftl/configurable.hpp b/components/common/cpp/include/ftl/configurable.hpp
index 244b548285b92902e645436891b04acb182ffd08..219bd171edf4b7997cc35f97fa7787a0c4bb4559 100644
--- a/components/common/cpp/include/ftl/configurable.hpp
+++ b/components/common/cpp/include/ftl/configurable.hpp
@@ -104,7 +104,7 @@ class Configurable {
protected:
nlohmann::json *config_;
- virtual void inject(const std::string name, nlohmann::json &value) {}
+ virtual void inject(const std::string &name, nlohmann::json &value) {}
private:
std::map<std::string, std::list<std::function<void(const config::Event&)>>> observers_;
diff --git a/components/net/cpp/include/ftl/net_configurable.hpp b/components/net/cpp/include/ftl/net_configurable.hpp
index 3a31ecc27675f2e37b72bc42ee83da0dee5b7e4c..bdd21c4700e7664cff996672585054c30b6e8e6a 100644
--- a/components/net/cpp/include/ftl/net_configurable.hpp
+++ b/components/net/cpp/include/ftl/net_configurable.hpp
@@ -13,7 +13,7 @@ namespace ftl {
~NetConfigurable();
protected:
- void inject(const std::string name, nlohmann::json &value);
+ void inject(const std::string &name, nlohmann::json &value);
private:
ftl::UUID peer;
diff --git a/components/net/cpp/src/net_configurable.cpp b/components/net/cpp/src/net_configurable.cpp
index edcec65b8c4ba452472e30efd29b672b7a6e2c55..be98cf7edbb221c61e9732a9e59d8546f1185644 100644
--- a/components/net/cpp/src/net_configurable.cpp
+++ b/components/net/cpp/src/net_configurable.cpp
@@ -7,6 +7,6 @@ ftl::NetConfigurable::NetConfigurable(ftl::UUID peer, const std::string &suri, f
ftl::NetConfigurable::~NetConfigurable(){}
-void ftl::NetConfigurable::inject(const std::string name, nlohmann::json &value) {
+void ftl::NetConfigurable::inject(const std::string &name, nlohmann::json &value) {
ctrl.set(peer, suri + std::string("/") + name, value);
}
diff --git a/components/operators/CMakeLists.txt b/components/operators/CMakeLists.txt
index 76e67de0209533bead6c6b065eb197e4eba3a858..ae282147dbbd3038d3dbe78981835748de0a81a3 100644
--- a/components/operators/CMakeLists.txt
+++ b/components/operators/CMakeLists.txt
@@ -17,7 +17,8 @@ set(OPERSRC
)
if (HAVE_OPTFLOW)
- list(APPEND OPERSRC src/nvopticalflow.cpp
+ list(APPEND OPERSRC
+ src/nvopticalflow.cpp
src/disparity/optflow_smoothing.cu
src/disparity/optflow_smoothing.cpp)
endif()
diff --git a/components/operators/include/ftl/operators/filling.hpp b/components/operators/include/ftl/operators/filling.hpp
index 1c06074e060dc7ed329829e7894b271f07256773..2825dde5f6840658d7c49dd374792e80952e9391 100644
--- a/components/operators/include/ftl/operators/filling.hpp
+++ b/components/operators/include/ftl/operators/filling.hpp
@@ -21,6 +21,17 @@ class ScanFieldFill : public ftl::operators::Operator {
};
+class CrossSupportFill : public ftl::operators::Operator {
+ public:
+ explicit CrossSupportFill(ftl::Configurable*);
+ ~CrossSupportFill();
+
+ 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;
+
+};
+
}
}
diff --git a/components/operators/include/ftl/operators/operator.hpp b/components/operators/include/ftl/operators/operator.hpp
index 04f542b3b1b088b53adfdb79cb012d9fb3e89e9c..729092dd10a96418e6ec4b3cf0073f82238b455d 100644
--- a/components/operators/include/ftl/operators/operator.hpp
+++ b/components/operators/include/ftl/operators/operator.hpp
@@ -56,7 +56,7 @@ class Operator {
namespace detail {
struct ConstructionHelperBase {
- ConstructionHelperBase(ftl::Configurable *cfg) : config(cfg) {}
+ explicit ConstructionHelperBase(ftl::Configurable *cfg) : config(cfg) {}
virtual ~ConstructionHelperBase() {}
virtual ftl::operators::Operator *make()=0;
@@ -65,7 +65,7 @@ struct ConstructionHelperBase {
template <typename T>
struct ConstructionHelper : public ConstructionHelperBase {
- ConstructionHelper(ftl::Configurable *cfg) : ConstructionHelperBase(cfg) {}
+ explicit ConstructionHelper(ftl::Configurable *cfg) : ConstructionHelperBase(cfg) {}
~ConstructionHelper() {}
ftl::operators::Operator *make() override {
return new T(config);
diff --git a/components/operators/src/filling.cpp b/components/operators/src/filling.cpp
index c0e6bc997cbead6c480b7dedc3f8661a36b84681..b1e68b4d7e3c1dd79fc70e7b7879a820ec708ec6 100644
--- a/components/operators/src/filling.cpp
+++ b/components/operators/src/filling.cpp
@@ -3,6 +3,7 @@
#include "filling_cuda.hpp"
using ftl::operators::ScanFieldFill;
+using ftl::operators::CrossSupportFill;
using ftl::codecs::Channel;
ScanFieldFill::ScanFieldFill(ftl::Configurable *cfg) : ftl::operators::Operator(cfg) {
@@ -26,3 +27,26 @@ bool ScanFieldFill::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd
return true;
}
+
+
+CrossSupportFill::CrossSupportFill(ftl::Configurable *cfg) : ftl::operators::Operator(cfg) {
+
+}
+
+CrossSupportFill::~CrossSupportFill() {
+
+}
+
+bool CrossSupportFill::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::Source *s, cudaStream_t stream) {
+
+ /*ftl::cuda::filling_csr(
+ in.createTexture<uchar4>(Channel::Colour2),
+ in.createTexture<float4>(Channel::Normals),
+ in.createTexture<float>(Channel::Depth),
+ in.createTexture<float>(Channel::Depth),
+ in.createTexture<uchar4>(Channel::Colour),
+ s->parameters(), 0
+ );*/
+
+ return true;
+}
diff --git a/components/operators/src/filling.cu b/components/operators/src/filling.cu
index e0b6cf978bc2b974324f8f88501db61ff6993f79..7780530ff69ab178c46803a38f05de85d037f033 100644
--- a/components/operators/src/filling.cu
+++ b/components/operators/src/filling.cu
@@ -1,5 +1,6 @@
#include "filling_cuda.hpp"
+#define T_PER_BLOCK 8
#define WARP_SIZE 32
#define MAX_EDGES 128
@@ -128,3 +129,73 @@ void ftl::cuda::scan_field_fill(
scan_field_fill_kernel<<<gridSize, blockSize, 0, stream>>>(depth_in, depth_out, smoothing, thresh, camera);
cudaSafeCall( cudaGetLastError() );
}
+
+
+//===== Cross Support Region Filling ===========================================
+
+__global__ void filling_csr_kernel(
+ TextureObject<uchar4> region,
+ TextureObject<float4> normals_in,
+ TextureObject<float> depth_in, // Virtual depth map
+ TextureObject<float> depth_out, // Accumulated output
+ TextureObject<uchar4> colour_in,
+ ftl::rgbd::Camera camera) {
+
+ const int x = blockIdx.x*blockDim.x + threadIdx.x;
+ const int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+ if (x < 0 || y < 0 || x >= depth_in.width() || y >= depth_in.height()) return;
+
+ float3 aX = make_float3(0.0f,0.0f,0.0f);
+ float3 nX = make_float3(0.0f,0.0f,0.0f);
+ float contrib = 0.0f;
+
+ float d0 = depth_in.tex2D(x, y);
+ //depth_out(x,y) = d0;
+ //normals_out(x,y) = normals_in(x,y);
+ if (d0 < camera.minDepth || d0 > camera.maxDepth) return;
+ float3 X = camera.screenToCam((int)(x),(int)(y),d0);
+
+ uchar4 c0 = colour_in.tex2D(x, y);
+
+ // Neighbourhood
+ uchar4 base = region.tex2D(x,y);
+
+ for (int v=-base.z; v<=base.w; ++v) {
+ uchar4 baseY = region.tex2D(x,y+v);
+
+ for (int u=-baseY.x; u<=baseY.y; ++u) {
+ const float d = depth_in.tex2D(x+u, y+v);
+ if (d < camera.minDepth || d > camera.maxDepth) continue;
+
+ // Point and normal of neighbour
+ const float3 Xi = camera.screenToCam((int)(x)+u,(int)(y)+v,d);
+ const float3 Ni = make_float3(normals_in.tex2D((int)(x)+u, (int)(y)+v));
+
+
+ }
+ }
+
+
+}
+
+void ftl::cuda::filling_csr(
+ ftl::cuda::TextureObject<uchar4> ®ion,
+ ftl::cuda::TextureObject<float4> &normals_in,
+ ftl::cuda::TextureObject<float> &depth_in,
+ ftl::cuda::TextureObject<float> &depth_out,
+ ftl::cuda::TextureObject<uchar4> &colour_in,
+ const ftl::rgbd::Camera &camera,
+ cudaStream_t stream) {
+
+ const dim3 gridSize((depth_out.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth_out.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
+ const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
+
+ filling_csr_kernel<<<gridSize, blockSize, 0, stream>>>(region, normals_in, depth_in, depth_out, colour_in, camera);
+ cudaSafeCall( cudaGetLastError() );
+
+
+ #ifdef _DEBUG
+ cudaSafeCall(cudaDeviceSynchronize());
+ #endif
+}
diff --git a/components/operators/src/filling_cuda.hpp b/components/operators/src/filling_cuda.hpp
index c89423fe2b3984694d16c0ef2c5d69735fe5c712..e47781ac7fdd3656809de496dad21d1bb3a48862 100644
--- a/components/operators/src/filling_cuda.hpp
+++ b/components/operators/src/filling_cuda.hpp
@@ -15,6 +15,15 @@ void scan_field_fill(
const ftl::rgbd::Camera &camera,
cudaStream_t stream);
+void filling_csr(
+ ftl::cuda::TextureObject<uchar4> ®ion,
+ ftl::cuda::TextureObject<float4> &normals_in,
+ ftl::cuda::TextureObject<float> &depth_in,
+ ftl::cuda::TextureObject<float> &depth_out,
+ ftl::cuda::TextureObject<uchar4> &colour_in,
+ const ftl::rgbd::Camera &camera,
+ cudaStream_t stream);
+
}
}
diff --git a/components/operators/src/mls.cu b/components/operators/src/mls.cu
index e189bf0e502845a18f565d3df539f8560930b8e1..d7770f1345a9545b67e3534ebaf637ab42f73370 100644
--- a/components/operators/src/mls.cu
+++ b/components/operators/src/mls.cu
@@ -214,6 +214,14 @@ void ftl::cuda::colour_mls_smooth(
// ===== Colour MLS using cross support region =================================
+__device__ inline int segmentID(int u, int v) {
+ if (u < 0 && v < 0) return 0x001;
+ if (u > 0 && v < 0) return 0x002;
+ if (u > 0 && v > 0) return 0x004;
+ if (u < 0 && v > 0) return 0x008;
+ return 0;
+}
+
/*
* Smooth depth map using Moving Least Squares. This version uses colour
* similarity weights to adjust the spatial smoothing factor. It is naive in
@@ -221,6 +229,7 @@ void ftl::cuda::colour_mls_smooth(
* level of single pixels can be subject to noise. Colour noise should first
* be removed from the image.
*/
+ template <bool FILLING>
__global__ void colour_mls_smooth_csr_kernel(
TextureObject<uchar4> region,
TextureObject<float4> normals_in,
@@ -244,14 +253,21 @@ void ftl::cuda::colour_mls_smooth(
float d0 = depth_in.tex2D(x, y);
depth_out(x,y) = d0;
normals_out(x,y) = normals_in(x,y);
- if (d0 < camera.minDepth || d0 > camera.maxDepth) return;
+ if (d0 < camera.minDepth || d0 > camera.maxDepth) {
+ if(FILLING) d0 = 0.0f;
+ else return;
+ }
float3 X = camera.screenToCam((int)(x),(int)(y),d0);
uchar4 c0 = colour_in.tex2D(x, y);
// Neighbourhood
uchar4 base = region.tex2D(x,y);
+ int segment_check = 0;
+ // TODO: Does using a fixed loop size with range test work better?
+ // Or with warp per pixel version, this would be less of a problem...
+ // TODO: Do a separate vote fill step?
for (int v=-base.z; v<=base.w; ++v) {
uchar4 baseY = region.tex2D(x,y+v);
@@ -268,12 +284,17 @@ void ftl::cuda::colour_mls_smooth(
const uchar4 c = colour_in.tex2D(x+u, y+v);
const float cw = ftl::cuda::colourWeighting(c0,c,colour_smoothing);
+ // Allow missing point to borrow z value
+ // TODO: This is a bad choice of Z! Perhaps try histogram vote approach
+ if (FILLING && d0 == 0.0f) X = camera.screenToCam((int)(x),(int)(y),Xi.z);
+
// Gauss approx weighting function using point distance
const float w = ftl::cuda::spatialWeighting(X,Xi,smoothing*cw);
aX += Xi*w;
nX += Ni*w;
contrib += w;
+ if (FILLING && w > 0.0f && v > -base.z+1 && v < base.w-1 && u > -baseY.x+1 && u < baseY.y-1) segment_check |= segmentID(u,v);
}
}
@@ -282,6 +303,11 @@ void ftl::cuda::colour_mls_smooth(
nX /= contrib; // Weighted average normal
aX /= contrib; // Weighted average point (centroid)
+ if (FILLING && d0 == 0.0f) {
+ if (__popc(segment_check) < 3) return;
+ X = camera.screenToCam((int)(x),(int)(y),aX.z);
+ }
+
// Signed-Distance Field function
float fX = nX.x * (X.x - aX.x) + nX.y * (X.y - aX.y) + nX.z * (X.z - aX.z);
@@ -306,13 +332,18 @@ void ftl::cuda::colour_mls_smooth_csr(
ftl::cuda::TextureObject<uchar4> &colour_in,
float smoothing,
float colour_smoothing,
+ bool filling,
const ftl::rgbd::Camera &camera,
cudaStream_t stream) {
const dim3 gridSize((depth_out.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (depth_out.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
- colour_mls_smooth_csr_kernel<<<gridSize, blockSize, 0, stream>>>(region, normals_in, normals_out, depth_in, depth_out, colour_in, smoothing, colour_smoothing, camera);
+ if (filling) {
+ colour_mls_smooth_csr_kernel<true><<<gridSize, blockSize, 0, stream>>>(region, normals_in, normals_out, depth_in, depth_out, colour_in, smoothing, colour_smoothing, camera);
+ } else {
+ colour_mls_smooth_csr_kernel<false><<<gridSize, blockSize, 0, stream>>>(region, normals_in, normals_out, depth_in, depth_out, colour_in, smoothing, colour_smoothing, camera);
+ }
cudaSafeCall( cudaGetLastError() );
diff --git a/components/operators/src/segmentation.cpp b/components/operators/src/segmentation.cpp
index 24405005a90efa68093d9f746ada0f5c070b8f5f..a4bb9b0f7ebb717ee40cbf42015b45629d01ae36 100644
--- a/components/operators/src/segmentation.cpp
+++ b/components/operators/src/segmentation.cpp
@@ -14,13 +14,25 @@ CrossSupport::~CrossSupport() {
}
bool CrossSupport::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::Source *s, cudaStream_t stream) {
- ftl::cuda::support_region(
- in.createTexture<uchar4>(Channel::Colour),
- out.createTexture<uchar4>(Channel::Colour2, ftl::rgbd::Format<uchar4>(in.get<cv::cuda::GpuMat>(Channel::Colour).size())),
- config()->value("tau", 5),
- config()->value("v_max", 5),
- config()->value("h_max", 5), 0
- );
+ bool use_depth = config()->value("depth_region", false);
+
+ if (use_depth) {
+ ftl::cuda::support_region(
+ in.createTexture<float>(Channel::Depth),
+ 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
+ );
+ } //else {
+ ftl::cuda::support_region(
+ in.createTexture<uchar4>(Channel::Colour),
+ 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
+ );
+ //}
return true;
}
@@ -37,11 +49,48 @@ VisCrossSupport::~VisCrossSupport() {
}
bool VisCrossSupport::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::Source *s, cudaStream_t stream) {
- ftl::cuda::vis_support_region(
- in.createTexture<uchar4>(Channel::Colour),
- in.createTexture<uchar4>(Channel::Colour2),
- 0
- );
+ bool show_depth = false;
+ if (in.hasChannel(Channel::Support2) && config()->value("show_depth_support", false)) {
+ show_depth = true;
+ }
+
+ bool show_bad = config()->value("show_bad", false) && in.hasChannel(Channel::Support2);
+
+ if (show_bad) {
+ ftl::cuda::vis_bad_region(
+ in.createTexture<uchar4>(Channel::Colour),
+ in.createTexture<float>(Channel::Depth),
+ in.createTexture<uchar4>(Channel::Support1),
+ in.createTexture<uchar4>(Channel::Support2),
+ 0
+ );
+ } else {
+ ftl::cuda::vis_support_region(
+ in.createTexture<uchar4>(Channel::Colour),
+ in.createTexture<uchar4>(Channel::Support1),
+ make_uchar4(0,0,255,0),
+ make_uchar4(255,0,0,0),
+ config()->value("offset_x", 0),
+ config()->value("offset_y", 0),
+ config()->value("spacing_x", 50),
+ config()->value("spacing_y", 50),
+ 0
+ );
+
+ if (show_depth) {
+ ftl::cuda::vis_support_region(
+ in.createTexture<uchar4>(Channel::Colour),
+ in.createTexture<uchar4>(Channel::Support2),
+ make_uchar4(0,0,255,0),
+ make_uchar4(0,255,0,0),
+ config()->value("offset_x", 0),
+ config()->value("offset_y", 0),
+ config()->value("spacing_x", 50),
+ config()->value("spacing_y", 50),
+ 0
+ );
+ }
+ }
return true;
}
\ No newline at end of file
diff --git a/components/operators/src/segmentation.cu b/components/operators/src/segmentation.cu
index c5c337edecd79173c327ae7816d892893b269096..baeda1023b99610396e7388b2ec2b5bae3c3cdbf 100644
--- a/components/operators/src/segmentation.cu
+++ b/components/operators/src/segmentation.cu
@@ -4,65 +4,93 @@
using ftl::cuda::TextureObject;
+template <typename T>
+__device__ inline float cross(T p1, T p2);
-__device__ inline int cross(uchar4 p1, uchar4 p2) {
+template <>
+__device__ inline float cross<uchar4>(uchar4 p1, uchar4 p2) {
return max(max(__sad(p1.x,p2.x,0),__sad(p1.y,p2.y,0)), __sad(p1.z,p2.z,0));
}
-__device__ uchar4 calculate_support_region(const TextureObject<uchar4> &img, int x, int y, int tau, int v_max, int h_max) {
+template <>
+__device__ inline float cross<float>(float p1, float p2) {
+ return fabs(p1-p2);
+}
+
+template <typename T>
+__device__ uchar4 calculate_support_region(const TextureObject<T> &img, int x, int y, float tau, int v_max, int h_max) {
int x_min = max(0, x - h_max);
int x_max = min(img.width()-1, x + h_max);
int y_min = max(0, y - v_max);
int y_max = min(img.height()-1, y + v_max);
- uchar4 result = make_uchar4(x - x_min, x_max - x, y - y_min, y_max - y);
+ uchar4 result = make_uchar4(0, 0, 0, 0);
- uchar4 colour = img.tex2D(x,y);
+ T colour = img.tex2D(x,y);
+ T prev_colour = colour;
- for (int u=x-1; u >= x_min; --u) {
- if (cross(colour, img.tex2D(u,y)) > tau) {
- result.x = x - u;
+ int u;
+ for (u=x-1; u >= x_min; --u) {
+ T next_colour = img.tex2D(u,y);
+ if (cross(prev_colour, next_colour) > tau) {
+ result.x = x - u - 1;
break;
- }
- }
-
- for (int u=x+1; u <= x_max; ++u) {
- if (cross(colour, img.tex2D(u,y)) > tau) {
- result.y = u - x;
+ }
+ prev_colour = next_colour;
+ }
+ if (u < x_min) result.x = x - x_min;
+
+ prev_colour = colour;
+ for (u=x+1; u <= x_max; ++u) {
+ T next_colour = img.tex2D(u,y);
+ if (cross(prev_colour, next_colour) > tau) {
+ result.y = u - x - 1;
break;
- }
- }
-
- for (int v=y-1; v >= y_min; --v) {
- if (cross(colour, img.tex2D(x,v)) > tau) {
- result.z = y - v;
+ }
+ prev_colour = next_colour;
+ }
+ if (u > x_max) result.y = x_max - x;
+
+ int v;
+ prev_colour = colour;
+ for (v=y-1; v >= y_min; --v) {
+ T next_colour = img.tex2D(x,v);
+ if (cross(prev_colour, next_colour) > tau) {
+ result.z = y - v - 1;
break;
- }
- }
+ }
+ prev_colour = next_colour;
+ }
+ if (v < y_min) result.z = y - y_min;
- for (int v=y+1; v <= y_max; ++v) {
- if (cross(colour, img.tex2D(x,v)) > tau) {
- result.w = v - y;
+ prev_colour = colour;
+ for (v=y+1; v <= y_max; ++v) {
+ T next_colour = img.tex2D(x,v);
+ if (cross(prev_colour, next_colour) > tau) {
+ result.w = v - y - 1;
break;
- }
- }
+ }
+ prev_colour = next_colour;
+ }
+ if (v > y_max) result.w = y_max - y;
return result;
}
-__global__ void support_region_kernel(TextureObject<uchar4> colour, TextureObject<uchar4> region, int tau, int v_max, int h_max) {
+template <typename T>
+__global__ void support_region_kernel(TextureObject<T> img, TextureObject<uchar4> region, float tau, int v_max, int h_max) {
const int x = blockIdx.x*blockDim.x + threadIdx.x;
const int y = blockIdx.y*blockDim.y + threadIdx.y;
- if (x < 0 || y < 0 || x >= colour.width() || y >= colour.height()) return;
+ if (x < 0 || y < 0 || x >= img.width() || y >= img.height()) return;
- region(x,y) = calculate_support_region(colour, x, y, tau, v_max, h_max);
+ region(x,y) = calculate_support_region(img, x, y, tau, v_max, h_max);
}
void ftl::cuda::support_region(
ftl::cuda::TextureObject<uchar4> &colour,
ftl::cuda::TextureObject<uchar4> ®ion,
- int tau,
+ float tau,
int v_max,
int h_max,
cudaStream_t stream) {
@@ -79,16 +107,40 @@ void ftl::cuda::support_region(
#endif
}
-__global__ void vis_support_region_kernel(TextureObject<uchar4> colour, TextureObject<uchar4> region) {
+void ftl::cuda::support_region(
+ ftl::cuda::TextureObject<float> &depth,
+ ftl::cuda::TextureObject<uchar4> ®ion,
+ float tau,
+ int v_max,
+ int h_max,
+ cudaStream_t stream) {
+
+ const dim3 gridSize((region.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (region.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
+ const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
+
+ support_region_kernel<<<gridSize, blockSize, 0, stream>>>(depth, region, tau, v_max, h_max);
+ cudaSafeCall( cudaGetLastError() );
+
+
+ #ifdef _DEBUG
+ cudaSafeCall(cudaDeviceSynchronize());
+ #endif
+}
+
+__global__ void vis_support_region_kernel(TextureObject<uchar4> colour, TextureObject<uchar4> region, uchar4 bcolour, uchar4 acolour,
+ int ox, int oy, int dx, int dy) {
const int x = blockIdx.x*blockDim.x + threadIdx.x;
const int y = blockIdx.y*blockDim.y + threadIdx.y;
if (x < 0 || y < 0 || x >= colour.width() || y >= colour.height()) return;
// Grid pattern
- if (x % 50 != 0 || y % 50 != 0) return;
+ if (x % dx != ox || y % dy != oy) return;
uchar4 base = region.tex2D(x,y);
+
+ // Edge pattern
+ //if (base.x != 1) return;
for (int v=-base.z; v<=base.w; ++v) {
uchar4 baseY = region.tex2D(x,y+v);
@@ -96,20 +148,87 @@ __global__ void vis_support_region_kernel(TextureObject<uchar4> colour, TextureO
for (int u=-baseY.x; u<=baseY.y; ++u) {
if (x+u < 0 || y+v < 0 || x+u >= colour.width() || y+v >= colour.height()) continue;
uchar4 col = colour.tex2D(x+u, y+v);
- colour(x+u, y+v) = make_uchar4(255,col.y,col.z,0);
+ colour(x+u, y+v) = (u==0 || v == 0) ?
+ make_uchar4(max(bcolour.x, col.x), max(bcolour.y, col.y), max(bcolour.z, col.z), 0) :
+ make_uchar4(max(acolour.x, col.x), max(acolour.y, col.y), max(acolour.z, col.z), 0);
}
}
}
void ftl::cuda::vis_support_region(
ftl::cuda::TextureObject<uchar4> &colour,
- ftl::cuda::TextureObject<uchar4> ®ion,
+ ftl::cuda::TextureObject<uchar4> ®ion,
+ uchar4 bar_colour,
+ uchar4 area_colour,
+ int ox, int oy, int dx, int dy,
+ cudaStream_t stream) {
+
+ const dim3 gridSize((region.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (region.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
+ const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
+
+ vis_support_region_kernel<<<gridSize, blockSize, 0, stream>>>(
+ colour,
+ region,
+ bar_colour,
+ area_colour,
+ ox,oy,dx,dy
+ );
+ cudaSafeCall( cudaGetLastError() );
+
+
+ #ifdef _DEBUG
+ cudaSafeCall(cudaDeviceSynchronize());
+ #endif
+}
+
+// ===== Vis bad edges =========================================================
+
+__global__ void vis_bad_region_kernel(
+ TextureObject<uchar4> colour,
+ TextureObject<float> depth,
+ TextureObject<uchar4> region,
+ TextureObject<uchar4> dregion) {
+ const int x = blockIdx.x*blockDim.x + threadIdx.x;
+ const int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+ if (x < 0 || y < 0 || x >= colour.width() || y >= colour.height()) return;
+
+ // Grid pattern
+ //if (x % 50 != 0 || y % 50 != 0) return;
+
+ uchar4 base = region.tex2D(x,y);
+ uchar4 baseD = dregion.tex2D(x,y);
+ uchar4 col = colour.tex2D(x,y);
+ float d = depth.tex2D(x,y);
+
+ if (baseD.x > base.x && baseD.y < base.y) {
+ uchar4 baseR = region.tex2D(x+baseD.y+1, y);
+ float dR = depth.tex2D(x+baseD.y+1, y);
+ //if (x+baseD.y+1-baseR.x <= x) {
+ if (d > 0.0f && d < 30.0f && (dR <= 0.0f || dR >= 30.0f)) {
+ colour(x,y) = make_uchar4(col.x, col.y, 255, 0);
+ depth(x,y) = 0.0f;
+ }
+ //}
+ }
+}
+
+void ftl::cuda::vis_bad_region(
+ ftl::cuda::TextureObject<uchar4> &colour,
+ ftl::cuda::TextureObject<float> &depth,
+ ftl::cuda::TextureObject<uchar4> ®ion,
+ ftl::cuda::TextureObject<uchar4> &dregion,
cudaStream_t stream) {
const dim3 gridSize((region.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (region.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
- vis_support_region_kernel<<<gridSize, blockSize, 0, stream>>>(colour, region);
+ vis_bad_region_kernel<<<gridSize, blockSize, 0, stream>>>(
+ colour,
+ depth,
+ region,
+ dregion
+ );
cudaSafeCall( cudaGetLastError() );
diff --git a/components/operators/src/segmentation_cuda.hpp b/components/operators/src/segmentation_cuda.hpp
index 566b9d033a77935798e294f6126ede330963be4c..5ef55abb10a0d61aa83f900761b26352a3ae4a87 100644
--- a/components/operators/src/segmentation_cuda.hpp
+++ b/components/operators/src/segmentation_cuda.hpp
@@ -9,14 +9,31 @@ namespace cuda {
void support_region(
ftl::cuda::TextureObject<uchar4> &colour,
ftl::cuda::TextureObject<uchar4> ®ion,
- int tau, int v_max, int h_max,
+ float tau, int v_max, int h_max,
+ cudaStream_t stream);
+
+void support_region(
+ ftl::cuda::TextureObject<float> &depth,
+ ftl::cuda::TextureObject<uchar4> ®ion,
+ float tau, int v_max, int h_max,
cudaStream_t stream);
void vis_support_region(
ftl::cuda::TextureObject<uchar4> &colour,
ftl::cuda::TextureObject<uchar4> ®ion,
+ uchar4 bar_colour,
+ uchar4 area_colour,
+ int ox, int oy, int dx, int dy,
cudaStream_t stream);
+void vis_bad_region(
+ ftl::cuda::TextureObject<uchar4> &colour,
+ ftl::cuda::TextureObject<float> &depth,
+ ftl::cuda::TextureObject<uchar4> ®ion,
+ ftl::cuda::TextureObject<uchar4> &dregion,
+ cudaStream_t stream);
+
+
}
}
diff --git a/components/operators/src/smoothing.cpp b/components/operators/src/smoothing.cpp
index c23ba0c5f3711ab1e044e4dcaa162b32de562367..a9041087ad2a3432e04957bfc983c8d29a2ae525 100644
--- a/components/operators/src/smoothing.cpp
+++ b/components/operators/src/smoothing.cpp
@@ -184,6 +184,7 @@ bool ColourMLS::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::So
int iters = config()->value("mls_iterations", 10);
int radius = config()->value("mls_radius",3);
bool crosssup = config()->value("cross_support", false);
+ bool filling = config()->value("filling", false);
if (!in.hasChannel(Channel::Normals)) {
LOG(ERROR) << "Required normals channel missing for MLS";
@@ -207,7 +208,7 @@ bool ColourMLS::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::So
);
} else {
ftl::cuda::colour_mls_smooth_csr(
- in.createTexture<uchar4>(Channel::Colour2),
+ in.createTexture<uchar4>(Channel::Support1),
in.createTexture<float4>(Channel::Normals),
in.createTexture<float4>(Channel::Points, ftl::rgbd::Format<float4>(in.get<cv::cuda::GpuMat>(Channel::Depth).size())),
in.createTexture<float>(Channel::Depth),
@@ -215,6 +216,7 @@ bool ColourMLS::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::So
in.createTexture<uchar4>(Channel::Colour),
thresh,
col_smooth,
+ filling,
s->parameters(),
0
);
diff --git a/components/operators/src/smoothing_cuda.hpp b/components/operators/src/smoothing_cuda.hpp
index 729c962218c48d204c175f6a338bc797878d70d0..a92a49260a48dcabe115768b347dfca5cc2283ab 100644
--- a/components/operators/src/smoothing_cuda.hpp
+++ b/components/operators/src/smoothing_cuda.hpp
@@ -38,6 +38,7 @@ void colour_mls_smooth_csr(
ftl::cuda::TextureObject<uchar4> &colour_in,
float smoothing,
float colour_smoothing,
+ bool filling,
const ftl::rgbd::Camera &camera,
cudaStream_t stream);
diff --git a/components/rgbd-sources/src/sources/ftlfile/file_source.cpp b/components/rgbd-sources/src/sources/ftlfile/file_source.cpp
index 9597bde796f80ac67f38a5a569190001e13d7065..0962c1886dc199e50530343c0d01edf4e74e37f0 100644
--- a/components/rgbd-sources/src/sources/ftlfile/file_source.cpp
+++ b/components/rgbd-sources/src/sources/ftlfile/file_source.cpp
@@ -30,6 +30,13 @@ FileSource::FileSource(ftl::rgbd::Source *s, ftl::rgbd::Player *r, int sid) : ft
realtime_ = host_->value("realtime", true);
timestamp_ = r->getStartTime();
sourceid_ = sid;
+ freeze_ = host_->value("freeze", false);
+ have_frozen_ = false;
+
+ host_->on("freeze", [this](const ftl::config::Event &e) {
+ have_frozen_ = false;
+ freeze_ = host_->value("freeze", false);
+ });
r->onPacket(sid, [this](const ftl::codecs::StreamPacket &spkt, ftl::codecs::Packet &pkt) {
host_->notifyRaw(spkt, pkt);
@@ -133,18 +140,28 @@ bool FileSource::capture(int64_t ts) {
}
bool FileSource::retrieve() {
- if (!reader_->read(timestamp_)) {
+ if (!have_frozen_ && !reader_->read(timestamp_)) {
cache_write_ = -1;
}
return true;
}
void FileSource::swap() {
+ if (have_frozen_) return;
cache_read_ = cache_write_;
cache_write_ = (cache_write_ == 0) ? 1 : 0;
}
bool FileSource::compute(int n, int b) {
+ // Freeze frame requires a copy to be made each time...
+ if (have_frozen_) {
+ cv::cuda::GpuMat t1, t2;
+ if (!rgb_.empty()) rgb_.copyTo(t1);
+ if (!depth_.empty()) depth_.copyTo(t2);
+ host_->notify(timestamp_, t1, t2);
+ return true;
+ }
+
if (cache_read_ < 0) return false;
if (cache_[cache_read_].size() == 0) return false;
@@ -191,7 +208,11 @@ bool FileSource::compute(int n, int b) {
if (rgb_.empty() || depth_.empty()) return false;
// Inform about a decoded frame pair
- host_->notify(timestamp_, rgb_, depth_);
+ if (freeze_ && !have_frozen_) {
+ have_frozen_ = true;
+ } else {
+ host_->notify(timestamp_, rgb_, depth_);
+ }
return true;
}
diff --git a/components/rgbd-sources/src/sources/ftlfile/file_source.hpp b/components/rgbd-sources/src/sources/ftlfile/file_source.hpp
index 2d59b68cb1e07d10aade07b33780af2497736a12..f8a26da37802e0aa2aae4681b420e6061c230fd2 100644
--- a/components/rgbd-sources/src/sources/ftlfile/file_source.hpp
+++ b/components/rgbd-sources/src/sources/ftlfile/file_source.hpp
@@ -43,6 +43,8 @@ class FileSource : public detail::Source {
ftl::codecs::Decoder *decoders_[2];
bool realtime_;
+ bool freeze_;
+ bool have_frozen_;
void _processCalibration(ftl::codecs::Packet &pkt);
void _processPose(ftl::codecs::Packet &pkt);