Skip to content
Snippets Groups Projects
Commit 12363ff0 authored by Nicolas Pope's avatar Nicolas Pope
Browse files

Correct bug in garbage collect kernel

parent 4bcb6e02
No related branches found
No related tags found
1 merge request!57Resolves #40 by colouring registration errors
This commit is part of merge request !57. Comments created here will be created in the context of that merge request.
...@@ -137,7 +137,7 @@ void ftl::cuda::garbageCollectIdentify(HashData& hashData, const HashParams& has ...@@ -137,7 +137,7 @@ void ftl::cuda::garbageCollectIdentify(HashData& hashData, const HashParams& has
__global__ void garbageCollectFreeKernel(HashData hashData) { __global__ void garbageCollectFreeKernel(HashData hashData) {
// Stride over all allocated blocks // 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 if (hashData.d_hashDecision[bi] != 0) { //decision to delete the hash entry
......
...@@ -245,7 +245,7 @@ __global__ void integrateRegistrationKernel(HashData hashData, HashParams hashPa ...@@ -245,7 +245,7 @@ __global__ void integrateRegistrationKernel(HashData hashData, HashParams hashPa
out.color.x = v1.color.x*redshift; out.color.x = v1.color.x*redshift;
out.color.y = v1.color.y*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.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); out.weight = min(c_hashParams.m_integrationWeightMax, (unsigned int)v0.weight + (unsigned int)v1.weight);
......
...@@ -149,7 +149,7 @@ void SceneRep::integrate() { ...@@ -149,7 +149,7 @@ void SceneRep::integrate() {
//volumetrically integrate the depth data into the depth SDFBlocks //volumetrically integrate the depth data into the depth SDFBlocks
_integrateDepthMap(cam.gpu, cam.params); _integrateDepthMap(cam.gpu, cam.params);
//_garbageCollect(cam.gpu); //_garbageCollect();
m_numIntegratedFrames++; m_numIntegratedFrames++;
} }
...@@ -205,6 +205,7 @@ void SceneRep::nextFrame() { ...@@ -205,6 +205,7 @@ void SceneRep::nextFrame() {
_destroy(); _destroy();
_create(_parametersFromConfig()); _create(_parametersFromConfig());
} else { } else {
//ftl::cuda::compactifyAllocated(m_hashData, m_hashParams, integ_stream_);
if (reg_mode_) ftl::cuda::clearVoxels(m_hashData, m_hashParams); if (reg_mode_) ftl::cuda::clearVoxels(m_hashData, m_hashParams);
else ftl::cuda::starveVoxels(m_hashData, m_hashParams); else ftl::cuda::starveVoxels(m_hashData, m_hashParams);
m_numIntegratedFrames = 0; m_numIntegratedFrames = 0;
......
0% Loading or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment