First attempt at new splat

components/renderers/cpp/src/splat_render.cpp

@@ -94,13 +94,30 @@ void Splatter::renderChannel(
 		ftl::cuda::dibr_merge(
 			f.createTexture<float4>(Channel::Points),
 			f.createTexture<float4>(Channel::Normals),
-			temp_.getTexture<int>(Channel::Depth),
+			temp_.createTexture<int>(Channel::Depth2),
 			params, backcull_, stream
 		);
 
 		//LOG(INFO) << "DIBR DONE";
 	}
 
+	// Now splat the points
+	for (size_t i=0; i < scene_->frames.size(); ++i) {
+		auto &f = scene_->frames[i];
+		auto *s = scene_->sources[i];
+
+		auto pose = MatrixConversion::toCUDA(s->getPose().cast<float>().inverse()); //.inverse());
+
+		ftl:cuda::splat(
+			f.getTexture<float4>(Channel::Points),
+			f.getTexture<float4>(Channel::Normals),
+			temp_.getTexture<int>(Channel::Depth2),
+			temp_.getTexture<int>(Channel::Depth),
+			s->parameters(),
+			pose, params, stream
+		);
+	}
+
 	// TODO: Add the depth splatting step..
 
 	temp_.createTexture<float4>(Channel::Colour);

components/renderers/cpp/src/splatter.cu

@@ -4,6 +4,7 @@
 #include <ftl/cuda_common.hpp>
 
 #include <ftl/cuda/weighting.hpp>
+#include <ftl/cuda/intersections.hpp>
 
 #define T_PER_BLOCK 8
 #define UPSAMPLE_FACTOR 1.8f
@@ -13,6 +14,13 @@
 #define MAX_ITERATIONS 32  // Note: Must be multiple of 32
 #define SPATIAL_SMOOTHING 0.005f
 
+#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 1.0f		// For colour contribution
+
+#define ACCUM_DIAMETER 8
+
 using ftl::cuda::TextureObject;
 using ftl::render::SplatParams;
 
@@ -72,17 +80,103 @@ void ftl::cuda::dibr_merge(TextureObject<float4> &points, TextureObject<float4>
 
 //==============================================================================
 
-__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
+ */
+ template <int SEARCH_DIAMETER>
+ __global__ void splat_kernel(
+        TextureObject<float4> points,       // Original 3D points
+        TextureObject<float4> normals,
+        TextureObject<int> depth_in,        // Virtual depth map
+        TextureObject<int> depth_out,   // Accumulated output
+        ftl::rgbd::Camera camera,
+        float4x4 pose_inv,
+        SplatParams params) {
+        
+    //const ftl::voxhash::DepthCameraCUDA &camera = c_cameras[cam];
+
+    const int tid = (threadIdx.x + threadIdx.y * blockDim.x);
+    //const int warp = tid / WARP_SIZE;
+    const int x = (blockIdx.x*blockDim.x + threadIdx.x) / WARP_SIZE;
+    const int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+    const float3 origin = params.m_viewMatrixInverse * make_float3(0.0f);
+    float3 ray = params.m_viewMatrixInverse.getFloat3x3() * params.camera.screenToCam(x,y,1.0f);
+    //ray = ray / length(ray);
+
+    //float depth = 0.0f;
+    //float contrib = 0.0f;
+
+    // Each thread in warp takes an upsample point and updates corresponding depth buffer.
+    const int lane = tid % WARP_SIZE;
+    for (int i=lane; i<SEARCH_DIAMETER*SEARCH_DIAMETER; i+=WARP_SIZE) {
+        const float u = (i % SEARCH_DIAMETER) - (SEARCH_DIAMETER / 2);
+        const float v = (i / SEARCH_DIAMETER) - (SEARCH_DIAMETER / 2);
+
+        // 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);
+
+        if (d < params.camera.minDepth || d > params.camera.maxDepth) continue;
+
+        const float3 camPos = params.camera.screenToCam((int)(x+u),(int)(y+v),d);
+        const float3 worldPos = params.m_viewMatrixInverse * camPos;
+        const float3 camPos2 = pose_inv * worldPos;
+        const uint2 screenPos = camera.camToScreen<uint2>(camPos2);
+
+        if (screenPos.x < points.width() && screenPos.y < points.height()) {
+            // Can now read points, normals and colours from source cam
+
+            // What is contribution of our current point at this pixel?
+            const float3 p = make_float3(points.tex2D((int)screenPos.x, (int)screenPos.y));
+            const float weight = ftl::cuda::spatialWeighting(worldPos, p, 0.01f);  //1.0f*(camPos2.z/camera.fx)
+            if (weight <= 0.0f) continue;
+
+            const float3 n = make_float3(normals.tex2D((int)screenPos.x, (int)screenPos.y));
+
+            // Does the ray intersect plane of splat?
+            float t;
+            if (ftl::cuda::intersectPlane(n, p, origin, ray, t)) {
+                //depth += t * weight;
+                //contrib += weight;
+                atomicMin(&depth_out(x,y), t * 1000.0f);
+            }
+        }
     }
 
+    /*depth = warpSum(depth);
+    contrib = warpSum(contrib);
+    if (lane == 0) {
+        depth_out(x, y) = 
+    }*/
+}
 
-#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 1.0f		// For colour contribution
+void ftl::cuda::splat(
+        TextureObject<float4> &points,       // Original 3D points
+        TextureObject<float4> &normals,
+        TextureObject<int> &depth_in,        // Virtual depth map
+        TextureObject<int> &depth_out,
+        const ftl::rgbd::Camera &camera, const float4x4 &pose_inv,
+        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);
 
-#define ACCUM_DIAMETER 8
+    splat_kernel<8><<<gridSize, blockSize, 0, stream>>>(
+        points,
+        normals,
+        depth_in,
+        depth_out,
+        camera,
+        pose_inv,
+        params
+    );
+    cudaSafeCall( cudaGetLastError() );
+}
+
+//==============================================================================
+
+__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

@@ -14,6 +14,14 @@ namespace cuda {
 		bool culling,
 		cudaStream_t stream);
 
+	void splat(
+        ftl::cuda::TextureObject<float4> &points,       // Original 3D points
+        ftl::cuda::TextureObject<float4> &normals,
+        ftl::cuda::TextureObject<int> &depth_in,        // Virtual depth map
+        ftl::cuda::TextureObject<int> &depth_out,
+        const ftl::rgbd::Camera &camera, const float4x4 &pose_inv,
+        const ftl::render::SplatParams &params, cudaStream_t stream);
+
 	void dibr_attribute(
 		ftl::cuda::TextureObject<uchar4> &in,	// Original colour image
 		ftl::cuda::TextureObject<float4> &points,		// Original 3D points