From c86b420d0af79305b68af0a212b2b295590e3c89 Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Fri, 9 Aug 2019 15:05:21 +0300
Subject: [PATCH] WIP Variable neighbor points
---
applications/reconstruct/src/dibr.cu | 10 +++++++---
1 file changed, 7 insertions(+), 3 deletions(-)
diff --git a/applications/reconstruct/src/dibr.cu b/applications/reconstruct/src/dibr.cu
index 97157c56f..88d480872 100644
--- a/applications/reconstruct/src/dibr.cu
+++ b/applications/reconstruct/src/dibr.cu
@@ -346,7 +346,8 @@ __global__ void OLD_dibr_visibility_kernel(TextureObject<int> depth, int cam, Sp
}
#define NEIGHBOR_RADIUS_2 5
-#define MAX_NEIGHBORS_2 ((NEIGHBOR_RADIUS_2*2+1)*(NEIGHBOR_RADIUS_2*2+1))
+#define NEIGHBOR_WINDOW ((NEIGHBOR_RADIUS_2*2+1)*(NEIGHBOR_RADIUS_2*2+1))
+#define MAX_NEIGHBORS_2 20
/*
@@ -360,6 +361,7 @@ __global__ void OLD_dibr_visibility_kernel(TextureObject<int> depth, int cam, Sp
__shared__ float3 neighborhood_cache[2*T_PER_BLOCK][MAX_NEIGHBORS_2];
__shared__ int minimum[2*T_PER_BLOCK];
__shared__ int maximum[2*T_PER_BLOCK];
+ __shared__ int nidx[2*T_PER_BLOCK];
const int warp = threadIdx.x / WARP_SIZE + threadIdx.y*2;
const int x = (blockIdx.x*blockDim.x + threadIdx.x) / WARP_SIZE;
@@ -393,17 +395,19 @@ __global__ void OLD_dibr_visibility_kernel(TextureObject<int> depth, int cam, Sp
if (lane == 0) {
minimum[warp] = 100000000;
maximum[warp] = -100000000;
+ nidx[warp] = 0;
}
__syncwarp();
- for (int i=lane; i<MAX_NEIGHBORS_2; i+=WARP_SIZE) {
+ 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);
- neighborhood_cache[warp][i] = point;
if (length(point - camPos) <= 0.04f) {
+ 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);
}
--
GitLab