DOCA GPUNetIO 如何统一 NVIDIA 软件栈中的 GPU 发起网络

GPU应用对网络和数据传输的需求日益增长，要求这些操作像GPU主导的一级操作一样高效，而非由主机驱动的服务。当CPU介入每一次网络事务时，它会成为关键路径上的瓶颈，增加延迟，并限制分布式应用在实时响应方面的效率。

NVIDIA DOCA GPUNetIO是一个面向实时数据包处理和数据传输的GPU中心化网络SDK层，直接解决了这一问题。它整合了GPUDirect RDMA、GPUDirect Async Kernel-Initiated (GDA-KI) 和 GDRCopy等技术，使CUDA内核能够直接驱动以太网、RDMA、Verbs和DMA操作，同时将CPU排除在应用关键路径之外。

具体而言，DOCA GPUNetIO提供以下两方面功能：

• 控制路径上的CPU函数，用于导出通过DOCA Ethernet、DOCA Verbs、DOCA DMA和DOCA CommChannel创建的GPU内存传输对象
• 数据路径上的GPU CUDA函数，允许创建能够操作上述导出传输对象的CUDA内核

自NVIDIA首次通过DOCA GPUNetIO引入GPU中心化数据包处理以来，该框架已显著成熟。它从一种强大的消除CPU关键路径瓶颈的方案，演变为一个在NVSHMEM、NCCL、Aerial 5G SDK、UCX/NIXL、NVQLink与Holoscan Sensor Bridge operator、Holoscan Advanced Network Operator、DeepEP/HybridEP及其他组件中集成的统一GDA-KI基础平台。

本文介绍新增内容：开源库、统一架构，以及这些集成在实际工作中如何运作。

为何统一的GPUNetIO框架至关重要

在GPUNetIO成为共享基础之前，每个通信库都独立构建了基于GDA-KI风格的GPU发起RDMA实现。这些实现各自独立运行，有着不同的假设、代码和维护成本，彼此之间没有任何共享。

将不同的GDA-KI风格实现纳入GPUNetIO框架带来了巨大优势：各库不再需要单独构建和维护基于GDA-KI的RDMA通信路径，而是可以汇聚到一个统一的GPUNetIO GDA-KI实现上。所有方都可以在同一位置使用、改进和扩展该实现。

这建立了一个共享的投入点，用于功能开发、优化和Bug修复，减少了SDK间的工程重复，并使得为一个框架开发的创新能更快地惠及其他框架。

换言之，GPUNetIO成为多个通信库共同构建的通用GDA-KI基础，而非一组并行且部分重叠的独立实现。这种统一方式对于NIXL/UCX、NCCL和NVSHMEM等更高级别的库尤为宝贵。这些库正协作丰富GPUNetIO的功能，并共同受益于整个技术栈的改进。从生态系统角度看，这意味着更少的代码重复、更少的行为碎片化，以及为长期演进提供更强劲的支持。

GPUNetIO：开源版本与SDK版本

NVIDIA 以两种紧密相关的形式提供 GPUNetIO，因为生态既需要广度，也需要开放性。完整的 DOCA SDK 版本是超集：它就是 DOCA Programming Guide 中记载的 GPUNetIO 实现，覆盖更广的 DOCA 技术栈，包括但不限于 Verbs、Ethernet、DMA 和 Comm Channel 集成，同时保留了 GPUNetIO 的核心模型——把控制路径尽量靠近 GPU，让 CPU 退出应用的关键路径。

与此同时，NVIDIA 还发布了一个开源的 GPUNetIO 项目，这是一个更轻量、专注于 RDMA-Verbs 的实现，面向那些希望在网络传输层集成方式上保持完全开放的框架。目前，开源版本是规模较小的、面向 Verbs 的子集，而 DOCA SDK 版本是超集，提供更丰富的能力，比如更完整的 RDMA 支持、Ethernet 和 DMA。

这种拆分并不是要制造两套分道扬镳的软件栈，而是为 GPU 发起的 RDMA 通信提供一个共同的开源基础，同时在有 DOCA SDK 的环境中保留通往更高级功能的路径。实际上，开源实现可以在运行时检测 DOCA SDK 是否存在，如果检测到，就通过 dlopen 调用部分闭源的 DOCA SDK 函数；如果 SDK 不存在，则继续使用开源实现运行。

这正是其关键的架构优势：各通信库不必各自构建和维护私有的 GDAKI 式 RDMA 底层设施，而是把 GPUNetIO 作为共享实现点，供多个框架和库共同采用、加固和改进。

在 CUDA 设备侧，Verbs 路径上的差距被有意控制在很小的范围内：开源版和 SDK 版都暴露面向设备的 API，因此即便主机侧实现从轻量的开源路径扩展到更完整的 DOCA SDK 功能集，GPU 编程模型也基本保持一致。

编程模型

本节介绍适用于所有 GPUNetIO 应用的核心概念与编程模型要素。

CPU 控制路径

使用 DOCA GPUNetIO 的应用需要一个初始的主机侧 CPU 配置阶段来执行控制路径，随后进入第二阶段，数据路径在 GPU 上通过 CUDA kernel 运行。

控制路径的工作流程通常遵循以下步骤：

• 初始化并配置 GPU 和网络设备，分配所需内存。
• 创建网络传输对象。例如，DOCA Verbs 或 DOCA Ethernet 在 CPU 上通过 mlx5dv 创建网络队列对象。
• GPUNetIO 的 CPU 函数将网络队列对象的相关元素导出到存储在 GPU 内存中的描述符，并提供该描述符的 GPU 地址。
• 应用启动 CUDA kernel，将网络队列的 GPU 描述符作为输入参数之一传入。

为了简化 DOCA Verbs 的部分操作，GPUNetIO 引入了一组高层函数（例如 doca_gpu_verbs_create_qp_hl），将创建和连接 RDMA QP 所需的步骤进行精简。这些函数在 SDK 和开源版本中均有提供。

控制路径完成后，数据路径即可在 GPU 上启动。此时，一个或多个 CUDA 内核会调用 GPUNetIO 的 CUDA 函数，操作已导出至 GPU 内存中的传输对象，从而发送或接收网络流量。

GPU 数据路径

用于 GPU 通信的通用 CUDA 内核结构通常包含以下步骤：

• 提交 WQE：一个或多个 CUDA 线程向网络队列对象提交工作队列条目（WQE），如 RDMA Write、RDMA Read 或以太网 Send/Recv。
• 敲门铃（Ring doorbell）：一个或多个 CUDA 线程通过向网卡寄存器写入数据，通知其有新的 WQE 就绪。
• 轮询 CQE（可选）：一个或多个 CUDA 线程等待完成队列条目（CQE），以确认 WQE 已成功完成。

API：高层与低层

针对 Ethernet 和 RDMA Verbs 传输，DOCA GPUNetIO 提供高层和低层两种不同级别的 API。

高级 API 实现了复杂和复合操作的封装。以 Verbs 侧为例，doca_gpu_dev_verbs_put_signal() 提供了一个预实现的线程安全组合操作，它将 RDMA Write WQE 与 RDMA Atomic Fetch & Add WQE 结合，可按线程或按 warp 粒度执行（warp 内所有线程协作完成操作，投递 WARP_SIZE 个 RDMA Write，最后在末尾投递一个 RDMA Atomic）。该函数负责处理同一网络队列上其他 CUDA 线程并发提交 WQE 的情况，并响应当铃信号以避免竞争条件。

类似地，在以太网侧，另一个示例是 doca_gpu_dev_eth_txq_send()，它提供线程安全的组合操作，即在网络队列上投递以太网 Send WQE 后响应当铃信号，支持线程、warp 或 block 粒度。接收侧也适用相同机制。

相比之下，低级 API 提供基础构件块，应用程序可用其创建自定义复合操作，如投递不同操作码的 WQE、响应当铃信号、轮询完成队列（CQE）等待完成事件。这些 API 不具备线程安全性，因此应用程序有责任妥善同步对同一网络队列对象的并发访问（如果存在）。

响应当铃信号

一旦 WQE 被投递到网络队列，必须通知网卡以便其执行这些操作。这一步称为“响应当铃信号”。在 GPUNetIO 中，指的是 CUDA 内核通知网卡有新工作准备好执行。

DOCA GPUNetIO 提供几种响应当铃信号的方式：

• 普通当铃信号：这是标准模式，NIC 寄存器通过 MMIO 映射到 CUDA 内存空间，允许 CUDA 线程直接写入。
• BlueFlame：遵循与普通当铃信号相同的基本模型，但此时整个 WQE 而非仅是通知被写入 NIC 寄存器。BlueFlame 通常用于对延迟敏感且具有少量网络队列的应用中。
• CPU 辅助门铃：在此模式下，NIC 寄存器映射到 CPU 内存而非 GPU 内存。GPU 将通知写入共享内存区域，由 CPU 线程进行轮询。一旦 CPU 线程检测到 GPU 更新，便会敲响 NIC 门铃。该模式主要用于在没有直接 GPU 到 NIC 连接的系统中启用 GDA-KI，例如 DGX Spark。

示例

为了方便 API 探索和性能基准测试，GPUNetIO 提供了多种参考实现。开发者可以直接在项目仓库中访问开源示例，而更全面的SDK 样本和一个基于以太网的专用应用程序则作为完整 DOCA SDK 的一部分分发。此外，在最近发布的版本中，这些样本已集成到GitHub 上，以便更轻松地获取。

以下章节将展示具体的代码片段，说明如何利用高层和低层原语，通过不同的编程语义实现相似的逻辑。

以太网示例

使用高层 API 可以实现 GPU 发起的以太网包生成器，如下面的代码片段所示。具体来说，高层函数 doca_gpu_dev_eth_txq_send 便于多个 CUDA 块并发地向共享发送队列提交 WQE，内部处理必要的同步以确保线程安全操作，无需自定义应用程序级锁。此外，它还抽象了队列结构的细粒度管理，例如跟踪 WQE 和 CQE 的具体索引。

__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;

	/* 只由 block 中的最后一个线程请求 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);

		/* 使用 BLOCK 范围的 send 内部已包含 __syncthreads */

		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 block 恰好分配一个 Send Queue，那么高层 doca_gpu_dev_eth_txq_send 内部的同步开销就变得没有必要了。在这种情况下，可以用如下模式来精简逻辑：

__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;
	/* 为简化代码，每个线程始终发送相同的缓冲区 */
	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)) {
			/* 敲击网卡门铃 */
			doca_gpu_dev_eth_txq_submit(txq, wqe_idx + 1);

			/* 轮询最后一次发送的完成状态 */
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 示例

GPUNetIO Verbs API 遵循类似的范式。例如，开发者可以利用高层原语实现以性能为导向的应用程序，通过以下 put 功能在 GPU 驱动的通信中复现 ib_write_bw 的逻辑：

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 块的带宽测试，手动管理每一个操作：

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 Device API

NCCL GIN（GPU-Initiated Networking，GPU 发起网络通信）为 CUDA 内核提供设备端通信抽象层，暴露出 put、get、signal、wait 和 flush 等原语。自 NCCL 2.27 起，GIN 已集成 GDA-KI 作为后端，利用开源的 GPUNetIO Verbs 路径。这种架构使 NCCL 算法和设备 API 应用能够直接从 GPU 驱动 RDMA 操作，同时在 GIN 接口之下屏蔽了底层队列管理和任务提交的复杂性。

工作流始于通信器初始化阶段，此时 CPU 执行控制路径，负责创建 RDMA 队列对、注册内存，并将传输描述符导出到 GPU 内存。配置完成后，数据路径完全移交至 GPU。GIN 操作会被映射为 GPUNetIO Verbs 调用，由后者管理 WQE 准备、门铃敲击以及完成轮询。这种模式成功地将 CPU 从每次网络事务的应用关键路径中移除。

通过将通信语义与传输实现解耦，NCCL 可以专注于集合通信算法，而由 GPUNetIO 提供经过生产验证的 GDA-KI 底层管道。NCCL 利用 GIN 级别的控制来管理线程、CTA 或 GPU 级别的资源共享，从而实现请求聚合和优化的队列管理。这些能力与 GPUNetIO 的优势高度契合，既能提供低延迟、高吞吐的路径，又不会让 RDMA 实现变得碎片化。最终，GPUNetIO 成为共享基础，让性能优化和 NIC 支持只需开发一次，就能在整个 GPU 通信生态中通用。

如果想了解 NCCL GIN 接口的功能实现，开发者可以访问 NCCL tests GitHub 仓库查看参考代码。

NVSHMEM

NVSHMEM 是一种基于 PGAS 的编程模型，为 GPU 集群提供可扩展的点对点和集合通信原语，如 put、get、原子操作、barrier 和规约。它让多 GPU 应用能够在 CUDA kernel 内实现细粒度的 GPU 间数据传输和同步，显著提升应用工作负载的强扩展性能。

在网络通信方面，NVSHMEM 根据可用的网络能力和应用需求提供了多种传输方式。NVSHMEM 提供 IBGDA 传输机制，该机制通过 MLX5 直接动词（direct verbs）实现了 GDA-KI。

随着开源 GPUNetIO 库的发布，NVSHMEM 3.7 引入了基于 GPUNetIO 的新传输层。依托这一通用实现层，NVSHMEM 得以享受 GPUNetIO 的优化以及 DOCA SDK 提供的高级网卡特性。新的 GPUNetIO 传输层同时支持 GDA-KI 通信和主机发起的通信（即 CPU 通信），并基于 RDMA 实现。在主机发起通信的场景下，GPUNetIO 仅用于控制路径以配置网络元素，而数据路径则由 CPU 通过 NVSHMEM 原生函数执行。用户可以通过设置 NVSHMEM_REMOTE_TRANSPORT=gpunetio 选择 GPUNetIO 传输层，并通过 NVSHMEM_GPUNETIO_ENABLE_GDAKI=1 启用 GDA-KI。

NVSHMEM 中 GPUNetIO 传输层的 GDA-KI 部分遵循 IBGDA 传输层的结构，但将大段底层代码替换为对 GPUNetIO 库的调用。例如，如上文“敲响门铃（Ring the doorbell）”一节所述，门铃的构建与敲响代码完全外部化到了 GPUNetIO 库中。在实现与 IBGDA 相同功能和性能的同时，GPUNetIO 传输层大幅降低了传输实现复杂度，彰显了整合 GDA-KI 通用功能模块的价值。

根据通信模式的不同，启用 GDA-KI 可能会显著提升小消息吞吐量。以下实验使用了 NVSHMEM 性能测试套件中的 shmem_put_bw 基准测试，利用 nvshmem_double_put_nbi 处理不同大小的消息。实验中每个 CTA（协作线程阵列，cooperative thread array）使用 1 个线程和 1 个 QP，并增加 CTA 的数量。下面两个图展示了 GPUNetIO 传输层达到的带宽，其中第一张图展示的是禁用 GDA-KI（即使用 CPU 通信）的结果，第二张图展示的是使用 GDA-KI 的结果。

如上所述，如图 5 所示，当 NVSHMEM 使用 CPU 通信（控制路径配置为 GPUNetIO 函数，而数据路径为 NVSHMEM 原生实现）时，由于 CPU 代理瓶颈，随着 CTA 和 QP 数量扩展，小消息的带宽会受限。而在图 6 中，基准测试执行时启用了 GDA-KI 功能的 GPUNetIO 后端（数据路径由 GPUNetIO CUDA 函数在 GPU 上执行）。

实验证实，使用 GDA-KI 消除代理瓶颈后，CTA 和 QP 的扩展性表现更好。GDA-KI 版本在更小的消息大小下就能达到更高的峰值带宽，而且比非 GDA-KI 版本更早进入峰值区间。此前对比 IBRC 和 IBGDA 带宽性能的实验也显示了类似结果。

总体来看，GPUNetIO 在保留 IBGDA 和 IBDEVX 功能与性能的同时，大幅降低了 NVSHMEM 的实现复杂度和维护成本。

NVQLink

NVIDIA NVQLink 是 NVIDIA 的低延迟参考架构，用于将量子处理器控制系统与 GPU 计算连接起来，支持通过与其 NVIDIA CUDA-Q 紧密集成，实现量子纠错（QEC）、前馈及自适应校准等实时量子-经典工作流。Holoscan Sensor Bridge（HSB） 通过采用 RoCE 与轻量级 FPGA IP 以及主机端控制软件，为这一方法提供了网络基础。NVQLink 则将该基于 HSB 的模型推向极低延迟区间，使其适用于实时 QPU 控制，而不仅仅是传统的传感器数据摄入流水线。

NVQLink 使用的 HSB 网络算子名为 GPU RoCE Transceiver。它基于 DOCA GPUNetIO 和 DOCA Verbs 库构建，旨在尽可能降低流量转发延迟。在独立模式下，该算子提供一条流量转发执行路径，用于网络一致性检查和性能分析。

该算子启动一个常驻 CUDA 内核，利用 DOCA GPUNetIO CUDA 函数持续从远程对等节点接收数据包并将其发回，全程无需 CPU 干预。在此配置中，远程对等节点是一块 FPGA，用于测量数据包从发出到收回的往返时间。

为了测量 HSB 2.7.0 GPU RoCE Transceiver 算子的网络转发延迟，我们使用了配备 Blackwell GPU 和 ConnectX-7 网卡的 NVIDIA IGX Thor 平台。

另一侧，我们通过 OSFP 线缆背对背直连，用一块 Xilinx RFSoC 4×2 开发板上的 FPGA 负责收发包并测量往返延迟。

以 FPGA 上测得的发包到收包的时间计算，最小延迟约为 2.6 微秒，中位数延迟为 2.7 微秒。

开始使用 DOCA GPUNetIO

• DOCA SDK GPUNetIO 编程指南 — 完整 SDK 文档
• GPUNetIO 开源仓库 — 示例与开源实现
• GitHub 上的 DOCA SDK 示例 — 参考实现
• NCCL 测试 — GIN device API — GIN 参考代码
• Holoscan Sensor Bridge GPU RoCE Transceiver — NVQLink 算子