Resolve neighbour buffer loop

applications/reconstruct/src/dibr.cu

@@ -13,7 +13,7 @@
 #define WARP_SIZE 32
 #define DEPTH_THRESHOLD 0.05f
 #define UPSAMPLE_MAX 60
-#define MAX_ITERATIONS 64  // Note: Must be multiple of 32
+#define MAX_ITERATIONS 32  // Note: Must be multiple of 32
 #define SPATIAL_SMOOTHING 0.005f
 
 using ftl::cuda::TextureObject;
@@ -457,7 +457,9 @@ __device__ inline float warpMin(float e) {
 		// If it is close enough...
 		if (point.z > params.camera.m_sensorDepthWorldMin && point.z < params.camera.m_sensorDepthWorldMax && length(point - minPos) <= 0.02f) {
 			// Append to neighbour list
-			unsigned int idx = atomicInc(&nidx[warp], MAX_NEIGHBORS_2-1);
+			//unsigned int idx = atomicInc(&nidx[warp], MAX_NEIGHBORS_2-1);
+			unsigned int idx = atomicAdd(&nidx[warp], 1);
+			if (idx >= MAX_NEIGHBORS_2) break;
 			neighborhood_cache[warp][idx] = point;
 			atomicMax(&maximum[warp], point.z*1000.0f);
 		}
@@ -475,6 +477,20 @@ __device__ inline float warpMin(float e) {
 	if (maxDepth <= params.camera.m_sensorDepthWorldMin) return;
 	//if (y == 200) printf("interval: %f\n", maxDepth);
 
+	// If all samples say same depth, then agree and return
+	// TODO: Check this is valid, since small energies should be removed...
+	/*if (fabs(minDepth - maxDepth) < 0.0001f) {
+		if (lane == 0) {
+			const unsigned int cx = x;
+			const unsigned int cy = y;
+			if (minDepth < params.camera.m_sensorDepthWorldMax && cx < depth.width() && cy < depth.height()) {
+				// Transform estimated point to virtual cam space and output z
+				atomicMin(&depth(cx,cy), minDepth * 1000.0f);
+			}
+		}
+		return;
+	}*/
+
 
 	float maxenergy = -1.0f;
 	float bestdepth = 0.0f;
@@ -482,7 +498,7 @@ __device__ inline float warpMin(float e) {
 	// Search for best or threshold energy
 	for (int k=lane; k<MAX_ITERATIONS; k+=WARP_SIZE) {
 		const float3 nearest = params.camera.kinectDepthToSkeleton(x,y,minDepth+float(k)*interval);
-		const float myenergy = ftl::cuda::mls_point_energy<MAX_NEIGHBORS_2>(neighborhood_cache[warp], nearest, nidx[warp], SPATIAL_SMOOTHING);
+		const float myenergy = ftl::cuda::mls_point_energy<MAX_NEIGHBORS_2>(neighborhood_cache[warp], nearest, min(nidx[warp], MAX_NEIGHBORS_2), SPATIAL_SMOOTHING);
 		const float newenergy = warpMax(max(myenergy, maxenergy));
 		bestdepth = (myenergy == newenergy) ? nearest.z : (newenergy > maxenergy) ? 0.0f : bestdepth;
 		maxenergy = newenergy;