From 91984f819fdbef89e7c7e4c0785857c9b6b2ee70 Mon Sep 17 00:00:00 2001 From: Taekjin LEE Date: Mon, 16 Dec 2024 09:04:53 +0900 Subject: [PATCH] feat(autoware_lidar_centerpoint): process front voxels first (#9608) * feat: optimize voxel indexing in preprocess_kernel.cu Signed-off-by: Taekjin LEE * fix: remove redundant index check Signed-off-by: Taekjin LEE * fix: modify voxel index for better memory access Signed-off-by: Taekjin LEE --------- Signed-off-by: Taekjin LEE --- .../lib/preprocess/preprocess_kernel.cu | 19 +++++++++++-------- 1 file changed, 11 insertions(+), 8 deletions(-) diff --git a/perception/autoware_lidar_centerpoint/lib/preprocess/preprocess_kernel.cu b/perception/autoware_lidar_centerpoint/lib/preprocess/preprocess_kernel.cu index 502ad04223ce9..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); @@ -185,12 +185,14 @@ __global__ void generateBaseFeatures_kernel( unsigned int * mask, float * voxels, int grid_y_size, int grid_x_size, int max_voxel_size, unsigned int * pillar_num, float * voxel_features, float * voxel_num, int * voxel_idxs) { - unsigned int voxel_idx = blockIdx.x * blockDim.x + threadIdx.x; - unsigned 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; + // exchange x and y to process in a row-major order + // flip x axis direction to process front to back + unsigned int voxel_idx_inverted = blockIdx.y * blockDim.y + threadIdx.y; + unsigned int voxel_idy = blockIdx.x * blockDim.x + threadIdx.x; + 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_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; @@ -220,9 +222,10 @@ cudaError_t generateBaseFeatures_launch( unsigned int * pillar_num, float * voxel_features, float * voxel_num, int * voxel_idxs, cudaStream_t stream) { + // exchange x and y to process in a row-major order dim3 threads = {32, 32}; dim3 blocks = { - (grid_x_size + threads.x - 1) / threads.x, (grid_y_size + threads.y - 1) / threads.y}; + (grid_y_size + threads.x - 1) / threads.x, (grid_x_size + threads.y - 1) / threads.y}; generateBaseFeatures_kernel<<>>( mask, voxels, grid_y_size, grid_x_size, max_voxel_size, pillar_num, voxel_features, voxel_num,