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!