WIP Gather upsample

applications/reconstruct/src/dibr.cu

@@ -39,12 +39,69 @@ __device__ inline bool isStable(const float3 &previous, const float3 &estimate,
         fabs(previous.z - estimate.z) <= psize;
 }
 
+// ===== PASS 1 : Gather & Upsample (Depth) ====================================
+
 /*
- * Pass 1: Determine depth buffer with enough accuracy for a visibility test in pass 2.
+ * Pass 1: Directly render raw points from all cameras, but upsample the points
+ * if their spacing is within smoothing threshold but greater than their pixel
+ * size in the original image.
+ */
+ __global__ void dibr_merge_upsample_kernel(TextureObject<int> depth, int cam, SplatParams params) {
+	const ftl::voxhash::DepthCameraCUDA &camera = c_cameras[cam];
+
+	const int x = (blockIdx.x*blockDim.x + threadIdx.x) / WARP_SIZE;
+	const int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+	const float3 worldPos = make_float3(tex2D<float4>(camera.points, x, y));
+	//const float3 normal = make_float3(tex2D<float4>(camera.normal, x, y));
+	if (worldPos.x == MINF) return;
+    const float r = (camera.poseInverse * worldPos).z / camera.params.fx;
+
+	// Get virtual camera ray for splat centre and backface cull if possible
+	//const float3 rayOrigin = params.m_viewMatrixInverse * make_float3(0.0f,0.0f,0.0f);
+	//const float3 rayDir = normalize(params.m_viewMatrixInverse * params.camera.kinectDepthToSkeleton(x,y,1.0f) - rayOrigin);
+	//if (dot(rayDir, normal) > 0.0f) return;
+
+    // Find the virtual screen position of current point
+	const float3 camPos = params.m_viewMatrix * worldPos;
+	if (camPos.z < params.camera.m_sensorDepthWorldMin) return;
+	if (camPos.z > params.camera.m_sensorDepthWorldMax) return;
+
+	const int upsample = 5; //min(UPSAMPLE_MAX, int((2.0f*r) * params.camera.fx / camPos.z));
+	const float interval = 1.0f / float(upsample / 2);
+
+            
+    // TODO:(Nick) Check depth buffer and don't do anything if already hidden?
+
+	// Each thread in warp takes an upsample point and updates corresponding depth buffer.
+	const int lane = threadIdx.x % WARP_SIZE;
+	for (int i=lane; i<upsample*upsample; i+=WARP_SIZE) {
+		const float u = (i % upsample) - (upsample / 2);
+		const float v = (i / upsample) - (upsample / 2);
+
+        // Make an initial estimate of the points location
+		// Use centroid depth as estimate...?
+		const float3 point = params.m_viewMatrix * ftl::cuda::upsampled_point(camera.points, make_float2(x+u*interval, y+v*interval));
+		const float d = point.z;
+
+		const uint2 screenPos = params.camera.cameraToKinectScreen(point);
+		const unsigned int cx = screenPos.x+u;
+        const unsigned int cy = screenPos.y+v;
+		if (d > params.camera.m_sensorDepthWorldMin && d < params.camera.m_sensorDepthWorldMax && cx < depth.width() && cy < depth.height()) {
+			// Transform estimated point to virtual cam space and output z
+			atomicMin(&depth(cx,cy), d * 1000.0f);
+		}
+	}
+}
+
+// ===== PASS 2 : Splat Visible Surface ========================================
+
+/*
+ * Pass 2: Determine depth buffer with enough accuracy for a visibility test in pass 2.
  * These values are also used as the actual surface estimate during rendering so should
  * at least be plane or sphere fitted if not MLS smoothed onto the actual surface.
  */
-__global__ void dibr_visibility_kernel(TextureObject<int> depth, int cam, SplatParams params) {
+__global__ void OLD_dibr_visibility_kernel(TextureObject<int> depth, int cam, SplatParams params) {
 	const ftl::voxhash::DepthCameraCUDA &camera = c_cameras[cam];
 
 	const int x = (blockIdx.x*blockDim.x + threadIdx.x) / WARP_SIZE;
@@ -158,13 +215,13 @@ __global__ void dibr_visibility_kernel(TextureObject<int> depth, int cam, SplatP
 	}
 }
 
-// ------ Alternative for pass 1: principle surfaces ---------------------------
+// ------ Alternative for pass 2: principle surfaces ---------------------------
 
 #define NEIGHBOR_RADIUS 1
 #define MAX_NEIGHBORS ((NEIGHBOR_RADIUS*2+1)*(NEIGHBOR_RADIUS*2+1))
 
 /*
- * Pass 1: Determine depth buffer with enough accuracy for a visibility test in pass 2.
+ * Pass 2: Determine depth buffer with enough accuracy for a visibility test in pass 2.
  * These values are also used as the actual surface estimate during rendering so should
  * at least be plane or sphere fitted if not MLS smoothed onto the actual surface.
  */
@@ -447,16 +504,20 @@ void ftl::cuda::dibr(const TextureObject<int> &depth_out,
 	cudaSafeCall(cudaDeviceSynchronize());
 #endif
 
-    int i=3;
-    // Pass 1, merge a depth map from each camera.
+	int i=3;
+	// Pass 1, gather and upsample depth maps
 	for (int i=0; i<numcams; ++i)
-        dibr_visibility_principal_kernel<<<sgridSize, sblockSize, 0, stream>>>(depth_out, i, params);
+		dibr_merge_upsample_kernel<<<sgridSize, sblockSize, 0, stream>>>(depth_out, i, params);
+
+    // Pass 2, merge a depth map from each camera.
+	//for (int i=0; i<numcams; ++i)
+    //    dibr_visibility_principal_kernel<<<sgridSize, sblockSize, 0, stream>>>(depth_out, i, params);
 
-    // Pass 2, accumulate all point contributions to pixels
+    // Pass 3, accumulate all point contributions to pixels
     for (int i=0; i<numcams; ++i)
         dibr_attribute_contrib_kernel<<<sgridSize, sblockSize, 0, stream>>>(depth_out, tmp_colour, normal_out, confidence_out, i, params);
 
-    // Pass 3, normalise contributions
+    // Pass 4, normalise contributions
     dibr_normalise_kernel<<<gridSize, blockSize, 0, stream>>>(tmp_colour, colour_out, normal_out, confidence_out);
 
 	cudaSafeCall( cudaGetLastError() );

applications/reconstruct/src/splat_params.hpp

@@ -8,8 +8,10 @@
 namespace ftl {
 namespace render {
 
-static const uint kShowBlockBorders = 0x0001;
-static const uint kNoSplatting = 0x0002;
+static const uint kShowBlockBorders = 0x00000001;  // Deprecated: from voxels system
+static const uint kNoSplatting = 0x00000002;
+static const uint kNoUpsampling = 0x00000004;
+static const uint kNoTexturing = 0x00000008;
 
 struct __align__(16) SplatParams {
 	float4x4 m_viewMatrix;