Almost working attribute accum

components/renderers/cpp/src/splat_render.cpp

@@ -33,7 +33,7 @@ bool Splatter::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out, cuda
 
 	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::Contribution, 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));
@@ -81,6 +81,8 @@ bool Splatter::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out, cuda
 
 	//LOG(INFO) << "Render ready: " << camera.width << "," << camera.height;
 
+	temp_.createTexture<int>(Channel::Depth);
+
 	// Render each camera into virtual view
 	for (size_t i=0; i<scene_->frames.size(); ++i) {
 		auto &f = scene_->frames[i];
@@ -104,13 +106,47 @@ bool Splatter::render(ftl::rgbd::VirtualSource *src, ftl::rgbd::Frame &out, cuda
 
 		ftl::cuda::dibr_merge(
 			f.createTexture<float4>(Channel::Points),
-			temp_.createTexture<int>(Channel::Depth),
+			temp_.getTexture<int>(Channel::Depth),
 			params, stream
 		);
 
 		//LOG(INFO) << "DIBR DONE";
 	}
 
+	temp_.createTexture<float4>(Channel::Colour);
+	temp_.createTexture<float>(Channel::Contribution);
+
+	// Accumulate attribute contributions for each pixel
+	for (auto &f : scene_->frames) {
+		// Convert colour from BGR to BGRA if needed
+		if (f.get<GpuMat>(Channel::Colour).type() == CV_8UC3) {
+			// Convert to 4 channel colour
+			auto &col = f.get<GpuMat>(Channel::Colour);
+			GpuMat tmp(col.size(), CV_8UC4);
+			cv::cuda::swap(col, tmp);
+			cv::cuda::cvtColor(tmp,col, cv::COLOR_BGR2BGRA);
+		}
+	
+		ftl::cuda::dibr_attribute(
+			f.createTexture<uchar4>(Channel::Colour),
+			f.createTexture<float4>(Channel::Points),
+			temp_.getTexture<int>(Channel::Depth),
+			temp_.getTexture<float4>(Channel::Colour),
+			temp_.getTexture<float>(Channel::Contribution),
+			params, stream
+		);
+	}
+
+	// Normalise attribute contributions
+	//for (auto &f : scene_->frames) {
+		ftl::cuda::dibr_normalise(
+			temp_.createTexture<float4>(Channel::Colour),
+			out.createTexture<uchar4>(Channel::Colour),
+			temp_.createTexture<float>(Channel::Contribution),
+			stream
+		);
+	//}
+
 		//ftl::cuda::dibr(depth1_, colour1_, normal1_, depth2_, colour_tmp_, depth3_, scene_->cameraCount(), params, stream);
 
 		// Step 1: Put all points into virtual view to gather them

components/renderers/cpp/src/splatter.cu

@@ -51,6 +51,8 @@ void ftl::cuda::dibr_merge(TextureObject<float4> &points, TextureObject<int> &de
     cudaSafeCall( cudaGetLastError() );
 }
 
+//==============================================================================
+
 __device__ inline float4 make_float4(const uchar4 &c) {
     return make_float4(c.x,c.y,c.z,c.w);
 }
@@ -135,25 +137,54 @@ __global__ void dibr_attribute_contrib_kernel(
 	}
 }
 
+void ftl::cuda::dibr_attribute(
+        TextureObject<uchar4> &colour_in,    // Original colour image
+        TextureObject<float4> &points,       // Original 3D points
+        TextureObject<int> &depth_in,        // Virtual depth map
+        TextureObject<float4> &colour_out,   // Accumulated output
+        //TextureObject<float4> normal_out,
+        TextureObject<float> &contrib_out,
+        SplatParams &params, cudaStream_t stream) {
+    const dim3 gridSize((depth_in.width() + 2 - 1)/2, (depth_in.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
+	const dim3 blockSize(2*WARP_SIZE, T_PER_BLOCK);
+
+    dibr_attribute_contrib_kernel<<<gridSize, blockSize, 0, stream>>>(
+        colour_in,
+        points,
+        depth_in,
+        colour_out,
+        contrib_out,
+        params
+    );
+    cudaSafeCall( cudaGetLastError() );
+}
 
-
+//==============================================================================
 
 __global__ void dibr_normalise_kernel(
-    TextureObject<float4> colour_in,
-    TextureObject<uchar4> colour_out,
-    //TextureObject<float4> normals,
-    TextureObject<float> contribs) {
-const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
-const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
-
-if (x < colour_in.width() && y < colour_in.height()) {
-    const float4 colour = colour_in.tex2D((int)x,(int)y);
-    //const float4 normal = normals.tex2D((int)x,(int)y);
-    const float contrib = contribs.tex2D((int)x,(int)y);
-
-    if (contrib > 0.0f) {
-        colour_out(x,y) = make_uchar4(colour.x / contrib, colour.y / contrib, colour.z / contrib, 0);
-        //normals(x,y) = normal / contrib;
+        TextureObject<float4> colour_in,
+        TextureObject<uchar4> colour_out,
+        //TextureObject<float4> normals,
+        TextureObject<float> contribs) {
+    const unsigned int x = blockIdx.x*blockDim.x + threadIdx.x;
+    const unsigned int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+    if (x < colour_in.width() && y < colour_in.height()) {
+        const float4 colour = colour_in.tex2D((int)x,(int)y);
+        //const float4 normal = normals.tex2D((int)x,(int)y);
+        const float contrib = contribs.tex2D((int)x,(int)y);
+
+        if (contrib > 0.0f) {
+            colour_out(x,y) = make_uchar4(colour.x / contrib, colour.y / contrib, colour.z / contrib, 0);
+            //normals(x,y) = normal / contrib;
+        }
     }
 }
+
+void ftl::cuda::dibr_normalise(TextureObject<float4> &colour_in, TextureObject<uchar4> &colour_out, TextureObject<float> &contribs, cudaStream_t stream) {
+    const dim3 gridSize((colour_in.width() + T_PER_BLOCK - 1)/T_PER_BLOCK, (colour_in.height() + T_PER_BLOCK - 1)/T_PER_BLOCK);
+    const dim3 blockSize(T_PER_BLOCK, T_PER_BLOCK);
+
+    dibr_normalise_kernel<<<gridSize, blockSize, 0, stream>>>(colour_in, colour_out, contribs);
+    cudaSafeCall( cudaGetLastError() );
 }

components/renderers/cpp/src/splatter_cuda.hpp

@@ -7,6 +7,21 @@
 namespace ftl {
 namespace cuda {
     void dibr_merge(ftl::cuda::TextureObject<float4> &points, ftl::cuda::TextureObject<int> &depth, ftl::render::SplatParams params, cudaStream_t stream);
+
+    void dibr_attribute(
+        ftl::cuda::TextureObject<uchar4> &colour_in,    // Original colour image
+        ftl::cuda::TextureObject<float4> &points,       // Original 3D points
+        ftl::cuda::TextureObject<int> &depth_in,        // Virtual depth map
+        ftl::cuda::TextureObject<float4> &colour_out,   // Accumulated output
+        //TextureObject<float4> normal_out,
+        ftl::cuda::TextureObject<float> &contrib_out,
+        ftl::render::SplatParams &params, cudaStream_t stream);
+
+    void dibr_normalise(
+        ftl::cuda::TextureObject<float4> &colour_in,
+        ftl::cuda::TextureObject<uchar4> &colour_out,
+        ftl::cuda::TextureObject<float> &contribs,
+        cudaStream_t stream);
 }
 }
 

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

@@ -17,9 +17,10 @@ enum struct Channel : int {
     Disparity = 3,
     Depth2 = 3,
     Deviation = 4,
-    Normals,            // 32FC4
-    Points,             // 32FC4
-    Confidence,         // 32F
+    Normals = 5,        // 32FC4
+    Points = 6,         // 32FC4
+    Confidence = 7,     // 32F
+    Contribution = 7,   // 32F
     Flow,               // 32F
     Energy,             // 32F
     LeftGray,