From 102e516f22abb90f5184ff696b7f66cf163ccc01 Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Tue, 3 Nov 2020 13:53:25 +0200
Subject: [PATCH] Add some MLS comments
---
.../include/ftl/operators/cuda/mls_cuda.hpp | 58 ++++++++++++++-----
.../operators/src/surface/mls/image_basic.cu | 2 +
2 files changed, 45 insertions(+), 15 deletions(-)
diff --git a/components/operators/include/ftl/operators/cuda/mls_cuda.hpp b/components/operators/include/ftl/operators/cuda/mls_cuda.hpp
index bb34b9a42..9790595d6 100644
--- a/components/operators/include/ftl/operators/cuda/mls_cuda.hpp
+++ b/components/operators/include/ftl/operators/cuda/mls_cuda.hpp
@@ -7,24 +7,52 @@
namespace ftl {
namespace cuda {
+/**
+ * Basic Moving Least Squares smoothing on a single depth map. Outputs boths
+ * normals and depth values. This is a single iteration of an algorithm that
+ * should iterate 2-3 times. Normals can use some basic estimation as initial
+ * input since they are smoothed by MLS.
+ *
+ * @param normals_in 4 channel 16-bit float (half).
+ * @param normals_out 4 channel 16-bit float (half).
+ * @param depth_in 1 channel 32-bit float
+ * @param depth_out 1 channel 32-bit float
+ * @param smoothing Gaussian radius in depth units
+ * @param radius Window radius for smoothing
+ * @param camera Camera intrinsics
+ * @param stream Optional CUDA stream
+ */
void mls_smooth(
- const cv::cuda::GpuMat &normals_in,
- cv::cuda::GpuMat &normals_out,
- const cv::cuda::GpuMat &depth_in,
- cv::cuda::GpuMat &depth_out,
- float smoothing,
- int radius,
- const ftl::rgbd::Camera &camera,
- cudaStream_t stream);
+ const cv::cuda::GpuMat &normals_in,
+ cv::cuda::GpuMat &normals_out,
+ const cv::cuda::GpuMat &depth_in,
+ cv::cuda::GpuMat &depth_out,
+ float smoothing,
+ int radius,
+ const ftl::rgbd::Camera &camera,
+ cudaStream_t stream);
+/**
+ * Basic Moving Least Squares smoothing on a single depth map. Outputs just the
+ * smoothed normals. This is a single iteration of an algorithm that should
+ * iterate 2-3 times.
+ *
+ * @param normals_in 4 channel 16-bit float (half).
+ * @param normals_out 4 channel 16-bit float (half).
+ * @param depth_in 1 channel 32-bit float
+ * @param smoothing Gaussian radius in depth units
+ * @param radius Window radius for smoothing
+ * @param camera Camera intrinsics
+ * @param stream Optional CUDA stream
+ */
void mls_smooth(
- const cv::cuda::GpuMat &normals_in,
- cv::cuda::GpuMat &normals_out,
- const cv::cuda::GpuMat &depth_in,
- float smoothing,
- int radius,
- const ftl::rgbd::Camera &camera,
- cudaStream_t stream);
+ const cv::cuda::GpuMat &normals_in,
+ cv::cuda::GpuMat &normals_out,
+ const cv::cuda::GpuMat &depth_in,
+ float smoothing,
+ int radius,
+ const ftl::rgbd::Camera &camera,
+ cudaStream_t stream);
}
}
diff --git a/components/operators/src/surface/mls/image_basic.cu b/components/operators/src/surface/mls/image_basic.cu
index a34d0b13a..17d528391 100644
--- a/components/operators/src/surface/mls/image_basic.cu
+++ b/components/operators/src/surface/mls/image_basic.cu
@@ -115,6 +115,8 @@ void ftl::cuda::mls_smooth(
const dim3 gridSize((depth_in.cols + THREADS_X - 1)/THREADS_X, (depth_in.rows + THREADS_Y - 1)/THREADS_Y);
const dim3 blockSize(THREADS_X, THREADS_Y);
+ normals_out.create(normals_in.size(), CV_16FC4);
+
switch (radius) {
case 5: mls_smooth_kernel<5><<<gridSize, blockSize, 0, stream>>>(normals_in.ptr<half4>(), normals_out.ptr<half4>(), depth_in.ptr<float>(), nullptr, normals_in.step1()/4, depth_in.step1(), smoothing, camera); break;
case 4: mls_smooth_kernel<4><<<gridSize, blockSize, 0, stream>>>(normals_in.ptr<half4>(), normals_out.ptr<half4>(), depth_in.ptr<float>(), nullptr, normals_in.step1()/4, depth_in.step1(), smoothing, camera); break;
--
GitLab