Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

revert(autoware_lidar_centerpoint): process front voxels first #1743

Merged
merged 1 commit into from
Dec 26, 2024
Merged

revert(autoware_lidar_centerpoint): process front voxels first #1743

Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
16 changes: 7 additions & 9 deletions perception/lidar_centerpoint/lib/preprocess/preprocess_kernel.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);

Expand Down Expand Up @@ -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;
Expand Down Expand Up @@ -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<<<blocks, threads, 0, stream>>>(
mask, voxels, grid_y_size, grid_x_size, max_voxel_size, pillar_num, voxel_features, voxel_num,
Expand Down
Loading