From 0f1dacb5ce1dfc12051453eb01e161b9e4c82a23 Mon Sep 17 00:00:00 2001
From: Nicolas Pope <nwpope@utu.fi>
Date: Sun, 30 Jun 2019 09:35:39 +0300
Subject: [PATCH] Add stream to starve voxels
---
applications/reconstruct/src/garbage.cu | 13 ++++++-------
applications/reconstruct/src/garbage.hpp | 2 +-
applications/reconstruct/src/voxel_scene.cpp | 2 +-
3 files changed, 8 insertions(+), 9 deletions(-)
diff --git a/applications/reconstruct/src/garbage.cu b/applications/reconstruct/src/garbage.cu
index beb7027b9..b6226e588 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 477db1ae8..5d1d7574d 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 a81680490..e274f068c 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;
}
}
--
GitLab