diff --git a/src/preprocess_kernels.cu b/src/preprocess_kernels.cu index 030c683..5228d19 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}; @@ -194,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];