Tutorial - Subscribe to the ZED images with CUDA zero-copy

This tutorial explains how to receive the ZED images directly in GPU memory with the CUDA zero-copy transport of ROS 2 Lyrical, and how to use them in your own accelerated ROS 2 node. It is the Isaac™ ROS 5 replacement for the NITROS subscriber tutorial of Isaac™ ROS 3.

The example creates a ROS 2 component that subscribes to a ZED image topic twice: first accepting only standard CPU data, then accepting the CUDA buffers published by the ZED node. It measures the latency, the frame rate, and the CPU and GPU load of both methods, and compares them.

Read the CUDA zero-copy transport page first for an introduction to how the transport works.

Set up the example

1. Set up the CUDA zero-copy transport in your Isaac™ ROS 5 environment by following the Set up the CUDA zero-copy transport section: install the cuda_buffer package, build the cuda_buffer_backend plugin, and build the ZED ROS 2 Wrapper.

2. If the ZED ROS 2 Examples are not in your workspace yet, clone them in the src folder:

   cd ${ISAAC_ROS_WS}/src
   git clone https://github.com/stereolabs/zed-ros2-examples.git

3. Build the example:

   cd ${ISAAC_ROS_WS}
   colcon build --symlink-install --cmake-args -DCMAKE_BUILD_TYPE=Release \
     --packages-up-to zed_wrapper zed_isaac_ros_cuda_img_sub
   source install/local_setup.bash

Subscribe to CUDA buffers in a ROS 2 node

The example node is a ROS 2 component implemented in C++, in the folder isaac_ros/isaac_ros_5/zed_isaac_ros_cuda_img_sub/src/component of the ZED ROS 2 Examples repository. The component is the class ZedCudaImgSubComponent, which inherits from rclcpp::Node.

Accept the CUDA buffer backend

The subscription is a standard sensor_msgs/msg/Image subscription. The acceptable_buffer_backends option declares which buffer backends the node can handle:

void ZedCudaImgSubComponent::create_subscriber(Phase phase)
{
[...]
  rclcpp::SubscriptionOptions sub_opts;
  sub_opts.acceptable_buffer_backends = (phase == Phase::CUDA) ? "cuda" : "cpu";
  _sub = this->create_subscription<sensor_msgs::msg::Image>(
    "image", _defaultQoS,
    std::bind(&ZedCudaImgSubComponent::sub_callback, this, std::placeholders::_1),
    sub_opts);
}

• "cpu" is the ROS 2 default: the node always receives CPU data.
• "cuda" accepts the CUDA buffers. CPU data is always accepted too, so the node also works when the zero-copy transport cannot be used, for example when the ZED node runs on another host.

The subscription QoS _defaultQoS keeps only the latest message (history depth 1). Use the same depth in the nodes that receive CUDA buffers.

Read the GPU memory of the message

Because the backend is negotiated for each connection, always check the backend of each received message with get_backend_type() before accessing its data. The function process_on_gpu() gets a GPU pointer to the image in both cases:

#include <cuda_buffer/cuda_buffer_api.hpp>
[...]
  if (img.data.get_backend_type() == "cuda") {
    try {
      // Zero-copy: the read handle waits for the GPU work of the publisher on
      // `_stream` and gives read-only access to the GPU memory of the message.
      // Use the pointer only while the handle is alive.
      cuda_buffer_backend::ReadHandle rh =
        cuda_buffer_backend::from_input_buffer(img.data, _stream);
      return process(rh.get_ptr());
    } catch (const std::exception & e) {
      RCLCPP_ERROR(get_logger(), "CUDA buffer error: %s", e.what());
      return false;
    }
  }
  // CPU data: the buffer behaves like a std::vector<uint8_t> and must be
  // uploaded to GPU memory before being processed
  [...]
  cudaMemcpyAsync(_devBuffer, img.data.data(), size, cudaMemcpyHostToDevice, _stream);
  return process(_devBuffer);

• cuda_buffer_backend::from_input_buffer() returns a ReadHandle. Its constructor makes the CUDA stream wait for the GPU work of the ZED node, so the data is complete when your kernels run on the same stream. Its destructor records the end of your work, so the ZED node does not reuse the memory too early.
• rh.get_ptr() returns a const uint8_t * device pointer: the memory belongs to the ZED node and is read-only.
• On a CPU message, img.data behaves like a std::vector<uint8_t>. Calling data() or the [] operator on a CUDA message throws an exception.

The process() lambda stands for the GPU processing of a real application: it copies the central pixel of the image back to CPU memory. Replace it with your CUDA kernels, or pass the device pointer to a CUDA-accelerated library.

Link your component to the cuda_buffer package in the CMakeLists.txt file:

find_package(cuda_buffer REQUIRED)
find_package(CUDAToolkit REQUIRED)
[...]
target_link_libraries(${PROJECT_NAME}_component
  [...]
  cuda_buffer::cuda_buffer
  CUDA::cudart
)

Then declare the dependency in the package.xml file:

<depend>cuda_buffer</depend>

How the example works

The node runs two benchmark phases on the same topic:

1. CPU: the subscription accepts only CPU data. The middleware copies the ZED image from GPU to CPU memory, and the node uploads it back to GPU memory to process it, as an accelerated node does with standard messages.
2. CUDA: the subscription accepts the cuda backend. The node processes the GPU memory of the ZED node directly, without copies.

Where the CUDA zero-copy transport makes the difference

Figure 1: the two benchmark phases of the example. Green arrows carry data in GPU memory, orange arrows mark the copies, the serialization, and the upload of the full image.

The numbers below match the badges of Figure 1.

1. CUDA phase (cross-process). The middleware sends only a small descriptor of the GPU memory where the ZED SDK wrote the image. The example node maps the same memory and reads it through a ReadHandle: the image is neither copied nor serialized.
2. CPU phase, copy to CPU memory. The subscription does not accept the cuda backend, so the publisher middleware copies the full image from GPU to CPU memory, 3.7 MB for each HD720 bgra8 image.
3. CPU phase, serialization and transfer. The middleware serializes the full image, transfers it to the example process, and deserializes it.
4. CPU phase, upload. The example node uploads the image back to GPU memory with cudaMemcpyAsync before processing it.

The bottom of Figure 1 shows the latency reduction measured with the example (see Fast DDS and Zenoh): the CUDA zero-copy transport gives its best with large images and with the Zenoh middleware. The subscribers that read CPU data, such as RViz2, ros2 bag, and Python nodes, keep following the CPU path and are not affected.

Each phase receives the number of messages set by the parameter benchmark.tot_samples, and measures:

• The latency: the time from the publishing of the message to the availability of the image on the GPU of the subscriber.
• The period between two messages, and the corresponding average frame rate.
• The CPU load and the GPU load of the whole system while the node is receiving the messages.

At the end of the CUDA phase, the node prints the results, compares the two phases, and shuts down. The results are also saved in the CSV file set by the parameter benchmark.csv_log_file.

The node also counts the CPU messages received in the CUDA phase: a value different from 0 means that the zero-copy transport is not available, and the node logs a warning with the possible causes.

Advanced configuration of the example

The folder config contains the file example_params.yaml with the parameters of the example.

The section general overrides the configuration of the ZED ROS 2 Wrapper node:

• grab_resolution: The native camera grab resolution. ‘HD1200’, ‘HD2K’, ‘HD1080’, ‘HD720’, ‘SVGA’, ‘VGA’, ‘AUTO’.
• pub_resolution: The resolution used for image and depth map publishing. ‘NATIVE’ to use the same general.grab_resolution; ‘CUSTOM’ to apply the general.pub_downscale_factor downscale factor.
• pub_downscale_factor: Downscale factor applied to the images before publishing, when pub_resolution is ‘CUSTOM’.
• pub_frame_rate: Publishing frequency of the images and of the depth data. Set to 0 for no limit.
• buffer_backend: cuda to publish the images with the CUDA zero-copy transport, cpu to always publish CPU data.

The section benchmark configures the benchmark:

• tot_samples: Number of samples to process for each benchmark phase, from 1 to 1000, 300 in example_params.yaml.
• cpu_gpu_load_period: Period in milliseconds to retrieve CPU and GPU load statistics, from 10 to 1000.
• cpu_gpu_load_avg_wnd_size: Number of samples to use to compute the average CPU and GPU load, from 1 to 100.
• csv_log_file: If not empty, log the benchmark results in a CSV file with the specified name.

A value out of range or of a wrong type is ignored: the node logs a warning with the valid range and uses its built-in default value.

The section debug is used to debug the nodes:

• debug_cuda: Enable the debug information of the example, for example the value of the central pixel of each image.
• use_pub_timestamps: If true, the ZED node sets the timestamps of the messages to the system time just before publishing, to measure the communication latency. If false, the latency is measured from the time the data is grabbed from the camera.

Run the example

Make sure you built the example and the ZED ROS 2 Wrapper as described in Set up the example.

Launch the example, replacing <camera_model> with the model of your camera:

ros2 launch zed_isaac_ros_cuda_img_sub zed_cuda_img_sub_example.launch.py camera_model:=<camera_model>

By default, the example subscribes to the rectified color image /zed_isaac/zed/rgb/color/rect/image. Use the topic_name argument to subscribe to another ZED image topic, for example the depth map:

ros2 launch zed_isaac_ros_cuda_img_sub zed_cuda_img_sub_example.launch.py \
  camera_model:=<camera_model> \
  topic_name:=/zed_isaac/zed/depth/depth_registered

To run the example without a camera, use the svo_path argument to replay an SVO recording. Use the camera_model of the camera that recorded the file. The SVO file is not replayed in a loop: for the benchmark to complete, it must contain more than two times benchmark.tot_samples frames, 600 frames by default, because a few frames are lost when the subscriber connects and when it changes phase. Record at least 30 seconds at 30 FPS:

ros2 launch zed_isaac_ros_cuda_img_sub zed_cuda_img_sub_example.launch.py \
  camera_model:=<camera_model> \
  svo_path:=<full_path_to_svo_file>

Example output

We executed the example with the camera at HD720 @ 60 FPS. While the example is running, it prints a line for each received message, with the buffer backend of the message:

[...] [zed_isaac.zed_cuda_img_sub_component]:  *** [CPU] Sample 150/300 - backend: cpu - FPS: 60.5 - Latency: 0.0013 sec - CPU: 2.8% - GPU: 5.0%
[...]
[...] [zed_isaac.zed_cuda_img_sub_component]:  *** [CUDA] Sample 150/300 - backend: cuda - FPS: 59.8 - Latency: 0.0007 sec - CPU: 1.4% - GPU: 5.0%

At the end of the benchmark, the example prints the results. This is the content of the CSV log file of the same run:

BENCHMARK, SAMPLES, UNITS, AVG, MIN, MAX, STDDEV
CPU Subscriber Period, 299, sec, 0.0168656, 0.00342222, 0.0332761, 0.00297642
CPU Subscriber Latency, 300, sec, 0.00181679, 0.00124442, 0.0143422, 0.00182357
CPU Subscriber CPU Load, 300, %, 2.9741, 0, 11.094, 1.81834
CPU Subscriber GPU Load, 300, %, 5.35067, 0, 7, 1.48754
CUDA Subscriber Period, 299, sec, 0.0168303, 0.00918486, 0.0567544, 0.00257081
CUDA Subscriber Latency, 300, sec, 0.000986213, 0.00065735, 0.0416681, 0.00236553
CUDA Subscriber CPU Load, 300, %, 3.1569, 0, 10.8403, 2.26447
CUDA Subscriber GPU Load, 300, %, 5.39733, 5, 7, 0.781106
 * CPU Subscriber Latency vs CUDA Subscriber Latency
   * Avg: 0.00181679 sec vs 0.000986213 sec
     - CPU Subscriber Latency is higher by 0.000830572 sec (84.2183%)
 * CPU Subscriber Period vs CUDA Subscriber Period
   * Avg: 0.0168656 sec vs 0.0168303 sec
     - CPU Subscriber Period is higher by 3.52887e-05 sec (0.209674%)
 * CPU Subscriber CPU Load vs CUDA Subscriber CPU Load
   * Avg: 2.9741 % vs 3.1569 %
     - CUDA Subscriber CPU Load is higher by 0.182795 % (6.14622%)
 * CPU Subscriber GPU Load vs CUDA Subscriber GPU Load
   * Avg: 5.35067 % vs 5.39733 %
     - CUDA Subscriber GPU Load is higher by 0.0466667 % (0.872165%)

If you set the CSV log file name without a path, the file is saved in the folder where you launched the example. Each run appends its results to the file.

With the CUDA zero-copy transport, the image is available on the GPU of the subscriber with a lower latency: the copy from GPU to CPU memory by the publisher middleware, the serialization of the image to send it to the process of the subscriber, and the upload back to GPU memory in the subscriber are no longer required. In this run, both methods kept the camera frame rate. The CPU and GPU loads are measured on the whole system, including the ZED SDK, so they vary between runs: compare the two phases of the same run.

The gain grows with the size of the data. With the same setup at HD1080 @ 30 FPS, the average latency dropped from 3.27 - 4.32 ms to 1.32 - 1.35 ms (65% lower). The results also depend on the camera, the resolution, and the load of the system.

Fast DDS and Zenoh

The CUDA zero-copy transport works with both the rmw_fastrtps_cpp (Fast DDS) and the rmw_zenoh_cpp (Zenoh) middleware: see Use Zenoh with Isaac™ ROS 5. We measured both middlewares with the Bare Metal installation, in two runs of each configuration. Zenoh used the recommended setting for large messages, with the default configuration otherwise (shared memory disabled). Fast DDS used its default configuration. These are the average latencies of the example:

With both middlewares, all the images were received as CUDA buffers in the CUDA phase, and the system CPU and GPU loads were similar. The CUDA phase kept the camera frame rate in all the runs. With Zenoh at HD1080, the CPU phase received 28.5 FPS instead of 30.

• With CPU data, Fast DDS has the lowest latency. Prefer it when most of the subscribers of the ZED topics read CPU data, such as RViz2, ros2 bag, or Python nodes.
• With the CUDA zero-copy transport, Zenoh has the lowest latency, so it benefits the most from the transport. Consider it when the subscribers of the ZED topics are mainly accelerated nodes, or when you need its network features, such as the discovery without UDP multicast.

The point cloud tutorial shows the same comparison with point clouds of 1.8 MB and 14.7 MB.

The launch file explained

The launch file is located in the folder launch and is named zed_cuda_img_sub_example.launch.py. It starts two ROS 2 component containers, the ZED node, and the example node.

Retrieve the path to the configuration file:

    example_params_file = os.path.join(
        get_package_share_directory('zed_isaac_ros_cuda_img_sub'),
        'config',
        'example_params.yaml'
    )

Create two ROS 2 component containers, /zed_isaac/zed_container for the ZED node and /zed_isaac/cuda_sub_container for the example node:

    container_name = 'zed_container'
    sub_container_name = 'cuda_sub_container'
    for name in [container_name, sub_container_name]:
        [...]
        actions.append(ComposableNodeContainer(
            name=name,
            namespace=namespace_val,
            package='rclcpp_components',
            executable='component_container',
            arguments=['--executor-type', 'multi-threaded',
                       '--ros-args', '--log-level', 'info',
                       '--log-level', 'serialize_buffer_with_endpoint:=warn',
                       '--log-level', 'deserialize_buffer_with_endpoint:=warn'],
            output='screen',
        ))

The serialize_buffer_with_endpoint and deserialize_buffer_with_endpoint loggers of the ROS 2 buffer typesupport print several INFO lines for each CUDA buffer message, thousands of lines per minute. The containers raise their level to WARN to keep the log readable.

The two nodes run in different processes, and the cuda buffer backend shares the ZED images between them through CUDA IPC, without copies. Do not load the nodes that receive the CUDA buffers in the container of the ZED node: see the warning about the queue depth and the processes.

Call the ZED ROS 2 Wrapper launch file to start the ZED node in its container, passing the camera model, the namespace, and the parameters file:

    zed_wrapper_launch = IncludeLaunchDescription(
        launch_description_source=PythonLaunchDescriptionSource([
            get_package_share_directory('zed_wrapper'),
            '/launch/zed_camera.launch.py'
        ]),
        launch_arguments={
            'camera_model': camera_model,
            'container_name': container_name,
            'namespace': namespace_val,
            'enable_ipc': 'false',
            'ros_params_override_path': example_params_file
        }.items()
    )

Load the example node in its own container, remapping its image topic to the topic to subscribe to:

    cuda_sub_node = ComposableNode(
        package='zed_isaac_ros_cuda_img_sub',
        plugin='stereolabs::ZedCudaImgSubComponent',
        name='zed_cuda_img_sub_component',
        namespace=namespace_val,
        remappings=[
            ('image', topic_name)
        ],
        parameters=[
            example_params_file
        ]
    )
    [...]
    load_sub_node = LoadComposableNodes(
        composable_node_descriptions=[cuda_sub_node],
        target_container=namespace_val + '/' + sub_container_name
    )

Conclusion

This tutorial showed how to receive the ZED images directly in GPU memory with the CUDA zero-copy transport of ROS 2 Lyrical: subscribe to the standard ZED topics, accept the cuda buffer backend, check the backend of each message, and read the GPU memory through a ReadHandle.

Compared to standard CPU data, the CUDA zero-copy transport removes the copies between GPU and CPU memory and the serialization of the images, which lowers the latency. The same node keeps working with CPU data when the zero-copy transport is not available.

To receive the ZED point cloud in GPU memory and read the coordinates of its points, follow the point cloud CUDA zero-copy tutorial.