From 15bc458bc363a699579c8dd9823e4beec1d90015 Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Sun, 11 Aug 2019 11:11:06 +0300
Subject: [PATCH] Resolve neighbour buffer loop
---
applications/reconstruct/src/dibr.cu | 22 +++++++++++++++++++---
1 file changed, 19 insertions(+), 3 deletions(-)
diff --git a/applications/reconstruct/src/dibr.cu b/applications/reconstruct/src/dibr.cu
index 52dd31965..ebf53b770 100644
--- a/applications/reconstruct/src/dibr.cu
+++ b/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;
--
GitLab