From e9a16f08240f30f59ccc3fe1b1a8e7372e03960d Mon Sep 17 00:00:00 2001 From: thinker1065948474 Date: Wed, 21 Sep 2022 22:04:01 +0800 Subject: [PATCH 1/2] Fix memory access out of bounds --- src/preprocess_kernels.cu | 2 ++ 1 file changed, 2 insertions(+) diff --git a/src/preprocess_kernels.cu b/src/preprocess_kernels.cu index 030c683..cc2d6ba 100644 --- a/src/preprocess_kernels.cu +++ b/src/preprocess_kernels.cu @@ -146,6 +146,8 @@ __global__ void generateBaseFeatures_kernel(unsigned int *mask, float *voxels, unsigned int current_pillarId = 0; current_pillarId = atomicAdd(pillar_num, 1); + if( current_pillarId >= MAX_VOXELS ) return; + voxel_num[current_pillarId] = count; uint4 idx = {0, 0, voxel_idy, voxel_idx}; From 647a9a7709d97abd8d409e44ab39425a2f9789ca Mon Sep 17 00:00:00 2001 From: thinker1065948474 Date: Thu, 22 Sep 2022 21:50:41 +0800 Subject: [PATCH 2/2] Fixe memory access out of bounds --- src/preprocess_kernels.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/preprocess_kernels.cu b/src/preprocess_kernels.cu index cc2d6ba..5228d19 100644 --- a/src/preprocess_kernels.cu +++ b/src/preprocess_kernels.cu @@ -196,10 +196,10 @@ __global__ void generateFeatures_kernel(float* voxel_features, int pillar_idx = blockIdx.x * WARPS_PER_BLOCK + threadIdx.x/WARP_SIZE; int point_idx = threadIdx.x % WARP_SIZE; - int pillar_idx_inBlock = threadIdx.x/32; + int pillar_idx_inBlock = threadIdx.x/WARP_SIZE; unsigned int num_pillars = params[0]; - if (pillar_idx >= num_pillars) return; + if (pillar_idx >= num_pillars || pillar_idx >= MAX_VOXELS) return; __shared__ float4 pillarSM[WARPS_PER_BLOCK][WARP_SIZE]; __shared__ float4 pillarSumSM[WARPS_PER_BLOCK];