autoware_cuda_pointcloud_preprocessor

Purpose

The pointcloud preprocessing implemented in autoware_pointcloud_preprocessor has been thoroughly tested in autoware. However, the latency it introduces does not scale well with modern LiDAR devices due to the high number of points they introduce.

To alleviate this issue, this package reimplements most of the pipeline presented in autoware_pointcloud_preprocessor leveraging the use of GPGPUs. In particular, this package makes use of CUDA to provide accelerated versions of the already established implementations, while also maintaining compatibility with normal ROS nodes/topics.

Inner-workings / Algorithms

A detailed description of each filter’s algorithm is available in the following links.

(Optional) Future extensions / Unimplemented parts

The subsample filters implemented in autoware_pointcloud_preprocessor will have similar counterparts in this package.

Changelog for package autoware_cuda_pointcloud_preprocessor

0.53.0 (2026-09-29)

• Merge remote-tracking branch 'origin/main' into prepare-0.53.0-changelog

• fix(cuda_pointcloud_preprocessor): link CUDA utilities in test (#13435) The test compiles CUDA utility headers with the host compiler. Link the exported target so CUDA 13 CCCL headers are available.

• perf(autoware_cuda_pointcloud_preprocessor): process points in ring order instead of an organized layout (#13412)

• test(cuda_pointcloud_preprocessor): cover capacity bounds behavior (#13137)

  * perf(cuda_pointcloud_preprocessor): avoid runtime thrust allocations Thrust's algorithm calls allocate temporary device memory on every frame. Replace them with explicit kernels and CUB calls running on storage that is allocated once at startup:

  • thrust_stream::fill / fill_n become a fill kernel launched on the stream
  • thrust::inclusive_scan becomes cub::DeviceScan::InclusiveSum

  - thrust_stream::count becomes cub::DeviceReduce::Sum over a transform iterator, accumulating the three diagnostic counters into one device buffer that is copied back in a single transfer The CUB calls share a single scratch workspace, sized at startup to the largest requirement among the sort, the scan and the reductions. The twist struct counts become locals, since they are only used within process(). The counting iterator is thrust::transform_iterator rather than cub::TransformInputIterator, since CCCL 3.0 (CUDA 13, used on jazzy) removed the latter along with its header.

  * refactor(cuda_pointcloud_preprocessor): name the processing stat indices Hoist the device_processing_stats_ layout into class-scope constants so that the buffer sizing in initializeBuffers() and the readback in process() share one definition instead of repeating a bare 3. Co-Authored-By: Claude Opus 5 (1M context) <<noreply@anthropic.com>>

  * refactor(cuda_pointcloud_preprocessor): clarify queue update flow The twist and IMU queue bookkeeping was spread over the insertion callbacks and a pair of bounding helpers, so capacity was enforced after insertion and the pruning rules were duplicated per queue. Move the bookkeeping into detail:: helpers in a new queue_bounds.hpp and give each queue one update path:

  - prepare_queue_update() prunes queue entries older than the current pointcloud, drops and counts incoming messages beyond the free capacity, and sorts the remainder, so capacity is enforced before insertion

  - twistCallback()/imuCallback() become insertTwistMessage()/ insertImuMessage(), which now only insert in stamp order

  - boundTwistQueue()/boundImuQueue() are gone, since nothing is inserted beyond capacity any more Backward time jump handling moves to detail::is_backward_time_jump() and now tolerates jumps of up to one second: the queue is cleared only when its oldest entry is more than one second ahead of the incoming stamp, rather than on any backward step. The old per-entry pop for queues spanning more than a second is dropped, since prepare_queue_update() already prunes everything older than the current pointcloud.

  • refactor(cuda_pointcloud_preprocessor): address queue_bounds review comments

  - Convert builtin_interfaces/Time to nanoseconds from its own fields, which drops the rclcpp dependency from queue_bounds.hpp entirely. A negative [sec]{.title-ref} throws rather than wrapping, which is what rclcpp::Time did too

  - Document and assert that prune_old_queue_entries() takes a stamp-sorted queue and leaves it sorted

  - Document what prepare_queue_update() does and the state it leaves its two containers in Co-Authored-By: Claude Opus 5 (1M context) <<noreply@anthropic.com>>

  * test(cuda_pointcloud_preprocessor): cover capacity bounds behavior Add a gtest suite for the behavior introduced by the preceding commits in this stack:

  - CudaPointcloudPreprocessor accepts positive capacities and rejects non-positive ones

  - points exactly at the ring and points-per-ring capacities are processed without reporting an overflow, and exceeding either reports one

  • input clouds longer than max_input_point_count are truncated

  - detail::prepare_queue_update() prunes, sorts, bounds and counts drops for both the twist and IMU queues, and rejects an already over-capacity queue

File truncated at 100 lines see the full file