From 9de0033a11562470e9eb574fb41cd473cbabb514 Mon Sep 17 00:00:00 2001 From: Taekjin LEE Date: Thu, 26 Dec 2024 09:43:12 +0900 Subject: [PATCH] Revert "Merge pull request #1688 from technolojin/feat/centerpoint/voxel-fill-from-front" This reverts commit 3326ba5dae3d9f22f78cbe9c6891e595ea31de25, reversing changes made to 16569fc00f58ed0fa45bee7c0bb6d66b3872c66f. --- .../lib/preprocess/preprocess_kernel.cu | 16 +++++++--------- 1 file changed, 7 insertions(+), 9 deletions(-) diff --git a/perception/lidar_centerpoint/lib/preprocess/preprocess_kernel.cu b/perception/lidar_centerpoint/lib/preprocess/preprocess_kernel.cu index b032835924c6f..01f62f22e892b 100644 --- a/perception/lidar_centerpoint/lib/preprocess/preprocess_kernel.cu +++ b/perception/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 = (grid_x_size - 1 - voxel_idx) * grid_y_size + voxel_idy; + unsigned int voxel_index = voxel_idy * grid_x_size + voxel_idx; unsigned int point_id = atomicAdd(&(mask[voxel_index]), 1); @@ -185,13 +185,12 @@ __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) { - // flip x and y to process in a row-major order - 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_idx = blockIdx.x * blockDim.x + threadIdx.x; + unsigned int voxel_idy = blockIdx.y * blockDim.y + threadIdx.y; - unsigned int voxel_index = voxel_idx_inverted * grid_y_size + voxel_idy; + 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 < MAX_POINT_IN_VOXEL_SIZE ? count : MAX_POINT_IN_VOXEL_SIZE; @@ -221,10 +220,9 @@ cudaError_t generateBaseFeatures_launch( unsigned int * pillar_num, float * voxel_features, float * voxel_num, int * voxel_idxs, cudaStream_t stream) { - // flip x and y to process in a row-major order dim3 threads = {32, 32}; dim3 blocks = { - (grid_y_size + threads.x - 1) / threads.x, (grid_x_size + threads.y - 1) / threads.y}; + (grid_x_size + threads.x - 1) / threads.x, (grid_y_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,