diff --git a/build_depends_humble.repos b/build_depends_humble.repos
index 0f35bd3d0e357..7d4fa266293d0 100644
--- a/build_depends_humble.repos
+++ b/build_depends_humble.repos
@@ -46,3 +46,11 @@ repositories:
     type: git
     url: https://github.com/tier4/glog.git
     version: v0.6.0_t4-ros
+  universe/external/cuda_blackboard:
+    type: git
+    url: https://github.com/autowarefoundation/cuda_blackboard.git
+    version: v0.1.1
+  universe/external/negotiated:
+    type: git
+    url: https://github.com/osrf/negotiated.git
+    version: eac198b55dcd052af5988f0f174902913c5f20e7
diff --git a/sensing/autoware_cuda_pointcloud_preprocessor/CMakeLists.txt b/sensing/autoware_cuda_pointcloud_preprocessor/CMakeLists.txt
new file mode 100644
index 0000000000000..3a77c27ce5f8b
--- /dev/null
+++ b/sensing/autoware_cuda_pointcloud_preprocessor/CMakeLists.txt
@@ -0,0 +1,97 @@
+cmake_minimum_required(VERSION 3.14)
+project(autoware_cuda_pointcloud_preprocessor)
+
+find_package(ament_cmake_auto REQUIRED)
+find_package(CUDA)
+
+if(NOT ${CUDA_FOUND})
+  message(WARNING "cuda was not found, so the autoware_cuda_pointcloud_preprocessor package will not be built.")
+  return()
+elseif(CMAKE_BUILD_TYPE STREQUAL "Debug")
+  set(CMAKE_CUDA_FLAGS ${CMAKE_CUDA_FLAGS} "-g -G")
+  set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS} "-g -G")
+endif()
+
+ament_auto_find_build_dependencies()
+
+# Default to C++17
+if(NOT CMAKE_CXX_STANDARD)
+    set(CMAKE_CXX_STANDARD 17)
+endif()
+
+if(CMAKE_COMPILER_IS_GNUCXX OR CMAKE_CXX_COMPILER_ID MATCHES "Clang")
+    add_compile_options(-Wall -Wextra -Wpedantic -Wunused-function)
+endif()
+
+if(BUILD_TESTING)
+    find_package(ament_lint_auto REQUIRED)
+    ament_lint_auto_find_test_dependencies()
+endif()
+
+include_directories(
+    include
+    SYSTEM
+    ${CUDA_INCLUDE_DIRS}
+)
+
+list(APPEND CUDA_NVCC_FLAGS "--expt-relaxed-constexpr -diag-suppress 20012") # cSpell: ignore expt
+
+cuda_add_library(cuda_pointcloud_preprocessor_lib SHARED
+  src/cuda_pointcloud_preprocessor/cuda_pointcloud_preprocessor.cu
+  src/cuda_pointcloud_preprocessor/common_kernels.cu
+  src/cuda_pointcloud_preprocessor/organize_kernels.cu
+  src/cuda_pointcloud_preprocessor/outlier_kernels.cu
+  src/cuda_pointcloud_preprocessor/undistort_kernels.cu
+)
+
+target_link_libraries(cuda_pointcloud_preprocessor_lib
+  ${autoware_pointcloud_preprocessor_TARGETS}
+)
+
+target_include_directories(cuda_pointcloud_preprocessor_lib SYSTEM PRIVATE
+  ${autoware_pointcloud_preprocessor_INCLUDE_DIRS}
+  ${autoware_point_types_INCLUDE_DIRS}
+  ${cuda_blackboard_INCLUDE_DIRS}
+  ${diagnostic_msgs_INCLUDE_DIRS}
+  ${geometry_msgs_INCLUDE_DIRS}
+  ${rclcpp_INCLUDE_DIRS}
+  ${rclcpp_components_INCLUDE_DIRS}
+  ${rcl_interfaces_INCLUDE_DIRS}
+  ${sensor_msgs_INCLUDE_DIRS}
+  ${tf2_INCLUDE_DIRS}
+  ${tf2_msgs_INCLUDE_DIRS}
+)
+
+
+# Targets
+ament_auto_add_library(cuda_pointcloud_preprocessor SHARED
+  src/cuda_pointcloud_preprocessor/cuda_pointcloud_preprocessor_node.cpp
+)
+
+target_link_libraries(cuda_pointcloud_preprocessor
+  ${CUDA_LIBRARIES}
+  cuda_pointcloud_preprocessor_lib
+)
+
+rclcpp_components_register_node(cuda_pointcloud_preprocessor
+  PLUGIN "autoware::cuda_pointcloud_preprocessor::CudaPointcloudPreprocessorNode"
+  EXECUTABLE cuda_pointcloud_preprocessor_node
+)
+
+install(DIRECTORY launch config
+  DESTINATION share/${PROJECT_NAME}
+)
+
+ament_auto_package()
+
+# Set ROS_DISTRO macros
+set(ROS_DISTRO $ENV{ROS_DISTRO})
+if(${ROS_DISTRO} STREQUAL "rolling")
+    add_compile_definitions(ROS_DISTRO_ROLLING)
+elseif(${ROS_DISTRO} STREQUAL "foxy")
+    add_compile_definitions(ROS_DISTRO_FOXY)
+elseif(${ROS_DISTRO} STREQUAL "galactic")
+    add_compile_definitions(ROS_DISTRO_GALACTIC)
+elseif(${ROS_DISTRO} STREQUAL "humble")
+    add_compile_definitions(ROS_DISTRO_HUMBLE)
+endif()
diff --git a/sensing/autoware_cuda_pointcloud_preprocessor/README.md b/sensing/autoware_cuda_pointcloud_preprocessor/README.md
new file mode 100644
index 0000000000000..0fa5ab7e7c65a
--- /dev/null
+++ b/sensing/autoware_cuda_pointcloud_preprocessor/README.md
@@ -0,0 +1,19 @@
+# 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. <!-- cSpell: ignore GPGPUs >
+
+## Inner-workings / Algorithms
+
+A detailed description of each filter's algorithm is available in the following links.
+
+| Filter Name                  | Description                                                                                                                                  | Detail                                       |
+| ---------------------------- | -------------------------------------------------------------------------------------------------------------------------------------------- | -------------------------------------------- |
+| cuda_pointcloud_preprocessor | Implements the cropping, distortion correction, and outlier filtering (ring-based) of the `autoware_pointcloud_preprocessor`'s CPU versions. | [link](docs/cuda-pointcloud-preprocessor.md) |
+
+## (Optional) Future extensions / Unimplemented parts
+
+The subsample filters implemented in `autoware_pointcloud_preprocessor` will have similar counterparts in this package.
diff --git a/sensing/autoware_cuda_pointcloud_preprocessor/config/cuda_pointcloud_preprocessor.param.yaml b/sensing/autoware_cuda_pointcloud_preprocessor/config/cuda_pointcloud_preprocessor.param.yaml
new file mode 100644
index 0000000000000..5a3d138484d0c
--- /dev/null
+++ b/sensing/autoware_cuda_pointcloud_preprocessor/config/cuda_pointcloud_preprocessor.param.yaml
@@ -0,0 +1,16 @@
+/**:
+  ros__parameters:
+
+    base_frame: base_link
+    use_imu: true
+    use_3d_distortion_correction: false
+    distance_ratio: 1.03
+    object_length_threshold: 0.1
+    num_points_threshold: 4
+
+    crop_box.min_x: [0.0]
+    crop_box.min_y: [0.0]
+    crop_box.min_z: [0.0]
+    crop_box.max_x: [0.0]
+    crop_box.max_y: [0.0]
+    crop_box.max_z: [0.0]
diff --git a/sensing/autoware_cuda_pointcloud_preprocessor/docs/cuda-pointcloud-preprocessor.md b/sensing/autoware_cuda_pointcloud_preprocessor/docs/cuda-pointcloud-preprocessor.md
new file mode 100644
index 0000000000000..2ba24e6a8fb52
--- /dev/null
+++ b/sensing/autoware_cuda_pointcloud_preprocessor/docs/cuda-pointcloud-preprocessor.md
@@ -0,0 +1,45 @@
+# cuda_pointcloud_preprocessor
+
+## Purpose
+
+This node implements all standard pointcloud preprocessing algorithms applied to a single LiDAR's pointcloud in CUDA.
+In particular, this node implements:
+
+- box cropping (ego-vehicle and ego-vehicle's mirrors)
+- distortion correction
+- ring-based outlier filtering
+
+## Inner-workings / Algorithms
+
+As this node reimplements the functionalities of the CPU-version algorithms, please have a look at the documentations of [crop-box](../../autoware_pointcloud_preprocessor/docs/crop-box-filter.md), [distortion correction](../../autoware_pointcloud_preprocessor/docs/distortion-corrector.md), and [ring-based outlier filter](../../autoware_pointcloud_preprocessor/docs/ring-outlier-filter.md) for more information about these algorithms.
+
+In addition to the individual algorithms previously mentioned, this node uses the `cuda_blackboard`, a cuda transport layer that enables a zero-copy mechanism between GPU and GPU memory for both input and output.
+
+## Inputs / Outputs
+
+### Input
+
+| Name                      | Type                                             | Description                               |
+| ------------------------- | ------------------------------------------------ | ----------------------------------------- |
+| `~/input/pointcloud`      | `sensor_msgs::msg::PointCloud2`                  | Input pointcloud's topic.                 |
+| `~/input/pointcloud/cuda` | `negotiated_interfaces/msg/NegotiatedTopicsInfo` | Input pointcloud's type negotiation topic |
+| `~/input/twist`           | `geometry_msgs::msg::TwistWithCovarianceStamped` | Topic of the twist information.           |
+| `~/input/imu`             | `sensor_msgs::msg::Imu`                          | Topic of the IMU data.                    |
+
+### Output
+
+| Name                       | Type                                             | Description                              |
+| -------------------------- | ------------------------------------------------ | ---------------------------------------- |
+| `~/output/pointcloud`      | `sensor_msgs::msg::PointCloud2`                  | Processed pointcloud's topic             |
+| `~/output/pointcloud/cuda` | `negotiated_interfaces/msg/NegotiatedTopicsInfo` | Processed pointcloud's negotiation topic |
+
+## Parameters
+
+### Core Parameters
+
+{{ json_to_markdown("sensing/autoware_cuda_pointcloud_preprocessor/schema/cuda_pointcloud_preprocessor.schema.schema.json") }}
+
+## Assumptions / Known limits
+
+- The CUDA implementations, while following the original CPU ones, will not offer the same numerical results, and small approximations were needed to maximize GPU usage.
+- This node expects that the input pointcloud follows the `autoware::point_types::PointXYZIRCAEDT` layout and the output pointcloud will use the `autoware::point_types::PointXYZIRC` layout defined in the `autoware_point_types` package.
diff --git a/sensing/autoware_cuda_pointcloud_preprocessor/include/autoware/cuda_pointcloud_preprocessor/common_kernels.hpp b/sensing/autoware_cuda_pointcloud_preprocessor/include/autoware/cuda_pointcloud_preprocessor/common_kernels.hpp
new file mode 100644
index 0000000000000..d450a206e9d59
--- /dev/null
+++ b/sensing/autoware_cuda_pointcloud_preprocessor/include/autoware/cuda_pointcloud_preprocessor/common_kernels.hpp
@@ -0,0 +1,45 @@
+// Copyright 2025 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.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef AUTOWARE__CUDA_POINTCLOUD_PREPROCESSOR__COMMON_KERNELS_HPP_
+#define AUTOWARE__CUDA_POINTCLOUD_PREPROCESSOR__COMMON_KERNELS_HPP_
+
+#include "autoware/cuda_pointcloud_preprocessor/point_types.hpp"
+#include "autoware/cuda_pointcloud_preprocessor/types.hpp"
+
+#include <cuda_runtime.h>
+
+namespace autoware::cuda_pointcloud_preprocessor
+{
+void transformPointsLaunch(
+  const InputPointType * input_points, InputPointType * output_points, int num_points,
+  TransformStruct transform, int threads_per_block, int blocks_per_grid, cudaStream_t & stream);
+
+void cropBoxLaunch(
+  InputPointType * d_points, std::uint32_t * output_mask, int num_points,
+  const CropBoxParameters * crop_box_parameters_ptr, int num_crop_boxes, int threads_per_block,
+  int blocks_per_grid, cudaStream_t & stream);
+
+void combineMasksLaunch(
+  const std::uint32_t * mask1, const std::uint32_t * mask2, int num_points,
+  std::uint32_t * output_mask, int threads_per_block, int blocks_per_grid, cudaStream_t & stream);
+
+void extractPointsLaunch(
+  InputPointType * input_points, std::uint32_t * masks, std::uint32_t * indices, int num_points,
+  OutputPointType * output_points, int threads_per_block, int blocks_per_grid,
+  cudaStream_t & stream);
+
+}  // namespace autoware::cuda_pointcloud_preprocessor
+
+#endif  // AUTOWARE__CUDA_POINTCLOUD_PREPROCESSOR__COMMON_KERNELS_HPP_
diff --git a/sensing/autoware_cuda_pointcloud_preprocessor/include/autoware/cuda_pointcloud_preprocessor/cuda_pointcloud_preprocessor.hpp b/sensing/autoware_cuda_pointcloud_preprocessor/include/autoware/cuda_pointcloud_preprocessor/cuda_pointcloud_preprocessor.hpp
new file mode 100644
index 0000000000000..bd4a3a1b20601
--- /dev/null
+++ b/sensing/autoware_cuda_pointcloud_preprocessor/include/autoware/cuda_pointcloud_preprocessor/cuda_pointcloud_preprocessor.hpp
@@ -0,0 +1,109 @@
+// Copyright 2025 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.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef AUTOWARE__CUDA_POINTCLOUD_PREPROCESSOR__CUDA_POINTCLOUD_PREPROCESSOR_HPP_
+#define AUTOWARE__CUDA_POINTCLOUD_PREPROCESSOR__CUDA_POINTCLOUD_PREPROCESSOR_HPP_
+
+#include "autoware/cuda_pointcloud_preprocessor/point_types.hpp"
+#include "autoware/cuda_pointcloud_preprocessor/types.hpp"
+
+#include <Eigen/Core>
+#include <Eigen/Geometry>
+#include <cuda_blackboard/cuda_pointcloud2.hpp>
+
+#include <geometry_msgs/msg/transform_stamped.hpp>
+#include <geometry_msgs/msg/twist_with_covariance_stamped.hpp>
+#include <geometry_msgs/msg/vector3_stamped.hpp>
+#include <sensor_msgs/msg/point_cloud2.hpp>
+
+#include <thrust/device_vector.h>
+
+#include <cstdint>
+#include <deque>
+#include <memory>
+#include <vector>
+
+namespace autoware::cuda_pointcloud_preprocessor
+{
+
+class CudaPointcloudPreprocessor
+{
+public:
+  enum class UndistortionType { Invalid, Undistortion2D, Undistortion3D };
+
+  CudaPointcloudPreprocessor();
+  ~CudaPointcloudPreprocessor() = default;
+
+  void setCropBoxParameters(const std::vector<CropBoxParameters> & crop_box_parameters);
+  void setRingOutlierFilterParameters(const RingOutlierFilterParameters & ring_outlier_parameters);
+  void setUndistortionType(const UndistortionType & undistortion_type);
+
+  void preallocateOutput();
+
+  std::unique_ptr<cuda_blackboard::CudaPointCloud2> process(
+    const sensor_msgs::msg::PointCloud2::ConstSharedPtr input_pointcloud_msg_ptr,
+    const geometry_msgs::msg::TransformStamped & transform_msg,
+    const std::deque<geometry_msgs::msg::TwistWithCovarianceStamped> & twist_queue,
+    const std::deque<geometry_msgs::msg::Vector3Stamped> & angular_velocity_queue,
+    const std::uint32_t first_point_rel_stamp_nsec);
+
+private:
+  void organizePointcloud();
+
+  CropBoxParameters self_crop_box_parameters_{};
+  CropBoxParameters mirror_crop_box_parameters_{};
+  RingOutlierFilterParameters ring_outlier_parameters_{};
+  UndistortionType undistortion_type_{UndistortionType::Invalid};
+
+  int num_rings_{};
+  int max_points_per_ring_{};
+  std::size_t num_raw_points_{};
+  std::size_t num_organized_points_{};
+
+  std::vector<sensor_msgs::msg::PointField> point_fields_{};
+  std::unique_ptr<cuda_blackboard::CudaPointCloud2> output_pointcloud_ptr_{};
+
+  cudaStream_t stream_;
+  int max_blocks_per_grid_{};
+  const int threads_per_block_{256};
+  cudaMemPool_t device_memory_pool_;
+
+  // Organizing buffers
+  thrust::device_vector<InputPointType> device_input_points_;
+  thrust::device_vector<InputPointType> device_organized_points_;
+  thrust::device_vector<std::int32_t> device_ring_index_;
+  thrust::device_vector<std::uint32_t> device_indexes_tensor_;
+  thrust::device_vector<std::uint32_t> device_sorted_indexes_tensor_;
+  thrust::device_vector<std::int32_t> device_segment_offsets_;
+  thrust::device_vector<std::int32_t> device_max_ring_;
+  thrust::device_vector<std::int32_t> device_max_points_per_ring_;
+
+  thrust::device_vector<std::uint8_t> device_sort_workspace_;
+  std::size_t sort_workspace_bytes_{0};
+
+  // Pointcloud preprocessing buffers
+  thrust::device_vector<InputPointType> device_transformed_points_{};
+  thrust::device_vector<OutputPointType> device_output_points_{};
+  thrust::device_vector<std::uint32_t> device_crop_mask_{};
+  thrust::device_vector<std::uint32_t> device_ring_outlier_mask_{};
+  thrust::device_vector<std::uint32_t> device_indices_{};
+  thrust::device_vector<TwistStruct2D> device_twist_2d_structs_{};
+  thrust::device_vector<TwistStruct3D> device_twist_3d_structs_{};
+  thrust::device_vector<CropBoxParameters> host_crop_box_structs_{};
+  thrust::device_vector<CropBoxParameters> device_crop_box_structs_{};
+};
+
+}  // namespace autoware::cuda_pointcloud_preprocessor
+
+#endif  // AUTOWARE__CUDA_POINTCLOUD_PREPROCESSOR__CUDA_POINTCLOUD_PREPROCESSOR_HPP_
diff --git a/sensing/autoware_cuda_pointcloud_preprocessor/include/autoware/cuda_pointcloud_preprocessor/cuda_pointcloud_preprocessor_node.hpp b/sensing/autoware_cuda_pointcloud_preprocessor/include/autoware/cuda_pointcloud_preprocessor/cuda_pointcloud_preprocessor_node.hpp
new file mode 100644
index 0000000000000..1c2b15b5de5d3
--- /dev/null
+++ b/sensing/autoware_cuda_pointcloud_preprocessor/include/autoware/cuda_pointcloud_preprocessor/cuda_pointcloud_preprocessor_node.hpp
@@ -0,0 +1,116 @@
+// Copyright 2025 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.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef AUTOWARE__CUDA_POINTCLOUD_PREPROCESSOR__CUDA_POINTCLOUD_PREPROCESSOR_NODE_HPP_
+#define AUTOWARE__CUDA_POINTCLOUD_PREPROCESSOR__CUDA_POINTCLOUD_PREPROCESSOR_NODE_HPP_
+
+#include "autoware/cuda_pointcloud_preprocessor/cuda_pointcloud_preprocessor.hpp"
+#include "autoware/cuda_pointcloud_preprocessor/point_types.hpp"
+
+#include <autoware/point_types/types.hpp>
+#include <autoware/universe_utils/ros/debug_publisher.hpp>
+#include <autoware/universe_utils/system/stop_watch.hpp>
+#include <cuda_blackboard/cuda_adaptation.hpp>
+#include <cuda_blackboard/cuda_blackboard_publisher.hpp>
+#include <cuda_blackboard/cuda_blackboard_subscriber.hpp>
+#include <cuda_blackboard/cuda_pointcloud2.hpp>
+#include <rclcpp/rclcpp.hpp>
+
+#include <geometry_msgs/msg/twist_with_covariance_stamped.hpp>
+#include <geometry_msgs/msg/vector3_stamped.hpp>
+#include <sensor_msgs/msg/imu.hpp>
+#include <sensor_msgs/msg/point_cloud2.hpp>
+
+#include <tf2/transform_datatypes.h>
+#include <tf2_ros/buffer.h>
+#include <tf2_ros/transform_listener.h>
+
+#include <deque>
+#include <memory>
+#include <string>
+
+/* *INDENT-OFF* */
+#define CHECK_OFFSET(structure1, structure2, field)             \
+  static_assert(                                                \
+    offsetof(structure1, field) == offsetof(structure2, field), \
+    "Offset of " #field " in " #structure1 " does not match expected offset.")
+/* *INDENT-ON* */
+
+namespace autoware::cuda_pointcloud_preprocessor
+{
+
+static_assert(sizeof(InputPointType) == sizeof(autoware::point_types::PointXYZIRCAEDT));
+static_assert(sizeof(OutputPointType) == sizeof(autoware::point_types::PointXYZIRC));
+
+CHECK_OFFSET(InputPointType, autoware::point_types::PointXYZIRCAEDT, x);
+CHECK_OFFSET(InputPointType, autoware::point_types::PointXYZIRCAEDT, y);
+CHECK_OFFSET(InputPointType, autoware::point_types::PointXYZIRCAEDT, z);
+CHECK_OFFSET(InputPointType, autoware::point_types::PointXYZIRCAEDT, intensity);
+CHECK_OFFSET(InputPointType, autoware::point_types::PointXYZIRCAEDT, return_type);
+CHECK_OFFSET(InputPointType, autoware::point_types::PointXYZIRCAEDT, channel);
+CHECK_OFFSET(InputPointType, autoware::point_types::PointXYZIRCAEDT, azimuth);
+CHECK_OFFSET(InputPointType, autoware::point_types::PointXYZIRCAEDT, elevation);
+CHECK_OFFSET(InputPointType, autoware::point_types::PointXYZIRCAEDT, distance);
+CHECK_OFFSET(InputPointType, autoware::point_types::PointXYZIRCAEDT, time_stamp);
+
+CHECK_OFFSET(OutputPointType, autoware::point_types::PointXYZIRCAEDT, x);
+CHECK_OFFSET(OutputPointType, autoware::point_types::PointXYZIRCAEDT, y);
+CHECK_OFFSET(OutputPointType, autoware::point_types::PointXYZIRCAEDT, z);
+CHECK_OFFSET(OutputPointType, autoware::point_types::PointXYZIRCAEDT, intensity);
+CHECK_OFFSET(OutputPointType, autoware::point_types::PointXYZIRCAEDT, return_type);
+CHECK_OFFSET(OutputPointType, autoware::point_types::PointXYZIRCAEDT, channel);
+
+class CudaPointcloudPreprocessorNode : public rclcpp::Node
+{
+public:
+  explicit CudaPointcloudPreprocessorNode(const rclcpp::NodeOptions & node_options);
+
+private:
+  bool getTransform(
+    const std::string & target_frame, const std::string & source_frame,
+    tf2::Transform * tf2_transform_ptr);
+
+  // Callback
+  void pointcloudCallback(
+    const sensor_msgs::msg::PointCloud2::ConstSharedPtr input_pointcloud_msg_ptr);
+  void cudaPointcloudCallback(
+    const std::shared_ptr<const cuda_blackboard::CudaPointCloud2> cuda_msg);
+  void twistCallback(
+    const geometry_msgs::msg::TwistWithCovarianceStamped::ConstSharedPtr pointcloud_msg);
+  void imuCallback(const sensor_msgs::msg::Imu::ConstSharedPtr imu_msg);
+
+  tf2_ros::Buffer tf2_buffer_;
+  tf2_ros::TransformListener tf2_listener_;
+
+  std::string base_frame_;
+  std::deque<geometry_msgs::msg::TwistWithCovarianceStamped> twist_queue_;
+  std::deque<geometry_msgs::msg::Vector3Stamped> angular_velocity_queue_;
+
+  // Subscribers
+  rclcpp::Subscription<geometry_msgs::msg::TwistWithCovarianceStamped>::SharedPtr twist_sub_{};
+  rclcpp::Subscription<sensor_msgs::msg::Imu>::SharedPtr imu_sub_{};
+  rclcpp::Subscription<sensor_msgs::msg::PointCloud2>::SharedPtr pointcloud_sub_{};
+
+  // CUDA pub
+  std::unique_ptr<cuda_blackboard::CudaBlackboardPublisher<cuda_blackboard::CudaPointCloud2>> pub_;
+
+  std::unique_ptr<CudaPointcloudPreprocessor> cuda_pointcloud_preprocessor_;
+
+  std::unique_ptr<autoware::universe_utils::StopWatch<std::chrono::milliseconds>> stop_watch_ptr_;
+  std::unique_ptr<autoware::universe_utils::DebugPublisher> debug_publisher_;
+};
+
+}  // namespace autoware::cuda_pointcloud_preprocessor
+
+#endif  // AUTOWARE__CUDA_POINTCLOUD_PREPROCESSOR__CUDA_POINTCLOUD_PREPROCESSOR_NODE_HPP_
diff --git a/sensing/autoware_cuda_pointcloud_preprocessor/include/autoware/cuda_pointcloud_preprocessor/memory.hpp b/sensing/autoware_cuda_pointcloud_preprocessor/include/autoware/cuda_pointcloud_preprocessor/memory.hpp
new file mode 100644
index 0000000000000..9fb24b8a29c3f
--- /dev/null
+++ b/sensing/autoware_cuda_pointcloud_preprocessor/include/autoware/cuda_pointcloud_preprocessor/memory.hpp
@@ -0,0 +1,91 @@
+// Copyright 2025 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.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef AUTOWARE__CUDA_POINTCLOUD_PREPROCESSOR__MEMORY_HPP_
+#define AUTOWARE__CUDA_POINTCLOUD_PREPROCESSOR__MEMORY_HPP_
+
+#include "autoware/cuda_pointcloud_preprocessor/point_types.hpp"
+
+#include <sensor_msgs/msg/point_cloud2.hpp>
+
+#include <vector>
+
+namespace autoware::cuda_pointcloud_preprocessor
+{
+
+// TODO(knzo25): this checks already exists within the pointcloud_preprocessor. After this PR is
+// merged I will move it somewhere all classes that require checks can access it
+bool is_data_layout_compatible_with_point_xyzircaedt(
+  const std::vector<sensor_msgs::msg::PointField> & fields)
+{
+  if (fields.size() != 10) {
+    return false;
+  }
+  bool same_layout = true;
+  const auto & field_x = fields.at(0);
+  same_layout &= field_x.name == "x";
+  same_layout &= field_x.offset == offsetof(InputPointType, x);
+  same_layout &= field_x.datatype == sensor_msgs::msg::PointField::FLOAT32;
+  same_layout &= field_x.count == 1;
+  const auto & field_y = fields.at(1);
+  same_layout &= field_y.name == "y";
+  same_layout &= field_y.offset == offsetof(InputPointType, y);
+  same_layout &= field_y.datatype == sensor_msgs::msg::PointField::FLOAT32;
+  same_layout &= field_y.count == 1;
+  const auto & field_z = fields.at(2);
+  same_layout &= field_z.name == "z";
+  same_layout &= field_z.offset == offsetof(InputPointType, z);
+  same_layout &= field_z.datatype == sensor_msgs::msg::PointField::FLOAT32;
+  same_layout &= field_z.count == 1;
+  const auto & field_intensity = fields.at(3);
+  same_layout &= field_intensity.name == "intensity";
+  same_layout &= field_intensity.offset == offsetof(InputPointType, intensity);
+  same_layout &= field_intensity.datatype == sensor_msgs::msg::PointField::UINT8;
+  same_layout &= field_intensity.count == 1;
+  const auto & field_return_type = fields.at(4);
+  same_layout &= field_return_type.name == "return_type";
+  same_layout &= field_return_type.offset == offsetof(InputPointType, return_type);
+  same_layout &= field_return_type.datatype == sensor_msgs::msg::PointField::UINT8;
+  same_layout &= field_return_type.count == 1;
+  const auto & field_ring = fields.at(5);
+  same_layout &= field_ring.name == "channel";
+  same_layout &= field_ring.offset == offsetof(InputPointType, channel);
+  same_layout &= field_ring.datatype == sensor_msgs::msg::PointField::UINT16;
+  same_layout &= field_ring.count == 1;
+  const auto & field_azimuth = fields.at(6);
+  same_layout &= field_azimuth.name == "azimuth";
+  same_layout &= field_azimuth.offset == offsetof(InputPointType, azimuth);
+  same_layout &= field_azimuth.datatype == sensor_msgs::msg::PointField::FLOAT32;
+  same_layout &= field_azimuth.count == 1;
+  const auto & field_elevation = fields.at(7);
+  same_layout &= field_elevation.name == "elevation";
+  same_layout &= field_elevation.offset == offsetof(InputPointType, elevation);
+  same_layout &= field_elevation.datatype == sensor_msgs::msg::PointField::FLOAT32;
+  same_layout &= field_elevation.count == 1;
+  const auto & field_distance = fields.at(8);
+  same_layout &= field_distance.name == "distance";
+  same_layout &= field_distance.offset == offsetof(InputPointType, distance);
+  same_layout &= field_distance.datatype == sensor_msgs::msg::PointField::FLOAT32;
+  same_layout &= field_distance.count == 1;
+  const auto & field_time_stamp = fields.at(9);
+  same_layout &= field_time_stamp.name == "time_stamp";
+  same_layout &= field_time_stamp.offset == offsetof(InputPointType, time_stamp);
+  same_layout &= field_time_stamp.datatype == sensor_msgs::msg::PointField::UINT32;
+  same_layout &= field_time_stamp.count == 1;
+  return same_layout;
+}
+
+}  // namespace autoware::cuda_pointcloud_preprocessor
+
+#endif  // AUTOWARE__CUDA_POINTCLOUD_PREPROCESSOR__MEMORY_HPP_
diff --git a/sensing/autoware_cuda_pointcloud_preprocessor/include/autoware/cuda_pointcloud_preprocessor/organize_kernels.hpp b/sensing/autoware_cuda_pointcloud_preprocessor/include/autoware/cuda_pointcloud_preprocessor/organize_kernels.hpp
new file mode 100644
index 0000000000000..0fe52bfe76409
--- /dev/null
+++ b/sensing/autoware_cuda_pointcloud_preprocessor/include/autoware/cuda_pointcloud_preprocessor/organize_kernels.hpp
@@ -0,0 +1,41 @@
+// Copyright 2025 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.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef AUTOWARE__CUDA_POINTCLOUD_PREPROCESSOR__ORGANIZE_KERNELS_HPP_
+#define AUTOWARE__CUDA_POINTCLOUD_PREPROCESSOR__ORGANIZE_KERNELS_HPP_
+
+#include "autoware/cuda_pointcloud_preprocessor/point_types.hpp"
+
+#include <cuda_runtime.h>
+
+namespace autoware::cuda_pointcloud_preprocessor
+{
+
+std::size_t querySortWorkspace(
+  int num_items, int num_segments, int * offsets_device, std::uint32_t * keys_in_device,
+  std::uint32_t * keys_out_device);
+
+void organizeLaunch(
+  const InputPointType * input_points, std::uint32_t * index_tensor, std::int32_t * ring_indexes,
+  std::int32_t initial_max_rings, std::int32_t * output_max_rings,
+  std::int32_t initial_max_points_per_ring, std::int32_t * output_max_points_per_ring,
+  int num_points, int threads_per_block, int blocks_per_grid, cudaStream_t & stream);
+
+void gatherLaunch(
+  const InputPointType * input_points, const std::uint32_t * index_tensor,
+  InputPointType * output_points, int num_rings, int max_points_per_ring, int threads_per_block,
+  int blocks_per_grid, cudaStream_t & stream);
+}  // namespace autoware::cuda_pointcloud_preprocessor
+
+#endif  // AUTOWARE__CUDA_POINTCLOUD_PREPROCESSOR__ORGANIZE_KERNELS_HPP_
diff --git a/sensing/autoware_cuda_pointcloud_preprocessor/include/autoware/cuda_pointcloud_preprocessor/outlier_kernels.hpp b/sensing/autoware_cuda_pointcloud_preprocessor/include/autoware/cuda_pointcloud_preprocessor/outlier_kernels.hpp
new file mode 100644
index 0000000000000..23d46458edce3
--- /dev/null
+++ b/sensing/autoware_cuda_pointcloud_preprocessor/include/autoware/cuda_pointcloud_preprocessor/outlier_kernels.hpp
@@ -0,0 +1,31 @@
+// Copyright 2025 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.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef AUTOWARE__CUDA_POINTCLOUD_PREPROCESSOR__OUTLIER_KERNELS_HPP_
+#define AUTOWARE__CUDA_POINTCLOUD_PREPROCESSOR__OUTLIER_KERNELS_HPP_
+
+#include "autoware/cuda_pointcloud_preprocessor/point_types.hpp"
+
+#include <cuda_runtime.h>
+
+namespace autoware::cuda_pointcloud_preprocessor
+{
+void ringOutlierFilterLaunch(
+  const InputPointType * points, std::uint32_t * output_mask, int num_rings,
+  int max_points_per_ring, float distance_ratio, float object_length_threshold_squared,
+  int num_points_threshold, int threads_per_block, int blocks_per_grid, cudaStream_t & stream);
+
+}  // namespace autoware::cuda_pointcloud_preprocessor
+
+#endif  // AUTOWARE__CUDA_POINTCLOUD_PREPROCESSOR__OUTLIER_KERNELS_HPP_
diff --git a/sensing/autoware_cuda_pointcloud_preprocessor/include/autoware/cuda_pointcloud_preprocessor/point_types.hpp b/sensing/autoware_cuda_pointcloud_preprocessor/include/autoware/cuda_pointcloud_preprocessor/point_types.hpp
new file mode 100644
index 0000000000000..e92fde50d2782
--- /dev/null
+++ b/sensing/autoware_cuda_pointcloud_preprocessor/include/autoware/cuda_pointcloud_preprocessor/point_types.hpp
@@ -0,0 +1,49 @@
+// Copyright 2025 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.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+#ifndef AUTOWARE__CUDA_POINTCLOUD_PREPROCESSOR__POINT_TYPES_HPP_
+#define AUTOWARE__CUDA_POINTCLOUD_PREPROCESSOR__POINT_TYPES_HPP_
+
+#include <cstdint>
+
+namespace autoware::cuda_pointcloud_preprocessor
+{
+
+// Note: We can not use PCL nor uniform initialization here because of thrust
+struct OutputPointType
+{
+  float x;
+  float y;
+  float z;
+  std::uint8_t intensity;
+  std::uint8_t return_type;
+  std::uint16_t channel;
+};
+
+struct InputPointType
+{
+  float x;
+  float y;
+  float z;
+  std::uint8_t intensity;
+  std::uint8_t return_type;
+  std::uint16_t channel;
+  float azimuth;
+  float elevation;
+  float distance;
+  std::uint32_t time_stamp;
+};
+
+}  // namespace autoware::cuda_pointcloud_preprocessor
+
+#endif  // AUTOWARE__CUDA_POINTCLOUD_PREPROCESSOR__POINT_TYPES_HPP_
diff --git a/sensing/autoware_cuda_pointcloud_preprocessor/include/autoware/cuda_pointcloud_preprocessor/types.hpp b/sensing/autoware_cuda_pointcloud_preprocessor/include/autoware/cuda_pointcloud_preprocessor/types.hpp
new file mode 100644
index 0000000000000..dbe62b4e82f47
--- /dev/null
+++ b/sensing/autoware_cuda_pointcloud_preprocessor/include/autoware/cuda_pointcloud_preprocessor/types.hpp
@@ -0,0 +1,106 @@
+// Copyright 2025 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.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef AUTOWARE__CUDA_POINTCLOUD_PREPROCESSOR__TYPES_HPP_
+#define AUTOWARE__CUDA_POINTCLOUD_PREPROCESSOR__TYPES_HPP_
+
+#include <cuda_runtime.h>
+
+#include <cmath>
+#include <cstddef>
+#include <cstdint>
+
+namespace autoware::cuda_pointcloud_preprocessor
+{
+
+struct TwistStruct2D
+{
+  float cum_x;
+  float cum_y;
+  float cum_theta;
+  float cum_cos_theta;
+  float cum_sin_theta;
+  std::uint32_t last_stamp_nsec;  // relative to the start of the pointcloud
+  std::uint32_t stamp_nsec;       // relative to the start of the pointcloud
+  float v_x;
+  float v_theta;
+};
+
+struct TwistStruct3D
+{
+  float cum_transform_buffer[16];
+  std::uint32_t last_stamp_nsec;  // relative to the start of the pointcloud
+  std::uint32_t stamp_nsec;       // relative to the start of the pointcloud
+  float v[3];
+  float w[3];
+};
+
+struct TransformStruct
+{
+  float x;
+  float y;
+  float z;
+  float m11;
+  float m12;
+  float m13;
+  float m21;
+  float m22;
+  float m23;
+  float m31;
+  float m32;
+  float m33;
+};
+
+struct CropBoxParameters
+{
+  float min_x;
+  float max_x;
+  float min_y;
+  float max_y;
+  float min_z;
+  float max_z;
+};
+
+struct RingOutlierFilterParameters
+{
+  float distance_ratio{std::nanf("")};
+  float object_length_threshold{std::nanf("")};
+  std::size_t num_points_threshold{0};
+};
+
+/* *INDENT-OFF* */
+template <typename T>
+class MemoryPoolAllocator
+{
+public:
+  using value_type = T;
+  explicit MemoryPoolAllocator(cudaMemPool_t pool) : m_pool(pool) {}
+
+  T * allocate(std::size_t n)
+  {
+    void * ptr = nullptr;
+    cudaMallocFromPoolAsync(&ptr, n * sizeof(T), m_pool, cudaStreamDefault);
+    return static_cast<T *>(ptr);
+  }
+
+  void deallocate(T * ptr, std::size_t) { cudaFreeAsync(ptr, cudaStreamDefault); }
+
+protected:
+  cudaMemPool_t m_pool;
+};
+/* *INDENT-ON* */
+
+}  // namespace autoware::cuda_pointcloud_preprocessor
+
+#endif  // AUTOWARE__CUDA_POINTCLOUD_PREPROCESSOR__TYPES_HPP_
diff --git a/sensing/autoware_cuda_pointcloud_preprocessor/include/autoware/cuda_pointcloud_preprocessor/undistort_kernels.hpp b/sensing/autoware_cuda_pointcloud_preprocessor/include/autoware/cuda_pointcloud_preprocessor/undistort_kernels.hpp
new file mode 100644
index 0000000000000..8b791f171285c
--- /dev/null
+++ b/sensing/autoware_cuda_pointcloud_preprocessor/include/autoware/cuda_pointcloud_preprocessor/undistort_kernels.hpp
@@ -0,0 +1,53 @@
+// Copyright 2025 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.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef AUTOWARE__CUDA_POINTCLOUD_PREPROCESSOR__UNDISTORT_KERNELS_HPP_
+#define AUTOWARE__CUDA_POINTCLOUD_PREPROCESSOR__UNDISTORT_KERNELS_HPP_
+
+#include "autoware/cuda_pointcloud_preprocessor/point_types.hpp"
+#include "autoware/cuda_pointcloud_preprocessor/types.hpp"
+
+#include <geometry_msgs/msg/twist_with_covariance_stamped.hpp>
+#include <geometry_msgs/msg/vector3_stamped.hpp>
+
+#include <cuda_runtime.h>
+#include <thrust/device_vector.h>
+
+#include <deque>
+
+namespace autoware::cuda_pointcloud_preprocessor
+{
+void undistort2DLaunch(
+  InputPointType * input_points, int num_points, TwistStruct2D * twist_structs, int num_twists,
+  int threads_per_block, int blocks_per_grid, cudaStream_t & stream);
+
+void undistort3DLaunch(
+  InputPointType * input_points, int num_points, TwistStruct3D * twist_structs, int num_twists,
+  int threads_per_block, int blocks_per_grid, cudaStream_t & stream);
+
+void setupTwist2DStructs(
+  const std::deque<geometry_msgs::msg::TwistWithCovarianceStamped> & twist_queue,
+  const std::deque<geometry_msgs::msg::Vector3Stamped> & angular_velocity_queue,
+  const std::uint64_t pointcloud_stamp_nsec, const std::uint32_t first_point_rel_stamp_nsec,
+  thrust::device_vector<TwistStruct2D> & device_twist_2d_structs, cudaStream_t & stream);
+
+void setupTwist3DStructs(
+  const std::deque<geometry_msgs::msg::TwistWithCovarianceStamped> & twist_queue,
+  const std::deque<geometry_msgs::msg::Vector3Stamped> & angular_velocity_queue,
+  const std::uint64_t pointcloud_stamp_nsec, const std::uint32_t first_point_rel_stamp_nsec,
+  thrust::device_vector<TwistStruct3D> & device_twist_3d_structs, cudaStream_t & stream);
+
+}  // namespace autoware::cuda_pointcloud_preprocessor
+
+#endif  // AUTOWARE__CUDA_POINTCLOUD_PREPROCESSOR__UNDISTORT_KERNELS_HPP_
diff --git a/sensing/autoware_cuda_pointcloud_preprocessor/launch/cuda_pointcloud_preprocessor.launch.xml b/sensing/autoware_cuda_pointcloud_preprocessor/launch/cuda_pointcloud_preprocessor.launch.xml
new file mode 100644
index 0000000000000..a1da651aa76e2
--- /dev/null
+++ b/sensing/autoware_cuda_pointcloud_preprocessor/launch/cuda_pointcloud_preprocessor.launch.xml
@@ -0,0 +1,16 @@
+<launch>
+  <arg name="input/pointcloud" default="/sensing/lidar/top/pointcloud_raw_ex"/>
+  <arg name="input/imu" default="/sensing/imu/imu_data"/>
+  <arg name="input/twist" default="/sensing/vehicle_velocity_converter/twist_with_covariance"/>
+  <arg name="output/pointcloud" default="/sensing/lidar/top/test"/>
+
+  <arg name="cuda_pointcloud_preprocessor_param_file" default="$(find-pkg-share autoware_cuda_pointcloud_preprocessor)/config/cuda_pointcloud_preprocessor.param.yaml"/>
+
+  <node pkg="autoware_cuda_pointcloud_preprocessor" exec="cuda_pointcloud_preprocessor_node" name="cuda_pointcloud_preprocessor" output="screen">
+    <remap from="~/input/pointcloud" to="$(var input/pointcloud)"/>
+    <remap from="~/input/imu" to="$(var input/imu)"/>
+    <remap from="~/input/twist" to="$(var input/twist)"/>
+    <remap from="~/output/pointcloud" to="$(var output/pointcloud)"/>
+    <param from="$(var cuda_pointcloud_preprocessor_param_file)"/>
+  </node>
+</launch>
diff --git a/sensing/autoware_cuda_pointcloud_preprocessor/package.xml b/sensing/autoware_cuda_pointcloud_preprocessor/package.xml
new file mode 100644
index 0000000000000..abbdb5ecb6b78
--- /dev/null
+++ b/sensing/autoware_cuda_pointcloud_preprocessor/package.xml
@@ -0,0 +1,33 @@
+<?xml version="1.0"?>
+<package format="3">
+  <name>autoware_cuda_pointcloud_preprocessor</name>
+  <version>0.1.0</version>
+  <description>autoware_cuda_pointcloud_preprocessor</description>
+  <maintainer email="kenzo.lobos@tier4.jp">Kenzo Lobos-Tsunekawa</maintainer>
+  <author email="kenzo.lobos@tier4.jp">Kenzo Lobos-Tsunekawa</author>
+  <license>Apache License 2.0</license>
+
+  <buildtool_depend>ament_cmake_auto</buildtool_depend>
+  <buildtool_depend>autoware_cmake</buildtool_depend>
+
+  <depend>autoware_cuda_dependency_meta</depend>
+  <depend>autoware_point_types</depend>
+  <depend>autoware_pointcloud_preprocessor</depend>
+  <depend>autoware_universe_utils</depend>
+  <depend>cuda_blackboard</depend>
+  <depend>diagnostic_msgs</depend>
+  <depend>rclcpp</depend>
+  <depend>rclcpp_components</depend>
+  <depend>sensor_msgs</depend>
+  <depend>tf2</depend>
+  <depend>tf2_msgs</depend>
+  <depend>tier4_debug_msgs</depend>
+
+  <test_depend>ament_clang_format</test_depend>
+  <test_depend>ament_lint_auto</test_depend>
+  <test_depend>ament_lint_common</test_depend>
+
+  <export>
+    <build_type>ament_cmake</build_type>
+  </export>
+</package>
diff --git a/sensing/autoware_cuda_pointcloud_preprocessor/schema/cuda_pointcloud_preprocessor.schema.json b/sensing/autoware_cuda_pointcloud_preprocessor/schema/cuda_pointcloud_preprocessor.schema.json
new file mode 100644
index 0000000000000..9d61b4932ad8a
--- /dev/null
+++ b/sensing/autoware_cuda_pointcloud_preprocessor/schema/cuda_pointcloud_preprocessor.schema.json
@@ -0,0 +1,104 @@
+{
+  "$schema": "http://json-schema.org/draft-07/schema#",
+  "title": "Parameters for the cuda_pointcloud_preprocessor",
+  "type": "object",
+  "definitions": {
+    "cuda_pointcloud_preprocessor": {
+      "type": "object",
+      "properties": {
+        "base_frame": {
+          "type": "string",
+          "description": "The undistortion algorithm is based on a base frame, which must be the same as the twist frame.",
+          "default": "base_link"
+        },
+        "use_imu": {
+          "type": "boolean",
+          "description": "Use IMU angular velocity, otherwise, use twist angular velocity.",
+          "default": "true"
+        },
+        "use_3d_distortion_correction": {
+          "type": "boolean",
+          "description": "Use 3d distortion correction algorithm, otherwise, use 2d distortion correction algorithm.",
+          "default": "false"
+        },
+        "distance_ratio": {
+          "type": "number",
+          "description": "distance ratio between consecutive points to determine is they belong to the same walk",
+          "default": "1.03",
+          "minimum": 0.0
+        },
+        "object_length_threshold": {
+          "type": "number",
+          "description": "minimum object length to be considered a valid cluster in meters",
+          "default": "0.1",
+          "minimum": 0.0
+        },
+        "num_points_threshold": {
+          "type": "integer",
+          "description": "threshold on the number of points in a walk to determine if it corresponds to a cluster",
+          "default": "4",
+          "minimum": 0
+        },
+        "crop_box.min_x": {
+          "type": "array",
+          "description": "An array in which each element is the minimum x value in a crop box",
+          "default": []
+        },
+        "crop_box.min_y": {
+          "type": "array",
+          "description": "An array in which each element is the minimum y value in a crop box",
+          "default": []
+        },
+        "crop_box.min_z": {
+          "type": "array",
+          "description": "An array in which each element is the minimum z value in a crop box",
+          "default": []
+        },
+        "crop_box.max_x": {
+          "type": "array",
+          "description": "An array in which each element is the maximum x value in a crop box",
+          "default": []
+        },
+        "crop_box.max_y": {
+          "type": "array",
+          "description": "An array in which each element is the maximum y value in a crop box",
+          "default": []
+        },
+        "crop_box.max_z": {
+          "type": "array",
+          "description": "An array in which each element is the maximum z value in a crop box",
+          "default": []
+        }
+      },
+      "required": [
+        "base_frame",
+        "use_imu",
+        "use_3d_distortion_correction",
+        "distance_ratio",
+        "object_length_threshold",
+        "num_points_threshold",
+        "crop_box.min_x",
+        "crop_box.min_y",
+        "crop_box.min_z",
+        "crop_box.max_x",
+        "crop_box.max_y",
+        "crop_box.max_z"
+      ],
+      "additionalProperties": false
+    }
+  },
+  "properties": {
+    "/**": {
+      "type": "object",
+      "properties": {
+        "ros__parameters": {
+          "$ref": "#/definitions/cuda_pointcloud_preprocessor"
+        }
+      },
+      "required": ["ros__parameters"],
+      "additionalProperties": false
+    }
+  },
+  "required": ["/**"],
+  "additionalProperties": false
+}
diff --git a/sensing/autoware_cuda_pointcloud_preprocessor/src/cuda_pointcloud_preprocessor/common_kernels.cu b/sensing/autoware_cuda_pointcloud_preprocessor/src/cuda_pointcloud_preprocessor/common_kernels.cu
new file mode 100644
index 0000000000000..57792e175ddfe
--- /dev/null
+++ b/sensing/autoware_cuda_pointcloud_preprocessor/src/cuda_pointcloud_preprocessor/common_kernels.cu
@@ -0,0 +1,129 @@
+// Copyright 2025 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.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "autoware/cuda_pointcloud_preprocessor/common_kernels.hpp"
+#include "autoware/cuda_pointcloud_preprocessor/point_types.hpp"
+#include "autoware/cuda_pointcloud_preprocessor/types.hpp"
+
+namespace autoware::cuda_pointcloud_preprocessor
+{
+__global__ void transformPointsKernel(
+  const InputPointType * __restrict__ input_points, InputPointType * output_points, int num_points,
+  TransformStruct transform)
+{
+  int idx = blockIdx.x * blockDim.x + threadIdx.x;
+  if (idx < num_points) {
+    output_points[idx] = input_points[idx];
+
+    const float x = input_points[idx].x;
+    const float y = input_points[idx].y;
+    const float z = input_points[idx].z;
+
+    output_points[idx].x = transform.m11 * x + transform.m12 * y + transform.m13 * z + transform.x;
+    output_points[idx].y = transform.m21 * x + transform.m22 * y + transform.m23 * z + transform.y;
+    output_points[idx].z = transform.m31 * x + transform.m32 * y + transform.m33 * z + transform.z;
+  }
+}
+
+__global__ void cropBoxKernel(
+  InputPointType * __restrict__ d_points, std::uint32_t * __restrict__ output_mask, int num_points,
+  const CropBoxParameters * __restrict__ crop_box_parameters_ptr, int num_crop_boxes)
+{
+  for (int idx = blockIdx.x * blockDim.x + threadIdx.x; idx < num_points;
+       idx += blockDim.x * gridDim.x) {
+    const float x = d_points[idx].x;
+    const float y = d_points[idx].y;
+    const float z = d_points[idx].z;
+
+    std::uint32_t mask = 1;
+
+    for (int i = 0; i < num_crop_boxes; i++) {
+      const CropBoxParameters & crop_box_parameters = crop_box_parameters_ptr[i];
+      const float & min_x = crop_box_parameters.min_x;
+      const float & min_y = crop_box_parameters.min_y;
+      const float & min_z = crop_box_parameters.min_z;
+      const float & max_x = crop_box_parameters.max_x;
+      const float & max_y = crop_box_parameters.max_y;
+      const float & max_z = crop_box_parameters.max_z;
+      mask &=
+        (x <= min_x || x >= max_x) || (y <= min_y || y >= max_y) || (z <= min_z || z >= max_z);
+    }
+
+    output_mask[idx] = mask;
+  }
+}
+
+__global__ void combineMasksKernel(
+  const std::uint32_t * __restrict__ mask1, const std::uint32_t * __restrict__ mask2,
+  int num_points, std::uint32_t * __restrict__ output_mask)
+{
+  int idx = blockIdx.x * blockDim.x + threadIdx.x;
+  if (idx < num_points) {
+    output_mask[idx] = mask1[idx] & mask2[idx];
+  }
+}
+
+__global__ void extractPointsKernel(
+  InputPointType * __restrict__ input_points, std::uint32_t * __restrict__ masks,
+  std::uint32_t * __restrict__ indices, int num_points,
+  OutputPointType * __restrict__ output_points)
+{
+  int idx = blockIdx.x * blockDim.x + threadIdx.x;
+  if (idx < num_points && masks[idx] == 1) {
+    InputPointType & input_point = input_points[idx];
+    OutputPointType & output_point = output_points[indices[idx] - 1];
+    output_point.x = input_point.x;
+    output_point.y = input_point.y;
+    output_point.z = input_point.z;
+    output_point.intensity = input_point.intensity;
+    output_point.return_type = input_point.return_type;
+    output_point.channel = input_point.channel;
+  }
+}
+
+void transformPointsLaunch(
+  const InputPointType * input_points, InputPointType * output_points, int num_points,
+  TransformStruct transform, int threads_per_block, int blocks_per_grid, cudaStream_t & stream)
+{
+  transformPointsKernel<<<blocks_per_grid, threads_per_block, 0, stream>>>(
+    input_points, output_points, num_points, transform);
+}
+
+void cropBoxLaunch(
+  InputPointType * d_points, std::uint32_t * output_mask, int num_points,
+  const CropBoxParameters * crop_box_parameters_ptr, int num_crop_boxes, int threads_per_block,
+  int blocks_per_grid, cudaStream_t & stream)
+{
+  cropBoxKernel<<<blocks_per_grid, threads_per_block, 0, stream>>>(
+    d_points, output_mask, num_points, crop_box_parameters_ptr, num_crop_boxes);
+}
+
+void combineMasksLaunch(
+  const std::uint32_t * mask1, const std::uint32_t * mask2, int num_points,
+  std::uint32_t * output_mask, int threads_per_block, int blocks_per_grid, cudaStream_t & stream)
+{
+  combineMasksKernel<<<blocks_per_grid, threads_per_block, 0, stream>>>(
+    mask1, mask2, num_points, output_mask);
+}
+
+void extractPointsLaunch(
+  InputPointType * input_points, std::uint32_t * masks, std::uint32_t * indices, int num_points,
+  OutputPointType * output_points, int threads_per_block, int blocks_per_grid,
+  cudaStream_t & stream)
+{
+  extractPointsKernel<<<blocks_per_grid, threads_per_block, 0, stream>>>(
+    input_points, masks, indices, num_points, output_points);
+}
+
+}  // namespace autoware::cuda_pointcloud_preprocessor
diff --git a/sensing/autoware_cuda_pointcloud_preprocessor/src/cuda_pointcloud_preprocessor/cuda_pointcloud_preprocessor.cu b/sensing/autoware_cuda_pointcloud_preprocessor/src/cuda_pointcloud_preprocessor/cuda_pointcloud_preprocessor.cu
new file mode 100644
index 0000000000000..aa9f19b8824f6
--- /dev/null
+++ b/sensing/autoware_cuda_pointcloud_preprocessor/src/cuda_pointcloud_preprocessor/cuda_pointcloud_preprocessor.cu
@@ -0,0 +1,454 @@
+// Copyright 2025 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.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "autoware/cuda_pointcloud_preprocessor/common_kernels.hpp"
+#include "autoware/cuda_pointcloud_preprocessor/cuda_pointcloud_preprocessor.hpp"
+#include "autoware/cuda_pointcloud_preprocessor/organize_kernels.hpp"
+#include "autoware/cuda_pointcloud_preprocessor/outlier_kernels.hpp"
+#include "autoware/cuda_pointcloud_preprocessor/point_types.hpp"
+#include "autoware/cuda_pointcloud_preprocessor/types.hpp"
+#include "autoware/cuda_pointcloud_preprocessor/undistort_kernels.hpp"
+
+#include <Eigen/Core>
+#include <Eigen/Dense>
+#include <Eigen/Geometry>
+#include <cub/cub.cuh>
+
+#include <cuda_runtime.h>
+#include <tf2/utils.h>
+#include <thrust/execution_policy.h>
+#include <thrust/reduce.h>
+#include <thrust/scan.h>
+
+namespace autoware::cuda_pointcloud_preprocessor
+{
+
+CudaPointcloudPreprocessor::CudaPointcloudPreprocessor()
+{
+  sensor_msgs::msg::PointField x_field, y_field, z_field, intensity_field, return_type_field,
+    channel_field, azimuth_field, elevation_field, distance_field, time_stamp_field;
+  x_field.name = "x";
+  x_field.offset = 0;
+  x_field.datatype = sensor_msgs::msg::PointField::FLOAT32;
+  x_field.count = 1;
+
+  y_field.name = "y";
+  y_field.offset = 4;
+  y_field.datatype = sensor_msgs::msg::PointField::FLOAT32;
+  y_field.count = 1;
+
+  z_field.name = "z";
+  z_field.offset = 8;
+  z_field.datatype = sensor_msgs::msg::PointField::FLOAT32;
+  z_field.count = 1;
+
+  intensity_field.name = "intensity";
+  intensity_field.offset = 12;
+  intensity_field.datatype = sensor_msgs::msg::PointField::UINT8;
+  intensity_field.count = 1;
+
+  return_type_field.name = "return_type";
+  return_type_field.offset = 13;
+  return_type_field.datatype = sensor_msgs::msg::PointField::UINT8;
+  return_type_field.count = 1;
+
+  channel_field.name = "channel";
+  channel_field.offset = 14;
+  channel_field.datatype = sensor_msgs::msg::PointField::UINT16;
+  channel_field.count = 1;
+
+  static_assert(sizeof(OutputPointType) == 16, "OutputPointType size is not 16 bytes");
+  static_assert(offsetof(OutputPointType, x) == 0);
+  static_assert(offsetof(OutputPointType, y) == 4);
+  static_assert(offsetof(OutputPointType, z) == 8);
+  static_assert(offsetof(OutputPointType, intensity) == 12);
+  static_assert(offsetof(OutputPointType, return_type) == 13);
+  static_assert(offsetof(OutputPointType, channel) == 14);
+
+  point_fields_.push_back(x_field);
+  point_fields_.push_back(y_field);
+  point_fields_.push_back(z_field);
+  point_fields_.push_back(intensity_field);
+  point_fields_.push_back(return_type_field);
+  point_fields_.push_back(channel_field);
+
+  cudaStreamCreate(&stream_);
+
+  int num_sm;
+  cudaDeviceGetAttribute(&num_sm, cudaDevAttrMultiProcessorCount, 0);
+  max_blocks_per_grid_ = 4 * num_sm;  // used for strided loops
+
+  cudaMemPoolProps pool_props;
+  memset(&pool_props, 0, sizeof(cudaMemPoolProps));
+  pool_props.allocType = cudaMemAllocationTypePinned;
+  pool_props.handleTypes = cudaMemHandleTypePosixFileDescriptor;
+
+  pool_props.location.type = cudaMemLocationTypeDevice;
+  cudaGetDevice(&(pool_props.location.id));
+
+  cudaMemPoolCreate(&device_memory_pool_, &pool_props);
+  MemoryPoolAllocator<TwistStruct2D> allocator_2d(device_memory_pool_);
+  MemoryPoolAllocator<TwistStruct3D> allocator_3d(device_memory_pool_);
+  MemoryPoolAllocator<std::int32_t> allocator_int32(device_memory_pool_);
+  MemoryPoolAllocator<std::uint32_t> allocator_uint32(device_memory_pool_);
+  MemoryPoolAllocator<std::uint8_t> allocator_uint8(device_memory_pool_);
+  MemoryPoolAllocator<InputPointType> allocator_points(device_memory_pool_);
+
+  device_twist_2d_structs_ =
+    thrust::device_vector<TwistStruct2D, MemoryPoolAllocator<TwistStruct2D>>(allocator_2d);
+  device_twist_3d_structs_ =
+    thrust::device_vector<TwistStruct3D, MemoryPoolAllocator<TwistStruct3D>>(allocator_3d);
+
+  num_rings_ = 1;
+  max_points_per_ring_ = 1;
+  num_organized_points_ = num_rings_ * max_points_per_ring_;
+  device_ring_index_ =
+    thrust::device_vector<std::int32_t, MemoryPoolAllocator<std::int32_t>>(allocator_int32);
+  device_ring_index_.resize(num_rings_);
+
+  device_indexes_tensor_ =
+    thrust::device_vector<std::uint32_t, MemoryPoolAllocator<std::uint32_t>>(allocator_uint32);
+  device_sorted_indexes_tensor_ =
+    thrust::device_vector<std::uint32_t, MemoryPoolAllocator<std::uint32_t>>(allocator_uint32);
+
+  device_indexes_tensor_.resize(num_organized_points_);
+  device_sorted_indexes_tensor_.resize(num_organized_points_);
+
+  device_segment_offsets_ =
+    thrust::device_vector<std::int32_t, MemoryPoolAllocator<std::int32_t>>(allocator_int32);
+  device_segment_offsets_.resize(num_rings_ + 1);
+  device_segment_offsets_[0] = 0;
+  device_segment_offsets_[1] = 1;
+
+  device_max_ring_ =
+    thrust::device_vector<std::int32_t, MemoryPoolAllocator<std::int32_t>>(allocator_int32);
+  device_max_ring_.resize(1);
+
+  device_max_points_per_ring_ =
+    thrust::device_vector<std::int32_t, MemoryPoolAllocator<std::int32_t>>(allocator_int32);
+  device_max_points_per_ring_.resize(1);
+
+  device_input_points_ =
+    thrust::device_vector<InputPointType, MemoryPoolAllocator<InputPointType>>(allocator_points);
+  device_organized_points_ =
+    thrust::device_vector<InputPointType, MemoryPoolAllocator<InputPointType>>(allocator_points);
+  device_organized_points_.resize(num_organized_points_);
+
+  cudaMemsetAsync(
+    thrust::raw_pointer_cast(device_max_ring_.data()), 0, sizeof(std::int32_t), stream_);
+  cudaMemsetAsync(
+    thrust::raw_pointer_cast(device_max_points_per_ring_.data()), 0, sizeof(std::int32_t), stream_);
+  cudaMemsetAsync(
+    thrust::raw_pointer_cast(device_indexes_tensor_.data()), 0x255, sizeof(std::uint32_t), stream_);
+
+  sort_workspace_bytes_ = querySortWorkspace(
+    num_rings_ * max_points_per_ring_, num_rings_,
+    thrust::raw_pointer_cast(device_segment_offsets_.data()),
+    thrust::raw_pointer_cast(device_indexes_tensor_.data()),
+    thrust::raw_pointer_cast(device_sorted_indexes_tensor_.data()));
+
+  device_sort_workspace_ =
+    thrust::device_vector<std::uint8_t, MemoryPoolAllocator<std::uint8_t>>(allocator_uint8);
+
+  device_transformed_points_.resize(num_organized_points_);
+  device_crop_mask_.resize(num_organized_points_);
+  device_ring_outlier_mask_.resize(num_organized_points_);
+  device_indices_.resize(num_organized_points_);
+
+  preallocateOutput();
+}
+
+void CudaPointcloudPreprocessor::setCropBoxParameters(
+  const std::vector<CropBoxParameters> & crop_box_parameters)
+{
+  host_crop_box_structs_ = crop_box_parameters;
+  device_crop_box_structs_ = host_crop_box_structs_;
+}
+
+void CudaPointcloudPreprocessor::setRingOutlierFilterParameters(
+  const RingOutlierFilterParameters & ring_outlier_parameters)
+{
+  ring_outlier_parameters_ = ring_outlier_parameters;
+}
+
+void CudaPointcloudPreprocessor::setUndistortionType(const UndistortionType & undistortion_type)
+{
+  if (undistortion_type == UndistortionType::Invalid) {
+    throw std::runtime_error("Invalid undistortion type");
+  }
+
+  undistortion_type_ = undistortion_type;
+}
+
+void CudaPointcloudPreprocessor::preallocateOutput()
+{
+  output_pointcloud_ptr_ = std::make_unique<cuda_blackboard::CudaPointCloud2>();
+  output_pointcloud_ptr_->data = cuda_blackboard::make_unique<std::uint8_t[]>(
+    num_rings_ * max_points_per_ring_ * sizeof(OutputPointType));
+}
+
+void CudaPointcloudPreprocessor::organizePointcloud()
+{
+  cudaMemsetAsync(
+    thrust::raw_pointer_cast(device_ring_index_.data()), 0, num_rings_ * sizeof(std::int32_t),
+    stream_);
+  cudaMemsetAsync(
+    thrust::raw_pointer_cast(device_indexes_tensor_.data()), 0xFF,
+    num_organized_points_ * sizeof(std::uint32_t), stream_);
+
+  if (num_raw_points_ == 0) {
+    return;
+  }
+
+  const int raw_points_blocks_per_grid =
+    (num_raw_points_ + threads_per_block_ - 1) / threads_per_block_;
+
+  organizeLaunch(
+    thrust::raw_pointer_cast(device_input_points_.data()),
+    thrust::raw_pointer_cast(device_indexes_tensor_.data()),
+    thrust::raw_pointer_cast(device_ring_index_.data()), num_rings_,
+    thrust::raw_pointer_cast(device_max_ring_.data()), max_points_per_ring_,
+    thrust::raw_pointer_cast(device_max_points_per_ring_.data()), num_raw_points_,
+    threads_per_block_, raw_points_blocks_per_grid, stream_);
+
+  std::int32_t max_ring_value;
+  std::int32_t max_points_per_ring;
+
+  cudaMemcpyAsync(
+    &max_ring_value, thrust::raw_pointer_cast(device_max_ring_.data()), sizeof(std::int32_t),
+    cudaMemcpyDeviceToHost, stream_);
+  cudaMemcpyAsync(
+    &max_points_per_ring, thrust::raw_pointer_cast(device_max_points_per_ring_.data()),
+    sizeof(std::int32_t), cudaMemcpyDeviceToHost, stream_);
+  cudaStreamSynchronize(stream_);
+
+  if (max_ring_value >= num_rings_ || max_points_per_ring > max_points_per_ring_) {
+    num_rings_ = max_ring_value + 1;
+    max_points_per_ring_ = std::max((max_points_per_ring + 511) / 512 * 512, 512);
+    num_organized_points_ = num_rings_ * max_points_per_ring_;
+
+    device_ring_index_.resize(num_rings_);
+    device_indexes_tensor_.resize(num_organized_points_);
+    device_sorted_indexes_tensor_.resize(num_organized_points_);
+    device_segment_offsets_.resize(num_rings_ + 1);
+    device_organized_points_.resize(num_organized_points_);
+
+    device_transformed_points_.resize(num_organized_points_);
+    device_crop_mask_.resize(num_organized_points_);
+    device_ring_outlier_mask_.resize(num_organized_points_);
+    device_indices_.resize(num_organized_points_);
+
+    preallocateOutput();
+
+    std::vector<std::int32_t> segment_offsets_host(num_rings_ + 1);
+    for (std::size_t i = 0; i < num_rings_ + 1; i++) {
+      segment_offsets_host[i] = i * max_points_per_ring_;
+    }
+
+    cudaMemcpyAsync(
+      thrust::raw_pointer_cast(device_segment_offsets_.data()), segment_offsets_host.data(),
+      (num_rings_ + 1) * sizeof(std::int32_t), cudaMemcpyHostToDevice, stream_);
+
+    cudaMemsetAsync(
+      thrust::raw_pointer_cast(device_ring_index_.data()), 0, num_rings_ * sizeof(std::int32_t),
+      stream_);
+    cudaMemsetAsync(
+      thrust::raw_pointer_cast(device_indexes_tensor_.data()), 0xFF,
+      num_organized_points_ * sizeof(std::int32_t), stream_);
+
+    sort_workspace_bytes_ = querySortWorkspace(
+      num_organized_points_, num_rings_, thrust::raw_pointer_cast(device_segment_offsets_.data()),
+      thrust::raw_pointer_cast(device_indexes_tensor_.data()),
+      thrust::raw_pointer_cast(device_sorted_indexes_tensor_.data()));
+    device_sort_workspace_.resize(sort_workspace_bytes_);
+
+    organizeLaunch(
+      thrust::raw_pointer_cast(device_input_points_.data()),
+      thrust::raw_pointer_cast(device_indexes_tensor_.data()),
+      thrust::raw_pointer_cast(device_ring_index_.data()), num_rings_,
+      thrust::raw_pointer_cast(device_max_ring_.data()), max_points_per_ring_,
+      thrust::raw_pointer_cast(device_max_points_per_ring_.data()), num_raw_points_,
+      threads_per_block_, raw_points_blocks_per_grid, stream_);
+  }
+
+  if (num_organized_points_ == num_rings_) {
+    cudaMemcpyAsync(
+      thrust::raw_pointer_cast(device_sorted_indexes_tensor_.data()),
+      thrust::raw_pointer_cast(device_indexes_tensor_.data()),
+      num_organized_points_ * sizeof(std::uint32_t), cudaMemcpyDeviceToDevice, stream_);
+  } else {
+    cub::DeviceSegmentedRadixSort::SortKeys(
+      reinterpret_cast<void *>(thrust::raw_pointer_cast(device_sort_workspace_.data())),
+      sort_workspace_bytes_, thrust::raw_pointer_cast(device_indexes_tensor_.data()),
+      thrust::raw_pointer_cast(device_sorted_indexes_tensor_.data()), num_organized_points_,
+      num_rings_, thrust::raw_pointer_cast(device_segment_offsets_.data()),
+      thrust::raw_pointer_cast(device_segment_offsets_.data()) + 1, 0, sizeof(std::uint32_t) * 8,
+      stream_);
+  }
+
+  const int organized_points_blocks_per_grid =
+    (num_organized_points_ + threads_per_block_ - 1) / threads_per_block_;
+
+  gatherLaunch(
+    thrust::raw_pointer_cast(device_input_points_.data()),
+    thrust::raw_pointer_cast(device_sorted_indexes_tensor_.data()),
+    thrust::raw_pointer_cast(device_organized_points_.data()), num_rings_, max_points_per_ring_,
+    threads_per_block_, organized_points_blocks_per_grid, stream_);
+}
+
+std::unique_ptr<cuda_blackboard::CudaPointCloud2> CudaPointcloudPreprocessor::process(
+  const sensor_msgs::msg::PointCloud2::ConstSharedPtr input_pointcloud_msg_ptr,
+  const geometry_msgs::msg::TransformStamped & transform_msg,
+  const std::deque<geometry_msgs::msg::TwistWithCovarianceStamped> & twist_queue,
+  const std::deque<geometry_msgs::msg::Vector3Stamped> & angular_velocity_queue,
+  const std::uint32_t first_point_rel_stamp_nsec)
+{
+  auto frame_id = input_pointcloud_msg_ptr->header.frame_id;
+  num_raw_points_ = input_pointcloud_msg_ptr->width * input_pointcloud_msg_ptr->height;
+  num_organized_points_ = num_rings_ * max_points_per_ring_;
+
+  if (num_raw_points_ > device_input_points_.size()) {
+    std::size_t new_capacity = (num_raw_points_ + 1024) / 1024 * 1024;
+    device_input_points_.resize(new_capacity);
+  }
+
+  cudaMemcpyAsync(
+    thrust::raw_pointer_cast(device_input_points_.data()), input_pointcloud_msg_ptr->data.data(),
+    num_raw_points_ * sizeof(InputPointType), cudaMemcpyHostToDevice, stream_);
+
+  cudaStreamSynchronize(stream_);
+
+  organizePointcloud();
+
+  tf2::Quaternion rotation_quaternion(
+    transform_msg.transform.rotation.x, transform_msg.transform.rotation.y,
+    transform_msg.transform.rotation.z, transform_msg.transform.rotation.w);
+  tf2::Matrix3x3 rotation_matrix;
+  rotation_matrix.setRotation(rotation_quaternion);
+
+  TransformStruct transform_struct;
+  transform_struct.x = static_cast<float>(transform_msg.transform.translation.x);
+  transform_struct.y = static_cast<float>(transform_msg.transform.translation.y);
+  transform_struct.z = static_cast<float>(transform_msg.transform.translation.z);
+  transform_struct.m11 = static_cast<float>(rotation_matrix[0][0]);
+  transform_struct.m12 = static_cast<float>(rotation_matrix[0][1]);
+  transform_struct.m13 = static_cast<float>(rotation_matrix[0][2]);
+  transform_struct.m21 = static_cast<float>(rotation_matrix[1][0]);
+  transform_struct.m22 = static_cast<float>(rotation_matrix[1][1]);
+  transform_struct.m23 = static_cast<float>(rotation_matrix[1][2]);
+  transform_struct.m31 = static_cast<float>(rotation_matrix[2][0]);
+  transform_struct.m32 = static_cast<float>(rotation_matrix[2][1]);
+  transform_struct.m33 = static_cast<float>(rotation_matrix[2][2]);
+
+  // Twist preprocessing
+  std::uint64_t pointcloud_stamp_nsec = 1'000'000'000 * input_pointcloud_msg_ptr->header.stamp.sec +
+                                        input_pointcloud_msg_ptr->header.stamp.nanosec;
+
+  if (undistortion_type_ == UndistortionType::Undistortion3D) {
+    setupTwist3DStructs(
+      twist_queue, angular_velocity_queue, pointcloud_stamp_nsec, first_point_rel_stamp_nsec,
+      device_twist_3d_structs_, stream_);
+  } else if (undistortion_type_ == UndistortionType::Undistortion2D) {
+    setupTwist2DStructs(
+      twist_queue, angular_velocity_queue, pointcloud_stamp_nsec, first_point_rel_stamp_nsec,
+      device_twist_2d_structs_, stream_);
+  } else {
+    throw std::runtime_error("Invalid undistortion type");
+  }
+
+  // Obtain raw pointers for the kernels
+  TwistStruct2D * device_twist_2d_structs =
+    thrust::raw_pointer_cast(device_twist_2d_structs_.data());
+  TwistStruct3D * device_twist_3d_structs =
+    thrust::raw_pointer_cast(device_twist_3d_structs_.data());
+  InputPointType * device_transformed_points =
+    thrust::raw_pointer_cast(device_transformed_points_.data());
+  std::uint32_t * device_crop_mask = thrust::raw_pointer_cast(device_crop_mask_.data());
+  std::uint32_t * device_ring_outlier_mask =
+    thrust::raw_pointer_cast(device_ring_outlier_mask_.data());
+  std::uint32_t * device_indices = thrust::raw_pointer_cast(device_indices_.data());
+
+  const int blocks_per_grid = (num_organized_points_ + threads_per_block_ - 1) / threads_per_block_;
+
+  transformPointsLaunch(
+    thrust::raw_pointer_cast(device_organized_points_.data()), device_transformed_points,
+    num_organized_points_, transform_struct, threads_per_block_, blocks_per_grid, stream_);
+
+  int crop_box_blocks_per_grid = std::min(blocks_per_grid, max_blocks_per_grid_);
+  if (host_crop_box_structs_.size() > 0) {
+    cropBoxLaunch(
+      device_transformed_points, device_crop_mask, num_organized_points_,
+      thrust::raw_pointer_cast(device_crop_box_structs_.data()), host_crop_box_structs_.size(),
+      crop_box_blocks_per_grid, threads_per_block_, stream_);
+  } else {
+    thrust::fill(thrust::device, device_crop_mask, device_crop_mask + num_organized_points_, 1);
+  }
+
+  if (
+    undistortion_type_ == UndistortionType::Undistortion3D && device_twist_3d_structs_.size() > 0) {
+    undistort3DLaunch(
+      device_transformed_points, num_organized_points_, device_twist_3d_structs,
+      device_twist_3d_structs_.size(), threads_per_block_, blocks_per_grid, stream_);
+  } else if (
+    undistortion_type_ == UndistortionType::Undistortion2D && device_twist_2d_structs_.size() > 0) {
+    undistort2DLaunch(
+      device_transformed_points, num_organized_points_, device_twist_2d_structs,
+      device_twist_2d_structs_.size(), threads_per_block_, blocks_per_grid, stream_);
+  }
+
+  ringOutlierFilterLaunch(
+    device_transformed_points, device_ring_outlier_mask, num_rings_, max_points_per_ring_,
+    ring_outlier_parameters_.distance_ratio,
+    ring_outlier_parameters_.object_length_threshold *
+      ring_outlier_parameters_.object_length_threshold,
+    ring_outlier_parameters_.num_points_threshold, threads_per_block_, blocks_per_grid, stream_);
+
+  combineMasksLaunch(
+    device_crop_mask, device_ring_outlier_mask, num_organized_points_, device_ring_outlier_mask,
+    threads_per_block_, blocks_per_grid, stream_);
+
+  thrust::inclusive_scan(
+    thrust::device, device_ring_outlier_mask, device_ring_outlier_mask + num_organized_points_,
+    device_indices);
+
+  std::uint32_t num_output_points;
+  cudaMemcpyAsync(
+    &num_output_points, device_indices + num_organized_points_ - 1, sizeof(std::uint32_t),
+    cudaMemcpyDeviceToHost, stream_);
+  cudaStreamSynchronize(stream_);
+
+  if (num_output_points > 0) {
+    extractPointsLaunch(
+      device_transformed_points, device_ring_outlier_mask, device_indices, num_organized_points_,
+      reinterpret_cast<OutputPointType *>(output_pointcloud_ptr_->data.get()), threads_per_block_,
+      blocks_per_grid, stream_);
+  }
+
+  cudaStreamSynchronize(stream_);
+
+  // Copy the transformed points back
+  output_pointcloud_ptr_->row_step = num_output_points * sizeof(OutputPointType);
+  output_pointcloud_ptr_->width = num_output_points;
+  output_pointcloud_ptr_->height = 1;
+
+  output_pointcloud_ptr_->fields = point_fields_;
+  output_pointcloud_ptr_->is_dense = true;
+  output_pointcloud_ptr_->is_bigendian = input_pointcloud_msg_ptr->is_bigendian;
+  output_pointcloud_ptr_->point_step = sizeof(OutputPointType);
+  output_pointcloud_ptr_->header.stamp = input_pointcloud_msg_ptr->header.stamp;
+
+  return std::move(output_pointcloud_ptr_);
+}
+
+}  // namespace autoware::cuda_pointcloud_preprocessor
diff --git a/sensing/autoware_cuda_pointcloud_preprocessor/src/cuda_pointcloud_preprocessor/cuda_pointcloud_preprocessor_node.cpp b/sensing/autoware_cuda_pointcloud_preprocessor/src/cuda_pointcloud_preprocessor/cuda_pointcloud_preprocessor_node.cpp
new file mode 100644
index 0000000000000..b206cddb3e203
--- /dev/null
+++ b/sensing/autoware_cuda_pointcloud_preprocessor/src/cuda_pointcloud_preprocessor/cuda_pointcloud_preprocessor_node.cpp
@@ -0,0 +1,301 @@
+// Copyright 2025 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.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "autoware/cuda_pointcloud_preprocessor/cuda_pointcloud_preprocessor_node.hpp"
+
+#include "autoware/cuda_pointcloud_preprocessor/memory.hpp"
+#include "autoware/cuda_pointcloud_preprocessor/point_types.hpp"
+
+#include <autoware/point_types/types.hpp>
+
+#include <geometry_msgs/msg/transform_stamped.hpp>
+#include <sensor_msgs/point_cloud2_iterator.hpp>
+#include <tf2_geometry_msgs/tf2_geometry_msgs.hpp>
+
+#include <cuda_runtime.h>
+#include <pcl/pcl_base.h>
+#include <pcl/point_cloud.h>
+#include <pcl/point_types.h>
+
+#include <chrono>
+#include <string>
+#include <vector>
+
+namespace autoware::cuda_pointcloud_preprocessor
+{
+using sensor_msgs::msg::PointCloud2;
+
+CudaPointcloudPreprocessorNode::CudaPointcloudPreprocessorNode(
+  const rclcpp::NodeOptions & node_options)
+: Node("cuda_pointcloud_preprocessor", node_options),
+  tf2_buffer_(this->get_clock()),
+  tf2_listener_(tf2_buffer_)
+{
+  using std::placeholders::_1;
+
+  // Parameters
+  base_frame_ = declare_parameter<std::string>("base_frame");
+
+  RingOutlierFilterParameters ring_outlier_filter_parameters;
+  ring_outlier_filter_parameters.distance_ratio = declare_parameter<float>("distance_ratio");
+  ring_outlier_filter_parameters.object_length_threshold =
+    declare_parameter<float>("object_length_threshold");
+  ring_outlier_filter_parameters.num_points_threshold =
+    declare_parameter<int>("num_points_threshold");
+
+  const auto crop_box_min_x_vector = declare_parameter<std::vector<double>>("crop_box.min_x");
+  const auto crop_box_min_y_vector = declare_parameter<std::vector<double>>("crop_box.min_y");
+  const auto crop_box_min_z_vector = declare_parameter<std::vector<double>>("crop_box.min_z");
+
+  const auto crop_box_max_x_vector = declare_parameter<std::vector<double>>("crop_box.max_x");
+  const auto crop_box_max_y_vector = declare_parameter<std::vector<double>>("crop_box.max_y");
+  const auto crop_box_max_z_vector = declare_parameter<std::vector<double>>("crop_box.max_z");
+
+  /* *INDENT-OFF* */
+  if (
+    crop_box_min_x_vector.size() != crop_box_min_y_vector.size() ||
+    crop_box_min_x_vector.size() != crop_box_min_z_vector.size() ||
+    crop_box_min_x_vector.size() != crop_box_max_x_vector.size() ||
+    crop_box_min_x_vector.size() != crop_box_max_y_vector.size() ||
+    crop_box_min_x_vector.size() != crop_box_max_z_vector.size()) {
+    throw std::runtime_error("Crop box parameters must have the same size");
+  }
+  /* *INDENT-ON* */
+
+  std::vector<CropBoxParameters> crop_box_parameters;
+
+  for (std::size_t i = 0; i < crop_box_min_x_vector.size(); i++) {
+    CropBoxParameters parameters;
+    parameters.min_x = crop_box_min_x_vector[i];
+    parameters.min_y = crop_box_min_y_vector[i];
+    parameters.min_z = crop_box_min_z_vector[i];
+    parameters.max_x = crop_box_max_x_vector[i];
+    parameters.max_y = crop_box_max_y_vector[i];
+    parameters.max_z = crop_box_max_z_vector[i];
+    crop_box_parameters.push_back(parameters);
+  }
+
+  bool use_3d_undistortion = declare_parameter<bool>("use_3d_distortion_correction");
+  bool use_imu = declare_parameter<bool>("use_imu");
+
+  // Publisher
+  /* *INDENT-OFF* */
+  pub_ =
+    std::make_unique<cuda_blackboard::CudaBlackboardPublisher<cuda_blackboard::CudaPointCloud2>>(
+      *this, "~/output/pointcloud");
+  /* *INDENT-ON* */
+
+  // Subscriber
+  rclcpp::SubscriptionOptions sub_options;
+  sub_options.use_intra_process_comm = rclcpp::IntraProcessSetting::Enable;
+
+  pointcloud_sub_ = this->create_subscription<sensor_msgs::msg::PointCloud2>(
+    "~/input/pointcloud", rclcpp::SensorDataQoS{}.keep_last(1),
+    std::bind(&CudaPointcloudPreprocessorNode::pointcloudCallback, this, std::placeholders::_1),
+    sub_options);
+  twist_sub_ = this->create_subscription<geometry_msgs::msg::TwistWithCovarianceStamped>(
+    "~/input/twist", 10,
+    std::bind(&CudaPointcloudPreprocessorNode::twistCallback, this, std::placeholders::_1),
+    sub_options);
+
+  if (use_imu) {
+    imu_sub_ = this->create_subscription<sensor_msgs::msg::Imu>(
+      "~/input/imu", 10,
+      std::bind(&CudaPointcloudPreprocessorNode::imuCallback, this, std::placeholders::_1),
+      sub_options);
+  }
+
+  /* *INDENT-OFF* */
+  CudaPointcloudPreprocessor::UndistortionType undistortion_type =
+    use_3d_undistortion ? CudaPointcloudPreprocessor::UndistortionType::Undistortion3D
+                        : CudaPointcloudPreprocessor::UndistortionType::Undistortion2D;
+  /* *INDENT-ON* */
+
+  cuda_pointcloud_preprocessor_ = std::make_unique<CudaPointcloudPreprocessor>();
+  cuda_pointcloud_preprocessor_->setRingOutlierFilterParameters(ring_outlier_filter_parameters);
+  cuda_pointcloud_preprocessor_->setCropBoxParameters(crop_box_parameters);
+  cuda_pointcloud_preprocessor_->setUndistortionType(undistortion_type);
+
+  // initialize debug tool
+  {
+    using autoware::universe_utils::DebugPublisher;
+    using autoware::universe_utils::StopWatch;
+
+    stop_watch_ptr_ = std::make_unique<StopWatch<std::chrono::milliseconds>>();
+    debug_publisher_ = std::make_unique<DebugPublisher>(this, "cuda_pointcloud_preprocessor");
+    stop_watch_ptr_->tic("processing_time");
+  }
+}
+
+bool CudaPointcloudPreprocessorNode::getTransform(
+  const std::string & target_frame, const std::string & source_frame,
+  tf2::Transform * tf2_transform_ptr)
+{
+  if (target_frame == source_frame) {
+    tf2_transform_ptr->setOrigin(tf2::Vector3(0.0, 0.0, 0.0));
+    tf2_transform_ptr->setRotation(tf2::Quaternion(0.0, 0.0, 0.0, 1.0));
+    return true;
+  }
+
+  try {
+    const auto transform_msg =
+      tf2_buffer_.lookupTransform(target_frame, source_frame, tf2::TimePointZero);
+    tf2::convert(transform_msg.transform, *tf2_transform_ptr);
+  } catch (const tf2::TransformException & ex) {
+    RCLCPP_WARN(get_logger(), "%s", ex.what());
+    RCLCPP_ERROR(
+      get_logger(), "Please publish TF %s to %s", target_frame.c_str(), source_frame.c_str());
+
+    tf2_transform_ptr->setOrigin(tf2::Vector3(0.0, 0.0, 0.0));
+    tf2_transform_ptr->setRotation(tf2::Quaternion(0.0, 0.0, 0.0, 1.0));
+    return false;
+  }
+  return true;
+}
+
+void CudaPointcloudPreprocessorNode::twistCallback(
+  const geometry_msgs::msg::TwistWithCovarianceStamped::ConstSharedPtr twist_msg_ptr)
+{
+  twist_queue_.push_back(*twist_msg_ptr);
+
+  while (!twist_queue_.empty()) {
+    // for replay rosbag
+    bool backwards_time_jump_detected =
+      rclcpp::Time(twist_queue_.front().header.stamp) > rclcpp::Time(twist_msg_ptr->header.stamp);
+    bool is_queue_longer_than_1s =
+      rclcpp::Time(twist_queue_.front().header.stamp) <
+      rclcpp::Time(twist_msg_ptr->header.stamp) - rclcpp::Duration::from_seconds(1.0);
+
+    if (backwards_time_jump_detected) {
+      twist_queue_.clear();
+    } else if (is_queue_longer_than_1s) {
+      twist_queue_.pop_front();
+    } else {
+      break;
+    }
+  }
+}
+
+void CudaPointcloudPreprocessorNode::imuCallback(
+  const sensor_msgs::msg::Imu::ConstSharedPtr imu_msg)
+{
+  tf2::Transform imu_to_base_tf2{};
+  getTransform(base_frame_, imu_msg->header.frame_id, &imu_to_base_tf2);
+  geometry_msgs::msg::TransformStamped imu_to_base_msg;
+  imu_to_base_msg.transform.rotation = tf2::toMsg(imu_to_base_tf2.getRotation());
+
+  geometry_msgs::msg::Vector3Stamped angular_velocity;
+  angular_velocity.vector = imu_msg->angular_velocity;
+
+  geometry_msgs::msg::Vector3Stamped transformed_angular_velocity;
+  tf2::doTransform(angular_velocity, transformed_angular_velocity, imu_to_base_msg);
+  transformed_angular_velocity.header = imu_msg->header;
+  angular_velocity_queue_.push_back(transformed_angular_velocity);
+
+  /* *INDENT-OFF* */
+  while (!angular_velocity_queue_.empty()) {
+    // for rosbag replay
+    bool backwards_time_jump_detected = rclcpp::Time(angular_velocity_queue_.front().header.stamp) >
+                                        rclcpp::Time(imu_msg->header.stamp);
+    bool is_queue_longer_than_1s =
+      rclcpp::Time(angular_velocity_queue_.front().header.stamp) <
+      rclcpp::Time(imu_msg->header.stamp) - rclcpp::Duration::from_seconds(1.0);
+
+    if (backwards_time_jump_detected) {
+      angular_velocity_queue_.clear();
+    } else if (is_queue_longer_than_1s) {
+      angular_velocity_queue_.pop_front();
+    } else {
+      break;
+    }
+  }
+  /* *INDENT-ON* */
+}
+
+void CudaPointcloudPreprocessorNode::pointcloudCallback(
+  const sensor_msgs::msg::PointCloud2::ConstSharedPtr input_pointcloud_msg_ptr)
+{
+  if (!is_data_layout_compatible_with_point_xyzircaedt(input_pointcloud_msg_ptr->fields)) {
+    RCLCPP_ERROR(
+      get_logger(), "Input pointcloud data layout is not compatible with PointXYZIRCAEDT");
+  }
+
+  static_assert(
+    sizeof(InputPointType) == sizeof(autoware::point_types::PointXYZIRCAEDT),
+    "PointStruct and PointXYZIRCAEDT must have the same size");
+
+  stop_watch_ptr_->toc("processing_time", true);
+
+  // Make sure that the first twist is newer than the first point
+  sensor_msgs::PointCloud2ConstIterator<std::uint32_t> iter_stamp(
+    *input_pointcloud_msg_ptr, "time_stamp");
+
+  /* *INDENT-OFF* */
+  auto num_points = input_pointcloud_msg_ptr->width * input_pointcloud_msg_ptr->height;
+  std::uint32_t first_point_rel_stamp = num_points > 0 ? *iter_stamp : 0;
+  double first_point_stamp = input_pointcloud_msg_ptr->header.stamp.sec +
+                             input_pointcloud_msg_ptr->header.stamp.nanosec * 1e-9 +
+                             first_point_rel_stamp * 1e-9;
+
+  while (twist_queue_.size() > 1 &&
+         rclcpp::Time(twist_queue_.front().header.stamp).seconds() < first_point_stamp) {
+    twist_queue_.pop_front();
+  }
+
+  while (angular_velocity_queue_.size() > 1 &&
+         rclcpp::Time(angular_velocity_queue_.front().header.stamp).seconds() < first_point_stamp) {
+    angular_velocity_queue_.pop_front();
+  }
+  /* *INDENT-ON* */
+
+  // Obtain the base link to input pointcloud transform
+  geometry_msgs::msg::TransformStamped transform_msg;
+
+  try {
+    transform_msg = tf2_buffer_.lookupTransform(
+      base_frame_, input_pointcloud_msg_ptr->header.frame_id, tf2::TimePointZero);
+  } catch (const tf2::TransformException & ex) {
+    RCLCPP_WARN(get_logger(), "%s", ex.what());
+    return;
+  }
+
+  auto output_pointcloud_ptr = cuda_pointcloud_preprocessor_->process(
+    input_pointcloud_msg_ptr, transform_msg, twist_queue_, angular_velocity_queue_,
+    first_point_rel_stamp);
+  output_pointcloud_ptr->header.frame_id = base_frame_;
+
+  // Publish
+  pub_->publish(std::move(output_pointcloud_ptr));
+
+  if (debug_publisher_) {
+    const double processing_time_ms = stop_watch_ptr_->toc("processing_time", true);
+    debug_publisher_->publish<tier4_debug_msgs::msg::Float64Stamped>(
+      "debug/processing_time_ms", processing_time_ms);
+
+    double now_stamp_seconds = rclcpp::Time(this->get_clock()->now()).seconds();
+    double cloud_stamp_seconds = rclcpp::Time(input_pointcloud_msg_ptr->header.stamp).seconds();
+
+    debug_publisher_->publish<tier4_debug_msgs::msg::Float64Stamped>(
+      "debug/latency_ms", 1000.f * (now_stamp_seconds - cloud_stamp_seconds));
+  }
+
+  // Preallocate for next iteration
+  cuda_pointcloud_preprocessor_->preallocateOutput();
+}
+
+}  // namespace autoware::cuda_pointcloud_preprocessor
+
+#include "rclcpp_components/register_node_macro.hpp"
+RCLCPP_COMPONENTS_REGISTER_NODE(
+  autoware::cuda_pointcloud_preprocessor::CudaPointcloudPreprocessorNode)
diff --git a/sensing/autoware_cuda_pointcloud_preprocessor/src/cuda_pointcloud_preprocessor/organize_kernels.cu b/sensing/autoware_cuda_pointcloud_preprocessor/src/cuda_pointcloud_preprocessor/organize_kernels.cu
new file mode 100644
index 0000000000000..ea74709cb1b54
--- /dev/null
+++ b/sensing/autoware_cuda_pointcloud_preprocessor/src/cuda_pointcloud_preprocessor/organize_kernels.cu
@@ -0,0 +1,108 @@
+// Copyright 2025 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.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "autoware/cuda_pointcloud_preprocessor/organize_kernels.hpp"
+#include "autoware/cuda_pointcloud_preprocessor/point_types.hpp"
+
+#include <cub/cub.cuh>
+
+#include <limits>
+
+namespace autoware::cuda_pointcloud_preprocessor
+{
+
+__global__ void organizeKernel(
+  const InputPointType * __restrict__ input_points, std::uint32_t * index_tensor,
+  std::int32_t * ring_indexes, std::int32_t initial_max_rings, std::int32_t * output_max_rings,
+  std::int32_t initial_max_points_per_ring, std::int32_t * output_max_points_per_ring,
+  int num_points)
+{
+  int idx = blockIdx.x * blockDim.x + threadIdx.x;
+  if (idx >= num_points) {
+    return;
+  }
+
+  auto ring = input_points[idx].channel;
+
+  if (ring >= initial_max_rings) {
+    atomicMax(output_max_rings, ring);
+    return;
+  }
+
+  int next_offset = atomicAdd(&ring_indexes[ring], 1);
+
+  if (next_offset >= initial_max_points_per_ring) {
+    atomicMax(output_max_points_per_ring, next_offset);
+    return;
+  }
+
+  index_tensor[ring * initial_max_points_per_ring + next_offset] = idx;
+}
+
+__global__ void gatherKernel(
+  const InputPointType * __restrict__ input_points, const std::uint32_t * __restrict__ index_tensor,
+  InputPointType * __restrict__ output_points, int num_rings, int max_points_per_ring)
+{
+  const int idx = blockIdx.x * blockDim.x + threadIdx.x;
+  if (idx >= num_rings * max_points_per_ring) {
+    return;
+  }
+
+  const int ring = idx / max_points_per_ring;
+  const int point = idx % max_points_per_ring;
+
+  const std::uint32_t input_idx = index_tensor[ring * max_points_per_ring + point];
+
+  if (input_idx < std::numeric_limits<std::uint32_t>::max()) {
+    output_points[ring * max_points_per_ring + point] = input_points[input_idx];
+  } else {
+    output_points[ring * max_points_per_ring + point].distance = 0.0f;
+  }
+}
+
+std::size_t querySortWorkspace(
+  int num_items, int num_segments, int * offsets_device, std::uint32_t * keys_in_device,
+  std::uint32_t * keys_out_device)
+{
+  // Determine temporary device storage requirements
+  void * temp_storage = nullptr;
+  size_t temp_storage_bytes = 0;
+  cub::DeviceSegmentedRadixSort::SortKeys(
+    temp_storage, temp_storage_bytes, keys_in_device, keys_out_device, num_items, num_segments,
+    offsets_device, offsets_device + 1);
+
+  return temp_storage_bytes;
+}
+
+void organizeLaunch(
+  const InputPointType * input_points, std::uint32_t * index_tensor, std::int32_t * ring_indexes,
+  std::int32_t initial_max_rings, std::int32_t * output_max_rings,
+  std::int32_t initial_max_points_per_ring, std::int32_t * output_max_points_per_ring,
+  int num_points, int threads_per_block, int blocks_per_grid, cudaStream_t & stream)
+{
+  organizeKernel<<<blocks_per_grid, threads_per_block, 0, stream>>>(
+    input_points, index_tensor, ring_indexes, initial_max_rings, output_max_rings,
+    initial_max_points_per_ring, output_max_points_per_ring, num_points);
+}
+
+void gatherLaunch(
+  const InputPointType * input_points, const std::uint32_t * index_tensor,
+  InputPointType * output_points, int num_rings, int max_points_per_ring, int threads_per_block,
+  int blocks_per_grid, cudaStream_t & stream)
+{
+  gatherKernel<<<blocks_per_grid, threads_per_block, 0, stream>>>(
+    input_points, index_tensor, output_points, num_rings, max_points_per_ring);
+}
+
+}  // namespace autoware::cuda_pointcloud_preprocessor
diff --git a/sensing/autoware_cuda_pointcloud_preprocessor/src/cuda_pointcloud_preprocessor/outlier_kernels.cu b/sensing/autoware_cuda_pointcloud_preprocessor/src/cuda_pointcloud_preprocessor/outlier_kernels.cu
new file mode 100644
index 0000000000000..8ce0634ed2a41
--- /dev/null
+++ b/sensing/autoware_cuda_pointcloud_preprocessor/src/cuda_pointcloud_preprocessor/outlier_kernels.cu
@@ -0,0 +1,86 @@
+// Copyright 2025 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.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "autoware/cuda_pointcloud_preprocessor/outlier_kernels.hpp"
+#include "autoware/cuda_pointcloud_preprocessor/point_types.hpp"
+
+namespace autoware::cuda_pointcloud_preprocessor
+{
+
+__global__ void ringOutlierFilterKernel(
+  const InputPointType * points, std::uint32_t * output_mask, int num_rings,
+  int max_points_per_ring, float distance_ratio, float object_length_threshold_squared,
+  int num_points_threshold)
+{
+  int idx = blockIdx.x * blockDim.x + threadIdx.x;
+  int j = idx / max_points_per_ring;
+  int i = idx % max_points_per_ring;
+
+  if (j >= num_rings || i >= max_points_per_ring) {
+    return;
+  }
+
+  int min_i = max(i - num_points_threshold, 0);
+  int max_i = min(i + num_points_threshold, max_points_per_ring);
+
+  int walk_size = 1;
+  int left_idx = min_i;
+  int right_idx = min_i + 1;
+
+  for (int k = min_i; k < max_i - 1; k++) {
+    const InputPointType & left_point = points[j * max_points_per_ring + k];
+    const InputPointType & right_point = points[j * max_points_per_ring + k + 1];
+
+    // Find biggest walk that passes through i
+    float azimuth_diff = right_point.azimuth - left_point.azimuth;
+    azimuth_diff = azimuth_diff < 0.f ? azimuth_diff + 2 * M_PI : azimuth_diff;
+
+    if (
+      max(left_point.distance, right_point.distance) <
+        min(left_point.distance, right_point.distance) * distance_ratio &&
+      azimuth_diff < 1.f * M_PI / 180.f) {
+      // Determined to be included in the same walk
+      walk_size++;
+      right_idx++;
+    } else if (k >= i) {
+      break;
+    } else {
+      walk_size = 1;
+      left_idx = k + 1;
+      right_idx = k + 2;  // this is safe since we break if k >= i
+    }
+  }
+
+  const InputPointType & left_point = points[j * max_points_per_ring + left_idx];
+  const InputPointType & right_point = points[j * max_points_per_ring + right_idx - 1];
+  const float x = left_point.x - right_point.x;
+  const float y = left_point.y - right_point.y;
+  const float z = left_point.z - right_point.z;
+
+  output_mask[j * max_points_per_ring + i] = static_cast<std::uint32_t>(
+    (walk_size > num_points_threshold) ||
+    (x * x + y * y + z * z >= object_length_threshold_squared));
+}
+
+void ringOutlierFilterLaunch(
+  const InputPointType * points, std::uint32_t * output_mask, int num_rings,
+  int max_points_per_ring, float distance_ratio, float object_length_threshold_squared,
+  int num_points_threshold, int threads_per_block, int blocks_per_grid, cudaStream_t & stream)
+{
+  ringOutlierFilterKernel<<<blocks_per_grid, threads_per_block, 0, stream>>>(
+    points, output_mask, num_rings, max_points_per_ring, distance_ratio,
+    object_length_threshold_squared, num_points_threshold);
+}
+
+}  // namespace autoware::cuda_pointcloud_preprocessor
diff --git a/sensing/autoware_cuda_pointcloud_preprocessor/src/cuda_pointcloud_preprocessor/undistort_kernels.cu b/sensing/autoware_cuda_pointcloud_preprocessor/src/cuda_pointcloud_preprocessor/undistort_kernels.cu
new file mode 100644
index 0000000000000..8457cad79590d
--- /dev/null
+++ b/sensing/autoware_cuda_pointcloud_preprocessor/src/cuda_pointcloud_preprocessor/undistort_kernels.cu
@@ -0,0 +1,334 @@
+// Copyright 2025 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.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "autoware/cuda_pointcloud_preprocessor/point_types.hpp"
+#include "autoware/cuda_pointcloud_preprocessor/types.hpp"
+#include "autoware/cuda_pointcloud_preprocessor/undistort_kernels.hpp"
+
+#include <Eigen/Core>
+#include <Eigen/Dense>
+#include <Eigen/Geometry>
+
+namespace autoware::cuda_pointcloud_preprocessor
+{
+
+__host__ __device__ Eigen::Matrix3f skewSymmetric(const Eigen::Vector3f & v)
+{
+  Eigen::Matrix3f m;
+  m << 0, -v.z(), v.y(), v.z(), 0, -v.x(), -v.y(), v.x(), 0;
+  return m;
+}
+
+__host__ __device__ Eigen::Matrix3f leftJacobianSO3(const Eigen::Vector3f & omega)
+{
+  double theta = omega.norm();
+  if (std::abs(theta) < 1e-6) {
+    return Eigen::Matrix3f::Identity();
+  }
+
+  Eigen::Matrix3f Omega = skewSymmetric(omega);
+
+  Eigen::Matrix3f Omega2 = Omega * Omega;
+  double theta2 = theta * theta;
+  double theta3 = theta2 * theta;
+
+  // Rodrigues' formula for Jacobian
+  return Eigen::Matrix3f::Identity() + ((1 - cos(theta)) / theta2) * Omega +
+         ((theta - sin(theta)) / theta3) * Omega2;
+}
+
+__host__ __device__ Eigen::Matrix4f transformationMatrixFromVelocity(
+  const Eigen::Vector3f & linear_velocity, const Eigen::Vector3f & angular_velocity, double dt)
+{
+  Eigen::Matrix3f R = Eigen::AngleAxisf(angular_velocity.norm() * dt, angular_velocity.normalized())
+                        .toRotationMatrix();
+  Eigen::Matrix3f J = leftJacobianSO3(angular_velocity * dt);
+
+  Eigen::Vector3f translation = J * (linear_velocity * dt);
+
+  Eigen::Matrix4f transformation = Eigen::Matrix4f::Identity();
+  transformation.block<3, 3>(0, 0) = R;
+  transformation.block<3, 1>(0, 3) = translation;
+
+  return transformation;
+}
+
+__global__ void undistort2DKernel(
+  InputPointType * input_points, int num_points, TwistStruct2D * twist_structs, int num_twists)
+{
+  int idx = blockIdx.x * blockDim.x + threadIdx.x;
+  if (idx < num_points) {
+    InputPointType & point = input_points[idx];
+
+    // The twist must always be newer than the point ! (or it was the last twist)
+    int twist_index = 0;
+    while (twist_index < num_twists && twist_structs[twist_index].stamp_nsec < point.time_stamp) {
+      twist_index++;
+    }
+
+    twist_index = min(twist_index, num_twists - 1);
+
+    TwistStruct2D twist = twist_structs[twist_index];
+    float x = twist.cum_x;
+    float y = twist.cum_y;
+    float theta = twist.cum_theta;
+
+    double dt_nsec =
+      point.time_stamp > twist.last_stamp_nsec ? point.time_stamp - twist.last_stamp_nsec : 0;
+    double dt = 1e-9 * (dt_nsec);
+
+    theta += twist.v_theta * dt;
+    float d = twist.v_x * dt;
+    x += d * cos(theta);
+    y += d * sin(theta);
+
+    float distorted_x = point.x;
+    float distorted_y = point.y;
+
+    point.x = distorted_x * cos(theta) - distorted_y * sin(theta) + x;
+    point.y = distorted_x * sin(theta) + distorted_y * cos(theta) + y;
+  }
+}
+
+__global__ void undistort3DKernel(
+  InputPointType * input_points, int num_points, TwistStruct3D * twist_structs, int num_twists)
+{
+  int idx = blockIdx.x * blockDim.x + threadIdx.x;
+  if (idx < num_points) {
+    InputPointType & point = input_points[idx];
+
+    // The twist must always be newer than the point ! (or it was the last twist)
+    int twist_index = 0;
+    while (twist_index < num_twists && twist_structs[twist_index].stamp_nsec < point.time_stamp) {
+      twist_index++;
+    }
+
+    twist_index = min(twist_index, num_twists - 1);
+
+    TwistStruct3D twist = twist_structs[twist_index];
+    Eigen::Map<Eigen::Matrix4f> cum_transform_buffer_map(twist.cum_transform_buffer);
+    Eigen::Map<Eigen::Vector3f> v_map(twist.v);
+    Eigen::Map<Eigen::Vector3f> w_map(twist.w);
+
+    double dt_nsec =
+      point.time_stamp > twist.last_stamp_nsec ? point.time_stamp - twist.last_stamp_nsec : 0;
+    double dt = 1e-9 * (dt_nsec);
+
+    Eigen::Matrix4f transform =
+      cum_transform_buffer_map * transformationMatrixFromVelocity(v_map, w_map, dt);
+    Eigen::Vector3f p(point.x, point.y, point.z);
+    p = transform.block<3, 3>(0, 0) * p + transform.block<3, 1>(0, 3);
+
+    point.x = p.x();
+    point.y = p.y();
+    point.z = p.z();
+  }
+}
+
+void undistort2DLaunch(
+  InputPointType * input_points, int num_points, TwistStruct2D * twist_structs, int num_twists,
+  int threads_per_block, int blocks_per_grid, cudaStream_t & stream)
+{
+  undistort2DKernel<<<blocks_per_grid, threads_per_block, 0, stream>>>(
+    input_points, num_points, twist_structs, num_twists);
+}
+
+void undistort3DLaunch(
+  InputPointType * input_points, int num_points, TwistStruct3D * twist_structs, int num_twists,
+  int threads_per_block, int blocks_per_grid, cudaStream_t & stream)
+{
+  undistort3DKernel<<<blocks_per_grid, threads_per_block, 0, stream>>>(
+    input_points, num_points, twist_structs, num_twists);
+}
+
+void setupTwist2DStructs(
+  const std::deque<geometry_msgs::msg::TwistWithCovarianceStamped> & twist_queue,
+  const std::deque<geometry_msgs::msg::Vector3Stamped> & angular_velocity_queue,
+  const std::uint64_t pointcloud_stamp_nsec, const std::uint32_t first_point_rel_stamp_nsec,
+  thrust::device_vector<TwistStruct2D> & device_twist_2d_structs, cudaStream_t & stream)
+{
+  std::vector<TwistStruct2D> host_twist_2d_structs;
+  host_twist_2d_structs.reserve(twist_queue.size() + angular_velocity_queue.size());
+
+  // Twist preprocessing
+  float cum_x = 0;
+  float cum_y = 0;
+  float cum_theta = 0;
+  // All time stamps from now on are in nsec from the "beginning of the pointcloud"
+  std::uint32_t last_stamp_nsec = first_point_rel_stamp_nsec;
+
+  std::size_t twist_index = 0;
+  std::size_t angular_velocity_index = 0;
+  float v_x{0.f}, v_theta{0.f};
+
+  for (; twist_index < twist_queue.size() ||
+         angular_velocity_index < angular_velocity_queue.size();) {
+    std::uint64_t twist_stamp, input_twist_global_stamp_nsec, angular_velocity_global_stamp_nsec;
+
+    if (twist_index < twist_queue.size()) {
+      input_twist_global_stamp_nsec =
+        1'000'000'000 * static_cast<std::uint64_t>(twist_queue[twist_index].header.stamp.sec) +
+        static_cast<std::uint64_t>(twist_queue[twist_index].header.stamp.nanosec);
+      v_x = twist_queue[twist_index].twist.twist.linear.x;
+    } else {
+      input_twist_global_stamp_nsec = std::numeric_limits<std::uint64_t>::max();
+    }
+
+    if (angular_velocity_index < angular_velocity_queue.size()) {
+      angular_velocity_global_stamp_nsec =
+        1'000'000'000 * static_cast<std::uint64_t>(
+                          angular_velocity_queue[angular_velocity_index].header.stamp.sec) +
+        static_cast<std::uint64_t>(
+          angular_velocity_queue[angular_velocity_index].header.stamp.nanosec);
+      v_theta = angular_velocity_queue[angular_velocity_index].vector.z;
+    } else {
+      angular_velocity_global_stamp_nsec = std::numeric_limits<std::uint64_t>::max();
+    }
+
+    if (input_twist_global_stamp_nsec < angular_velocity_global_stamp_nsec) {
+      twist_stamp = input_twist_global_stamp_nsec;
+      twist_index++;
+    } else if (input_twist_global_stamp_nsec > angular_velocity_global_stamp_nsec) {
+      twist_stamp = angular_velocity_global_stamp_nsec;
+      angular_velocity_index++;
+    } else {
+      twist_stamp = input_twist_global_stamp_nsec;
+      twist_index++;
+      angular_velocity_index++;
+    }
+
+    TwistStruct2D twist;
+    twist.cum_x = cum_x;
+    twist.cum_y = cum_y;
+    twist.cum_theta = cum_theta;
+
+    std::uint64_t twist_global_stamp_nsec = twist_stamp;
+    assert(twist_global_stamp_nsec > pointcloud_stamp_nsec);  // by construction
+    std::uint32_t twist_from_pointcloud_start_nsec =
+      twist_global_stamp_nsec - pointcloud_stamp_nsec;
+
+    twist.stamp_nsec = twist_from_pointcloud_start_nsec;
+    twist.v_x = v_x;
+    twist.v_theta = v_theta;
+    twist.last_stamp_nsec = last_stamp_nsec;
+    host_twist_2d_structs.push_back(twist);
+
+    double dt_seconds = 1e-9 * (twist.stamp_nsec - last_stamp_nsec);
+    last_stamp_nsec = twist.stamp_nsec;
+    cum_theta += v_theta * dt_seconds;
+    float d = twist.v_x * dt_seconds;
+    cum_x += d * cos(cum_theta);
+    cum_y += d * sin(cum_theta);
+  }
+
+  // Copy to device
+  device_twist_2d_structs.resize(host_twist_2d_structs.size());
+  cudaMemcpyAsync(
+    thrust::raw_pointer_cast(device_twist_2d_structs.data()), host_twist_2d_structs.data(),
+    host_twist_2d_structs.size() * sizeof(TwistStruct2D), cudaMemcpyHostToDevice, stream);
+}
+
+void setupTwist3DStructs(
+  const std::deque<geometry_msgs::msg::TwistWithCovarianceStamped> & twist_queue,
+  const std::deque<geometry_msgs::msg::Vector3Stamped> & angular_velocity_queue,
+  const std::uint64_t pointcloud_stamp_nsec, const std::uint32_t first_point_rel_stamp_nsec,
+  thrust::device_vector<TwistStruct3D> & device_twist_3d_structs, cudaStream_t & stream)
+{
+  std::vector<TwistStruct3D> host_twist_3d_structs;
+  host_twist_3d_structs.reserve(twist_queue.size() + angular_velocity_queue.size());
+
+  Eigen::Matrix4f cum_transform = Eigen::Matrix4f::Identity();
+  Eigen::Vector3f v = Eigen::Vector3f::Zero();
+  Eigen::Vector3f w = Eigen::Vector3f::Zero();
+
+  // All time stamps from now on are in nsec from the "beginning of the pointcloud"
+  std::uint32_t last_stamp_nsec = first_point_rel_stamp_nsec;
+
+  std::size_t twist_index = 0;
+  std::size_t angular_velocity_index = 0;
+
+  for (; twist_index < twist_queue.size() ||
+         angular_velocity_index < angular_velocity_queue.size();) {
+    std::uint64_t twist_stamp, input_twist_global_stamp_nsec, angular_velocity_global_stamp_nsec;
+
+    if (twist_index < twist_queue.size()) {
+      input_twist_global_stamp_nsec =
+        1'000'000'000 * static_cast<std::uint64_t>(twist_queue[twist_index].header.stamp.sec) +
+        static_cast<std::uint64_t>(twist_queue[twist_index].header.stamp.nanosec);
+      v.x() = twist_queue[twist_index].twist.twist.linear.x;
+      v.y() = twist_queue[twist_index].twist.twist.linear.y;
+      v.z() = twist_queue[twist_index].twist.twist.linear.z;
+    } else {
+      input_twist_global_stamp_nsec = std::numeric_limits<uint64_t>::max();
+      v = Eigen::Vector3f::Zero();
+    }
+
+    if (angular_velocity_index < angular_velocity_queue.size()) {
+      angular_velocity_global_stamp_nsec =
+        1'000'000'000 * static_cast<std::uint64_t>(
+                          angular_velocity_queue[angular_velocity_index].header.stamp.sec) +
+        static_cast<std::uint64_t>(
+          angular_velocity_queue[angular_velocity_index].header.stamp.nanosec);
+      w.x() = angular_velocity_queue[angular_velocity_index].vector.x;
+      w.y() = angular_velocity_queue[angular_velocity_index].vector.y;
+      w.z() = angular_velocity_queue[angular_velocity_index].vector.z;
+    } else {
+      angular_velocity_global_stamp_nsec = std::numeric_limits<std::uint64_t>::max();
+      w = Eigen::Vector3f::Zero();
+    }
+
+    if (input_twist_global_stamp_nsec < angular_velocity_global_stamp_nsec) {
+      twist_stamp = input_twist_global_stamp_nsec;
+      twist_index++;
+    } else if (input_twist_global_stamp_nsec > angular_velocity_global_stamp_nsec) {
+      twist_stamp = angular_velocity_global_stamp_nsec;
+      angular_velocity_index++;
+    } else {
+      twist_index++;
+      angular_velocity_index++;
+    }
+
+    TwistStruct3D twist;
+
+    Eigen::Map<Eigen::Matrix4f> cum_transform_buffer_map(twist.cum_transform_buffer);
+    Eigen::Map<Eigen::Vector3f> v_map(twist.v);
+    Eigen::Map<Eigen::Vector3f> w_map(twist.w);
+    cum_transform_buffer_map = cum_transform;
+
+    std::uint64_t twist_global_stamp_nsec = twist_stamp;
+    assert(twist_global_stamp_nsec > pointcloud_stamp_nsec);  // by construction
+    std::uint32_t twist_from_pointcloud_start_nsec =
+      twist_global_stamp_nsec - pointcloud_stamp_nsec;
+
+    twist.stamp_nsec = twist_from_pointcloud_start_nsec;
+    v_map = v;
+    w_map = w;
+    twist.last_stamp_nsec = last_stamp_nsec;
+    host_twist_3d_structs.push_back(twist);
+
+    double dt_seconds = 1e-9 * (twist.stamp_nsec - last_stamp_nsec);
+    last_stamp_nsec = twist.stamp_nsec;
+
+    auto delta_transform = transformationMatrixFromVelocity(v, w, dt_seconds);
+    cum_transform = cum_transform * delta_transform;
+  }
+
+  // Copy to device
+  device_twist_3d_structs.resize(host_twist_3d_structs.size());
+  cudaMemcpyAsync(
+    thrust::raw_pointer_cast(device_twist_3d_structs.data()), host_twist_3d_structs.data(),
+    host_twist_3d_structs.size() * sizeof(TwistStruct3D), cudaMemcpyHostToDevice, stream);
+}
+
+}  // namespace autoware::cuda_pointcloud_preprocessor