Ongoing #133 improvements

applications/reconstruct/src/ilw/ilw.cu

@@ -16,6 +16,14 @@ __device__ inline float warpMin(float e) {
 	return e;
 }
 
+__device__ inline float warpSum(float e) {
+	for (int i = WARP_SIZE/2; i > 0; i /= 2) {
+		const float other = __shfl_xor_sync(FULL_MASK, e, i, WARP_SIZE);
+		e += other;
+	}
+	return e;
+}
+
 #define COR_WIN_RADIUS 17
 #define COR_WIN_SIZE (COR_WIN_RADIUS * COR_WIN_RADIUS)
 
@@ -41,8 +49,9 @@ __global__ void correspondence_energy_vector_kernel(
     const uint2 screen2 = cam2.camToScreen<uint2>(camPos2);
     
     float bestcost = 1.1f;
-    float nextbest = 1.0f;
-    float3 bestpoint;
+    float avgcost = 0.0f;
+	float3 bestpoint;
+	int count = 0;
 
     // Project to p2 using cam2
     // Each thread takes a possible correspondence and calculates a weighting
@@ -55,6 +64,8 @@ __global__ void correspondence_energy_vector_kernel(
         const uchar4 colour2 = c2.tex2D(screen2.x+u, screen2.y+v);
 		if (world2.x == MINF) continue;
 
+		++count;
+
         // Determine degree of correspondence
         float cost = 1.0f - ftl::cuda::spatialWeighting(world1, world2, 0.04f);
         cost *= 1.0f - ftl::cuda::colourWeighting(colour1, colour2, 50.0f);
@@ -62,15 +73,16 @@ __global__ void correspondence_energy_vector_kernel(
 
         if (cost < bestcost) {
             bestpoint = world2;
-            nextbest = bestcost;
+            avgcost += cost;
             bestcost = cost;
         }
     }
 
+	count = warpSum(count);
     const float mincost = warpMin(bestcost);
-    bool best = mincost == bestcost;
-    bestcost = (best) ? nextbest : bestcost;
-    const float confidence = mincost / warpMin(bestcost);
+	bool best = mincost == bestcost;
+	avgcost = warpSum(avgcost) / count;
+    const float confidence = (avgcost - mincost);
 
     if (best && mincost < 1.0f) {
         vout(x,y) = vout.tex2D(x, y) + make_float4(
@@ -79,11 +91,9 @@ __global__ void correspondence_energy_vector_kernel(
             (bestpoint.z - world1.z),
 			mincost);
 			
-		eout(x,y) = max(eout(x,y), length(bestpoint-world1) * 10.0f);
-		//eout(x,y) = max(eout(x,y), (1.0f - mincost) * 7.0f); //confidence * 5.0f;
-
-		// FIXME: This needs to be summed across all frames
-        //eout(x,y) = max(eout(x, y), confidence * 7.0f);
+		//eout(x,y) = max(eout(x,y), length(bestpoint-world1) * 10.0f);
+		//eout(x,y) = max(eout(x,y), (1.0f - mincost) * 7.0f);
+        eout(x,y) = max(eout(x, y), confidence * 7.0f);
     } else if (mincost >= 1.0f && lane == 0) {
         //vout(x,y) = make_float4(0.0f);
         //eout(x,y) = 0.0f;