From 167b69d3cdb678b1fa7ae7aedf96da9b33f0848d Mon Sep 17 00:00:00 2001 From: Taekjin LEE Date: Wed, 11 Dec 2024 17:49:10 +0900 Subject: [PATCH] fix: modify voxel index for better memory access Signed-off-by: Taekjin LEE --- .../lib/preprocess/preprocess_kernel.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/perception/autoware_lidar_centerpoint/lib/preprocess/preprocess_kernel.cu b/perception/autoware_lidar_centerpoint/lib/preprocess/preprocess_kernel.cu index 2d1f77fe6dac6..f300411a44aad 100644 --- a/perception/autoware_lidar_centerpoint/lib/preprocess/preprocess_kernel.cu +++ b/perception/autoware_lidar_centerpoint/lib/preprocess/preprocess_kernel.cu @@ -148,7 +148,7 @@ __global__ void generateVoxels_random_kernel( int voxel_idx = floorf((point.x - min_x_range) / pillar_x_size); int voxel_idy = floorf((point.y - min_y_range) / pillar_y_size); - unsigned int voxel_index = voxel_idy * grid_x_size + voxel_idx; + unsigned int voxel_index = (grid_x_size - 1 - voxel_idx) * grid_y_size + voxel_idy; unsigned int point_id = atomicAdd(&(mask[voxel_index]), 1); @@ -192,7 +192,7 @@ __global__ void generateBaseFeatures_kernel( if (voxel_idx_inverted >= grid_x_size || voxel_idy >= grid_y_size) return; unsigned int voxel_idx = grid_x_size - 1 - voxel_idx_inverted; - unsigned int voxel_index = voxel_idy * grid_x_size + voxel_idx; + unsigned int voxel_index = voxel_idx_inverted * grid_y_size + voxel_idy; unsigned int count = mask[voxel_index]; if (!(count > 0)) return; count = count < MAX_POINT_IN_VOXEL_SIZE ? count : MAX_POINT_IN_VOXEL_SIZE;