Add full 3D MLS and carving

components/common/cpp/include/ftl/cuda_matrix_util.hpp

@@ -750,7 +750,7 @@ public:
 	}
 
 	//! returns the 3x3 part of the matrix
-	inline __device__ __host__ float3x3 getFloat3x3() {
+	inline __device__ __host__ float3x3 getFloat3x3() const {
 		float3x3 ret;
 		ret.m11 = m11;	ret.m12 = m12;	ret.m13 = m13;
 		ret.m21 = m21;	ret.m22 = m22;	ret.m23 = m23;
@@ -1086,7 +1086,7 @@ public:
 
 
 	//! returns the 3x3 part of the matrix
-	inline __device__ __host__ float3x3 getFloat3x3() {
+	inline __device__ __host__ float3x3 getFloat3x3() const {
 		float3x3 ret;
 		ret.m11 = m11;	ret.m12 = m12;	ret.m13 = m13;
 		ret.m21 = m21;	ret.m22 = m22;	ret.m23 = m23;

components/operators/CMakeLists.txt

 set(OPERSRC
-	src/smoothing.cpp
-	src/smoothing.cu
-	src/mls.cu
-	src/smoothchan.cu
+	src/surface/smoothing.cpp
+	src/surface/smoothing.cu
+	src/surface/mls/image_basic.cu
+	src/surface/mls.cu
+	src/analysis/local/smoothchan.cu
 	src/operator.cpp
-	src/colours.cpp
-	src/normals.cpp
-	src/filling.cpp
-	src/filling.cu
+	src/misc/colours.cpp
+	src/analysis/local/normals.cpp
+	src/surface/filling.cpp
+	src/surface/filling.cu
 	src/disparity/libstereo.cpp
 	src/disparity/disp2depth.cu
 	src/disparity/disparity_to_depth.cpp
 	src/disparity/bilateral_filter.cpp
 	src/disparity/opencv/disparity_bilateral_filter.cpp
 	src/disparity/opencv/disparity_bilateral_filter.cu
-	src/segmentation.cu
-	src/segmentation.cpp
-	src/mask.cu
-	src/mask.cpp
-	src/antialiasing.cpp
-	src/antialiasing.cu
+	src/analysis/segmentation/segmentation.cu
+	src/analysis/segmentation/segmentation.cpp
+	src/analysis/local/mask.cu
+	src/analysis/local/mask.cpp
+	src/misc/antialiasing.cpp
+	src/misc/antialiasing.cu
 	src/fusion/mvmls.cpp
 	src/fusion/correspondence.cu
 	src/fusion/correspondence_depth.cu
 	src/fusion/correspondence_util.cu
 	src/fusion/mls_aggr.cu
-	src/clipping.cpp
-	src/depth.cpp
-	src/detectandtrack.cpp
-	src/aruco.cpp
-	src/weighting.cpp
-	src/weighting.cu
-	src/poser.cpp
-	src/gt_analysis.cpp
-	src/gt_analysis.cu
+	src/fusion/smoothing/mls_multi_weighted.cu
+	src/fusion/carving/carver.cu
+	src/fusion/fusion.cpp
+	src/misc/clipping.cpp
+	src/disparity/depth.cpp
+	src/analysis/tracking/detectandtrack.cpp
+	src/analysis/tracking/aruco.cpp
+	src/analysis/local/weighting.cpp
+	src/analysis/local/weighting.cu
+	src/misc/poser.cpp
+	src/analysis/evaluation/gt_analysis.cpp
+	src/analysis/evaluation/gt_analysis.cu
 )
 
 if (HAVE_LIBSGM)
@@ -42,8 +46,8 @@ endif (HAVE_LIBSGM)
 
 if (HAVE_OPTFLOW)
 	list(APPEND OPERSRC
-		src/nvopticalflow.cpp
-		src/opticalflow.cu
+		src/analysis/tracking/nvopticalflow.cpp
+		src/analysis/tracking/opticalflow.cu
 		src/disparity/optflow_smoothing.cu
 		src/disparity/optflow_smoothing.cpp)
 endif()

components/operators/include/ftl/cuda/fixed.hpp

@@ -3,7 +3,7 @@
 
 template <int FRAC>
 __device__ inline float fixed2float(short v) {
-    return v / (1 << FRAC);
+    return float(v) / float(1 << FRAC);
 }
 
 template <int FRAC>

components/operators/include/ftl/operators/cuda/carver.hpp

+#ifndef _FTL_CUDA_CARVER_HPP_
+#define _FTL_CUDA_CARVER_HPP_
+
+#include <ftl/cuda_common.hpp>
+#include <ftl/cuda_matrix_util.hpp>
+#include <ftl/rgbd/camera.hpp>
+
+namespace ftl {
+namespace cuda {
+
+/**
+ * Carve `in` using `ref` as visibility reference.
+ */
+void depth_carve(
+	cv::cuda::GpuMat &in,
+	const cv::cuda::GpuMat &ref,
+	//const cv::cuda::GpuMat &in_colour,
+	//const cv::cuda::GpuMat &ref_colour,
+	//cv::cuda::GpuMat &colour_scale,
+	const float4x4 &transform,
+	const ftl::rgbd::Camera &incam,
+	const ftl::rgbd::Camera &refcam,
+	cudaStream_t stream);
+
+void apply_colour_scaling(
+	const cv::cuda::GpuMat &scale,
+	cv::cuda::GpuMat &colour,
+	int radius,
+	cudaStream_t stream);
+
+}
+}
+
+#endif
\ No newline at end of file

components/operators/include/ftl/operators/cuda/mls/multi_intensity.hpp

+#ifndef _FTL_CUDA_MLS_MULTIINTENSITY_HPP_
+#define _FTL_CUDA_MLS_MULTIINTENSITY_HPP_
+
+#include <ftl/cuda_common.hpp>
+#include <ftl/rgbd/camera.hpp>
+#include <ftl/cuda_matrix_util.hpp>
+
+namespace ftl {
+namespace cuda {
+
+// TODO: 4th channel of normals could be used as curvature estimate? Then can
+// adaptively adjust the smoothing amount by the degree of curvature. This
+// curvature value must be low for noise data, is only high if a consistent
+// high curvature is detected over a radius. Other option is to also or instead
+// use the curvature value as a feature, to only smooth against points with
+// similar curvature?
+
+/**
+ * For a particular viewpoint, use a set of other viewpoints to estimate a
+ * smooth surface for the focus viewpoint. This version of MLS uses absolute
+ * difference of some 8-bit value as a weight, this can be raw intensity or
+ * a local contrast measure. This `prime`, `gather`, `adjust` cycle should
+ * be iterated 2-3 times, but perhaps better to do one iteration of each view
+ * first before repeating a view.
+ * 
+ * Note: Have one instance per frameset because the internal buffers are
+ * allocated based upon frame image size.
+ */
+class MLSMultiIntensity
+{
+public:
+	MLSMultiIntensity(int radius);
+	~MLSMultiIntensity();
+
+	void prime(
+		const cv::cuda::GpuMat &depth_prime,
+		const cv::cuda::GpuMat &intensity_prime,
+		const ftl::rgbd::Camera &cam_prime,
+		const float4x4 &pose_prime,
+		cudaStream_t stream
+	);
+
+	void gatherPrime(float smoothing, cudaStream_t stream);
+
+	void gather(
+		const cv::cuda::GpuMat &depth_src,
+		const cv::cuda::GpuMat &normals_src,
+		const cv::cuda::GpuMat &intensity_src,
+		const ftl::rgbd::Camera &cam_src,
+		const float4x4 &pose_src,
+		float smoothing,
+		cudaStream_t stream
+	);
+
+	void adjust(
+		cv::cuda::GpuMat &depth_out,
+		cv::cuda::GpuMat &normals_out,
+		cudaStream_t stream
+	);
+
+private:
+	cv::cuda::GpuMat depth_prime_;
+	cv::cuda::GpuMat intensity_prime_;
+	ftl::rgbd::Camera cam_prime_;
+	float4x4 pose_prime_;
+	int radius_;
+
+	cv::cuda::GpuMat normal_accum_;
+	cv::cuda::GpuMat centroid_accum_;
+	cv::cuda::GpuMat weight_accum_;
+};
+
+}
+}
+
+#endif
\ No newline at end of file

components/operators/include/ftl/operators/cuda/mls_cuda.hpp

+#ifndef _FTL_CUDA_MLS_HPP_
+#define _FTL_CUDA_MLS_HPP_
+
+#include <ftl/rgbd/camera.hpp>
+#include <ftl/cuda_common.hpp>
+
+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);
+
+/**
+ * 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);
+
+}
+}
+
+#endif  // _FTL_CUDA_MLS_HPP_

components/operators/src/smoothing_cuda.hpp → components/operators/include/ftl/operators/cuda/smoothing_cuda.hpp

@@ -7,16 +7,6 @@
 namespace ftl {
 namespace cuda {
 
-void mls_smooth(
-		ftl::cuda::TextureObject<half4> &normals_in,
-		ftl::cuda::TextureObject<half4> &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);
-
 void colour_mls_smooth(
 		ftl::cuda::TextureObject<half4> &normals_in,
 		ftl::cuda::TextureObject<half4> &normals_out,

components/operators/include/ftl/operators/fusion.hpp

+#ifndef _FTL_OPERATORS_FUSION_HPP_
+#define _FTL_OPERATORS_FUSION_HPP_
+
+#include <ftl/operators/operator.hpp>
+#include <ftl/operators/cuda/mls/multi_intensity.hpp>
+#include <vector>
+
+namespace ftl {
+namespace operators {
+
+class Fusion : public ftl::operators::Operator {
+	public:
+	Fusion(ftl::operators::Graph *g, ftl::Configurable*);
+	~Fusion();
+
+	inline Operator::Type type() const override { return Operator::Type::ManyToMany; }
+
+	bool apply(ftl::rgbd::FrameSet &in, ftl::rgbd::FrameSet &out, cudaStream_t stream) override;
+
+	private:
+	ftl::cuda::MLSMultiIntensity mls_;
+	std::vector<cv::cuda::GpuMat> weights_;
+	cv::cuda::GpuMat temp_;
+};
+
+}
+}
+
+#endif

components/operators/src/smoothchan.cu → components/operators/src/analysis/local/smoothchan.cu

-#include "smoothing_cuda.hpp"
+#include <ftl/operators/cuda/smoothing_cuda.hpp>
 
 #include <ftl/cuda/weighting.hpp>