diff --git a/applications/reconstruct/src/garbage.cu b/applications/reconstruct/src/garbage.cu index beb7027b98162b2f7b79761f3f27d0fca81d006b..b6226e58826b4ba0ef594a87a8201562209df32d 100644 --- a/applications/reconstruct/src/garbage.cu +++ b/applications/reconstruct/src/garbage.cu @@ -7,27 +7,26 @@ using namespace ftl::voxhash; #define NUM_CUDA_BLOCKS 10000 __global__ void starveVoxelsKernel(HashData hashData) { + int ptr; // Stride over all allocated blocks for (int bi=blockIdx.x; bi<*hashData.d_hashCompactifiedCounter; bi+=NUM_CUDA_BLOCKS) { - const HashEntry& entry = hashData.d_hashCompactified[bi]; - - //is typically exectued only every n'th frame - int weight = hashData.d_SDFBlocks[entry.ptr + threadIdx.x].weight; + ptr = hashData.d_hashCompactified[bi].ptr; + int weight = hashData.d_SDFBlocks[ptr + threadIdx.x].weight; weight = max(0, weight-2); - hashData.d_SDFBlocks[entry.ptr + threadIdx.x].weight = weight; //CHECK Remove to totally clear previous frame (Nick) + hashData.d_SDFBlocks[ptr + threadIdx.x].weight = weight; //CHECK Remove to totally clear previous frame (Nick) } } -void ftl::cuda::starveVoxels(HashData& hashData, const HashParams& hashParams) { +void ftl::cuda::starveVoxels(HashData& hashData, const HashParams& hashParams, cudaStream_t stream) { const unsigned int threadsPerBlock = SDF_BLOCK_SIZE*SDF_BLOCK_SIZE*SDF_BLOCK_SIZE; const dim3 gridSize(NUM_CUDA_BLOCKS, 1); const dim3 blockSize(threadsPerBlock, 1); //if (hashParams.m_numOccupiedBlocks > 0) { - starveVoxelsKernel << <gridSize, blockSize >> >(hashData); + starveVoxelsKernel << <gridSize, blockSize, 0, stream >> >(hashData); //} #ifdef _DEBUG cudaSafeCall(cudaDeviceSynchronize()); diff --git a/applications/reconstruct/src/garbage.hpp b/applications/reconstruct/src/garbage.hpp index 477db1ae831e60568902475b6afeeeda3b959563..5d1d7574d252b40da18008da39f1bf89a7d667fb 100644 --- a/applications/reconstruct/src/garbage.hpp +++ b/applications/reconstruct/src/garbage.hpp @@ -5,7 +5,7 @@ namespace ftl { namespace cuda { void clearVoxels(ftl::voxhash::HashData& hashData, const ftl::voxhash::HashParams& hashParams); -void starveVoxels(ftl::voxhash::HashData& hashData, const ftl::voxhash::HashParams& hashParams); +void starveVoxels(ftl::voxhash::HashData& hashData, const ftl::voxhash::HashParams& hashParams, cudaStream_t stream); void garbageCollectIdentify(ftl::voxhash::HashData& hashData, const ftl::voxhash::HashParams& hashParams, cudaStream_t stream); void garbageCollectFree(ftl::voxhash::HashData& hashData, const ftl::voxhash::HashParams& hashParams, cudaStream_t stream); diff --git a/applications/reconstruct/src/voxel_scene.cpp b/applications/reconstruct/src/voxel_scene.cpp index a816804905e4e29fec81d3fdd5e7b89ee32eed99..e274f068c6d80cf0cc9d7149291016181e377608 100644 --- a/applications/reconstruct/src/voxel_scene.cpp +++ b/applications/reconstruct/src/voxel_scene.cpp @@ -238,7 +238,7 @@ void SceneRep::nextFrame() { } 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); + else ftl::cuda::starveVoxels(m_hashData, m_hashParams, integ_stream_); m_numIntegratedFrames = 0; } }