Render colours during splat

components/renderers/cpp/src/splat_render.cpp

@@ -115,6 +115,7 @@ void Splatter::renderChannel(
 			temp_.getTexture<float>(Channel::Contribution),
 			params, stream
 		);
+	}
 
 	// Normalise attribute contributions
 	ftl::cuda::dibr_normalise(
@@ -123,16 +124,53 @@ void Splatter::renderChannel(
 		temp_.createTexture<float>(Channel::Contribution),
 		stream
 	);
+
+	temp_.get<GpuMat>(Channel::Colour).setTo(cv::Scalar(0.0f,0.0f,0.0f,0.0f), cvstream);
+	temp_.get<GpuMat>(Channel::Contribution).setTo(cv::Scalar(0.0f), cvstream);
+
+	// Create colours first
+	for (auto &f : scene_->frames) {
+		// Convert colour from BGR to BGRA if needed
+		if (f.get<GpuMat>(channel).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::Left),
+			f.createTexture<float4>(Channel::Points),
+			temp_.getTexture<int>(Channel::Depth2),
+			temp_.getTexture<float4>(Channel::Colour),
+			temp_.getTexture<float>(Channel::Contribution),
+			params, stream
+		);
 	}
 
+	// Normalise attribute contributions
+	ftl::cuda::dibr_normalise(
+		temp_.createTexture<float4>(Channel::Colour),
+		temp_.createTexture<uchar4>(Channel::Colour2),
+		temp_.createTexture<float>(Channel::Contribution),
+		stream
+	);
+
+	out.get<GpuMat>(Channel::Left).setTo(cv::Scalar(0,0,0,0), cvstream);
+
 	// Now splat the points
 	ftl:cuda::splat(
 		out.getTexture<float4>(Channel::Normals),
+		temp_.getTexture<uchar4>(Channel::Colour2),
 		temp_.getTexture<int>(Channel::Depth2),
-		temp_.getTexture<int>(Channel::Depth),
+		out.createTexture<float>(Channel::Depth),
+		out.createTexture<uchar4>(Channel::Left),
 		params, stream
 	);
 
+	return;
+
 	temp_.get<GpuMat>(Channel::Colour).setTo(cv::Scalar(0.0f,0.0f,0.0f,0.0f), cvstream);
 	temp_.get<GpuMat>(Channel::Contribution).setTo(cv::Scalar(0.0f), cvstream);
 

components/renderers/cpp/src/splatter.cu

@@ -98,6 +98,10 @@ __device__ inline float warpSum(float e) {
 	return e;
 }
 
+__device__ inline float4 make_float4(const uchar4 &c) {
+    return make_float4(c.x,c.y,c.z,c.w);
+}
+
 /*
  * Pass 1b: Expand splats to full size and merge
  */
@@ -105,8 +109,10 @@ __device__ inline float warpSum(float e) {
  __global__ void splat_kernel(
         //TextureObject<float4> points,       // Original 3D points
         TextureObject<float4> normals,
+        TextureObject<uchar4> colour_in,
         TextureObject<int> depth_in,        // Virtual depth map
-        TextureObject<int> depth_out,   // Accumulated output
+        TextureObject<float> depth_out,   // Accumulated output
+        TextureObject<uchar4> colour_out,
         //ftl::rgbd::Camera camera,
         //float4x4 pose_inv,
         SplatParams params) {
@@ -130,7 +136,13 @@ __device__ inline float warpSum(float e) {
     //float contrib = 0.0f;
     float depth = 1000.0f;
 
-    float2 results[(SEARCH_DIAMETER*SEARCH_DIAMETER) / WARP_SIZE];
+    struct Result {
+        float weight;
+        float depth;
+        uchar4 colour;
+    };
+
+    Result results[(SEARCH_DIAMETER*SEARCH_DIAMETER) / WARP_SIZE];
 
     // Each thread in warp takes an upsample point and updates corresponding depth buffer.
     const int lane = tid % WARP_SIZE;
@@ -138,7 +150,7 @@ __device__ inline float warpSum(float e) {
         const float u = (i % SEARCH_DIAMETER) - (SEARCH_DIAMETER / 2);
         const float v = (i / SEARCH_DIAMETER) - (SEARCH_DIAMETER / 2);
 
-        results[i/WARP_SIZE] = make_float2(0.0f, 0.0f);
+        results[i/WARP_SIZE] = {0.0f, 0.0f, make_uchar4(0,0,0,0)};
 
         // Use the depth buffer to determine this pixels 3D position in camera space
         const float d = ((float)depth_in.tex2D(x+u, y+v)/1000.0f);
@@ -176,7 +188,7 @@ __device__ inline float warpSum(float e) {
                 //depth += t * weight;
                 //contrib += weight;
                 depth = min(depth, t);
-                results[i/WARP_SIZE] = make_float2(t, weight);
+                results[i/WARP_SIZE] = {weight, t, colour_in.tex2D((int)x, (int)y)}; //make_float2(t, weight);
                 //atomicMin(&depth_out(x,y), (int)(depth * 1000.0f));
             }
         //}
@@ -186,37 +198,46 @@ __device__ inline float warpSum(float e) {
 
     float adepth = 0.0f;
     float contrib = 0.0f;
+    float4 colour = make_float4(0.0f);
 
     // Loop over results array
     for (int i=0; i<(SEARCH_DIAMETER*SEARCH_DIAMETER) / WARP_SIZE; ++i) {
-        if (results[i].x - depth < 0.04f) {
-            adepth += results[i].x * results[i].y;
-            contrib += results[i].y;
+        if (results[i].depth - depth < 0.04f) {
+            adepth += results[i].depth * results[i].weight;
+            colour += make_float4(results[i].colour) * results[i].weight;
+            contrib += results[i].weight;
         }
     }
 
     // Sum all attributes and contributions
     adepth = warpSum(adepth);
+    colour.x = warpSum(colour.x);
+    colour.y = warpSum(colour.y);
+    colour.z = warpSum(colour.z);
     contrib = warpSum(contrib);
 
     if (lane == 0 && contrib > 0.0f) {
-        depth_out(x,y) = (int)(adepth / contrib * 1000.0f);
+        depth_out(x,y) = adepth / contrib;
+        colour_out(x,y) = make_uchar4(colour.x / contrib, colour.y / contrib, colour.z / contrib, 255.0f);
     }
 }
 
 void ftl::cuda::splat(
         TextureObject<float4> &normals,
+        TextureObject<uchar4> &colour_in,
         TextureObject<int> &depth_in,        // Virtual depth map
-        TextureObject<int> &depth_out,
-        //TextureObject<uchar4> &colour_out,
+        TextureObject<float> &depth_out,
+        TextureObject<uchar4> &colour_out,
         const 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);
 
     splat_kernel<8><<<gridSize, blockSize, 0, stream>>>(
         normals,
+        colour_in,
         depth_in,
         depth_out,
+        colour_out,
         params
     );
     cudaSafeCall( cudaGetLastError() );
@@ -224,10 +245,6 @@ void ftl::cuda::splat(
 
 //==============================================================================
 
-__device__ inline float4 make_float4(const uchar4 &c) {
-    return make_float4(c.x,c.y,c.z,c.w);
-}
-
 /*
  * Pass 2: Accumulate attribute contributions if the points pass a visibility test.
  */

components/renderers/cpp/src/splatter_cuda.hpp

@@ -16,8 +16,10 @@ namespace cuda {
 
 	void splat(
         ftl::cuda::TextureObject<float4> &normals,
+		ftl::cuda::TextureObject<uchar4> &colour_in,
         ftl::cuda::TextureObject<int> &depth_in,        // Virtual depth map
-        ftl::cuda::TextureObject<int> &depth_out,
+        ftl::cuda::TextureObject<float> &depth_out,
+		ftl::cuda::TextureObject<uchar4> &colour_out,
         const ftl::render::SplatParams &params, cudaStream_t stream);
 
 	void dibr_attribute(