diff --git a/applications/reconstruct/src/dibr.cu b/applications/reconstruct/src/dibr.cu index 08c73219df472ba526792f41ced599e1dd1e181c..2d64ddd151f09a289f8eeef411b170b32c7fd469 100644 --- a/applications/reconstruct/src/dibr.cu +++ b/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);