Implements #146 upsampling option

applications/reconstruct/src/dibr.cu

@@ -383,20 +383,16 @@ __global__ void OLD_dibr_visibility_kernel(TextureObject<int> depth, int cam, Sp
 
 #define FULL_MASK 0xffffffff
 
-__device__ inline float warpMax(float energy) {
-	//for (int i = WARP_SIZE/2; i > 0; i /= 2) {
-	float e = energy;
-	for (int i = 1; i < 32; i *= 2) {
+__device__ inline float warpMax(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 = max(e, other);
 	}
 	return e;
 }
 
-__device__ inline float warpMin(float energy) {
-	//for (int i = WARP_SIZE/2; i > 0; i /= 2) {
-	float e = energy;
-	for (int i = 1; i < 32; i *= 2) {
+__device__ inline float warpMin(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 = min(e, other);
 	}
@@ -431,9 +427,7 @@ __device__ inline float warpMin(float energy) {
 
 	__syncwarp();
 
-	// Preload valid neighbour points from within a window
-	// TODO: Should be nearest neighbours, which it currently isn't if there are
-	// more than MAX_NEIGHBOUR_2 points.
+	// Search for a valid minimum neighbour
 	for (int i=lane; i<NEIGHBOR_WINDOW; i+=WARP_SIZE) {
 		const int u = (i % (2*NEIGHBOR_RADIUS_2+1)) - NEIGHBOR_RADIUS_2;
 		const int v = (i / (2*NEIGHBOR_RADIUS_2+1)) - NEIGHBOR_RADIUS_2;
@@ -442,17 +436,38 @@ __device__ inline float warpMin(float energy) {
 
 		// If it is close enough...
 		if (point.z > params.camera.m_sensorDepthWorldMin && point.z < params.camera.m_sensorDepthWorldMax && length(point - camPos) <= 0.02f) {
+			atomicMin(&minimum[warp], point.z*1000.0f);
+		}
+	}
+
+	__syncwarp();
+
+	const float minDepth = float(minimum[warp])/1000.0f;
+	
+	// Preload valid neighbour points from within a window. A point is valid
+	// if it is within a specific distance of the minimum.
+	// Also calculate the maximum at the same time.
+	const float3 minPos = params.camera.kinectDepthToSkeleton(x, y, minDepth);
+
+	for (int i=lane; i<NEIGHBOR_WINDOW; i+=WARP_SIZE) {
+		const int u = (i % (2*NEIGHBOR_RADIUS_2+1)) - NEIGHBOR_RADIUS_2;
+		const int v = (i / (2*NEIGHBOR_RADIUS_2+1)) - NEIGHBOR_RADIUS_2;
+		const float3 point = params.camera.kinectDepthToSkeleton(x+u, y+v, float(point_in.tex2D(x+u, y+v)) / 1000.0f);
+
+		// 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);
 			neighborhood_cache[warp][idx] = point;
-			atomicMin(&minimum[warp], point.z*1000.0f);
 			atomicMax(&maximum[warp], point.z*1000.0f);
 		}
 	}
 
 	__syncwarp();
-	
-	const float minDepth = float(minimum[warp])/1000.0f;
+
+	// FIXME: What if minDepth fails energy test, an alternate min is needed.
+	// Perhaps a second pass can be used?
+
 	const float maxDepth = float(maximum[warp])/1000.0f;
 	const float interval = (maxDepth - minDepth) / float(MAX_ITERATIONS);