From de27e64c5803e0a45cae5278ffebc03adcad89f1 Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Sun, 11 Aug 2019 09:23:23 +0300
Subject: [PATCH] Fix incorrect colour map when splatting
---
applications/reconstruct/src/dibr.cu | 8 ++++++--
applications/reconstruct/src/splat_render.cpp | 4 ++--
2 files changed, 8 insertions(+), 4 deletions(-)
diff --git a/applications/reconstruct/src/dibr.cu b/applications/reconstruct/src/dibr.cu
index 2d64ddd15..bdd00de44 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 039d4ab99..aad8fb818 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);
--
GitLab