In the preprocessing module, why calculate the pillar index value twice？

[leidahhh]

global void generateBaseFeatures_kernel(unsigned int mask, float voxels, int grid_y_size, int grid_x_size, unsigned int pillar_num, float voxel_features, float voxel_num_points, float coords) { int voxel_idx = blockIdx.x blockDim.x + threadIdx.x; int voxel_idy = blockIdx.y blockDim.y + threadIdx.y;

if(voxel_idx >= grid_x_size ||voxel_idy >= grid_y_size) return;

unsigned int voxel_index = voxel_idy * grid_x_size

• voxel_idx; unsigned int count = mask[voxel_index]; if( !(count>0) ) return; count = count<POINTS_PER_VOXEL?count:POINTS_PER_VOXEL;

  unsigned int current_pillarId = 0; current_pillarId = atomicAdd(pillar_num+4, 1);

  voxel_num_points[current_pillarId] = count;

  float4 coord = {0.0, 0, (float)voxel_idy, (float)voxel_idx}; ((float4*)coords)[current_pillarId] = coord;

  for (int i=0; i<count; i++){ int inIndex = voxel_indexPOINTS_PER_VOXEL + i; int outIndex = current_pillarIdPOINTS_PER_VOXEL + i; ((float4)voxel_features)[outIndex] = ((float4)voxels)[inIndex]; }

  // clear buffer for next infer //mask[voxel_index] = 0; atomicExch(mask + voxel_index, 0); }

In the code block, the voxel_idx is calculated once according to the thread index value, why is it calculated again later current_pillarId and is there any difference between the two index values. I would very appreciate if could reply my confusions!