feat(autoware_lidar_centerpoint): process front voxels first (#9608)

* feat: optimize voxel indexing in preprocess_kernel.cu Signed-off-by: Taekjin LEE <[email protected]> * fix: remove redundant index check Signed-off-by: Taekjin LEE <[email protected]> * fix: modify voxel index for better memory access Signed-off-by: Taekjin LEE <[email protected]> --------- Signed-off-by: Taekjin LEE <[email protected]>
autowarefoundation · Dec 16, 2024 · 91984f8 · 91984f8
1 parent 5ced41e
commit 91984f8
Showing 1 changed file with 11 additions and 8 deletions.
diff --git 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<<<blocks, threads, 0, stream>>>(
     mask, voxels, grid_y_size, grid_x_size, max_voxel_size, pillar_num, voxel_features, voxel_num,