diff --git a/applications/reconstruct/src/dibr.cu b/applications/reconstruct/src/dibr.cu index 2d64ddd151f09a289f8eeef411b170b32c7fd469..bdd00de4412146a1979f2dbcd2d0f664d7db525a 100644 --- a/applications/reconstruct/src/dibr.cu +++ b/applications/reconstruct/src/dibr.cu @@ -518,6 +518,8 @@ __global__ void dibr_attribute_contrib_kernel( const ftl::voxhash::DepthCameraCUDA &camera = c_cameras[cam]; + const int tid = (threadIdx.x + threadIdx.y * blockDim.x); + const int warp = tid / WARP_SIZE; const int x = (blockIdx.x*blockDim.x + threadIdx.x) / WARP_SIZE; const int y = blockIdx.y*blockDim.y + threadIdx.y; @@ -546,7 +548,7 @@ __global__ void dibr_attribute_contrib_kernel( const float4 normal = tex2D<float4>(camera.normal, x, y); // Each thread in warp takes an upsample point and updates corresponding depth buffer. - const int lane = threadIdx.x % WARP_SIZE; + const int lane = tid % 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); @@ -559,7 +561,9 @@ __global__ void dibr_attribute_contrib_kernel( const float weight = ftl::cuda::spatialWeighting(length(nearest - camPos), SPATIAL_SMOOTHING); if (screenPos.x+u < colour_out.width() && screenPos.y+v < colour_out.height() && weight > 0.0f) { // TODO: Use confidence threshold here const float4 wcolour = colour * weight; - const float4 wnormal = normal * weight; + const float4 wnormal = normal * weight; + + //printf("Z %f\n", d); // Add this points contribution to the pixel buffer atomicAdd((float*)&colour_out(screenPos.x+u, screenPos.y+v), wcolour.x); diff --git a/applications/reconstruct/src/splat_render.cpp b/applications/reconstruct/src/splat_render.cpp index 039d4ab99a4de7add09797b1cf0b35d7775994dc..aad8fb818215f1925ff0f159fc85d2b70b1d127d 100644 --- a/applications/reconstruct/src/splat_render.cpp +++ b/applications/reconstruct/src/splat_render.cpp @@ -90,7 +90,7 @@ void Splatter::render(ftl::rgbd::Source *src, cudaStream_t stream) { if (src->value("splatting", false)) { //ftl::cuda::splat_points(depth1_, colour1_, normal1_, depth2_, colour2_, params, stream); ftl::cuda::int_to_float(depth1_, depth2_, 1.0f / 1000.0f, stream); - src->writeFrames(colour2_, depth2_, stream); + src->writeFrames(colour1_, depth2_, stream); } else { ftl::cuda::int_to_float(depth1_, depth2_, 1.0f / 1000.0f, stream); src->writeFrames(colour1_, depth2_, stream); @@ -108,7 +108,7 @@ void Splatter::render(ftl::rgbd::Source *src, cudaStream_t stream) { } else { if (src->value("splatting", false)) { //ftl::cuda::splat_points(depth1_, colour1_, normal1_, depth2_, colour2_, params, stream); - src->writeFrames(colour2_, depth2_, stream); + src->writeFrames(colour1_, depth2_, stream); } else { ftl::cuda::int_to_float(depth1_, depth2_, 1.0f / 1000.0f, stream); src->writeFrames(colour1_, depth2_, stream);