diff --git a/applications/reconstruct/src/dibr.cu b/applications/reconstruct/src/dibr.cu
index acb86216a7eecd6e10f3ebf3a3b999964b56cf09..8c9607da2bf94fd22198f103e26175ab941c1a53 100644
--- a/applications/reconstruct/src/dibr.cu
+++ b/applications/reconstruct/src/dibr.cu
@@ -39,12 +39,69 @@ __device__ inline bool isStable(const float3 &previous, const float3 &estimate,
fabs(previous.z - estimate.z) <= psize;
}
+// ===== PASS 1 : Gather & Upsample (Depth) ====================================
+
/*
- * Pass 1: Determine depth buffer with enough accuracy for a visibility test in pass 2.
+ * Pass 1: Directly render raw points from all cameras, but upsample the points
+ * if their spacing is within smoothing threshold but greater than their pixel
+ * size in the original image.
+ */
+ __global__ void dibr_merge_upsample_kernel(TextureObject<int> depth, int cam, SplatParams params) {
+ const ftl::voxhash::DepthCameraCUDA &camera = c_cameras[cam];
+
+ const int x = (blockIdx.x*blockDim.x + threadIdx.x) / WARP_SIZE;
+ const int y = blockIdx.y*blockDim.y + threadIdx.y;
+
+ const float3 worldPos = make_float3(tex2D<float4>(camera.points, x, y));
+ //const float3 normal = make_float3(tex2D<float4>(camera.normal, x, y));
+ if (worldPos.x == MINF) return;
+ const float r = (camera.poseInverse * worldPos).z / camera.params.fx;
+
+ // Get virtual camera ray for splat centre and backface cull if possible
+ //const float3 rayOrigin = params.m_viewMatrixInverse * make_float3(0.0f,0.0f,0.0f);
+ //const float3 rayDir = normalize(params.m_viewMatrixInverse * params.camera.kinectDepthToSkeleton(x,y,1.0f) - rayOrigin);
+ //if (dot(rayDir, normal) > 0.0f) return;
+
+ // Find the virtual screen position of current point
+ const float3 camPos = params.m_viewMatrix * worldPos;
+ if (camPos.z < params.camera.m_sensorDepthWorldMin) return;
+ if (camPos.z > params.camera.m_sensorDepthWorldMax) return;
+
+ const int upsample = 5; //min(UPSAMPLE_MAX, int((2.0f*r) * params.camera.fx / camPos.z));
+ const float interval = 1.0f / float(upsample / 2);
+
+
+ // TODO:(Nick) Check depth buffer and don't do anything if already hidden?
+
+ // Each thread in warp takes an upsample point and updates corresponding depth buffer.
+ const int lane = threadIdx.x % WARP_SIZE;
+ for (int i=lane; i<upsample*upsample; i+=WARP_SIZE) {
+ const float u = (i % upsample) - (upsample / 2);
+ const float v = (i / upsample) - (upsample / 2);
+
+ // Make an initial estimate of the points location
+ // Use centroid depth as estimate...?
+ const float3 point = params.m_viewMatrix * ftl::cuda::upsampled_point(camera.points, make_float2(x+u*interval, y+v*interval));
+ const float d = point.z;
+
+ const uint2 screenPos = params.camera.cameraToKinectScreen(point);
+ const unsigned int cx = screenPos.x+u;
+ const unsigned int cy = screenPos.y+v;
+ if (d > params.camera.m_sensorDepthWorldMin && d < params.camera.m_sensorDepthWorldMax && cx < depth.width() && cy < depth.height()) {
+ // Transform estimated point to virtual cam space and output z
+ atomicMin(&depth(cx,cy), d * 1000.0f);
+ }
+ }
+}
+
+// ===== PASS 2 : Splat Visible Surface ========================================
+
+/*
+ * Pass 2: Determine depth buffer with enough accuracy for a visibility test in pass 2.
* These values are also used as the actual surface estimate during rendering so should
* at least be plane or sphere fitted if not MLS smoothed onto the actual surface.
*/
-__global__ void dibr_visibility_kernel(TextureObject<int> depth, int cam, SplatParams params) {
+__global__ void OLD_dibr_visibility_kernel(TextureObject<int> depth, int cam, SplatParams params) {
const ftl::voxhash::DepthCameraCUDA &camera = c_cameras[cam];
const int x = (blockIdx.x*blockDim.x + threadIdx.x) / WARP_SIZE;
@@ -158,13 +215,13 @@ __global__ void dibr_visibility_kernel(TextureObject<int> depth, int cam, SplatP
}
}
-// ------ Alternative for pass 1: principle surfaces ---------------------------
+// ------ Alternative for pass 2: principle surfaces ---------------------------
#define NEIGHBOR_RADIUS 1
#define MAX_NEIGHBORS ((NEIGHBOR_RADIUS*2+1)*(NEIGHBOR_RADIUS*2+1))
/*
- * Pass 1: Determine depth buffer with enough accuracy for a visibility test in pass 2.
+ * Pass 2: Determine depth buffer with enough accuracy for a visibility test in pass 2.
* These values are also used as the actual surface estimate during rendering so should
* at least be plane or sphere fitted if not MLS smoothed onto the actual surface.
*/
@@ -447,16 +504,20 @@ void ftl::cuda::dibr(const TextureObject<int> &depth_out,
cudaSafeCall(cudaDeviceSynchronize());
#endif
- int i=3;
- // Pass 1, merge a depth map from each camera.
+ int i=3;
+ // Pass 1, gather and upsample depth maps
for (int i=0; i<numcams; ++i)
- dibr_visibility_principal_kernel<<<sgridSize, sblockSize, 0, stream>>>(depth_out, i, params);
+ dibr_merge_upsample_kernel<<<sgridSize, sblockSize, 0, stream>>>(depth_out, i, params);
+
+ // Pass 2, merge a depth map from each camera.
+ //for (int i=0; i<numcams; ++i)
+ // dibr_visibility_principal_kernel<<<sgridSize, sblockSize, 0, stream>>>(depth_out, i, params);
- // Pass 2, accumulate all point contributions to pixels
+ // Pass 3, accumulate all point contributions to pixels
for (int i=0; i<numcams; ++i)
dibr_attribute_contrib_kernel<<<sgridSize, sblockSize, 0, stream>>>(depth_out, tmp_colour, normal_out, confidence_out, i, params);
- // Pass 3, normalise contributions
+ // Pass 4, normalise contributions
dibr_normalise_kernel<<<gridSize, blockSize, 0, stream>>>(tmp_colour, colour_out, normal_out, confidence_out);
cudaSafeCall( cudaGetLastError() );
diff --git a/applications/reconstruct/src/splat_params.hpp b/applications/reconstruct/src/splat_params.hpp
index cb2d3febda3364a3407264711620a25f9dc116a1..8ae5bf345e4e0d348c414cc1ce7bb52190d5ffff 100644
--- a/applications/reconstruct/src/splat_params.hpp
+++ b/applications/reconstruct/src/splat_params.hpp
@@ -8,8 +8,10 @@
namespace ftl {
namespace render {
-static const uint kShowBlockBorders = 0x0001;
-static const uint kNoSplatting = 0x0002;
+static const uint kShowBlockBorders = 0x00000001; // Deprecated: from voxels system
+static const uint kNoSplatting = 0x00000002;
+static const uint kNoUpsampling = 0x00000004;
+static const uint kNoTexturing = 0x00000008;
struct __align__(16) SplatParams {
float4x4 m_viewMatrix;