diff --git a/applications/reconstruct/src/main.cpp b/applications/reconstruct/src/main.cpp
index 71f790dba0a7f5b8d158019e2a5ac08066bc839f..d8dd66b59323d5ddf5aa0f136f4c25f83213441a 100644
--- a/applications/reconstruct/src/main.cpp
+++ b/applications/reconstruct/src/main.cpp
@@ -254,8 +254,8 @@ 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::ColourMLS>("mls"); // Perform MLS (using smoothing channel)
+ pipeline1->append<ftl::operators::SmoothChannel>("smoothing"); // Generate a smoothing channel
+ pipeline1->append<ftl::operators::SimpleMLS>("mls"); // Perform MLS (using smoothing channel)
// Alignment
diff --git a/components/operators/CMakeLists.txt b/components/operators/CMakeLists.txt
index 58e21f43eb785a24d4e7aa77681192b5c5d77d44..c93a87f6693334c30e28039387db45554ebdb0d2 100644
--- a/components/operators/CMakeLists.txt
+++ b/components/operators/CMakeLists.txt
@@ -1,6 +1,8 @@
add_library(ftloperators
src/smoothing.cpp
src/smoothing.cu
+ src/mls.cu
+ src/smoothchan.cu
src/operator.cpp
src/colours.cpp
src/normals.cpp
diff --git a/components/operators/include/ftl/operators/smoothing.hpp b/components/operators/include/ftl/operators/smoothing.hpp
index 46bf082bede957b4a32984fe33c467e0fe0be690..a35053d2a97be851ca9afe75938a57b357465a29 100644
--- a/components/operators/include/ftl/operators/smoothing.hpp
+++ b/components/operators/include/ftl/operators/smoothing.hpp
@@ -63,7 +63,7 @@ class SimpleMLS : public ftl::operators::Operator {
/**
* Perform Moving Least Squares smoothing with a smoothing amount determined
- * by a simple colour similarity weighting.
+ * by a simple colour similarity weighting. In practice this is too naive.
*/
class ColourMLS : public ftl::operators::Operator {
public:
@@ -84,10 +84,10 @@ class ColourMLS : public ftl::operators::Operator {
* encompass hundreds of pixels with a smoothing factor approaching a meter, or
* it can be only a few or even no pixels with a zero smoothing factor.
*/
-class DynamicColourMLS : public ftl::operators::Operator {
+class AdaptiveMLS : public ftl::operators::Operator {
public:
- explicit DynamicColourMLS(ftl::Configurable*);
- ~DynamicColourMLS();
+ explicit AdaptiveMLS(ftl::Configurable*);
+ ~AdaptiveMLS();
inline Operator::Type type() const override { return Operator::Type::OneToOne; }
diff --git a/components/operators/src/mls.cu b/components/operators/src/mls.cu
new file mode 100644
index 0000000000000000000000000000000000000000..7e408b829fa39b15ad687ab8246886be3da3ee38
--- /dev/null
+++ b/components/operators/src/mls.cu
@@ -0,0 +1,208 @@
+#include "smoothing_cuda.hpp"
+
+#include <ftl/cuda/weighting.hpp>
+
+using ftl::cuda::TextureObject;
+
+#define T_PER_BLOCK 8
+
+// ===== MLS Smooth ============================================================
+
+/*
+ * Smooth depth map using Moving Least Squares
+ */
+ template <int SEARCH_RADIUS>
+ __global__ void mls_smooth_kernel(
+ TextureObject<float4> normals_in,
+ TextureObject<float4> normals_out,
+ TextureObject<float> depth_in, // Virtual depth map
+ TextureObject<float> depth_out, // Accumulated output
+ float 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;
+ if (d0 < camera.minDepth || d0 > camera.maxDepth) return;
+ float3 X = camera.screenToCam((int)(x),(int)(y),d0);
+
+ // Neighbourhood
+ for (int v=-SEARCH_RADIUS; v<=SEARCH_RADIUS; ++v) {
+ for (int u=-SEARCH_RADIUS; u<=SEARCH_RADIUS; ++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));
+
+ // Gauss approx weighting function using point distance
+ const float w = ftl::cuda::spatialWeighting(X,Xi,smoothing);
+
+ aX += Xi*w;
+ nX += Ni*w;
+ contrib += w;
+ }
+ }
+
+ 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::mls_smooth(
+ ftl::cuda::TextureObject<float4> &normals_in,
+ ftl::cuda::TextureObject<float4> &normals_out,
+ ftl::cuda::TextureObject<float> &depth_in,
+ ftl::cuda::TextureObject<float> &depth_out,
+ float smoothing,
+ int radius,
+ 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);
+
+ switch (radius) {
+ case 5: mls_smooth_kernel<5><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, smoothing, camera); break;
+ case 4: mls_smooth_kernel<4><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, smoothing, camera); break;
+ case 3: mls_smooth_kernel<3><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, smoothing, camera); break;
+ case 2: mls_smooth_kernel<2><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, smoothing, camera); break;
+ case 1: mls_smooth_kernel<1><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, smoothing, camera); break;
+ }
+ cudaSafeCall( cudaGetLastError() );
+
+
+ #ifdef _DEBUG
+ cudaSafeCall(cudaDeviceSynchronize());
+ #endif
+}
+
+
+// ===== Colour MLS Smooth =====================================================
+
+/*
+ * 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.
+ */
+ template <int SEARCH_RADIUS>
+ __global__ void colour_mls_smooth_kernel(
+ 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;
+ 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
+ for (int v=-SEARCH_RADIUS; v<=SEARCH_RADIUS; ++v) {
+ for (int u=-SEARCH_RADIUS; u<=SEARCH_RADIUS; ++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));
+
+ 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;
+ }
+ }
+
+ 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(
+ 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,
+ int radius,
+ 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);
+
+ switch (radius) {
+ case 5: colour_mls_smooth_kernel<5><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, colour_in, smoothing, colour_smoothing, camera); break;
+ case 4: colour_mls_smooth_kernel<4><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, colour_in, smoothing, colour_smoothing, camera); break;
+ case 3: colour_mls_smooth_kernel<3><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, colour_in, smoothing, colour_smoothing, camera); break;
+ case 2: colour_mls_smooth_kernel<2><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, colour_in, smoothing, colour_smoothing, camera); break;
+ case 1: colour_mls_smooth_kernel<1><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, colour_in, smoothing, colour_smoothing, camera); break;
+ }
+ cudaSafeCall( cudaGetLastError() );
+
+
+ #ifdef _DEBUG
+ cudaSafeCall(cudaDeviceSynchronize());
+ #endif
+}
+
diff --git a/components/operators/src/smoothchan.cu b/components/operators/src/smoothchan.cu
new file mode 100644
index 0000000000000000000000000000000000000000..963fc85ff3bac4d66263ce333882c8584ff6b748
--- /dev/null
+++ b/components/operators/src/smoothchan.cu
@@ -0,0 +1,68 @@
+#include "smoothing_cuda.hpp"
+
+#include <ftl/cuda/weighting.hpp>
+
+using ftl::cuda::TextureObject;
+
+#define T_PER_BLOCK 8
+
+template <int RADIUS>
+__global__ void smooth_chan_kernel(
+ ftl::cuda::TextureObject<uchar4> colour_in,
+ ftl::cuda::TextureObject<float> depth_in,
+ ftl::cuda::TextureObject<float> smoothing_out,
+ ftl::rgbd::Camera camera,
+ float alpha) {
+
+ const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
+ const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+ if (x < smoothing_out.width() && y < smoothing_out.height()) {
+ const int ix = (float)x / (float)smoothing_out.width() * (float)colour_in.width();
+ const int iy = (float)x / (float)smoothing_out.height() * (float)colour_in.height();
+
+ // A distance has already been found
+ if (smoothing_out(x,y) < 1.0f) return;
+
+ float mindist = 1.0f;
+
+ const uchar4 c0 = colour_in.tex2D(ix, iy);
+ const float3 pos = camera.screenToCam(ix, iy, depth_in.tex2D(ix, iy));
+
+ for (int v=-RADIUS; v<=RADIUS; ++v) {
+ #pragma unroll
+ for (int u=-RADIUS; u<=RADIUS; ++u) {
+ const uchar4 c = colour_in.tex2D(ix+u, iy+v);
+ const float3 posN = camera.screenToCam(ix+u, iy+v, depth_in.tex2D(ix+u, iy+v));
+
+ if (ftl::cuda::colourDistance(c, c0) >= alpha) mindist = min(mindist, length(pos - posN));
+ }
+ }
+
+ smoothing_out(x,y) = min(mindist, 1.0f);
+ }
+}
+
+void ftl::cuda::smooth_channel(
+ ftl::cuda::TextureObject<uchar4> &colour_in,
+ ftl::cuda::TextureObject<float> &depth_in,
+ ftl::cuda::TextureObject<float> &smoothing_out,
+ const ftl::rgbd::Camera &camera,
+ float alpha,
+ int radius,
+ cudaStream_t stream) {
+
+ const dim3 gridSize((smoothing_out.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (smoothing_out.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
+ const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
+
+
+ smooth_chan_kernel<1><<<gridSize, blockSize, 0, stream>>>(colour_in, depth_in, smoothing_out, camera, alpha);
+
+
+ cudaSafeCall( cudaGetLastError() );
+
+
+ #ifdef _DEBUG
+ cudaSafeCall(cudaDeviceSynchronize());
+ #endif
+}
diff --git a/components/operators/src/smoothing.cpp b/components/operators/src/smoothing.cpp
index 26fed4d17c65bf20faaa8240dff5bda860aea9ed..284658e8633362ebd7a17096e6cf874de47e5b9e 100644
--- a/components/operators/src/smoothing.cpp
+++ b/components/operators/src/smoothing.cpp
@@ -70,7 +70,23 @@ SmoothChannel::~SmoothChannel() {
}
bool SmoothChannel::apply(ftl::rgbd::Frame &in, ftl::rgbd::Frame &out, ftl::rgbd::Source *s, cudaStream_t stream) {
-
+ int radius = config()->value("radius", 1);
+ float threshold = config()->value("threshold", 30.0f);
+
+ // Clear to max smoothing
+ out.get<GpuMat>(Channel::Smoothing).setTo(cv::Scalar(1.0f));
+
+ // Reduce to nearest
+ ftl::cuda::smooth_channel(
+ in.createTexture<uchar4>(Channel::Colour),
+ in.createTexture<float>(Channel::Depth),
+ out.createTexture<float>(Channel::Smoothing),
+ s->parameters(),
+ threshold,
+ radius,
+ stream
+ );
+
return true;
}
diff --git a/components/operators/src/smoothing.cu b/components/operators/src/smoothing.cu
index 288e1ed2977f591520a2296cc7857f7bd5c372a0..a2b7377fba011243c08cf5347f3dea4800b8ab51 100644
--- a/components/operators/src/smoothing.cu
+++ b/components/operators/src/smoothing.cu
@@ -6,202 +6,6 @@ using ftl::cuda::TextureObject;
#define T_PER_BLOCK 8
-// ===== MLS Smooth ============================================================
-
-/*
- * Smooth depth map using Moving Least Squares
- */
- template <int SEARCH_RADIUS>
- __global__ void mls_smooth_kernel(
- TextureObject<float4> normals_in,
- TextureObject<float4> normals_out,
- TextureObject<float> depth_in, // Virtual depth map
- TextureObject<float> depth_out, // Accumulated output
- float 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;
- if (d0 < camera.minDepth || d0 > camera.maxDepth) return;
- float3 X = camera.screenToCam((int)(x),(int)(y),d0);
-
- // Neighbourhood
- for (int v=-SEARCH_RADIUS; v<=SEARCH_RADIUS; ++v) {
- for (int u=-SEARCH_RADIUS; u<=SEARCH_RADIUS; ++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));
-
- // Gauss approx weighting function using point distance
- const float w = ftl::cuda::spatialWeighting(X,Xi,smoothing);
-
- aX += Xi*w;
- nX += Ni*w;
- contrib += w;
- }
- }
-
- 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::mls_smooth(
- ftl::cuda::TextureObject<float4> &normals_in,
- ftl::cuda::TextureObject<float4> &normals_out,
- ftl::cuda::TextureObject<float> &depth_in,
- ftl::cuda::TextureObject<float> &depth_out,
- float smoothing,
- int radius,
- 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);
-
- switch (radius) {
- case 5: mls_smooth_kernel<5><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, smoothing, camera); break;
- case 4: mls_smooth_kernel<4><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, smoothing, camera); break;
- case 3: mls_smooth_kernel<3><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, smoothing, camera); break;
- case 2: mls_smooth_kernel<2><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, smoothing, camera); break;
- case 1: mls_smooth_kernel<1><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, smoothing, camera); break;
- }
- cudaSafeCall( cudaGetLastError() );
-
-
- #ifdef _DEBUG
- cudaSafeCall(cudaDeviceSynchronize());
- #endif
-}
-
-
-// ===== Colour MLS Smooth =====================================================
-
-/*
- * Smooth depth map using Moving Least Squares
- */
- template <int SEARCH_RADIUS>
- __global__ void colour_mls_smooth_kernel(
- 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;
- 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
- for (int v=-SEARCH_RADIUS; v<=SEARCH_RADIUS; ++v) {
- for (int u=-SEARCH_RADIUS; u<=SEARCH_RADIUS; ++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));
-
- 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;
- }
- }
-
- 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(
- 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,
- int radius,
- 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);
-
- switch (radius) {
- case 5: colour_mls_smooth_kernel<5><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, colour_in, smoothing, colour_smoothing, camera); break;
- case 4: colour_mls_smooth_kernel<4><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, colour_in, smoothing, colour_smoothing, camera); break;
- case 3: colour_mls_smooth_kernel<3><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, colour_in, smoothing, colour_smoothing, camera); break;
- case 2: colour_mls_smooth_kernel<2><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, colour_in, smoothing, colour_smoothing, camera); break;
- case 1: colour_mls_smooth_kernel<1><<<gridSize, blockSize, 0, stream>>>(normals_in, normals_out, depth_in, depth_out, colour_in, smoothing, colour_smoothing, camera); break;
- }
- cudaSafeCall( cudaGetLastError() );
-
-
- #ifdef _DEBUG
- cudaSafeCall(cudaDeviceSynchronize());
- #endif
-}
-
// ===== Colour-based Smooth ===================================================
template <int RADIUS>
diff --git a/components/operators/src/smoothing_cuda.hpp b/components/operators/src/smoothing_cuda.hpp
index f7f87f0508f9e5d16c58d0d089bb589d396ef9c8..dae25e221b1121bbfd02bab732b5c6ba5f465aff 100644
--- a/components/operators/src/smoothing_cuda.hpp
+++ b/components/operators/src/smoothing_cuda.hpp
@@ -29,6 +29,15 @@ void colour_mls_smooth(
const ftl::rgbd::Camera &camera,
cudaStream_t stream);
+void smooth_channel(
+ ftl::cuda::TextureObject<uchar4> &colour_in,
+ ftl::cuda::TextureObject<float> &depth_in,
+ ftl::cuda::TextureObject<float> &smoothing_out,
+ const ftl::rgbd::Camera &camera,
+ float alpha,
+ int radius,
+ cudaStream_t stream);
+
void depth_smooth(
ftl::cuda::TextureObject<float> &depth_in,
ftl::cuda::TextureObject<uchar4> &colour_in,
diff --git a/components/renderers/cpp/include/ftl/cuda/weighting.hpp b/components/renderers/cpp/include/ftl/cuda/weighting.hpp
index 15d3dbcec387f97d8ffe60690bdb5c1fda2c098c..bffff673d2009ac698c3c604ec565a2aa1a89901 100644
--- a/components/renderers/cpp/include/ftl/cuda/weighting.hpp
+++ b/components/renderers/cpp/include/ftl/cuda/weighting.hpp
@@ -25,6 +25,11 @@ __device__ inline float spatialWeighting(const float3 &a, const float3 &b, float
return rh*rh*rh*rh;
}
+__device__ inline float colourDistance(uchar4 a, uchar4 b) {
+ const float3 delta = make_float3((float)a.x - (float)b.x, (float)a.y - (float)b.y, (float)a.z - (float)b.z);
+ return length(delta);
+}
+
/*
* Colour weighting as suggested in:
* C. Kuster et al. Spatio-Temporal Geometry Fusion for Multiple Hybrid Cameras using Moving Least Squares Surfaces. 2014.