diff --git a/applications/reconstruct/src/main.cpp b/applications/reconstruct/src/main.cpp
index 73445da8c9afded93b510954212b9100fb5ee5b4..457ee985ed636a3cff8ea2a457103b120c1ed8e7 100644
--- a/applications/reconstruct/src/main.cpp
+++ b/applications/reconstruct/src/main.cpp
@@ -34,6 +34,7 @@
#include <ftl/operators/colours.hpp>
#include <ftl/operators/normals.hpp>
#include <ftl/operators/filling.hpp>
+#include <ftl/operators/segmentation.hpp>
#include <ftl/cuda/normals.hpp>
#include <ftl/registration.hpp>
@@ -165,7 +166,7 @@ static void run(ftl::Configurable *root) {
}
auto configproxy = ConfigProxy(net);
- configproxy.add(root, "source/disparity", "disparity");
+ //configproxy.add(root, "source/disparity", "disparity");
// Create scene transform, intended for axis aligning the walls and floor
Eigen::Matrix4d transform;
@@ -300,9 +301,11 @@ static void run(ftl::Configurable *root) {
pipeline1->append<ftl::operators::ColourChannels>("colour"); // Convert BGR to BGRA
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::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::ColourMLS>("mls"); // Perform MLS (using smoothing channel)
+ pipeline1->append<ftl::operators::VisCrossSupport>("viscross")->set("enabled", false);
// Alignment
diff --git a/components/operators/CMakeLists.txt b/components/operators/CMakeLists.txt
index 54da0eb4be9284b27ac9d5f9c338a7fa98eb6987..15dfb638e89e0ed37916d46b674d6471b245df09 100644
--- a/components/operators/CMakeLists.txt
+++ b/components/operators/CMakeLists.txt
@@ -8,6 +8,8 @@ add_library(ftloperators
src/normals.cpp
src/filling.cpp
src/filling.cu
+ src/segmentation.cu
+ src/segmentation.cpp
src/nvopticalflow.cpp
)
diff --git a/components/operators/include/ftl/operators/segmentation.hpp b/components/operators/include/ftl/operators/segmentation.hpp
new file mode 100644
index 0000000000000000000000000000000000000000..f905ce96966d0f68171ee94d209b6d888581005a
--- /dev/null
+++ b/components/operators/include/ftl/operators/segmentation.hpp
@@ -0,0 +1,40 @@
+#ifndef _FTL_OPERATORS_SEGMENTATION_HPP_
+#define _FTL_OPERATORS_SEGMENTATION_HPP_
+
+#include <ftl/operators/operator.hpp>
+
+namespace ftl {
+namespace operators {
+
+/**
+ * Generate the cross support regions channel.
+ */
+class CrossSupport : public ftl::operators::Operator {
+ public:
+ explicit CrossSupport(ftl::Configurable*);
+ ~CrossSupport();
+
+ 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;
+
+};
+
+/**
+ * Visualise the cross support regions channel.
+ */
+class VisCrossSupport : public ftl::operators::Operator {
+ public:
+ explicit VisCrossSupport(ftl::Configurable*);
+ ~VisCrossSupport();
+
+ 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_SEGMENTATION_HPP_
diff --git a/components/operators/src/mls.cu b/components/operators/src/mls.cu
index 9353734e6fb5a462fc8681783666f2e91ec011f9..e189bf0e502845a18f565d3df539f8560930b8e1 100644
--- a/components/operators/src/mls.cu
+++ b/components/operators/src/mls.cu
@@ -212,6 +212,117 @@ void ftl::cuda::colour_mls_smooth(
}
+// ===== Colour MLS using cross support region =================================
+
+/*
+ * Smooth depth map using Moving Least Squares. This version uses colour
+ * similarity weights to adjust the spatial smoothing factor. It is naive in
+ * that similar colours may exist on both sides of an edge and colours at the
+ * level of single pixels can be subject to noise. Colour noise should first
+ * be removed from the image.
+ */
+ __global__ void colour_mls_smooth_csr_kernel(
+ TextureObject<uchar4> region,
+ TextureObject<float4> normals_in,
+ TextureObject<float4> normals_out,
+ TextureObject<float> depth_in, // Virtual depth map
+ TextureObject<float> depth_out, // Accumulated output
+ TextureObject<uchar4> colour_in,
+ float smoothing,
+ float colour_smoothing,
+ 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));
+
+ if (Ni.x+Ni.y+Ni.z == 0.0f) continue;
+
+ const uchar4 c = colour_in.tex2D(x+u, y+v);
+ const float cw = ftl::cuda::colourWeighting(c0,c,colour_smoothing);
+
+ // 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 (contrib == 0.0f) return;
+
+ nX /= contrib; // Weighted average normal
+ aX /= contrib; // Weighted average point (centroid)
+
+ // Signed-Distance Field function
+ float fX = nX.x * (X.x - aX.x) + nX.y * (X.y - aX.y) + nX.z * (X.z - aX.z);
+
+ // Calculate new point using SDF function to adjust depth (and position)
+ X = X - nX * fX;
+
+ //uint2 screen = camera.camToScreen<uint2>(X);
+
+ //if (screen.x < depth_out.width() && screen.y < depth_out.height()) {
+ // depth_out(screen.x,screen.y) = X.z;
+ //}
+ depth_out(x,y) = X.z;
+ normals_out(x,y) = make_float4(nX / length(nX), 0.0f);
+}
+
+void ftl::cuda::colour_mls_smooth_csr(
+ ftl::cuda::TextureObject<uchar4> ®ion,
+ ftl::cuda::TextureObject<float4> &normals_in,
+ ftl::cuda::TextureObject<float4> &normals_out,
+ ftl::cuda::TextureObject<float> &depth_in,
+ ftl::cuda::TextureObject<float> &depth_out,
+ ftl::cuda::TextureObject<uchar4> &colour_in,
+ float smoothing,
+ float colour_smoothing,
+ 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);
+
+ cudaSafeCall( cudaGetLastError() );
+
+
+ #ifdef _DEBUG
+ cudaSafeCall(cudaDeviceSynchronize());
+ #endif
+}
+
+
// ===== Adaptive MLS Smooth =====================================================
/*
diff --git a/components/operators/src/segmentation.cpp b/components/operators/src/segmentation.cpp
new file mode 100644
index 0000000000000000000000000000000000000000..24405005a90efa68093d9f746ada0f5c070b8f5f
--- /dev/null
+++ b/components/operators/src/segmentation.cpp
@@ -0,0 +1,47 @@
+#include <ftl/operators/segmentation.hpp>
+#include "segmentation_cuda.hpp"
+
+using ftl::operators::CrossSupport;
+using ftl::operators::VisCrossSupport;
+using ftl::codecs::Channel;
+
+CrossSupport::CrossSupport(ftl::Configurable *cfg) : ftl::operators::Operator(cfg) {
+
+}
+
+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
+ );
+
+ return true;
+}
+
+
+
+
+VisCrossSupport::VisCrossSupport(ftl::Configurable *cfg) : ftl::operators::Operator(cfg) {
+
+}
+
+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
+ );
+
+ return true;
+}
\ No newline at end of file
diff --git a/components/operators/src/segmentation.cu b/components/operators/src/segmentation.cu
new file mode 100644
index 0000000000000000000000000000000000000000..c5c337edecd79173c327ae7816d892893b269096
--- /dev/null
+++ b/components/operators/src/segmentation.cu
@@ -0,0 +1,119 @@
+#include "segmentation_cuda.hpp"
+
+#define T_PER_BLOCK 8
+
+using ftl::cuda::TextureObject;
+
+
+__device__ inline int cross(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) {
+ 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 colour = img.tex2D(x,y);
+
+ for (int u=x-1; u >= x_min; --u) {
+ if (cross(colour, img.tex2D(u,y)) > tau) {
+ result.x = x - u;
+ break;
+ }
+ }
+
+ for (int u=x+1; u <= x_max; ++u) {
+ if (cross(colour, img.tex2D(u,y)) > tau) {
+ result.y = u - x;
+ break;
+ }
+ }
+
+ for (int v=y-1; v >= y_min; --v) {
+ if (cross(colour, img.tex2D(x,v)) > tau) {
+ result.z = y - v;
+ break;
+ }
+ }
+
+ for (int v=y+1; v <= y_max; ++v) {
+ if (cross(colour, img.tex2D(x,v)) > tau) {
+ result.w = v - y;
+ break;
+ }
+ }
+
+ return result;
+}
+
+__global__ void support_region_kernel(TextureObject<uchar4> colour, TextureObject<uchar4> region, int 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;
+
+ region(x,y) = calculate_support_region(colour, x, y, tau, v_max, h_max);
+}
+
+void ftl::cuda::support_region(
+ ftl::cuda::TextureObject<uchar4> &colour,
+ ftl::cuda::TextureObject<uchar4> ®ion,
+ int 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>>>(colour, 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) {
+ 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);
+
+ 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) {
+ 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);
+ }
+ }
+}
+
+void ftl::cuda::vis_support_region(
+ ftl::cuda::TextureObject<uchar4> &colour,
+ ftl::cuda::TextureObject<uchar4> ®ion,
+ 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);
+ cudaSafeCall( cudaGetLastError() );
+
+
+ #ifdef _DEBUG
+ cudaSafeCall(cudaDeviceSynchronize());
+ #endif
+}
diff --git a/components/operators/src/segmentation_cuda.hpp b/components/operators/src/segmentation_cuda.hpp
new file mode 100644
index 0000000000000000000000000000000000000000..566b9d033a77935798e294f6126ede330963be4c
--- /dev/null
+++ b/components/operators/src/segmentation_cuda.hpp
@@ -0,0 +1,23 @@
+#ifndef _FTL_CUDA_SEGMENTATION_HPP_
+#define _FTL_CUDA_SEGMENTATION_HPP_
+
+#include <ftl/cuda_common.hpp>
+
+namespace ftl {
+namespace cuda {
+
+void support_region(
+ ftl::cuda::TextureObject<uchar4> &colour,
+ ftl::cuda::TextureObject<uchar4> ®ion,
+ int tau, int v_max, int h_max,
+ cudaStream_t stream);
+
+void vis_support_region(
+ ftl::cuda::TextureObject<uchar4> &colour,
+ ftl::cuda::TextureObject<uchar4> ®ion,
+ cudaStream_t stream);
+
+}
+}
+
+#endif
diff --git a/components/operators/src/smoothing.cpp b/components/operators/src/smoothing.cpp
index 8561e7a11ed9714d586b4dae9de7a0014ca40135..c23ba0c5f3711ab1e044e4dcaa162b32de562367 100644
--- a/components/operators/src/smoothing.cpp
+++ b/components/operators/src/smoothing.cpp
@@ -183,6 +183,7 @@ bool ColourMLS::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::So
float col_smooth = config()->value("mls_colour_smoothing", 30.0f);
int iters = config()->value("mls_iterations", 10);
int radius = config()->value("mls_radius",3);
+ bool crosssup = config()->value("cross_support", false);
if (!in.hasChannel(Channel::Normals)) {
LOG(ERROR) << "Required normals channel missing for MLS";
@@ -191,18 +192,33 @@ bool ColourMLS::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::So
// FIXME: Assume in and out are the same frame.
for (int i=0; i<iters; ++i) {
- ftl::cuda::colour_mls_smooth(
- 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),
- in.createTexture<float>(Channel::Depth2, ftl::rgbd::Format<float>(in.get<cv::cuda::GpuMat>(Channel::Depth).size())),
- in.createTexture<uchar4>(Channel::Colour),
- thresh,
- col_smooth,
- radius,
- s->parameters(),
- 0
- );
+ if (!crosssup) {
+ ftl::cuda::colour_mls_smooth(
+ 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),
+ in.createTexture<float>(Channel::Depth2, ftl::rgbd::Format<float>(in.get<cv::cuda::GpuMat>(Channel::Depth).size())),
+ in.createTexture<uchar4>(Channel::Colour),
+ thresh,
+ col_smooth,
+ radius,
+ s->parameters(),
+ 0
+ );
+ } else {
+ ftl::cuda::colour_mls_smooth_csr(
+ in.createTexture<uchar4>(Channel::Colour2),
+ 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),
+ in.createTexture<float>(Channel::Depth2, ftl::rgbd::Format<float>(in.get<cv::cuda::GpuMat>(Channel::Depth).size())),
+ in.createTexture<uchar4>(Channel::Colour),
+ thresh,
+ col_smooth,
+ s->parameters(),
+ 0
+ );
+ }
in.swapChannels(Channel::Depth, Channel::Depth2);
in.swapChannels(Channel::Normals, Channel::Points);
diff --git a/components/operators/src/smoothing_cuda.hpp b/components/operators/src/smoothing_cuda.hpp
index 53e8cf898c4e99875dc4a9881df58700ee31c127..729c962218c48d204c175f6a338bc797878d70d0 100644
--- a/components/operators/src/smoothing_cuda.hpp
+++ b/components/operators/src/smoothing_cuda.hpp
@@ -29,6 +29,18 @@ void colour_mls_smooth(
const ftl::rgbd::Camera &camera,
cudaStream_t stream);
+void colour_mls_smooth_csr(
+ ftl::cuda::TextureObject<uchar4> ®ion,
+ ftl::cuda::TextureObject<float4> &normals_in,
+ ftl::cuda::TextureObject<float4> &normals_out,
+ ftl::cuda::TextureObject<float> &depth_in,
+ ftl::cuda::TextureObject<float> &depth_out,
+ ftl::cuda::TextureObject<uchar4> &colour_in,
+ float smoothing,
+ float colour_smoothing,
+ const ftl::rgbd::Camera &camera,
+ cudaStream_t stream);
+
void adaptive_mls_smooth(
ftl::cuda::TextureObject<float4> &normals_in,
ftl::cuda::TextureObject<float4> &normals_out,