From 12363ff0ac9003800213cc4482da946ee5b8cf9b Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Sat, 29 Jun 2019 10:35:29 +0300
Subject: [PATCH] Correct bug in garbage collect kernel
---
applications/reconstruct/src/garbage.cu | 2 +-
applications/reconstruct/src/integrators.cu | 2 +-
applications/reconstruct/src/voxel_scene.cpp | 3 ++-
3 files changed, 4 insertions(+), 3 deletions(-)
diff --git a/applications/reconstruct/src/garbage.cu b/applications/reconstruct/src/garbage.cu
index fc02c414f..beb7027b9 100644
--- a/applications/reconstruct/src/garbage.cu
+++ b/applications/reconstruct/src/garbage.cu
@@ -137,7 +137,7 @@ void ftl::cuda::garbageCollectIdentify(HashData& hashData, const HashParams& has
__global__ void garbageCollectFreeKernel(HashData hashData) {
// Stride over all allocated blocks
- for (int bi=blockIdx.x*blockDim.x + threadIdx.x; bi<*hashData.d_hashCompactifiedCounter; bi+=NUM_CUDA_BLOCKS) {
+ for (int bi=blockIdx.x*blockDim.x + threadIdx.x; bi<*hashData.d_hashCompactifiedCounter; bi+=NUM_CUDA_BLOCKS*blockDim.x) {
if (hashData.d_hashDecision[bi] != 0) { //decision to delete the hash entry
diff --git a/applications/reconstruct/src/integrators.cu b/applications/reconstruct/src/integrators.cu
index 76f132c08..35bb6f6ac 100644
--- a/applications/reconstruct/src/integrators.cu
+++ b/applications/reconstruct/src/integrators.cu
@@ -245,7 +245,7 @@ __global__ void integrateRegistrationKernel(HashData hashData, HashParams hashPa
out.color.x = v1.color.x*redshift;
out.color.y = v1.color.y*redshift;
- out.color.z = v1.color.z*(1.0f / redshift);
+ out.color.z = v1.color.z; //*(1.0f / redshift);
out.sdf = (v0.sdf * (float)v0.weight + v1.sdf * (float)v1.weight) / ((float)v0.weight + (float)v1.weight);
out.weight = min(c_hashParams.m_integrationWeightMax, (unsigned int)v0.weight + (unsigned int)v1.weight);
diff --git a/applications/reconstruct/src/voxel_scene.cpp b/applications/reconstruct/src/voxel_scene.cpp
index 8705d660e..ccdee9e7d 100644
--- a/applications/reconstruct/src/voxel_scene.cpp
+++ b/applications/reconstruct/src/voxel_scene.cpp
@@ -149,7 +149,7 @@ void SceneRep::integrate() {
//volumetrically integrate the depth data into the depth SDFBlocks
_integrateDepthMap(cam.gpu, cam.params);
- //_garbageCollect(cam.gpu);
+ //_garbageCollect();
m_numIntegratedFrames++;
}
@@ -205,6 +205,7 @@ void SceneRep::nextFrame() {
_destroy();
_create(_parametersFromConfig());
} else {
+ //ftl::cuda::compactifyAllocated(m_hashData, m_hashParams, integ_stream_);
if (reg_mode_) ftl::cuda::clearVoxels(m_hashData, m_hashParams);
else ftl::cuda::starveVoxels(m_hashData, m_hashParams);
m_numIntegratedFrames = 0;
--
GitLab