Add some MLS comments

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);
 
 }
 }

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;