fix(hesai): address review comments from WIP PR

- Copyright year 2024 -> 2026 for new files
- Replace deprecated find_package(CUDA) with find_package(CUDAToolkit)
- Remove --expt-relaxed-constexpr flag (not needed)
- Remove unused per-packet kernel and launcher (dead code)
- Batch launcher returns bool; caller logs via NEBULA_LOG_STREAM
- Reorder CudaNebulaPoint fields for better memory packing
- Remove redundant is_multi_frame member; use n_frames > 1
- Make HesaiCudaDecoder destructor virtual
- Add int32_t range guarantee comment in angle corrector

Author: k1832

src/nebula_hesai/nebula_hesai/tests/hesai_cuda_decoder_test.cpp

@@ -1,4 +1,4 @@
-// Copyright 2024 TIER IV, Inc.
+// Copyright 2026 TIER IV, Inc.
 //
 // Licensed under the Apache License, Version 2.0 (the "License");
 // you may not use this file except in compliance with the License.

src/nebula_hesai/nebula_hesai_decoders/CMakeLists.txt

@@ -9,17 +9,12 @@ option(BUILD_CUDA "Build with CUDA support" ON)
 
 # CUDA Configuration
 if(BUILD_CUDA)
-  find_package(CUDA)
-  if(NOT ${CUDA_FOUND})
-    message(WARNING "CUDA not found, building without CUDA support")
+  find_package(CUDAToolkit QUIET)
+  if(NOT CUDAToolkit_FOUND)
+    message(WARNING "CUDA Toolkit not found, building without CUDA support")
     set(BUILD_CUDA OFF)
   else()
-    message(STATUS "CUDA found: ${CUDA_VERSION}")
-    message(STATUS "CUDA include dirs: ${CUDA_INCLUDE_DIRS}")
-
-    # CUDA compiler flags
-    list(APPEND CUDA_NVCC_FLAGS "--expt-relaxed-constexpr")
-    list(APPEND CUDA_NVCC_FLAGS "-diag-suppress 20012")
+    message(STATUS "CUDA Toolkit found: ${CUDAToolkit_VERSION}")
 
     # GPU architectures — set via -DCMAKE_CUDA_ARCHITECTURES on the command line,
     # or fall back to a default set covering Turing through Ada + PTX for forward
@@ -28,14 +23,6 @@ if(BUILD_CUDA)
     if(NOT DEFINED CMAKE_CUDA_ARCHITECTURES OR CMAKE_CUDA_ARCHITECTURES STREQUAL "")
       set(CMAKE_CUDA_ARCHITECTURES "75;86;87;89" CACHE STRING "CUDA architectures" FORCE)
     endif()
-    # Also emit PTX for the highest arch so the driver can JIT for newer GPUs
-    list(APPEND CUDA_NVCC_FLAGS "-gencode arch=compute_89,code=compute_89")
-
-    # Extended lambda support
-    list(APPEND CUDA_NVCC_FLAGS "--extended-lambda")
-
-    # Optimization flags
-    list(APPEND CUDA_NVCC_FLAGS "-O3")
 
     # Check if CUDA kernel source exists
     if(EXISTS "${CMAKE_CURRENT_SOURCE_DIR}/src/cuda/hesai_cuda_kernels.cu")
@@ -67,7 +54,7 @@ target_link_libraries(
 if(BUILD_CUDA AND CUDA_KERNELS_EXIST)
   enable_language(CUDA)
 
-  set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -O3 --expt-relaxed-constexpr")
+  set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -O3 -diag-suppress 20012")
 
   add_library(nebula_hesai_decoders_cuda SHARED
     src/cuda/hesai_cuda_kernels.cu
@@ -90,17 +77,20 @@ if(BUILD_CUDA AND CUDA_KERNELS_EXIST)
            $<INSTALL_INTERFACE:include>)
 
   target_link_libraries(nebula_hesai_decoders_cuda
-    PRIVATE cudart)
+    PRIVATE CUDA::cudart_static)
 
   target_compile_definitions(nebula_hesai_decoders_cuda PUBLIC NEBULA_CUDA_ENABLED)
   target_compile_definitions(nebula_hesai_decoders PUBLIC NEBULA_CUDA_ENABLED)
 
   # Add CUDA include directories to main library for header resolution
   target_include_directories(nebula_hesai_decoders
-    PUBLIC ${CUDA_INCLUDE_DIRS})
+    PUBLIC ${CUDAToolkit_INCLUDE_DIRS})
 
-  # Link CUDA library to main library
-  target_link_libraries(nebula_hesai_decoders PUBLIC nebula_hesai_decoders_cuda ${CUDA_LIBRARIES})
+  # Link CUDA library to main library (CUDA::cudart is PRIVATE so it won't
+  # propagate the imported target to downstream packages that lack CUDAToolkit)
+  target_link_libraries(nebula_hesai_decoders
+    PUBLIC nebula_hesai_decoders_cuda
+    PRIVATE CUDA::cudart)
 
   install(TARGETS nebula_hesai_decoders_cuda EXPORT export_nebula_hesai_decoders_cuda)
   ament_export_targets(export_nebula_hesai_decoders_cuda)

src/nebula_hesai/nebula_hesai_decoders/include/nebula_hesai_decoders/cuda/hesai_cuda_decoder.hpp

@@ -1,4 +1,4 @@
-// Copyright 2024 TIER IV, Inc.
+// Copyright 2026 TIER IV, Inc.
 //
 // Licensed under the Apache License, Version 2.0 (the "License");
 // you may not use this file except in compliance with the License.
@@ -32,8 +32,8 @@ struct CudaNebulaPoint
   float elevation;
   float intensity;
   uint8_t return_type;
-  uint16_t channel;
   uint8_t in_current_scan;  // 1 = belongs to current scan, 0 = belongs to output/next scan
+  uint16_t channel;
   uint32_t entry_id;        // Block group ID for batched processing (used for sorting & filtering)
 };
 
@@ -89,15 +89,14 @@ struct CudaDecoderConfig
   // Multi-frame support for sensors like AT128 (has 4 mirror frames)
   uint32_t n_frames;  // Number of frames (1 for single-frame, 4 for AT128)
   CudaFrameAngleInfo frame_angles[MAX_CUDA_FRAMES];  // Per-frame angle boundaries
-  bool is_multi_frame;  // True if sensor has multiple frames (uses frame_angles)
 };
 
 /// @brief Main CUDA decoder class for Hesai LiDAR
 class HesaiCudaDecoder
 {
 public:
   HesaiCudaDecoder();
-  ~HesaiCudaDecoder();
+  virtual ~HesaiCudaDecoder();
 
   /// @brief Initialize decoder with maximum points and channels
   bool initialize(size_t max_points, uint32_t n_channels);
@@ -126,20 +125,9 @@ class HesaiCudaDecoder
 
 extern "C" {
 
-/// @brief Launch kernel to decode a single Hesai packet
-void launch_decode_hesai_packet(
-    const uint16_t* d_distances,
-    const uint8_t* d_reflectivities,
-    const nebula::drivers::cuda::CudaAngleCorrectionData* d_angle_lut,
-    const nebula::drivers::cuda::CudaDecoderConfig& config,
-    nebula::drivers::cuda::CudaNebulaPoint* d_points,
-    uint32_t* d_count,
-    uint32_t n_azimuths,
-    uint32_t raw_azimuth,
-    cudaStream_t stream);
-
 /// @brief Launch batched kernel to decode entire scan
-void launch_decode_hesai_scan_batch(
+/// @return true on success, false on CUDA error
+bool launch_decode_hesai_scan_batch(
     const uint16_t* d_distances_batch,
     const uint8_t* d_reflectivities_batch,
     const uint32_t* d_raw_azimuths,

src/nebula_hesai/nebula_hesai_decoders/include/nebula_hesai_decoders/decoders/angle_corrector_calibration_based.hpp

@@ -155,7 +155,9 @@ class AngleCorrectorCalibrationBased
   [[nodiscard]] std::tuple<uint32_t, uint32_t, uint32_t, uint32_t>
   get_cuda_raw_angles(double fov_start_deg, double fov_end_deg, double cut_angle_deg) const
   {
-    // Find min/max azimuth offset in raw units
+    // Find min/max azimuth offset in raw units.
+    // All values are guaranteed to fit in int32_t: angles are in [0, 360) degrees and
+    // AngleUnit is at most 100 (0.01-degree resolution), so raw values stay within ±36000.
     auto round_away_from_zero = [](float value) -> int32_t {
       return (value < 0) ? static_cast<int32_t>(std::floor(value))
                          : static_cast<int32_t>(std::ceil(value));

src/nebula_hesai/nebula_hesai_decoders/include/nebula_hesai_decoders/decoders/hesai_decoder.hpp

@@ -26,19 +26,8 @@
 #ifdef NEBULA_CUDA_ENABLED
 #include "nebula_hesai_decoders/cuda/hesai_cuda_decoder.hpp"
 
-// C-linkage kernel launcher declarations
-extern "C" void launch_decode_hesai_packet(
-  const uint16_t * d_distances,
-  const uint8_t * d_reflectivities,
-  const nebula::drivers::cuda::CudaAngleCorrectionData * d_angle_lut,
-  const nebula::drivers::cuda::CudaDecoderConfig & config,
-  nebula::drivers::cuda::CudaNebulaPoint * d_points,
-  uint32_t * d_count,
-  uint32_t n_azimuths,
-  uint32_t raw_azimuth,
-  cudaStream_t stream);
-
-extern "C" void launch_decode_hesai_scan_batch(
+// C-linkage kernel launcher declaration (defined in hesai_cuda_kernels.cu)
+extern "C" bool launch_decode_hesai_scan_batch(
   const uint16_t * d_distances_batch,
   const uint8_t * d_reflectivities_batch,
   const uint32_t * d_raw_azimuths,
@@ -328,7 +317,7 @@ class HesaiDecoder : public HesaiScanDecoder
     config.max_output_points = n_entries * n_channels * max_returns;
 
     if constexpr (SensorT::uses_calibration_based_angles) {
-      config.is_multi_frame = false;
+      // Single-frame sensor (calibration-based)
       config.n_frames = 1;
       config.timestamp_reset_angle_raw = cuda_timestamp_reset_angle_raw_;
       config.emit_angle_raw = cuda_emit_angle_raw_;
@@ -337,7 +326,7 @@ class HesaiDecoder : public HesaiScanDecoder
       config.frame_angles[0].timestamp_reset = cuda_timestamp_reset_angle_raw_;
       config.frame_angles[0].scan_emit = cuda_emit_angle_raw_;
     } else {
-      config.is_multi_frame = true;
+      // Multi-frame sensor (correction-based, e.g. AT128)
       config.n_frames = static_cast<uint32_t>(angle_corrector_.get_n_frames());
       config.timestamp_reset_angle_raw = 0;
       config.emit_angle_raw = 0;
@@ -465,11 +454,14 @@ class HesaiDecoder : public HesaiScanDecoder
       d_points_, 0, sparse_buffer_size * sizeof(cuda::CudaNebulaPoint), cuda_stream_);
 
     // Launch batched kernel
-    launch_decode_hesai_scan_batch(
+    bool kernel_ok = launch_decode_hesai_scan_batch(
       gpu_scan_buffer_.d_distances_batch, gpu_scan_buffer_.d_reflectivities_batch,
       gpu_scan_buffer_.d_raw_azimuths, gpu_scan_buffer_.d_n_returns,
       gpu_scan_buffer_.d_last_azimuths, cuda_decoder_->get_angle_lut(), config, d_points_,
       d_count_, cuda_n_azimuths_, n_entries, cuda_stream_);
+    if (!kernel_ok) {
+      NEBULA_LOG_STREAM(logger_->error, "CUDA batched kernel launch failed");
+    }
 
     cudaStreamSynchronize(cuda_stream_);
 

src/nebula_hesai/nebula_hesai_decoders/src/cuda/hesai_cuda_kernels.cu

@@ -1,4 +1,4 @@
-// Copyright 2024 TIER IV, Inc.
+// Copyright 2026 TIER IV, Inc.
 //
 // Licensed under the Apache License, Version 2.0 (the "License");
 // you may not use this file except in compliance with the License.
@@ -71,68 +71,6 @@ __device__ __forceinline__ bool cuda_is_inside_overlap_multiframe(
   return false;
 }
 
-// CUDA kernel for decoding a single Hesai LiDAR packet
-__global__ void decode_hesai_packet_kernel(
-  const uint16_t * __restrict__ distances,
-  const uint8_t * __restrict__ reflectivities,
-  const CudaAngleCorrectionData * __restrict__ angle_lut,
-  const CudaDecoderConfig config,
-  CudaNebulaPoint * __restrict__ output_points,
-  uint32_t * __restrict__ output_count,
-  uint32_t n_azimuths,
-  uint32_t raw_azimuth)
-{
-  const uint32_t channel_id = blockIdx.x * blockDim.x + threadIdx.x;
-  const uint32_t block_id = blockIdx.y;
-
-  if (channel_id >= config.n_channels || block_id >= config.n_blocks) {
-    return;
-  }
-
-  const uint32_t data_stride = config.data_stride > 0 ? config.data_stride : config.n_blocks;
-  const uint32_t data_idx = channel_id * data_stride + block_id;
-
-  const uint16_t raw_distance = distances[data_idx];
-  const uint8_t reflectivity = reflectivities[data_idx];
-
-  if (raw_distance == 0) {
-    return;
-  }
-
-  const float distance = static_cast<float>(raw_distance) * config.dis_unit;
-
-  if (distance < config.min_range || distance > config.max_range) {
-    return;
-  }
-
-  if (distance < config.sensor_min_range || distance > config.sensor_max_range) {
-    return;
-  }
-
-  const uint32_t azimuth_idx = (raw_azimuth / config.azimuth_scale) % n_azimuths;
-  const uint32_t lut_idx = azimuth_idx * config.n_channels + channel_id;
-  const CudaAngleCorrectionData angle_data = angle_lut[lut_idx];
-
-  const float xy_distance = distance * angle_data.cos_elevation;
-  const float x = xy_distance * angle_data.sin_azimuth;
-  const float y = xy_distance * angle_data.cos_azimuth;
-  const float z = distance * angle_data.sin_elevation;
-
-  const uint32_t output_idx = atomicAdd(output_count, 1);
-
-  CudaNebulaPoint & out_pt = output_points[output_idx];
-  out_pt.x = x;
-  out_pt.y = y;
-  out_pt.z = z;
-  out_pt.distance = distance;
-  out_pt.azimuth = angle_data.azimuth_rad;
-  out_pt.elevation = angle_data.elevation_rad;
-  out_pt.intensity = static_cast<float>(reflectivity);
-  out_pt.return_type = static_cast<uint8_t>(block_id);
-  out_pt.channel = static_cast<uint16_t>(channel_id);
-  out_pt.entry_id = config.entry_id;
-}
-
 /// @brief Batched kernel for processing an entire scan in one launch
 __global__ void decode_hesai_scan_batch_kernel(
   const uint16_t * __restrict__ d_distances_batch,
@@ -186,7 +124,7 @@ __global__ void decode_hesai_scan_batch_kernel(
   uint8_t in_current_scan = 1;
 
   bool is_in_overlap = false;
-  if (config.is_multi_frame) {
+  if (config.n_frames > 1) {
     is_in_overlap = cuda_is_inside_overlap_multiframe(
       last_azimuth, raw_azimuth, config.frame_angles, config.n_frames, config.n_azimuths_raw);
   } else {
@@ -336,37 +274,8 @@ bool HesaiCudaDecoder::upload_angle_corrections(
 
 }  // namespace nebula::drivers::cuda
 
-// C-linkage wrapper for per-packet kernel
-extern "C" void launch_decode_hesai_packet(
-  const uint16_t * d_distances,
-  const uint8_t * d_reflectivities,
-  const nebula::drivers::cuda::CudaAngleCorrectionData * d_angle_lut,
-  const nebula::drivers::cuda::CudaDecoderConfig & config,
-  nebula::drivers::cuda::CudaNebulaPoint * d_points,
-  uint32_t * d_count,
-  uint32_t n_azimuths,
-  uint32_t raw_azimuth,
-  cudaStream_t stream)
-{
-  const uint32_t threads_per_block = 128;
-  const uint32_t n_blocks_x = (config.n_channels + threads_per_block - 1) / threads_per_block;
-  const uint32_t n_blocks_y = config.n_blocks;
-
-  dim3 grid(n_blocks_x, n_blocks_y);
-  dim3 block(threads_per_block);
-
-  nebula::drivers::cuda::decode_hesai_packet_kernel<<<grid, block, 0, stream>>>(
-    d_distances, d_reflectivities, d_angle_lut, config,
-    d_points, d_count, n_azimuths, raw_azimuth);
-
-  cudaError_t err = cudaGetLastError();
-  if (err != cudaSuccess) {
-    fprintf(stderr, "CUDA kernel launch failed: %s\n", cudaGetErrorString(err));
-  }
-}
-
 // C-linkage wrapper for batched kernel
-extern "C" void launch_decode_hesai_scan_batch(
+extern "C" bool launch_decode_hesai_scan_batch(
   const uint16_t * d_distances_batch,
   const uint8_t * d_reflectivities_batch,
   const uint32_t * d_raw_azimuths,
@@ -391,8 +300,5 @@ extern "C" void launch_decode_hesai_scan_batch(
     d_distances_batch, d_reflectivities_batch, d_raw_azimuths, d_n_returns, d_last_azimuths,
     d_angle_lut, config, d_points, d_count, n_azimuths, n_packets);
 
-  cudaError_t err = cudaGetLastError();
-  if (err != cudaSuccess) {
-    fprintf(stderr, "CUDA batched kernel launch failed: %s\n", cudaGetErrorString(err));
-  }
+  return cudaGetLastError() == cudaSuccess;
 }