NVIDIA GPUNetIO：让 GPU 直接发起网络通信

GPU 应用越来越需要让网络通信和数据移动成为由 GPU 控制的一等操作，而不是由主机驱动的服务。当 CPU 介入每一次网络事务时，就会成为关键路径上的瓶颈，增加延迟，并限制分布式应用实时响应的效率。GPU applications increasingly need networking and data movement to behave like first-class GPU-controlled operations rather than host-driven services. When the CPU sits in the middle of every network transaction, it becomes a bottleneck on the critical path, adding latency and limiting how efficiently distributed applications can respond in real time.

NVIDIA DOCA GPUNetIO 是一个面向 GPU 的网络 SDK 层，用于实时数据包处理和数据传输，能够直接解决这一问题。它整合了 GPUDirect RDMA、GPUDirect Async Kernel-Initiated (GDA-KI) 和 GDRCopy 等技术，使 CUDA 内核能够直接驱动以太网、RDMA、Verbs 和 DMA 操作，同时让 CPU 脱离应用程序的关键路径。NVIDIA DOCA GPUNetIO, a GPU-centric networking SDK layer for real-time packet processing and data movement, addresses this directly. It brings together technologies such as GPUDirect RDMA, GPUDirect Async Kernel-Initiated (GDA-KI) and GDRCopy so CUDA kernels can directly drive Ethernet, RDMA, Verbs, and DMA operations while keeping the CPU out of the application critical path.

具体来说，DOCA GPUNetIO 同时提供：Specifically, DOCA GPUNetIO provides both:

• CPU 控制路径上的函数，用于导出通过 DOCA Ethernet、DOCA Verbs、DOCA DMA 和 DOCA CommChannel 创建的 GPU 内存传输对象。
  CPU functions on the control path to export on GPU memory transport objects created with DOCA Ethernet, DOCA Verbs, DOCA DMA, DOCA CommChannel
• 在数据路径上提供 GPU CUDA 函数，以便创建能够操作已导出的传输对象的 CUDA 内核。
  GPU CUDA functions on the data path to allow the creation of CUDA kernels that can manipulate the transport objects exported

自 NVIDIA 首次通过 DOCA GPUNetIO 引入以 GPU 为中心的数据包处理以来，该框架已显著成熟，从一种能够将 CPU 移出关键路径的强大方案，发展为统一的 GDA-KI 基础，并集成到 NVSHMEM、NCCL、Aerial 5G SDK、UCX/NIXL、配备 Holoscan Sensor Bridge 算子的 NVQLink、Holoscan Advanced Network Operator、DeepEP/HybridEP 等框架中。Since NVIDIA first introduced GPU-centric packet processing with DOCA GPUNetIO, the framework has matured significantly, evolving from a powerful way to remove the CPU from the critical path into a unified GDA-KI foundation integrated across NVSHMEM, NCCL, Aerial 5G SDK, UCX/NIXL, NVQLink with Holoscan Sensor Bridge operator , Holoscan Advanced Network Operator , DeepEP/HybridEP, and others.

本文将介绍最新动态：开源库、统一架构，以及这些集成在实际中的运作方式。This post covers what’s new: the open-source library, the unified architecture, and how these integrations work in practice.

为什么统一的 GPUNetIO 体系至关重要Why a unified GPUNetIO umbrella matters

在 GPUNetIO 成为共享基础之前，每个通信库都各自实现了一套 GDA-KI 风格的 GPU 发起型 RDMA。它们彼此独立运行，各有各的假设、代码和维护负担，互不共享任何内容。Before GPUNetIO became the shared foundation, every communication library had built its own separate implementation of GDA-KI-style, GPU-initiated RDMA. Each one worked independently, with its own assumptions, its own code, and its own maintenance burden. None of them shared anything with the others.

将不同的 GDA-KI 风格实现纳入 GPUNetIO 体系下，具有显著优势：各个库无需再各自构建和维护基于 GDA-KI 的 RDMA 通信路径，而是可以统一采用 GPUNetIO 的单一 GDA-KI 实现，并在同一处共同使用、改进和扩展它。Bringing the different GDA-KI-style implementations under the GPUNetIO umbrella provides a great advantage: instead of each library building and maintaining its own GDA-KI-based RDMA communication path, they can converge on a single GPUNetIO GDA-KI implementation that everyone can use, improve, and extend in one place.

这为功能开发、优化和错误修复提供了一个共享的投入点，减少了各 SDK 之间重复的工程工作，也让针对某一框架开发的创新能够更快地应用到其他框架中。That creates a shared investment point for features, optimizations, and bug fixes, reduces duplicated engineering across SDKs, and allows innovations developed for one framework to become available to others much more quickly.

换句话说，GPUNetIO 将成为通用的 GDA-KI 基础，让多个通信库都能在其上构建，而不再是彼此平行、部分功能重叠的一系列实现。这种统一方案对 NIXL/UCX、NCCL 和 NVSHMEM 等更高层级的库尤其有价值。它们正协同合作，持续完善 GPUNetIO，并从整个技术栈的改进中受益。从生态系统的角度来看，这意味着更少的重复代码、更少的行为分化，也为长期演进铺就了更稳固的道路。In other words, GPUNetIO becomes the common GDA-KI foundation that multiple communication libraries can build on, rather than a collection of parallel, partially overlapping implementations. This unified approach is especially valuable for higher-level libraries such as NIXL/UCX, NCCL, and NVSHMEM. They are collaborating to enrich GPUNetIO, all taking and benefit from improvements made across the stack. From an ecosystem perspective, that means less duplicated code, less fragmentation in behavior, and a much stronger path for long-term evolution.

GPUNetIO：开源版与 SDK版GPUNetIO: Open Source vs SDK

NVIDIA 以两种密切相关的形式提供 GPUNetIO，因为整个生态系统既需要广泛的适用性，也需要开放性。完整的 DOCA SDK 版本是功能超集：它是《DOCA 编程指南》中介绍的 GPUNetIO 实现，涵盖更广泛的 DOCA 技术栈，包括但不限于 Verbs、Ethernet、DMA 和 Comm Channel 集成，同时保留了 GPUNetIO 的核心模式：让控制路径更靠近 GPU，并将 CPU 移出应用程序的关键路径。NVIDIA ships GPUNetIO in two closely related forms because the ecosystem needs both breadth and openness. The full DOCA SDK version is the superset: it is the GPUNetIO implementation documented in the DOCA Programming Guide and it spans the broader DOCA stack, including (but not limited to) Verbs, Ethernet, DMA, and Comm Channel integration, while preserving the core GPUNetIO model of moving the control path closer to the GPU and removing the CPU from the application critical path.

与此同时，NVIDIA 还发布了开源的 GPUNetIO 项目。这是一个更轻量、专注于 RDMA Verbs 的实现，面向希望在网络传输的集成方式上保持完全开放的框架。目前，开源版本是较小的 Verbs 导向子集，而 DOCA SDK 版本则是功能更全面的超集，增加了更丰富的 RDMA 支持、以太网和 DMA 等能力。In parallel, NVIDIA also publishes an open-source GPUNetIO project as a lighter-weight, RDMA-Verbs-focused implementation designed for frameworks that want to remain fully open in how they integrate networking transports. Today, the open version is the smaller Verbs-oriented subset, while the DOCA SDK version is the superset that adds broader capabilities such as richer RDMA support, Ethernet, and DMA.

这种拆分并非为了创建两套彼此分化的软件栈，而是为了提供一个用于 GPU 发起的 RDMA 通信的通用开源基础，同时在具备 DOCA SDK 时保留使用更高级功能的途径。事实上，开源实现可以在运行时检测是否存在 DOCA SDK；如果存在，就通过 dlopen 调用选定的闭源 DOCA SDK 函数；如果不存在，则继续使用开源实现运行。That split is not about creating two divergent software stacks. It is about providing a common open-source foundation for GPU-initiated RDMA communications while still leaving a path to more advanced functionality when the DOCA SDK is present. Indeed, the open-source implementation can detect the presence of DOCA SDK at runtime and, when available, call selected closed-source DOCA SDK functions through dlopen; if the SDK is absent, it continues to operate using the open-source implementation.

这正是其关键的架构优势：多个框架和库无需各自构建并维护私有版本的 GDAKI 式 RDMA 基础设施，而是可以共同采用、加固并改进由 GPUNetIO 提供的共享实现。This is the key architectural advantage: instead of every communication library building and maintaining its own private version of GDAKI-style RDMA plumbing, GPUNetIO becomes the shared implementation point that multiple frameworks and libraries can adopt, harden, and improve together.

在 CUDA 设备端，Verbs 路径上的差异有意保持得很小：开源版和 SDK 版本都提供面向设备的 API，因此即使主机端实现从轻量级开源路径扩展到功能更全面的 DOCA SDK，GPU 编程模型仍基本保持一致。On the CUDA device side, the gap is intentionally small for the Verbs path: both the open and SDK variants expose a device-facing API, which keeps the GPU programming model largely aligned even as the host-side implementation scales from a lightweight open-source path to the broader DOCA SDK feature set.

编程模型Programming model

本节介绍适用于任何 GPUNetIO 应用程序的关键概念和编程模型要素。This section covers the key concepts and programming model elements that apply to any GPUNetIO application.

CPU 控制路径CPU control path

在使用 DOCA GPUNetIO 的应用中，首先需要在主机端通过 CPU 执行初始配置阶段，以运行控制路径；随后进入第二阶段，由 GPU 通过 CUDA 内核运行数据路径。In applications that use DOCA GPUNetIO, an initial host-side CPU configuration phase is required to execute the control path, followed by a second phase in which the data path runs on the GPU through CUDA kernels.

控制路径工作流通常遵循以下顺序：The control-path workflow generally follows this sequence:

• GPU 和网络设备已初始化并配置完成，所需内存也已分配。
  The GPU and network devices are initialized and configured, and the required memory is allocated.
• 网络传输对象被创建。例如，DOCA Verbs 或 DOCA Ethernet 会在 CPU 上通过 mlx5dv 创建网络队列对象。
  Network transport objects are created. For example, DOCA Verbs or DOCA Ethernet creates a network queue object via mlx5dv on the CPU.
• GPUNetIO CPU 函数会将网络队列对象的相关元素导出到存储在 GPU 内存中的描述符中，并提供该描述符的 GPU 地址。
  A GPUNetIO CPU function exports the relevant elements of the network queue object into a descriptor stored in GPU memory and provides a GPU address for it.
• 应用程序启动一个 CUDA 内核，并将网络队列的 GPU 描述符作为其中一个输入参数传入。
  The application launches a CUDA kernel, passing the GPU descriptor for the network queue as one of the input parameters.

为了简化 DOCA Verbs 的部分操作，引入了一组 GPUNetIO 高级函数（例如 doca_gpu_verbs_create_qp_hl），以精简创建和连接 RDMA QP 所需的步骤。这些函数在 SDK 版本和开源版本中均可使用。To simplify some of these operations for DOCA Verbs, a set of GPUNetIO high-level functions (for example, doca_gpu_verbs_create_qp_hl) has been introduced to condense the steps required to create and connect RDMA QPs. These functions are present in both SDK and open-source versions.

控制路径完成后，数据路径即可在 GPU 上启动，此时一个或多个 CUDA 内核会使用 GPUNetIO CUDA 函数，操作导出到 GPU 内存中的传输对象，以发送或接收流量。Once the control path is complete, the data path can start on the GPU, where one or more CUDA kernels use GPUNetIO CUDA functions to operate on the transport objects exported into GPU memory to send or receive traffic.

GPU 数据路径GPU data path

用于 GPU 通信的 CUDA 内核通常包含以下步骤：The generic structure of a CUDA kernel for GPU communication typically consists of the following steps:

• 提交 WQE：一个或多个 CUDA 线程向网络队列对象提交工作队列条目（WQE），例如 RDMA Write、RDMA Read 或 Ethernet Send/Recv。
  Post WQEs: one or more CUDA threads post Work Queue Entries (WQEs), such as RDMA Write, RDMA Read, or Ethernet Send/Recv, to the network queue object.
• 敲响门铃：一个或多个 CUDA 线程通过向网卡寄存器写入数据，通知网卡有新的 WQE 已就绪（也就是“敲响门铃”）。
  Ring the doorbell: one or more CUDA threads notify the network card that new WQEs are ready by writing to its registers (that is, “ringing the doorbell”).
• 轮询 CQE（可选）：一个或多个 CUDA 线程等待完成队列条目（CQE），以确认工作队列条目（WQE）已成功完成。
  Poll CQEs (optional): one or more CUDA threads wait for Completion Queue Entries (CQEs) to confirm that the WQEs have completed successfully.

API：高级接口与底层接口API: High-level vs low-level

对于以太网和 RDMA Verbs 传输，DOCA GPUNetIO 提供两种不同层级的 API：高级 API 和低级 API。For Ethernet and RDMA Verbs transports, DOCA GPUNetIO offers two different levels of API: high-level and low-level.

高级 API 提供复杂操作和复合操作的实现。例如，在 Verbs 侧，doca_gpu_dev_verbs_put_signal() 预先实现了一种线程安全的组合操作：将 RDMA 写工作队列条目（WQE）与 RDMA 原子 Fetch & Add WQE 组合起来，并可按线程范围或 warp 范围执行（warp 中的所有线程协同执行该操作，提交 WARP_SIZE 个 RDMA 写操作，并在最后仅提交一个 RDMA 原子操作）。该函数会处理同一网络队列上其他 CUDA 线程并发提交 WQE 的情况，并敲响网卡门铃，以避免竞态条件。The high-level API offers implementation of complex and composite operations. As an example, on the Verbs side, the doca_gpu_dev_verbs_put_signal() offers a pre-implemented thread-safe combination of combining an RDMA Write Work Queue Entry (WQE) with RDMA Atomic Fetch & Add WQE executed per-thread scope or per-warp scope (all threads in the warp cooperate to the operation posting WARP_SIZE RDMA Write and just a single RDMA Atomic at the end). The function takes care about concurrent submissions of WQEs from other CUDA threads on the same network queue and rings the network card doorbell to avoid race conditions.

类似地，在以太网侧，另一个示例是 doca_gpu_dev_eth_txq_send()：它以线程安全的方式，将在网络队列上提交一个 Ethernet Send WQE 与随后敲响 doorbell 结合起来，并支持线程、warp 或 block 级别的操作。接收侧也是如此。Similarly, on the Ethernet side, another example is the doca_gpu_dev_eth_txq_send(), offering a thread-safe combination of posting an Ethernet Send WQE on the network queue and then ringing the doorbell at thread, warp or block scope. The same applies to the receiving side.

相比之下，低级 API 提供了基本构件，应用程序可以利用这些构件创建自定义的复合操作，例如发布具有不同操作码的 WQE、敲响网卡门铃，以及轮询完成队列以等待完成队列条目（CQE）。这些 API 不是线程安全的，因此，如果多个线程并发访问同一个网络队列对象，应用程序有责任进行适当的同步。The low-level API, by contrast, offers basic building blocks an application can use to create its own customized composite operations like posting WQEs with different opcodes, ringing the network card doorbell, poll the Completion Queue waiting for Completion Queue Entries (CQE). These APIs are not thread safe, so it’s the application’s responsibility to properly synchronize the concurrent access (if present) to the same network queue object.

按门铃Ring the doorbell

一旦将 WQE 提交到网络队列，就必须通知网卡，以便它执行这些 WQE。这一步称为“敲门铃”。在 GPUNetIO 中，它指的是 CUDA 内核通知网卡有新的工作可以执行。Once WQEs have been posted to the network queue, the network card must be notified so it can execute them. This step is known as “ringing the doorbell.” In GPUNetIO, it refers to the CUDA kernel notifying the network card that new work is ready for execution.

DOCA GPUNetIO 提供多种敲响门铃的方式：DOCA GPUNetIO offers several ways to ring the doorbell:

• 常规门铃：这是标准模式，在此模式下，NIC 寄存器通过 MMIO 映射到 CUDA 内存空间，使 CUDA 线程能够直接向其写入数据。
  Regular doorbell: This is the canonical mode, in which the NIC registers are MMIO-mapped into the CUDA memory space, allowing CUDA threads to write to them directly.
• BlueFlame：它遵循与常规门铃相同的总体模型，但在这种情况下，写入网卡寄存器的是整个 WQE，而不仅仅是一条通知。BlueFlame 通常用于对延迟敏感且网络队列数量较少的应用程序。
  BlueFlame: This follows the same general model as the regular doorbell, but in this case the entire WQE, rather than just a notification, is written into the NIC registers. BlueFlame is typically used in latency-sensitive applications with a small number of network queues.
• CPU 辅助门铃：在此模式下，NIC 寄存器映射到 CPU 内存，而非 GPU 内存。GPU 会向一个共享内存区域写入通知，该区域由 CPU 线程轮询。CPU 线程检测到 GPU 更新后，就会触发 NIC 门铃。此模式主要用于在没有 GPU 到 NIC 直连的系统中启用 GDA-KI，例如 DGX Spark。
  CPU-assisted doorbell: In this mode, the NIC register is mapped into CPU memory rather than GPU memory. The GPU writes a notification to a shared memory area that is polled by a CPU thread. Once the CPU thread detects the GPU update, it rings the NIC doorbell. This mode is mainly useful to enable GDA-KI in systems without a direct GPU-to-NIC connection, such as DGX Spark.

示例Examples

为便于探索 API 和进行性能基准测试，GPUNetIO 提供了多种参考实现。开发者可以直接在项目代码仓库中获取开源示例，而更全面的 SDK 示例和专用的基于以太网的应用程序则作为完整 DOCA SDK 的一部分发布。此外，在近期版本中，这些示例也已集成到 GitHub，方便用户访问。To facilitate both API exploration and performance benchmarking, GPUNetIO offers various reference implementations. Developers can access the open-source examples directly within the project repository, while the more comprehensive SDK samples and a dedicated Ethernet-based application are distributed as part of the full DOCA SDK. Furthermore, these samples have been integrated into GitHub for easier accessibility in recent releases.

以下部分重点展示了特定代码片段，以说明如何利用高级和低级原语，通过不同的编程语义实现类似的逻辑。The following sections highlight specific code snippets to illustrate how high-level and low-level primitives can be leveraged to implement similar logic using different programming semantics.

以太网示例Ethernet example

可以使用如下代码片段所示的高级 API 实现由 GPU 发起的以太网数据包生成器。具体来说，高级函数 `doca_gpu_dev_eth_txq_send` 可将多个 CUDA block 中的 WQE 并发提交到共享 Send Queue，并在内部处理必要的同步，从而确保线程安全，无需应用程序自行加锁。此外，它还封装了对队列结构的细粒度管理，例如跟踪 WQE 和 CQE 的具体索引。A GPU-initiated Ethernet packet generator can be implemented using high-level API like in the code snippet below. Specifically, the high-level function doca_gpu_dev_eth_txq_send facilitates the concurrent submission of WQEs across multiple CUDA blocks to a shared Send Queue, handling necessary synchronization internally to ensure thread-safe operation without requiring custom application-level locks. Also, it abstracts away the granular management of queue structures, such as tracking the specific indices of WQEs and CQEs.

__global__ void send_packets(struct doca_gpu_eth_txq *txq, uint8_t *addr, const uint32_t mkey, const size_t size, uint32_t *exit_cond)
{
	enum doca_gpu_eth_send_flags flags = DOCA_GPUNETIO_ETH_SEND_FLAG_NONE;
	doca_gpu_dev_eth_ticket_t out_ticket;
	uint32_t num_completed = 0;

	/* Only the last thread in the block requests a CQE. */
	if (threadIdx.x == (blockDim.x - 1))
		flags = DOCA_GPUNETIO_ETH_SEND_FLAG_NOTIFY;

	while (DOCA_GPUNETIO_VOLATILE(*exit_cond) == 0) {
		doca_gpu_dev_eth_txq_send<DOCA_GPUNETIO_ETH_RESOURCE_SHARING_MODE_GPU,
				  DOCA_GPUNETIO_ETH_SYNC_SCOPE_GPU,
				  DOCA_GPUNETIO_ETH_NIC_HANDLER_AUTO,
				  DOCA_GPUNETIO_ETH_EXEC_SCOPE_BLOCK>(txq, addr, mkey, size, flags, &out_ticket);

		/* __syncthreads already present in send with BLOCK scope */

		if (threadIdx.x == 0)
			doca_gpu_dev_eth_txq_poll_completion<DOCA_GPUNETIO_ETH_CQ_POLL_LAST>(txq, 1, DOCA_GPUNETIO_ETH_WAIT_FLAG_B, &num_completed);
		
		__syncthreads();
	}
}

应用程序也可以利用底层原语，对其实现进行精确控制。例如，如果数据包生成器为每个 CUDA 块恰好分配一个 Send Queue，那么高层函数 doca_gpu_dev_eth_txq_send 的内部同步开销就不再必要。在这种情况下，可以使用以下模式简化逻辑：Applications can also utilize low-level primitives to achieve precise control over their implementation. For instance, if a packet generator allocates exactly one Send Queue per CUDA block, the internal synchronization overhead of the high-level doca_gpu_dev_eth_txq_send becomes unnecessary. In such scenarios, the logic can be streamlined using the following pattern:

__global__ void send_packets(struct doca_gpu_eth_txq *txq, uint8_t *start_addr, const uint32_t mkey, const size_t size, uint32_t *exit_cond)
{
	uint64_t wqe_idx = threadIdx.x, cqe_idx = 0;
	enum doca_gpu_eth_send_flags flags = DOCA_GPUNETIO_ETH_SEND_FLAG_NONE;
	struct doca_gpu_dev_eth_txq_wqe *wqe_ptr;
	/* For simplicity, every thread always sends the same buffer */
	uint64_t addr = ((uint64_t)start_addr) + (uint64_t)(size * threadIdx.x);

	if (threadIdx.x == (blockDim.x - 1))
		flags = DOCA_GPUNETIO_ETH_SEND_FLAG_NOTIFY;

	while (DOCA_GPUNETIO_VOLATILE(*exit_cond) == 0) {
		wqe_ptr = doca_gpu_dev_eth_txq_get_wqe_ptr(txq, wqe_idx);
		doca_gpu_dev_eth_txq_wqe_prepare_send(txq, wqe_ptr, wqe_idx, addr, mkey, size, flags);
		__syncthreads();

		if (threadIdx.x == (blockDim.x - 1)) {
			/* Ring network card doorbell */
			doca_gpu_dev_eth_txq_submit(txq, wqe_idx + 1);

			/* Poll for last send completion */
doca_gpu_dev_eth_txq_poll_completion_at<DOCA_GPUNETIO_ETH_RESOURCE_SHARING_MODE_GPU, DOCA_GPUNETIO_ETH_SYNC_SCOPE_CTA>(txq, cqe_idx, DOCA_GPUNETIO_ETH_WAIT_FLAG_B);
			cqe_idx++;
		}

		__syncthreads();
		wqe_idx += blockDim.x;
	}

动词示例Verbs example

GPUNetIO Verbs API 遵循类似的范式。例如，开发者可以利用高级原语，通过以下 put 功能实现由 GPU 驱动的通信，从而构建一个性能导向型应用，复现 ib_write_bw 的逻辑：The GPUNetIO Verbs API follows a similar paradigm. For instance, developers can leverage high-level primitives to implement a performance-oriented application that replicates the logic of ib_write_bw using GPU-driven communication through the following put functionality:

template <enum doca_gpu_dev_verbs_exec_scope scope>
__global__ void put_bw(struct doca_gpu_dev_verbs_qp *qp, uint32_t num_iters, uint32_t data_size, uint8_t *src_buf, uint32_t src_buf_mkey, uint8_t *dst_buf, uint32_t dst_buf_mkey) {
	doca_gpu_dev_verbs_ticket_t out_ticket;
	uint32_t lane_idx = doca_gpu_dev_verbs_get_lane_id();
	uint32_t tidx = threadIdx.x + (blockIdx.x * blockDim.x);

	for (uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x; idx < num_iters; idx += (blockDim.x * gridDim.x)) {
		doca_gpu_dev_verbs_put<DOCA_GPUNETIO_VERBS_RESOURCE_SHARING_MODE_GPU, DOCA_GPUNETIO_VERBS_NIC_HANDLER_AUTO, scope>(qp, 
doca_gpu_dev_verbs_addr{.addr = (uint64_t)(dst_buf + (data_size * tidx)), .key = (uint32_t)dst_buf_mkey},
doca_gpu_dev_verbs_addr{.addr = (uint64_t)(src_buf + (data_size * tidx)), .key = (uint32_t)src_buf_mkey},
data_size, &out_ticket);

		__syncthreads();
	}
}

应用程序还可以利用底层原语精确控制其实现方式，例如创建一个 CUDA 块专用带宽测试，手动管理每一项操作：Applications can also utilize low-level primitives to achieve precise control over their implementation, such as creating a CUDA block-specific bandwidth test that manages every individual operation manually:

global void write_bw(struct doca_gpu_dev_verbs_qp *qp, uint32_t num_iters, uint32_t size, uint8_t *src_buf, uint32_t src_buf_mkey, uint8_t *dst_buf, uint32_t dst_buf_mkey) {
	uint64_t wqe_idx;
	struct doca_gpu_dev_verbs_wqe *wqe_ptr;

	for (uint32_t idx = threadIdx.x; idx < num_iters; idx += blockDim.x) {
		wqe_idx = doca_gpu_dev_verbs_reserve_wq_slots(qp, 1);
		wqe_ptr = doca_gpu_dev_verbs_get_wqe_ptr(qp, wqe_idx);

		doca_gpu_dev_verbs_wqe_prepare_write(qp, wqe_ptr, wqe_idx, MLX5_OPCODE_RDMA_WRITE, DOCA_GPUNETIO_MLX5_WQE_CTRL_CQ_UPDATE, 0,
			(uint64_t)(dst_buf + (size * threadIdx.x)), dst_buf_mkey,
			(uint64_t)(src_buf + (size * threadIdx.x)), src_buf_mkey, size);
		__syncthreads();

		if (threadIdx.x == (blockDim.x - 1))
			doca_gpu_dev_verbs_submit<DOCA_GPUNETIO_VERBS_RESOURCE_SHARING_MODE_EXCLUSIVE>(qp, (wqe_idx + 1));
		__syncthreads();

		wqe_idx += blockDim.x;
	}

	if (threadIdx.x == (blockDim.x - 1))
		doca_gpu_dev_verbs_poll_cq_at(doca_gpu_dev_verbs_qp_get_cq_sq(qp), (wqe_idx - blockDim.x));

	__syncthreads();
}

用于 GIN 的 NCCL 设备 APINCCL Device API for GIN

NCCL GIN（GPU 发起的网络通信）为 CUDA 内核提供设备端通信抽象层，并暴露 put、get、signal、wait 和 flush 等原语。自 NCCL 2.27 起，GIN 已集成 GDA-KI 作为后端，并利用开源的 GPUNetIO Verbs 路径。该架构使 NCCL 算法和设备 API 应用能够直接从 GPU 发起 RDMA 操作，同时通过 GIN 接口屏蔽底层队列管理和工作提交的复杂性。NCCL GIN (GPU-Initiated Networking) delivers the device-side communication abstraction layer for CUDA kernels, exposing primitives such as put, get, signal, wait, and flush. Since NCCL 2.27, GIN has integrated GDA-KI as a backend leveraging the open-source GPUNetIO Verbs path. This architecture enables NCCL algorithms and device-API applications to drive RDMA operations directly from the GPU while abstracting away the low-level queue management and work-submission complexities beneath the GIN interface.

工作流始于 CPU 在通信器初始化期间执行控制路径，在此过程中处理 RDMA 队列对的创建、内存注册，以及将传输描述符导出到 GPU 内存。配置完成后，数据路径完全转移到 GPU。GIN 操作会映射为 GPUNetIO Verbs 调用，由其负责准备 WQE、敲响门铃并轮询完成状态。这种模式成功地将 CPU 从每一笔网络事务的应用关键路径中移除。The workflow begins with the CPU executing the control path during communicator initialization, where it handles RDMA queue pair creation, memory registration, and the export of transport descriptors to GPU memory. Once configured, the data path shifts entirely to the GPU. GIN operations are mapped into GPUNetIO Verbs calls, which manage WQE preparation, doorbell ringing, and completion polling. This model successfully removes the CPU from the application critical path for every network transaction.

通过将通信语义与传输实现解耦，NCCL 专注于集体通信算法，而 GPUNetIO 则提供经过加固的 GDA-KI 底层支持。NCCL 利用 GIN 级别的控制机制管理线程、CTA 或整个 GPU 范围内的资源共享，从而实现请求聚合和队列管理优化。这些能力与 GPUNetIO 的优势相得益彰，无需拆分 RDMA 实现即可实现低延迟、高吞吐量的数据传输路径。归根结底，GPUNetIO 构成了共享基础，让性能增强和 NIC 支持只需开发一次，便可应用于整个 GPU 通信生态系统。By decoupling communication semantics from transport implementation, NCCL focuses on collective algorithms while GPUNetIO provides the hardened GDA-KI plumbing. NCCL utilizes GIN-level controls to manage thread, CTA, or GPU-wide resource sharing, enabling request aggregation and optimized queue management. These capabilities align with GPUNetIO’s strengths, facilitating low-latency and high-throughput paths without fragmenting the RDMA implementation. Ultimately, GPUNetIO serves as the shared foundation where performance enhancements and NIC support can be developed once and utilized across the entire GPU communication ecosystem.

要探索 NCCL GIN 接口的功能实现，开发者可以访问 NCCL tests GitHub 仓库，查阅参考代码。To explore functional implementations of the NCCL GIN interface, developers can access the NCCL tests GitHub repository for reference code.

NVSHMEMNVSHMEM

NVSHMEM 是一种基于 PGAS 的编程模型，可在 GPU 集群中提供可扩展的点对点和集体通信原语，例如 put、get、原子操作、屏障和归约。它使多 GPU 应用能够在 CUDA 内核中实现细粒度的 GPU 间数据移动与同步，从而显著提升应用工作负载的强缩放能力。NVSHMEM is a PGAS-based programming model that provides scalable point-to-point and collective communication primitives, such as put, get, atomics, barriers, and reductions, in GPU clusters. It enables multi-GPU applications to implement fine-grained GPU-to-GPU data movement and synchronization within a CUDA kernel, providing significant gains in strong scaling of application workloads.

在网络通信方面，NVSHMEM 根据可用的网络能力和应用需求提供多种传输方式。NVSHMEM 提供 IBGDA 传输方式，使用 MLX5 direct verbs 实现 GDA-KI。For network communication, NVSHMEM offers several transports depending on the available network capabilities and application requirements. NVSHMEM provides an IBGDA transport, which implements GDA-KI using MLX5 direct verbs.

随着开源 GPUNetIO 库的发布，NVSHMEM 3.7 引入了基于 GPUNetIO 的新传输方式。通过基于这一通用实现层构建，NVSHMEM 可受益于 GPUNetIO 的优化，以及通过 DOCA SDK 提供的高级 NIC 功能。新的 GPUNetIO 传输方式支持通过 RDMA 进行 GDA-KI 通信和主机发起的通信（又称 CPU 通信）。在后一种场景中，GPUNetIO 仅用于控制路径，以配置网络组件；数据路径则通过 NVSHMEM 原生函数在 CPU 上执行。用户可通过设置 NVSHMEM_REMOTE_TRANSPORT=gpunetio 来选择 GPUNetIO 传输方式，并通过设置 NVSHMEM_GPUNETIO_ENABLE_GDAKI=1 来启用 GDA-KI。With the release of the open-source GPUNetIO library, NVSHMEM 3.7 introduced a new GPUNetIO-based transport. By building on this common implementation layer, NVSHMEM benefits from GPUNetIO optimizations and advanced NIC features available through the DOCA SDK. The new GPUNetIO transport supports both GDA-KI communication and host-initiated communication (a.k.a CPU communications) via RDMA (in this second scenario, GPUNetIO is only used on the control path to configure network elements while the data path is executed on the CPU via NVSHMEM native functions). Users can select the GPUNetIO transport by setting NVSHMEM_REMOTE_TRANSPORT=gpunetio and enable GDA-KI with NVSHMEM_GPUNETIO_ENABLE_GDAKI=1.

NVSHMEM 中 GPUNetIO 传输层的 GDA-KI 部分沿用了 IBGDA 传输层的结构，但将较大部分的底层代码替换为对 GPUNetIO 库的调用。例如，如上文“敲响门铃”一节所述，门铃的构造和敲响代码已完全交由 GPUNetIO 库实现。在实现与 IBGDA 相同的功能和性能的同时，GPUNetIO 传输层显著降低了传输层实现的复杂性，彰显了整合 GDA-KI 实现中通用功能的价值。The GDA-KI part of the GPUNetIO transport in NVSHMEM follows the structure of the IBGDA transport but replaces larger low-level code portions with calls to the GPUNetIO library. For example, as described in the “Ring the doorbell” section above, the doorbell construction and ringing code was entirely externalized into the GPUNetIO library. While achieving the same functionality and performance as IBGDA, the GPUNetIO transport reduces the complexity of the transport implementation significantly, showcasing the value of consolidating common functionality of the GDA-KI implementations.

根据通信模式的不同，启用 GDA-KI 可能会显著提高小消息尺寸下可达到的带宽。以下实验使用 NVSHMEM 性能测试套件中的 shmem_put_bw 基准测试，通过 nvshmem_double_put_nbi 发送不同消息尺寸的数据。实验使用 1 个线程和每个 CTA（协作线程阵列）1 个 QP，并逐步增加 CTA 的数量。下方两张图展示了 GPUNetIO 传输达到的带宽：第一张图为禁用 GDA-KI（即使用 CPU 通信）时的结果，第二张图为使用 GDA-KI 时的结果。Depending on the communication pattern, enabling GDA-KI may significantly enhance the achievable bandwidth for small message sizes. The following experiment uses the shmem_put_bw benchmark of the NVSHMEM performance test suite using nvshmem_double_put_nbi with different message sizes. The experiment uses 1 thread and 1 QP per CTA (cooperative thread array) with an increasing number of CTAs. The two plots below show the achieved bandwidth of the GPUNetIO transport with the first one showing the results with GDA-KI disabled, i.e., using CPU communications, and the second one showing the results using GDA-KI.

如上图 5 所示，当 NVSHMEM 使用 CPU 通信时（控制 CPU 路径配置为使用 GPUNetIO 函数，而 CPU 数据路径使用 NVSHMEM 原生实现），在扩展到更多 CTA 和 QP 时，由于 CPU 代理成为瓶颈，小消息的带宽会受到限制。而图 6 中则使用启用了 GDA-KI 特性的 GPUNetIO 后端执行相同的基准测试（数据路径由 GPU 使用 GPUNetIO CUDA 函数执行）。As shown in Figure 5, above, when NVSHMEM uses CPU communications (control CPU path is configured with GPUNetIO functions while CPU data path is NVSHMEM native), bandwidth is capped for small message sizes when scaling to more CTAs and QPs due to the CPU proxy bottleneck. In the Figure 6 instead, the same benchmark is executed with GPUNetIO backend enabling GDA-KI feature (data path executed by GPU with GPUNetIO CUDA functions).

该实验证实，使用 GDA-KI 移除代理瓶颈后，CTA 和 QP 的扩展表现更好。与非 GDA-KI 变体相比，GDA-KI 变体在较小的消息尺寸下即可达到更高的峰值带宽，并且在消息尺寸更小时便能达到峰值带宽。此前比较 IBRC 和 IBGDA 带宽性能的实验也得出了类似结果。The experiment confirms that the removed proxy bottleneck when using GDA-KI leads to a better CTA and QP scaling behavior. The GDA-KI variant can reach higher peak bandwidth for smaller message sizes and achieves peak bandwidth earlier in terms of message size than the non-GDA-KI variant. Previous experiments comparing the bandwidth performance of IBRC and IBGDA have shown similar results.

总体而言，GPUNetIO 在保留 IBGDA 和 IBDEVX 的功能与性能的同时，大幅降低了 NVSHMEM 的实现复杂度和维护成本。Overall, GPUNetIO preserves the functionality and performance of IBGDA and IBDEVX while substantially reducing NVSHMEM’s implementation complexity and maintenance costs.

NVQLinkNVQLink

NVIDIA NVQLink 是 NVIDIA 为连接量子处理器控制系统与 GPU 计算而设计的低延迟参考架构。它与 NVIDIA CUDA-Q 紧密集成，支持量子纠错（QEC）、前馈和自适应校准等实时量子—经典工作流。Holoscan Sensor Bridge（HSB）通过 RoCE、轻量级 FPGA IP 和主机端控制软件，为这一方案提供网络基础；而 NVQLink 则将这种基于 HSB 的模式推进到延迟更低的范围，适用于实时 QPU 控制，而非传统的传感器数据摄取流水线。NVIDIA NVQLink is NVIDIA’s low-latency reference architecture for connecting quantum processors control systems to GPU computing, enabling real-time quantum-classical workflows such as Quantum Error Correction (QEC), feed-forward, and adaptive calibration through tight integration with NVIDIA CUDA-Q. Holoscan Sensor Bridge (HSB) provides the networking foundation for this approach by using RoCE and lightweight FPGA IP plus host-side control software, while NVQLink pushes that HSB-based model into a much lower-latency regime suitable for live QPU control rather than conventional sensor-ingest pipelines.

NVQLink 使用的 HSB 网络算子称为 GPU RoCE Transceiver。它基于 DOCA GPUNetIO 和 DOCA Verbs 库构建，旨在尽可能降低流量转发延迟。在独立模式下，该算子为网络健全性检查和性能分析提供流量转发执行路径。The HSB network operator used by NVQLink is called GPU RoCE Transceiver. Built on the DOCA GPUNetIO and DOCA Verbs libraries, it is designed to minimize traffic-forwarding latency as much as possible. In standalone mode, the operator provides a traffic-forwarding execution path for network sanity checks and performance analysis.

操作员启动一个持久化 CUDA 内核，该内核使用 DOCA GPUNetIO CUDA 函数持续接收来自远程对等端的数据包，并在不需要 CPU 介入的情况下将其发回。在此设置中，远程对等端是一块 FPGA，用于测量数据包发送到接收返回之间经过的时间。The operator launches a persistent CUDA kernel that uses DOCA GPUNetIO CUDA functions to continuously receive packets from a remote peer and send them back without any CPU intervention. In this setup, the remote peer is an FPGA that measures the elapsed time between when a packet is transmitted and when it is received back.

为了测量 HSB 2.7.0 GPU RoCE Transceiver 算子的网络转发延迟，我们使用了搭载 Blackwell GPU 和 ConnectX-7 网卡的 NVIDIA IGX Thor 平台。To measure the network-forwarding latency of the HSB 2.7.0 GPU RoCE Transceiver operator, we used an NVIDIA IGX Thor platform with a Blackwell GPU and a ConnectX-7 network card.

在另一端，我们通过一根 OSFP 电缆进行背靠背连接，并使用 Xilinx RFSoC 4×2 开发板上的 FPGA 发送和接收数据，同时测量往返延迟。On the other side, connected back-to-back over an OSFP cable, we used an FPGA on a Xilinx RFSoC 4×2 board to send, receive, and measure round-trip latency.

测得的最低延迟（即数据包在 FPGA 处从发送到接收的时间）约为 2.6 微秒，中位延迟为 2.7 微秒。The minimum reported latency, measured as the time between packet transmission and reception at the FPGA, is approximately 2.6 microseconds while median latency is 2.7 microseconds.

开始使用 DOCA GPUNetIOGet started with DOCA GPUNetIO

• DOCA SDK GPUNetIO programming guide — full SDK documentation
• GPUNetIO 开源代码库——示例与开源实现
  Open-source GPUNetIO repository — examples and open-source implementation
• GitHub 上的 DOCA SDK 示例 — 参考实现
  DOCA SDK samples on GitHub — reference implementations
• NCCL 测试 — GIN 设备 API — GIN 参考代码
  NCCL tests — GIN device API — GIN reference code
• Holoscan 传感器桥接 GPU RoCE 收发器 — NVQLink 算子
  Holoscan Sensor Bridge GPU RoCE Transceiver — NVQLink operator