WIP fixing compile problems in refactor

applications/reconstruct/CMakeLists.txt

@@ -32,6 +32,6 @@ set_property(TARGET ftl-reconstruct PROPERTY CUDA_SEPARABLE_COMPILATION ON)
 endif()
 
 #target_include_directories(cv-node PUBLIC ${PROJECT_SOURCE_DIR}/include)
-target_link_libraries(ftl-reconstruct ftlcommon ftlrgbd Threads::Threads ${OpenCV_LIBS} ftlctrl ftlnet)
+target_link_libraries(ftl-reconstruct ftlcommon ftlrgbd Threads::Threads ${OpenCV_LIBS} ftlctrl ftlnet ftlrender)
 
 

components/renderers/cpp/include/ftl/cuda/weighting.hpp

+#ifndef _FTL_CUDA_WEIGHTING_HPP_
+#define _FTL_CUDA_WEIGHTING_HPP_
+
+namespace ftl {
+namespace cuda {
+
+/*
+ * Guennebaud, G.; Gross, M. Algebraic point set surfaces. ACMTransactions on Graphics Vol. 26, No. 3, Article No. 23, 2007.
+ * Used in: FusionMLS: Highly dynamic 3D reconstruction with consumer-grade RGB-D cameras
+ *     r = distance between points
+ *     h = smoothing parameter in meters (default 4cm)
+ */
+__device__ inline float spatialWeighting(float r, float h) {
+	if (r >= h) return 0.0f;
+	float rh = r / h;
+	rh = 1.0f - rh*rh;
+	return rh*rh*rh*rh;
+}
+
+}
+}
+
+#endif  // _FTL_CUDA_WEIGHTING_HPP_ 

components/renderers/cpp/include/ftl/render/splat_render.hpp

@@ -21,7 +21,7 @@ class Splatter : public ftl::render::Renderer {
 	explicit Splatter(nlohmann::json &config, const ftl::rgbd::FrameSet &fs);
 	~Splatter();
 
-	void render(ftl::rgbd::VirtualSource *src, cudaStream_t stream=0);
+	bool render(ftl::rgbd::VirtualSource *src, cudaStream_t stream=0) override;
 
 	//void setOutputDevice(int);
 

components/renderers/cpp/src/splat_render.cpp

@@ -14,22 +14,24 @@ Splatter::~Splatter() {
 
 }
 
-void Splatter::render(ftl::rgbd::VirtualSource *src, cudaStream_t stream) {
+bool Splatter::render(ftl::rgbd::VirtualSource *src, cudaStream_t stream) {
 	if (!src->isReady()) return;
 
 	const auto &camera = src->parameters();
 
 	//cudaSafeCall(cudaSetDevice(scene_->getCUDADevice()));
 
-	output_.create<cv::cuda::GpuMat>(Channel::Depth, Format<float>(camera.width, camera.height));
-	output_.create<cv::cuda::GpuMat>(Channel::Colour, Format<uchar4>(camera.width, camera.height));
+	output_.create<GpuMat>(Channel::Depth, Format<float>(camera.width, camera.height));
+	output_.create<GpuMat>(Channel::Colour, Format<uchar4>(camera.width, camera.height));
 
-	temp_.create<cv::cuda::GpuMat>(Channel::Colour, Format<float4>(camera.width, camera.height));
-	temp_.create<cv::cuda::GpuMat>(Channel::Colour2, Format<uchar4>(camera.width, camera.height));
-	temp_.create<cv::cuda::GpuMat>(Channel::Confidence, Format<float>(camera.width, camera.height));
-	temp_.create<cv::cuda::GpuMat>(Channel::Depth, Format<int>(camera.width, camera.height));
-	temp_.create<cv::cuda::GpuMat>(Channel::Depth2, Format<int>(camera.width, camera.height));
-	temp_.create<cv::cuda::GpuMat>(Channel::Normals, Format<float4>(camera.width, camera.height));
+	temp_.create<GpuMat>(Channel::Colour, Format<float4>(camera.width, camera.height));
+	temp_.create<GpuMat>(Channel::Colour2, Format<uchar4>(camera.width, camera.height));
+	temp_.create<GpuMat>(Channel::Confidence, Format<float>(camera.width, camera.height));
+	temp_.create<GpuMat>(Channel::Depth, Format<int>(camera.width, camera.height));
+	temp_.create<GpuMat>(Channel::Depth2, Format<int>(camera.width, camera.height));
+	temp_.create<GpuMat>(Channel::Normals, Format<float4>(camera.width, camera.height));
+
+	cv::cuda::Stream cvstream = cv::cuda::StreamAccessor::wrapStream(stream);
 
 	// Create buffers if they don't exist
 	/*if ((unsigned int)depth1_.width() != camera.width || (unsigned int)depth1_.height() != camera.height) {
@@ -90,10 +92,10 @@ void Splatter::render(ftl::rgbd::VirtualSource *src, cudaStream_t stream) {
 		//ftl::cuda::clear_depth(depth2_, stream);
 		//ftl::cuda::clear_colour(colour2_, stream);
 		
-		temp_.get<GpuMat>(Channel::Depth).setTo(cv::Scalar(0), stream);
-		temp_.get<GpuMat>(Channel::Depth2).setTo(cv::Scalar(0), stream);
-		output_.get<GpuMat>(Channel::Depth).setTo(cv::Scalar(0.0f), stream);
-		output_.get<GpuMat>(Channel::Colour).setTo(cv::Scalar(0,0,0), stream);
+		temp_.get<GpuMat>(Channel::Depth).setTo(cv::Scalar(0), cvstream);
+		temp_.get<GpuMat>(Channel::Depth2).setTo(cv::Scalar(0), cvstream);
+		output_.get<GpuMat>(Channel::Depth).setTo(cv::Scalar(0.0f), cvstream);
+		output_.get<GpuMat>(Channel::Colour).setTo(cv::Scalar(0,0,0), cvstream);
 
 		// Step 1: Put all points into virtual view to gather them
 		//ftl::cuda::dibr_raw(depth1_, scene_->cameraCount(), params, stream);
@@ -107,11 +109,11 @@ void Splatter::render(ftl::rgbd::VirtualSource *src, cudaStream_t stream) {
 				//ftl::cuda::splat_points(depth1_, colour1_, normal1_, depth2_, colour2_, params, stream);
 				//ftl::cuda::int_to_float(depth1_, depth2_, 1.0f / 1000.0f, stream);
 
-				temp_.get<GpuMat>(Channel::Depth).convertTo(output_.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 1000.0f, stream);
+				temp_.get<GpuMat>(Channel::Depth).convertTo(output_.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 1000.0f, cvstream);
 				//src->writeFrames(ts, colour1_, depth2_, stream);
 				src->write(scene_.timestamp, output_, stream);
 			} else {
-				temp_.get<GpuMat>(Channel::Depth).convertTo(output_.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 1000.0f, stream);
+				temp_.get<GpuMat>(Channel::Depth).convertTo(output_.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 1000.0f, cvstream);
 				//ftl::cuda::int_to_float(depth1_, depth2_, 1.0f / 1000.0f, stream);
 				//src->writeFrames(ts, colour1_, depth2_, stream);
 				src->write(scene_.timestamp, output_, stream);
@@ -145,7 +147,7 @@ void Splatter::render(ftl::rgbd::VirtualSource *src, cudaStream_t stream) {
 				src->write(scene_.timestamp, output_, stream);
 			} else {
 				//ftl::cuda::int_to_float(depth1_, depth2_, 1.0f / 1000.0f, stream);
-				temp_.get<GpuMat>(Channel::Depth).convertTo(output_.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 1000.0f, stream);
+				temp_.get<GpuMat>(Channel::Depth).convertTo(output_.get<GpuMat>(Channel::Depth), CV_32F, 1.0f / 1000.0f, cvstream);
 				//src->writeFrames(ts, colour1_, depth2_, stream);
 				src->write(scene_.timestamp, output_, stream);
 			}
@@ -156,6 +158,8 @@ void Splatter::render(ftl::rgbd::VirtualSource *src, cudaStream_t stream) {
 	//ftl::cuda::splat_points(depth1_, depth2_, params, stream);
 
 	// TODO: Second pass
+
+	return true;
 }
 
 void Splatter::setOutputDevice(int device) {

components/renderers/cpp/src/splatter.cu

@@ -2,6 +2,8 @@
 #include <ftl/rgbd/camera.hpp>
 #include <ftl/cuda_common.hpp>
 
+#include <ftl/cuda/weighting.hpp>
+
 #define T_PER_BLOCK 8
 #define UPSAMPLE_FACTOR 1.8f
 #define WARP_SIZE 32
@@ -44,6 +46,12 @@ __device__ inline float4 make_float4(const uchar4 &c) {
     return make_float4(c.x,c.y,c.z,c.w);
 }
 
+
+#define ENERGY_THRESHOLD 0.1f
+#define SMOOTHING_MULTIPLIER_A 10.0f	// For surface search
+#define SMOOTHING_MULTIPLIER_B 4.0f		// For z contribution
+#define SMOOTHING_MULTIPLIER_C 4.0f		// For colour contribution
+
 /*
  * Pass 2: Accumulate attribute contributions if the points pass a visibility test.
  */
@@ -100,7 +108,7 @@ __global__ void dibr_attribute_contrib_kernel(
         const float weight = ftl::cuda::spatialWeighting(length(nearest - camPos), SMOOTHING_MULTIPLIER_C*(nearest.z/params.camera.fx));
         if (screenPos.x+u < colour_out.width() && screenPos.y+v < colour_out.height() && weight > 0.0f) {  // TODO: Use confidence threshold here
             const float4 wcolour = colour * weight;
-			const float4 wnormal = normal * weight;
+			//const float4 wnormal = normal * weight;
 			
 			//printf("Z %f\n", d);
 

components/rgbd-sources/include/ftl/rgbd/source.hpp

@@ -55,7 +55,7 @@ class Source : public ftl::Configurable {
 	explicit Source(ftl::config::json_t &cfg);
 	Source(ftl::config::json_t &cfg, ftl::rgbd::SnapshotReader *);
 	Source(ftl::config::json_t &cfg, ftl::net::Universe *net);
-	~Source();
+	virtual ~Source();
 
 	public:
 	/**

components/rgbd-sources/include/ftl/rgbd/virtual.hpp

@@ -9,6 +9,7 @@ namespace rgbd {
 class VirtualSource : public ftl::rgbd::Source {
     public:
     explicit VirtualSource(ftl::config::json_t &cfg);
+	~VirtualSource();
 
     /**
 	 * Write frames into source buffers from an external renderer. Virtual