全文 - The Landscape of GPU-Centric Communication -NVSHMEM-GPUNetIO

最新推荐文章于 2026-09-20 16:54:43 发布

原创最新推荐文章于 2026-09-20 16:54:43 发布 · 401 阅读

12 ·

本内容遵循CC 4.0 BY-SA版权协议

GEO检测

GPU 中心通信全景

原文：arXiv:2409.09874v3，ACM Computing Surveys（CSUR）
CCS：网络·编程接口；计算方法论·并行编程语言；硬件·通信硬件、接口与存储；计算机系统组织·单指令多数据流（SIMD）

Ilyas Turimbetov（科奇大学，土耳其伊斯坦布尔，iturimbetov18@ku.edu.tr）、Mohammed Kefah Taha Issa（科奇大学，MISSA18@ku.edu.tr）、Doğan Sağbili（科奇大学，dsagbili17@ku.edu.tr）、Flavio Vella（特伦托大学，意大利，flavio.vella@unitn.it）、Daniele De Sensi（罗马第一大学，desensi@di.uniroma1.it）、Ismayil Ismayilov（科奇大学，iismayilov21@ku.edu.tr）

摘要

近年来，GPU 凭借其并行能力和高速内存带宽，已成为 HPC 与机器学习应用的首选加速器。虽然 GPU 能够大幅提升计算性能，但 GPU 之间的通信可能造成可扩展性瓶颈，尤其是随着每个节点和每个集群中 GPU 数量的增长。传统上，多 GPU 通信由 CPU 管理，但 GPU 中心通信（GPU-centric communication）的进步正在挑战 CPU 的这种主导地位：它减少了 CPU 的参与，赋予 GPU 在通信任务中更大的自主性，并试图解决多 GPU 通信与计算之间的语义不匹配问题。

本文对 GPU 中心通信进行了全面梳理，重点关注厂商机制和用户级库的支持。本文旨在厘清该领域的复杂性和多样化选项，定义术语，并对节点内和跨节点的现有方法进行分类。文章讨论了厂商在多 GPU 执行中提供的通信与内存管理机制，并综述了主要的通信库及其优势、挑战和性能洞见。随后，文章探讨了关键的研究范式、未来展望和开放性研究问题。通过对软硬件栈上 GPU 中心通信技术的深入描述，我们为研究人员、程序员、工程师和库设计者提供了如何充分利用多 GPU 系统的洞见。

关键词：GPU、通信、MPI、NVSHMEM、NCCL、RCCL、点对点通信、集合通信、GPUDirect 技术。

1. 引言

近年来，GPU 已成为 HPC 和机器学习领域众多应用的首选加速器。这一快速普及主要由 GPU 的大规模并行能力和高内存带宽驱动，意味着现代云和 HPC 算力的大部分现已集中于 GPU 集群。截至 2025 年 11 月，Top500 超级计算机前 10 名中有 9 台依赖 GPU 集群进行加速，这一趋势很可能会持续。唯一没有 GPU 的系统"富岳"（Fugaku）采用了高度向量化的 CPU 架构并结合高带宽内存（HBM）。

虽然使用大量 GPU 已被证明能显著加速计算，但 GPU 之间的通信会迅速成为可扩展性瓶颈。传统上，无论节点内还是跨节点的多 GPU 通信，一直都是 CPU 的职责。GPU 自诞生之初就被认为是能够提供大量算力、但在通信等辅助任务上天然依赖 CPU 的设备。在这种以 CPU 为中心的执行模型中，为 GPU 提供数据的中继例程对 GPU 的存在是无感知的。

过去十年中，多项广义上被称为"GPU 中心通信"的技术进步，试图挑战 CPU 在多 GPU 执行中的霸权。从高层来看，这些进步减少了 CPU 在执行关键路径中的参与，赋予 GPU 在发起和同步通信方面更大的自主性，并试图解决多 GPU 通信与计算之间的语义不匹配问题。

本文围绕厂商机制和用户级库支持，对 GPU 中心通信进行全面综述。本综述的目标是帮助消除潜在研究者进入该领域时普遍存在的困惑。我们希望帮助程序员、工程师、编程模型和库设计者理解现有选项的复杂性和多样性，因为 GPU 中心通信横跨非常广泛的方法谱系，包括专有的 GPU 间互连等硬件创新和软件机制。这些机制各有其独特的优势和挑战，因此何时何地应优先选用并不明确。文献中术语使用的不一致以及厂商产品之间的差异，使局面更加扑朔迷离。

本文结构如下：

第 2 节：定义术语，给出 GPU 中心通信的定义；对节点内和跨节点的现有方法进行分类和抽象，以消除混淆。
第 3 节：回顾历史，讨论厂商提供的用于实现通信与联网、在多 GPU 执行中跨设备管理内存的机制。这些机制是更高级 GPU 中心软件库的基本构件。
第 4 节：列举并比较节点内和跨节点场景下的主要通信库，讨论其优势与挑战，并借助现有的基准测试工作提供性能洞见。
第 5 节：讨论 GPU 中心通信背后的主要研究范式，展望该领域的未来，并提出开放性研究问题。

本综述中讨论的某些方法和技术天然与特定专有生态系统绑定。虽然现有文献和已部署系统的相当大一部分聚焦于 NVIDIA 技术，但我们力求给出平衡的视角：在信息可得的范围内，纳入 AMD 和 Intel 等其他厂商的相应机制。同时，某些能力仍为特定厂商所独有；我们明确标注这些厂商专有特性，以确保清晰和完整。

2. 术语与通信类型

我们可以将 GPU 中心通信宽泛地定义为：减少 CPU 在多 GPU 执行关键路径中参与的机制。这是一个非常宽泛的定义，涵盖了广泛的解决方案，既包括赋予 GPU 通信自主性的厂商级改进，也包括利用这些改进的用户级实现。为明确这一区别，我们在不同章节分别讨论。第 3 节聚焦于作为 NVIDIA CUDA 和 AMD ROCm 运行时原生组成部分提供的通信机制和原语；第 4 节讨论这些机制如何催生更高级的 GPU 中心通信库，并介绍 AMD、Intel 和 NVIDIA 的厂商软件支持，以及其他工业界和学术界的方案。

我们还指出节点内通信（intra-node）与跨节点通信（inter-node）之间的区别。单个 GPU 加速节点由一个共享内存主机和挂接的多个 GPU 卡组成。节点内通信时，任一给定 GPU 可由单个线程或进程控制，共享内存和地址空间。多节点系统包含多个这样的节点，每个 GPU 由不同进程控制，运行在不同节点上的进程之间不共享内存。通信格局因配置不同而变化，因为跨节点通信需要处理 GPU 与网卡（NIC）的交互以及跨进程通信。

2.1 节点内通信

尽管将通信方法划分为 GPU 侧和 CPU 侧是常用做法，对终端用户而言也往往足够，但这种划分并不总是有解释力和准确。为避免定义的模糊性，我们将通信方法划分为若干类型。这些类型基于通信过程中每项操作的执行者。在节点内场景下，我们定义通信发生所需的两个主要操作；跨节点场景则有四个操作。节点内场景下，一次通信调用的两个组成部分是：

API：定义程序员或库在哪里发起通信 API 调用。
数据路径：指明谁参与数据搬运，并展示相应的数据路径。

节点内通信机制分类示例及描绘数据路径的图示，见原文表 1 和图 1。

表 1 中的 ① 主机原生（host native） 通信方法在主机侧发起，不涉及设备间的直接点对点（P2P）访问。这涵盖所有在主机侧发起且禁用 P2P 访问的方法。而 ② 主机控制（host-controlled） 表明通信无需额外拷贝到主机内存，直接通过 PCIe、NVLink 或 Infinity Fabric 互连传输。GPUCCL（即 NCCL、RCCL、oneCCL 等）、GPU-aware MPI 和 *memcpy 操作具有主机侧 API，因此可能同时属于 ① 和 ②。

通过启用对等设备内存的直接访问，设备侧 API 将 CPU 从节点内通信的数据路径和控制路径中都移除了，如图 1 中 ③ 设备原生（device native） 所示。NVSHMEM、ROCSHMEM、Intel SHMEM 也提供主机侧 API，但需要 P2P 访问，因此其主机侧 API 属于 ②，设备侧 API 属于 ③。内核内 P2P 直接加载/存储（direct load/store）提供类似功能，属于 ③；但在禁用 P2P 访问时也能工作，此时属于 ④，数据路径回退到主机。

2.2 跨节点通信

跨节点场景更加多样，因为必须与网卡交互。每种方法的具体实现可能涉及复杂的数据路径和决策。为简化分类，我们区分跨节点通信的四个主要组成部分。除节点内场景使用的 API 和（通往网卡的）数据路径外，还新增两个涉及网卡交互的组成部分：

注册/构造消息：该步骤涉及数据包的构造及其在网卡上的注册。
触发通信：定义由谁来"按响网卡的门铃"（ring the doorbell）以发出数据传输。

表 2. 跨节点通信方法类型（加粗单元格表示发生变更或优化之处；D/H 表示设备侧和主机侧 API 调用均可属于该类型）

类型	API	注册	触发	数据路径	示例
① 主机原生	主机	主机	主机	经由主机（2 次拷贝）	GPUDirect 2.0 之前的 GPU-Aware MPI
② 固定主机原生（Pinned Host Native）	主机	主机	主机	经由主机（1 次拷贝）	GPU-Aware MPI 中的 GPUDirect 1.0、NCCL、RCCL、oneCCL
③ GPU RDMA	D/H	主机	主机	直接	GPU-Aware MPI 中的 GPUDirect RDMA、ROCm RDMA、NCCL、NVSHMEM、ROCSHMEM
④ GPU 触发	D/H	主机	设备	取决于 ③	GPU-Aware MPI 中的 GPUDirect Async、NCCL v2.28、NVSHMEM、ROCSHMEM
⑤ 设备原生	设备	设备	设备	取决于 ③	使用 IBGDA 的 NVSHMEM、GPUrdma

我们在分类跨节点通信时识别出五个主要类别。从 ① 到 ⑤（如表 2 和图 2 所示），通信调用的越来越多组成部分被移到了 GPU 侧，而数据传输路径上到达网卡所需的数据拷贝次数不断减少。多年来，通信方法与相应技术展示了对数据传输和通信控制两方面的优化，这将在第 3 节讨论。最初，只有 ① 这种完全 CPU 侧的通信方法可用；随后 ② 借助 GPU 与网卡之间共享的固定内存（pinned memory），消除了 CPU-GPU 与 CPU-NIC 缓冲区之间的额外拷贝，降低了 GPU-NIC 传输的延迟。之后，③ GPU RDMA（远程直接内存访问）使网卡能够经 PCIe 直接访问 GPU 内存，最小化了两者之间的数据路径。④ 代表 GPU 触发的通信技术，如 GPUDirect Async 和 GPU-TN：在 CPU 预先在网卡上准备好数据包的前提下，GPU 能够发起通信。⑤ 则将数据包准备和与网卡的交互也全部移到 GPU 上，使设备原生通信成为可能。

表 2 和图 2 给出的类型并未反映所有可能的组合，因为某些库基于配置和可用硬件，可能产生不同的数据路径与控制组合。例如，在没有 RDMA 技术的情况下，即便通信控制在 GPU 侧，数据路径仍会涉及主机内存。举例来说，Intel SHMEM 在主机上使用代理线程，通过标准 OpenSHMEM 库执行实际的 RDMA——尽管通信调用是在 GPU 上运行的设备内核内发起的。

3. 厂商机制

本节讨论厂商提供的机制：在多 GPU 执行中实现通信与联网、跨设备管理内存。这些机制由 GPU 编程模型运行时或扩展 API 提供。这些技术分为四类：内存管理器、GPUDirect 技术、硬件和库。原文图 3 总结了 NVIDIA 提供的技术及其时间线与可用性。

接下来，我们先介绍内存管理机制和 GPUDirect 技术，然后介绍作为先驱并最终使这些通信方法可行的硬件支持。这些技术构成了第 4 节讨论的更高级 GPU 中心库的骨干。

3.1 内存管理机制

3.1.1 页锁定/固定内存（Page-Locked / Pinned Memory）

默认情况下，主机上使用设备 malloc（如 cudaMalloc()、hipMalloc()）分配的内存是可分页的（pageable），GPU 无法访问。在可分页主机内存与设备内存之间执行传输时，GPU 运行时必须先将主机数据暂存到页锁定内存的临时缓冲区，然后再将数据从页锁定内存拷贝到 GPU。为避免"可分页 → 页锁定"内存拷贝，cudaMallocHost() 允许直接分配页锁定内存，跳过中间拷贝阶段。正因如此，页锁定内存也被称为零拷贝或固定内存。

固定内存以其在主机-设备传输中的高带宽和低延迟著称，通过实现全系统范围的直接访问，高效地协调 CPU-GPU 执行。它还与 GPUDirect RDMA 配合使用，以改进跨节点通信。然而，其物理内存锁定特性可能导致高内存占用，过度分配时可能影响系统性能。

3.1.2 统一虚拟寻址（UVA）

UVA 是一种内存管理技术，允许节点内所有 GPU 和 CPU 共享同一个统一虚拟地址空间。在 UVA 之前，主机 ↔ 设备和设备 ↔ 设备拷贝必须显式指定传输方向。借助 UVA，可以从指针值推断物理内存位置，从而降低了管理独立内存空间的开销，使库能够简化其接口。

3.1.3 进程间通信（IPC）

在早期的 GPU 运行时版本中，指针无法跨进程边界访问，因此 GPU 缓冲区之间的内存拷贝必须经由主机，形成瓶颈。为克服这一限制，IPC 使同一节点上的进程能够访问其他进程的设备缓冲区，而无需额外拷贝。借助 IPC，内存句柄（handle）被创建并通过标准 IPC 机制在进程间传递，从而获得比经由主机暂存拷贝更低的延迟。

3.1.4 统一虚拟内存（UVM）

UVM 允许通过 cudaMallocManaged() 调用分配托管内存（managed memory），创建一个节点内所有处理器均可访问的单一地址空间。UVM 的工作原理是将所请求的内存划分为驻留在 CPU 上的页。程序员无需显式拷贝即可在设备上访问内存。如果某次内存访问落在不在设备上的页中，UVM 驱动会触发缺页异常，自动将该页迁移到发起请求的设备。当总分页内存大小超过设备内存时，UVM 驱动还可以将页从某设备驱逐回主机内存。

UVM 在可编程性方面提供了若干好处。第一，程序员面对的是单一统一地址空间，可以像整块分配内存驻留在单个 GPU 上一样访问它。系统中发生的任何拷贝都是隐式的，对程序员隐藏。此外，UVM 允许内存超额订阅（oversubscription）：分配的内存可以超过所有多 GPU 设备内存的总和。这是可能的，因为大部分内存可以留在 CPU 上，在设备请求时才分页调入。

3.2 GPUDirect 技术

3.2.1 GPUDirect 1.0（网卡）

GPUDirect 1.0 允许 GPU 与网卡共享同一固定内存区域。在此之前，系统内存中供 GPU 和网卡使用的固定内存区域是相互独立的。由此推论，要跨节点传输 GPU 数据，GPU 先将数据拷贝到自己的固定内存区域，CPU 再将其拷贝到网卡的内存区域，只有在这之后网卡才能访问并将其发送到网络上（如图 2 ① 所示）。GPU → 网卡固定内存区域之间由 CPU 发起的中间拷贝增加了 CPU 开销和通信延迟。GPUDirect 1.0 引入了 GPU 与网卡共享的固定内存区域，从而避免了这次由 CPU 发起的中间拷贝。

3.2.2 GPUDirect 2.0（点对点）

随着 UVA 的引入，CUDA 4.0 增加了对同一节点内共享同一 PCIe 根复合体（root complex）的 GPU 之间直接点对点通信的支持。该功能被封装在名为 GPUDirect 2.0 或 GPUDirect P2P 的技术中。GPU 不再经由主机暂存数据，而是可以直接经 PCIe 访问彼此的内存，首次建立了直接的 GPU 到 GPU 数据路径。这些变化催生了两种新的通信机制：P2P DMA 拷贝（cudaMemcpy 调用会触发源 GPU 与目标 GPU 内存之间的直接 DMA 传输）和 P2P 直接加载/存储（GPU 通过解引用指向远程 GPU 缓冲区的指针直接访问数据）。GPUDirect P2P 在 NVLink（见 3.4 节）推出后也增加了对其的支持。

GPUDirect P2P 带来两大好处。它消除了冗余的 GPU ↔ CPU 拷贝和主机缓冲区（在传输经 CPU 暂存时是必需的）。此外，由于无需在主机上维护通信缓冲区，并提供了新的通信机制（P2P 直接加载/存储），GPUDirect P2P 提高了多 GPU 编程的便利性。

注意，P2P DMA 拷贝在没有 UVA 支持时也能工作。如果未启用 UVA，可以通过显式指定目标 GPU 的 cudaMemcpyPeer() 变体执行 P2P DMA 拷贝。然而，P2P 直接加载/存储在没有 UVA 时无法工作，因为直接访问远程 GPU 的指针以统一地址空间为前提。

3.2.3 GPUDirect RDMA

随着 CUDA 5.0 中 GPUDirect RDMA 的引入，NVIDIA GPU 之间的直接跨节点通信成为可能。GPUDirect RDMA 通过标准 PCIe 特性，在 GPU 与第三方设备之间建立直接通信通道。该技术将 GPU 内存段暴露在 PCIe 内存资源上，称为基址寄存器（BAR）区域，使网卡无需经由主机即可直接读写 GPU 内存。类似地，AMD 提供 ROCm RDMA（此前称为 ROCnRDMA）。GPUDirect RDMA 对数据路径进行了若干优化：消除了到主机内存的额外拷贝、降低了 GPU-NIC 交互的固有延迟、提高了带宽并减少了 CPU 开销。

3.2.4 GPUDirect Async

此前的 GPUDirect 技术聚焦于改进数据路径，而 GPUDirect Async 优化的是 GPU 与网卡之间的控制路径。它于 CUDA 8.0 引入，使 GPU 能够发起和同步网络传输，从而减少 CPU 在关键路径中的参与。GPUDirect Async 的工作方式是：CPU 预先注册消息，GPU 内核随后可通过按响网卡门铃来触发这些消息。因此，GPU 可以在通信被触发的同时继续执行，而不必像过去那样停下来等待 CPU 发起通信。

尽管 GPUDirect Async 在将控制路径从 CPU 移开方面取得了进展，但它并未将控制路径完全移交给 GPU，因为通信仍受限于内核启动边界。本质上，GPU 只能发起此前由 CPU 注册的消息。GPUDirect Async 的进一步改进作为 NVSHMEM 库中 IBGDA 传输层的一部分实现（见 4.3.1 节）。

3.3 GPUNetIO

GPUNetIO 是 NVIDIA 作为 DOCA（Datacenter-On-a-Chip Architecture，片上数据中心架构）一部分提出的技术方案。DOCA 是为开发 NVIDIA BlueField 数据处理单元（DPU）应用而设计的全栈软件框架。在非 RDMA 网络上，GPUNetIO 允许 GPU 发送、接收和处理网络数据包。在 RDMA 网络（RoCE 和 InfiniBand 均可）上，自 DOCA v2.7 起，GPUNetIO 允许 GPU 不仅在内核边界上、而且在内核执行的任意时刻执行 RDMA 发送和接收。简言之，GPUNetIO 允许 GPU 在完全无 CPU 干预的情况下与网卡交互。

在 RDMA 网络上，GPU 内核可以（以阻塞或非阻塞模式）等待 RDMA 接收操作完成。在非 RDMA 网络上，GPUNetIO 提供信号量，可在内核内显式用于与网卡同步数据包收发。信号量还可用于将 GPU 内核与 CPU 同步（当数据包处理在 CPU 与 GPU 之间拆分时），或与其他 CUDA 内核同步（当数据包处理跨多个内核拆分时）。

3.4 现代 GPU 中心互连

GPU 中心互连技术在节点内多个 GPU 之间提供高带宽、低延迟的通信，这一能力对高性能工作负载（尤其是 AI 训练和 HPC）至关重要。

NVLink 是 NVIDIA GPU 的专有互连技术。其设计旨在解决 PCIe 的带宽限制问题——PCIe 已被观察到是 GPU 加速应用中的传输瓶颈。表 3 给出了各代 NVLink 的规格。

表 3. NVLink 各代规格

代次	NVLink 链路数	单链路单向带宽（GB/s）	双向聚合总带宽（GB/s）	支持的架构
第一代	4	20	160	Pascal
第二代	6	25	300	Volta
第三代	12	25	600	Ampere
第四代	18	25	900	Hopper
第五代	18	50	1800	Blackwell

此外，NVLink 还曾用于连接 GPU 与 IBM Power8、Power9 CPU；随着 Grace Hopper 超级芯片的推出，NVLink 被用作芯片间（C2C）互连，双向带宽达 900 GB/s。随后，随着 Grace Blackwell 超级芯片的推出，NVLink-C2C 用于连接 Grace CPU 与 2 个 Blackwell GPU，双向总带宽达 3.6 TB/s。Blackwell 上的第五代 NVLink 每 GPU 提供 1.8 TB/s 双向吞吐，支持多达 576 个 GPU 之间的高速通信。

NVLink 的引入优化了 NVIDIA GPU 之间的带宽，使 P2P 通信成为节点内通信的可行机制，使数据路径大幅向 GPU 倾斜。NVLink 的一个缺点是它不能自路由：如果任意两个 GPU 之间没有直接的 NVLink 连接，通信必须经由中间 GPU 路由。这一限制由 NVSwitch 克服——这是一种背板技术，可以实现所有 GPU 之间的全互连。例如，DGX-2 节点由 16 个 V100 GPU 组成，通过 NVLink 和 NVSwitch 实现全互连。从第三代开始，NVSwitch 支持 SHARP，将 allreduce 操作卸载到 NVSwitch，使 allreduce 能够以全线速运行。

AMD 的替代方案是 Infinity Fabric/xGMI，用于现代 AMD GPU 加速器（如 AMD Instinct MI300X）。xGMI 在节点内 GPU 之间提供高带宽通信。该架构在 8 GPU 系统上支持 896 GB/s 的双向聚合 GPU 间带宽。与 NVSwitch 不同，AMD 当前的互连网格是扁平的，不依赖交换结构，这对扩展到每节点超过 8 个 GPU 形成了一些限制。最近，Ultra Accelerator Link（UALink）联盟成立，旨在开发一种更开放的共享内存加速器互连，兼容多种技术和厂商。

Intel 的 Xe-Link 架构是高带宽、全连接的节点内互连，用于 Aurora 超级计算机等系统，将 Intel Data Center GPU Max（Ponte Vecchio）设备连接成统一拓扑。在典型的 Aurora 节点中，六个 GPU 以全互连方式连接，Xe-Link 也可支持最多 8 路配置，其中每个 GPU 都有通往其他所有 GPU 的直接链路。这些链路支持加载/存储访问、拷贝引擎传输和远程原子操作，使 GPU 无需主机即可直接访问彼此内存。

3.5 厂商机制讨论

3.5.1 GPUDirect P2P 与直接加载/存储对编程的影响

GPUDirect P2P 的引入标志着多 GPU 执行范式的重大转变，使得在内核内使用加载和存储操作进行 GPU 间直接通信成为可能。

基于直接加载/存储的通信有若干好处。第一，它允许程序员将通信与计算内联。程序员不再依赖分离的通信模型和计算模型，而是可以在 GPU 内核内将两者结合。第二，直接加载/存储利用 GPU 提供的高度并行性，相比 DMA 拷贝可以实现更高的带宽和更低的延迟。第三，直接加载/存储可以通过 GPU 固有的延迟隐藏能力，隐式地实现通信与计算重叠。鉴于 GPU 提供的高度并行性以及现代互连不断提高的带宽，GPU 不仅能隐藏本地内存延迟，也能隐藏远程内存延迟。这对程序员而言是又一福音：实现重叠的方式从程序员通过流（stream）和事件（event）手动实现的软件方法，转变为自动的硬件重叠。由于通信/计算重叠的负担从程序员转移给了硬件，其另一层含义是：随着硬件隐藏内存延迟的能力增强，重叠效果也会随之改善。第四，直接加载/存储扩展了可通过多 GPU 加速的应用范围。传统上，具有细粒度通信模式的应用在多 GPU 系统上可扩展性很差，因为计算必须频繁中断和同步，以便 CPU 发起通信。借助内核内的直接加载/存储，GPU 能够很好地适应细粒度通信模式。最后，直接加载/存储允许在不离开 GPU 的情况下于内核内触发通信。这一方向与持久内核（persistent kernel）结合时尤其有前景——持久内核启动一次，通过在设备上运行内部循环跨多个工作迭代维持执行，从而最小化内核启动开销。事实上，Spector 等人在 Llama-70B 的张量并行推理上使用了直接加载/存储，并实现了一个"巨型内核"（megakernel）来执行异步通信，将其与计算和本地内存操作重叠。

尽管直接加载/存储带来诸多改进，也存在若干固有挑战。第一，一个根本性挑战是：通信和计算现在都需要大量 GPU 线程才能推进，因而争夺同一种有限资源。当通信以单独内核实现时，这一问题尤其严重。如果计算内核先启动，它可能独占所有 GPU 资源，使通信内核无法启动，实际上消除了任何重叠的可能。可以通过以更高优先级启动通信流来缓解此问题，使其总是被优先调度。注意，P2P DMA 拷贝没有这个问题，因为它们使用 GPU 的 DMA/拷贝引擎——一种物理上独立的资源——进行通信。第二，与单 GPU 内存访问类似，P2P 直接加载/存储对内存合并（coalescing）高度敏感：随机的非合并访问性能远差于合并访问。此类非合并直接读取可能暴露超出 GPU 调度器隐藏能力的远程内存延迟，最终导致执行停顿。类似地，零散的、亚缓存行粒度的非合并直接写入可能严重低利用互连带宽。

3.5.2 GPUDirect RDMA 的局限

GPUDirect RDMA 的一个重大局限是：内核运行期间，GPU 与网卡内存之间没有一致性保证。一致性只能通过将控制权交还 CPU——即销毁内核并启动新内核——来保证，因此通信被限制在内核边界。这也意味着将持久内核与 GPU 发起的跨节点通信结合，必然导致数据正确性问题。Chu 等人通过从网卡向 GPU 内存发起一次 PCIe 读取来绕过这一限制，该读取将网卡此前的写入冲刷到 GPU，保证内存序。自 11.3 版本起，CUDA 也提供 cudaDeviceFlushGPUDirectRDMAWrites() API，可用于类似地强制一致性。虽然有用，CUDA 仍依赖 CPU 来强制 GPU-NIC 一致性。另一方面，AMD 在持久内核的设备侧通信语境下明确修正了 GPU-NIC 一致性问题，并将所提议的修复集成到 ROC_SHMEM 中。我们将在 5.3 节"免 CPU 联网"语境下进一步讨论此问题。

3.5.3 在 GPU 中心通信中实现触发能力

在第 2 节介绍的类型 ③（GPUDirect/ROCn RDMA）中，CPU 仍负责系统的初始配置、数据传输准备和发起传输。第一阶段包括设置网络接口和加载 GPU 驱动；CPU 向支持 RDMA 的网卡注册 GPU 内存，使网卡能够直接访问 GPU 内存，绕开数据传输期间的中介 CPU 步骤。第二阶段，CPU 分配 GPU 内存缓冲区并确保正确对齐，这些缓冲区用于高效的 GPU 数据收发。然后，CPU 设置 GPU 流和事件，用于管理和排序数据传输并保证工作完成；流用于排队操作，确保按正确顺序执行。然而，对真正的低延迟应用而言，这一成本仍可能成为瓶颈，因为该机制依赖通过流的多个同步点。

第 2 节定义的类型 ④ 和 ⑤ 中的 GPU 触发通信，通过消除上述同步成本，促进将计算和通信控制路径都卸载到 GPU。这里，触发操作（triggered operations）扮演关键角色：它们是仅在满足特定条件时才执行的特殊任务。在流触发（ST）策略中，这些操作通过 GPU 控制处理器管理数据传输和同步，从而减少 CPU 参与。延迟执行（Deferred Execution）是另一个重要方面：CPU 创建具有延迟执行语义的命令描述符并将其追加到网卡命令队列；当 GPU 控制操作指定的条件满足时，这些描述符被执行。例如，HPE Slingshot 11 网卡支持此类延迟操作，包括当硬件计数器达到给定阈值时触发的收发通信——这通过启用特定的命令队列（如 Libfabric Deferred Work Queue）实现。GPU 控制处理器与网卡之间的同步通过特殊机制（例如 NVIDIA DOCA 中用 GPUNetIO 信号量实现）确保通信操作成功完成。

4. GPU 中心通信库

我们现在讨论近年来为简化多 GPU 编程而涌现的主要 GPU 中心通信库：GPU-aware MPI、GPU 中心集合通信库（GPUCCL）和 GPU 中心 OpenSHMEM（GPUSHMEM）。

4.1 GPU-Aware MPI

鉴于 MPI 是 HPC 事实上的通用语言，大量工作致力于使其与 GPU 编程模型互操作，最终形成了能够区分主机缓冲区和设备缓冲区的 GPU-Aware MPI 实现。在 GPU-Aware MPI 之前，所有多 GPU 通信都必须经主机暂存：源 GPU 产生一次 device → host 拷贝，目标 GPU 产生一次 host → device 拷贝。而使用 GPU-Aware MPI 实现时，程序员可以将设备缓冲区作为参数提供给 MPI 调用，使通信使用由 GPUDirect RDMA 或 ROCnRDMA 建立的直接 GPU 到 GPU 数据路径。在此过程中，GPU 感知消除了冗余的 host ↔ device 拷贝，并通过免去主机缓冲区简化了通信代码。

MVAPICH2 是第一个开始积极将 GPU 感知集成到其运行时中的 MPI 实现。早期的 MVAPICH2 工作引入了基本的 GPU 感知：将驻留在 GPU 上的缓冲区透明地经主机暂存，并通过流水线方案优化 host ↔ device 和 device ↔ device 传输。这些流水线方案因 UVA 而成为可能——UVA 使库无需依赖用户提示即可区分主机指针和设备指针。随后的 GPU 感知相比 GPU 无感知版本带来了性能提升。后续工作使用 CUDA IPC 优化节点内传输（此前必须经主机内存缓冲区暂存）。最终，在 rendezvous 协议上加入了对 GPUDirect RDMA 的支持，使传输绕开主机、消除冗余的 host ↔ device 拷贝。这降低了延迟；但由于当时的架构限制，带宽受限。后续工作在 eager 协议上加入 GPUDirect RDMA 支持，修正了带宽限制并进一步降低延迟。此外，还使用了一种新的回环（loopback）机制和早期版本的 GDRCopy 来消除昂贵的 host ↔ device cudaMemcpy。GDRCopy 允许将 GPU 内存映射到用户地址空间，针对小消息尺寸以最小开销进行了优化。另一项工作扩展了点对点 MPI 调用以支持 GPUDirect Async，允许 GPU 推进 CPU 排队的通信，从而优化控制路径。其他工作也越来越聚焦于为 MVAPICH2-GDR 添加 UVM 感知。

其他主流 MPI 实现也集成了 GPU 感知。Open MPI 在 1.7.0 版本引入了 CUDA 感知支持。使用其内部后端时，基于计算的集合操作（如 MPI_Allreduce）仍可能经主机暂存数据；使用 UCX 则可避免这一限制——UCX 支持 GPU 驻留数据的搬运和归约。Open MPI 还对小消息使用 GDRCopy，对节点内通信使用 CUDA IPC。GPU 感知通过 UCX 扩展到 AMD 设备（ROCm）。然而，Khorassani 等人为 MVAPICH2 提供了原生 ROCm 感知运行时，在 AMD GPU 集群上优于"Open MPI + UCX"。Open MPI 已集成统一集合通信（UCC）框架——UCX 生态的一个组件，旨在为高性能集合操作提供统一接口。UCC 构建于 UCX 传输层之上，利用其拓扑感知机制（如对 NVLink、PCIe 和共享内存层级的感知）为 GPU 缓冲区选择高效的集合算法。该集成使 Open MPI 能够自动使用最合适的 GPU 互连路径卸载或加速集合操作。

其他主流实现经历了类似演进。MPICH 在 3.4 版本通过其 CH4 通信层引入了 GPU 支持，该层管理设备缓冲区并选择高效通信路径。随后通过 ROCm 扩展了对 AMD GPU 的支持（自 4.0 版本起），Intel GPU 支持也在开发中。在 Cray 系统上，用户必须通过设置 MPICH_GPU_SUPPORT_ENABLED=1 显式启用 GPU 支持。MPICH 还集成了 GDRCopy、GPUDirect RDMA 和 IPC 以优化传输。MPICH 已在其 CH4 设备层中采用 UCX 作为网络模块，正在进行的工作探索集成 UCC 以启用优化的 GPU 集合操作。

4.2 GPU 中心集合通信库（GPUCCL）

随着深度学习模型越来越大，其计算需求 necessitates 将训练部署到多个 GPU 上。鉴于集合通信在深度学习训练中的普遍性，NVIDIA、AMD 和 Intel 都提供了针对各自 GPU 架构优化的高效集合通信库。为简明清晰起见，我们将这些厂商特定方案统称为 GPU 集合通信库（GPUCCL）。它们已被集成为多个最先进深度学习框架的通信后端，包括 PyTorch、TensorFlow、MXNet、Caffe、CNTK 和 Horovod。

MPI 中涉及计算的 GPU 感知集合操作（如 reduce/allreduce）的实现方式是：用 GPU 内核做局部归约，再由 CPU 发起 GPU 间的拷贝以完成聚合。这种方式会产生多次内核启动和通信调用延迟，并且需要主机上的中间缓冲区。GPUCCL 采取不同的方法：在单个内核中同时实现集合操作的通信与计算。接下来，我们按厂商方案支持的特性来突出它们之间的差异。

4.2.1 厂商集合通信库对比

原文表 4 列出了三大厂商支持的集合通信库的主要特性：NVIDIA 的 NCCL、AMD 的 RCCL 和 Intel 的 oneCCL（从集合原语、执行模型、加速器集成和传输语义等维度对比；缩写：B=broadcast、R=reduce、AR=allreduce、RS=reducescatter、AG/AGv=allgather 及变体、AA=all-to-all、P2P=点对点）。

除了 GPU 中心执行模型之外，NCCL 的独特之处还在于其集合算法设计哲学。NCCL 采用统一设计原则：所有集合操作——AllReduce、AllGather、Broadcast、ReduceScatter 和 AllToAll——都实现为少量可复用通信原语（send、recv、reduce、copy）的序列，应用于通信器创建时确定的逻辑拓扑。Hu 等人的研究表明，NCCL 将每个集合操作映射到环形或树形通信图（分别面向带宽主导或延迟主导的场景选择），并使用固定大小的块（chunk）对这些操作进行流水化，以最大化计算与通信之间的重叠。NCCL 算法设计的一个关键架构元素是**并行通信通道（channel）**的使用。每个通道对应集合算法的一个独立实例，处理消息的一个切片。通道允许 NCCL 并行运行多个环（或树）阶段，利用 GPU 线程块并行性，并在 NVLink 或网卡路径上实现带宽饱和。如 Hu 等人所观察，通道数量直接影响算法吞吐：通道过少会降低链路利用率，通道过多则会增加排队和同步开销。因此，通道是将集合操作分解为可并行子算法的结构性机制。

NVIDIA 提供 NCCL，AMD 则提供名为 RCCL（ROCm Collective Communication Library）的类似库，其镜像 NCCL 的 API 和编程抽象，使深度学习框架能在 ROCm 平台上兼容。尽管 API 几乎相同，但由于 AMD 独特的硬件设计，性能特征差异显著。RCCL 运行在复杂的多芯片（multi-die）拓扑上，如 MI250/MI300 加速器，其中通信要穿越多个 XCD 或 Infinity Fabric（IF）层级，各层带宽特性各异。理解这些异构带宽层级对高效集合算法构造至关重要，也促使 RCCL 更强调路由和环选择启发式。因此，RCCL 构造的逻辑环和树显式反映芯片单元（tile）级连接和链路不对称性；环段可能被加长、划分或重排，以最小化对低带宽 IF 路径的穿越。这种设计使集合执行与物理 GPU 结构的绑定比在基于 NVSwitch 的系统上更为紧密。

RCCL 也采用 NCCL 风格的并行通信通道来流水化集合执行，但通道策略受多 tile GPU 布局和 ROCm 分区模式（如 CPX/NPS4）影响。每个通道处理消息的一个切片，同时尝试将工作局部化到单个 XCD 内以减少跨 tile 流量，形成一种算法级的 tile 感知并行。Hidayetoglu 等人进一步观察到，RCCL 的延迟和扩展行为同时反映分块策略和多跳 IF 穿越的成本，使得通道配置对布局的敏感度高于同构的 NVIDIA 设计。

尽管存在这些架构和算法挑战，RCCL 实现了与 NCCL 相同的高层集合模式——环、树和层级算法——并复用 NCCL 的网络插件 ABI 以支持 InfiniBand 和 RoCE 传输。这一特性兼容性使得在底层传输允许时，可以跨厂商共享通信后端。然而，与配备 NVSwitch 交换结构的 NVIDIA 系统不同，当前 AMD 系统缺乏基于交换机的全互连，这使得拓扑感知的集合构造和 tile 级算法设计对于在大型多 GPU 节点上扩展 RCCL 尤为重要。

除 NCCL 和 RCCL 外，Intel 在 oneAPI 生态中为 CPU 和 GPU 集群提供 oneCCL（oneAPI Collective Communications Library）。与在 GPU 内核内执行集合操作的 NCCL/RCCL 不同，oneCCL 采用构建在 MPI 和 libfabric 之上的中间件导向设计，将传输选择和拓扑感知委托给底层通信栈。oneCCL 并非提供纯 GPU 驻留的集合引擎，而是暴露一组可移植的集合通信原语和 API，设计为在 CPU、集成 GPU 和独立 Xe 系列加速器上统一运行。其集合操作的实现反映了一种分层编排模型：GPU 内存通过 SYCL/Level Zero 设备队列管理，但集合操作的编排主要由 CPU 工作线程执行，这些线程代表 GPU 调度和推进操作。ATL（抽象传输层，oneCCL 架构中的模块化后端）将集合原语映射到 MPI 或 libfabric（OFI）等底层通信结构上。ATL 决定传输语义、消息推进和端点选择，而 Level Zero 或 SYCL 负责设备内和设备-主机数据搬运。该架构使 oneCCL 能将多种加速器类型集成到统一的集合命名空间中：CPU 进程和 GPU 进程可以参与同一次集合调用，由 oneCCL 协调跨主机内存、PCIe 和 Xe-Link 结构的数据搬运。其 API 确保集合原语的运行独立于底层缓冲区分配在 CPU RAM、GPU HBM 还是共享统一内存中。这在 Aurora 等系统上尤为重要——那里 CPU 与独立 PVC GPU 构成紧耦合的多加速器节点。

Kwack 等人的实证分析表明，oneCCL 在这些多样化端点之间统一了集合通信，传输层行为委托给 Slingshot 网络上的 MPI/libfabric。然而，oneCCL 的抽象对通信推进和 GPU 中心集合设计有影响。与集合内核完全在 GPU 上运行的 NCCL 和 RCCL 不同，oneCCL 的 GPU 参与仍是依赖性的、由主机驱动的。因此，集合原语通过命令队列和事件依赖运作，但 Level Zero 内核不能独立发起网络活动。其结果是，重叠和延迟特性取决于 CPU 工作线程调度和 ATL 后端的性能。

近期的 GPU-aware MPI 研究凸显了这一区别：它们表明当主机干预成为限制因素时，直接基于 Level Zero 和 IPC 的方法在 GPU-GPU 集合操作上可以优于 oneCCL。性能分析研究进一步表明，oneCCL 在层级化多 GPU 系统上呈现出与 NCCL/RCCL 不同的扩展模式，正是因为其集合原语通过混合 CPU/GPU 控制结构执行。

4.2.2 集合通信的进一步工作

尽管 GPUCCL 有显著优势，它也面临若干固有挑战，主要源于资源争用及其静态抽象模型。第一，使用 GPU 线程同时做通信和计算的根本问题在于：两者现在争夺同一种有限资源。此时，如果计算先于集合操作被调度，它可能独占所有 GPU 资源，实际上将集合操作串行化到计算之后。一种变通方法是以更高优先级的流启动集合操作，使其总是被优先调度。第二，GPUCCL 的执行模型针对规则集合模式定制——通信大小和数据类型作为主机参数静态表达。虽然适合深度学习，但这种设计缺乏更动态通信场景所需的灵活性。第三，NCCL 的设计往往将互连限制为单一传输模式（如线程拷贝而非 DMA 拷贝），其保守的同步方法也阻碍了细粒度双缓冲等高级优化技术的实现——后者能有效隐藏通信延迟。

为解决上述问题，已有多家集合通信库被提出；值得注意的是，许多由工业伙伴开发的方案现已开源、可直接使用。

Blink 是一个以最优链路利用率为明确目标的集合通信库。为此，Blink 检测底层拓扑，将拓扑建模为图，然后使用"打包生成树"（packing spanning trees）技术动态生成通信原语。结果表明，Blink 在图像分类任务上相比 NCCL2 将模型训练时间减少了 40%。
Aluminum（Dryden 等人）是一个用于大规模深度神经网络训练的 GPU 感知库。Aluminum 用树形算法扩展 NCCL，以避免其默认环形实现的延迟瓶颈，还增加了对非阻塞 NCCL allreduce 操作的支持。对 MPI，Aluminum 通过将单个 GPU 流与 MPI 通信器关联并仅对该流同步，绕过了 MPI-GPU 语义不匹配导致的强制同步。这些优化带来了优于 GPU-aware MPI 和基于 NCCL 实现的加速。
NCCLX（Meta 开发）为 NCCL 引入了一个透明加速层，专门针对深度学习工作负载。NCCLX 观察到，大多数数据并行训练作业在固定消息形状上反复调用相同的集合模式——主要是 AllReduce 和 AllGather。利用这种规律性，NCCLX 通过集合融合、路径合并和消除冗余同步与内存搬运阶段的专用内核流水线，在 NCCL 现有内核内构造优化执行路径。关键是，NCCLX 保持 NCCL 的公开 API，无需修改 PyTorch 或 TensorFlow 等框架，是标准深度学习训练栈的性能增强。
UCC（Unified Collective Communication） 是一个开源项目，提供高性能、可扩展集合操作的 API 和库实现，利用拓扑感知算法以及网内计算、DPU 卸载等技术。它依赖 UCX 点对点通信，以及 NCCL/RCCL、SHARP 等。UCC 提供一个可移植的、后端无关的框架，通过在基于 UCX 的实现、SHARP 卸载和 NCCL、RCCL 等厂商库之间动态选择，统一 CPU、GPU 和 DPU 上的集合操作。
MSCCL（Microsoft Collective Communication Library）引入了一个围绕设备驻留原语（put、signal、wait、flush）和用于算法合成的高级 DSL 构建的低层 GPU 中心通信底座，支持细粒度、融合的通信-计算内核，在 LLM 等新兴 AI 工作负载上优于传统 NCCL 路径。MSCCL++ 在小消息上分别比 NCCL 和 MSCCL 快达 2.8 倍和 1.6 倍，在大消息上快达 2.4 倍和 2.0 倍。
HiCCL（Hierarchical Collective Communication Library） 采用层级感知设计，将集合操作分解为通用原语（多播、归约、栅栏），并映射到多层加速器拓扑——包括 GPU tile、多 GPU 节点和多网卡集群——在 NVIDIA、AMD 和 Intel 系统上实现可移植性能。还有许多其他工作致力于将集合操作卸载统一到 MPI 运行时中。
MPI-xCCL 集成 NCCL、RCCL 和 HiCCL，创建混合执行模型：MPI 在有利时透明地将集合操作卸载到厂商库，同时保留 MPI 语义，并支持 GPUCCL 原生不支持的集合操作。

除可移植性和拓扑感知外，近期研究还开始沿能效和配置自动调优等新维度探索集合通信优化。PCCL 在 NCCL 之上引入了一个功耗感知的集合通信层，其依据的观察是：许多集合内核——尤其是 LLM 工作负载中的 AllReduce 和 AllGather——对频率不敏感，可以在大幅降低的 GPU 时钟频率下运行而无明显带宽损失。通过将细粒度 DVFS 管理集成到集合调用路径中，PCCL 为每个集合操作确定最低安全 GPU 频率，并在 NCCL 内核周围自动插入频率设置和恢复事件。

与以功耗为中心的优化互补，许多库致力于解决巨大配置空间（算法、协议、传输、通道、线程、块大小）的优化挑战——默认的 GPUCCL 成本模型对许多消息尺寸、GPU 拓扑和通信模式而言可能是次优的。AutoCCL 提供自动化在线调优框架：在训练早期迭代中对集合执行进行剖析，将参数分为实现级参数与资源分配参数。与以往的离线调优器不同，AutoCCL 直接考虑通信-计算均衡，在 PCIe 和 NVLink 系统上都频繁实现相比 NCCL 1.2–1.8 倍的带宽提升，LLM 工作负载的端到端迭代时间提升最高达 32%。该方法凸显了对能够根据硬件条件、工作负载模式和运行时干扰动态调整执行策略的集合运行时日益增长的需求。

4.3 GPU 中心 OpenSHMEM（GPUSHMEM）

GPU 中心 OpenSHMEM 运行时是 NVIDIA 的 NVSHMEM、AMD 的 ROC_SHMEM 和 Intel 的 Intel SHMEM 库。鉴于 NVSHMEM 出现最早，我们先讨论它，并在此过程中引入对三个库都基础的概念。注意，NVSHMEM 和 ROC_SHMEM 是为各自 GPU 生态从头设计的 GPU 中心 PGAS 运行时。两个库都遵循 OpenSHMEM 编程模型，但引入了自己的运行时结构和调用约定，缺少 OpenSHMEM 规范中的上下文（context）抽象，且其主机 API 仅作用于设备驻留的对称内存。相比之下，Intel SHMEM 严格遵循 OpenSHMEM 规范，通过基于 SYCL 的统一 API 同时支持主机和设备指针。

4.3.1 NVSHMEM

NVSHMEM 是 NVIDIA 面向 CUDA 设备的 OpenSHMEM 规范实现。它是一个分区全局地址空间（PGAS）库，提供高效的单边 put/get API，供进程访问远程数据对象。NVSHMEM 支持节点内和跨节点的 GPU 间点对点与集合通信。

NVSHMEM 基于对称堆（symmetric heap）概念工作。NVSHMEM 初始化期间，每个映射到 GPU 的进程——称为处理元素（PE）——使用 nvshmem_malloc() 保留一块 GPU 内存。在 NVSHMEM 中，所有内存分配必须以集合方式进行，即堆中所有对称内存区域必须大小相同、同时分配。要访问不同 PE 上的远程内存，给定 PE 需要对称内存的偏移量以及远程 PE 的 rank。

此外，NVSHMEM 提供同步一组 PE 的 API。这些 API 包括信号-等待机制（可作为点对点同步手段）和可充当全局屏障的集合同步调用。这一特性尤为重要，因为内核侧普遍缺乏全局屏障，通常由 CPU 扮演设备全局同步者的角色。设备在不终止内核执行的情况下跨设备高效同步的能力，是将控制平面转移到 GPU 的关键前提。

NVSHMEM 的一个显著特性是同时提供主机侧和设备侧 API。主机侧 API 暴露可选的流参数，可用于实现通信-计算重叠。对某些调用，GPU 侧变体提供三种粒度：线程、线程块和 warp。线程变体意味着调用由单个设备线程发起并由该线程执行；线程块和 warp 变体使用多个线程协作执行通信调用，应由相应线程块或 warp 中的所有线程调用。此前对主机侧与设备侧 API 的性能比较发现两者性能差异可忽略，主机侧 API 略优。该研究使用的是早期版本 NVSHMEM（0.3.0），此后 GPU 侧 API 性能已有改进。

自 2.7.0 版本起，NVSHMEM 引入了构建在 GPUDirect Async 之上的 InfiniBand GPUDirect Async（IBGDA）传输层。IBGDA 传输层允许 GPU 直接向网卡发起跨节点通信，完全绕过 CPU。没有 IBGDA 时，设备侧跨节点通信调用通过 CPU 上的代理线程执行，由代理线程触发相应的网卡操作。该代理线程消耗 CPU 资源，并在细粒度传输上形成达到网卡峰值吞吐的瓶颈。支持 IBGDA 的 NVSHMEM 与持久内核结合，能够将数据路径和控制路径完全转移到 GPU，标志着向完全自主多 GPU 执行的重大转变。然而，如 3.2.3 节所述，GPUDirect RDMA 仅在内核边界强制 GPU-NIC 内存一致性。这种对 CPU 内存一致性保证的固有依赖，是通往真正自主多 GPU 执行的潜在障碍。一种变通方法是使用回调机制：持久内核向 CPU 发信号，由 CPU 执行强制一致性的 API 调用（即 cudaDeviceFlushGPUDirectRDMAWrites()）。该方案的有效性尚不明确，有待进一步研究。在内核内强制 GPU-NIC 内存一致性由 ROC_SHMEM 支持，我们在下一节讨论。

近年来，NVSHMEM 已被集成为多个运行时的通信后端。PETSc 实现了 PetscSF——基于 NVSHMEM 的可扩展通信层，以补充其基于 MPI 的方法（后者与 CUDA 流语义配合不佳，阻碍内核启动流水线化）。Kokkos Remote Spaces 为 Kokkos 编程模型添加分布式内存支持，使用 NVSHMEM 作为通信后端；Kokkos 共轭梯度求解器的 NVSHMEM 实现优于 CUDA-aware MPI 实现，同时显著减少了代码量。Choi 等人使用持久内核和 NVSHMEM 实现了 CharminG——一个受 Charm++ 启发的完全 GPU 驻留运行时系统。Livermore Big Artificial Neural Network（LBANN）使用 NVSHMEM 实现了空间并行卷积，优于 MPI 和 Aluminum 实现。QUDA（格点 QCD 计算库）使用 NVSHMEM 和持久内核改进了 Dirac 算子的强扩展性。

NVSHMEM 也在运行时方法之外被用于实现性能提升。Chu 等人将 NVSHMEM 与持久内核结合，实现了最先进的 GPU 键值存储。Xie 等人使用 NVSHMEM 实现单节点多 GPU 稀疏三角求解器（SpTRSV），相比基于 UVM 的设计取得了良好的性能可扩展性。Ding 等人将持久内核与 NVSHMEM 结合，在单节点和多节点稀疏三角求解器上取得了令人印象深刻的性能。Atos 实现了持久内核和离散内核与基于 NVSHMEM 的通信，在节点内和跨节点的多 GPU BFS 上达到最先进的性能。Wang 等人提出 MGG：一种使用 NVSHMEM 细粒度通信、以 GPU 中心软件通信-计算流水线在多 GPU 系统上加速图神经网络（GNN）的系统设计。Ismayilov 等人使用持久内核和设备侧 NVSHMEM 实现了完全 GPU 侧的 Jacobi 2D/3D 和 CG 求解器，优于 CPU 控制的基线；他们保留部分线程块用于通信、其余用于计算——称为线程块专化（thread block specialization）技术——以实现显式的设备侧通信-计算重叠。Punniyamurthy 等人使用 ROC_SHMEM 和持久内核，在深度学习推荐模型中将嵌入操作与集合通信重叠。

4.3.2 ROC_SHMEM

ROC_SHMEM 是 AMD 面向 AMD GPU 的 OpenSHMEM 规范实现。ROC_SHMEM 提供两种通信后端：第一种称为 GPU-IB，在 GPU 上实现 InfiniBand，类似 NVSHMEM 的 IBGDA 传输层；第二种称为反向卸载（Reverse Offload，RO），使用主机侧代理线程将通信卸载到 CPU。GPU-IB 是默认后端，性能最佳。ROC_SHMEM 的工作方式与 NVSHMEM 几乎相同，提供类似的 API，但存在若干重要差异。

第一，如上一节所述，NVSHMEM 在从持久内核发起节点内通信时会遇到 GPU-NIC 内存一致性问题。ROC_SHMEM 则明确解决了该问题，保证使用持久内核时的正确性。Hamidouche 等人分析了源于 GPU 宽松内存模型的 GPU-NIC 内存不匹配问题，所提修改已集成到 ROC_SHMEM 中。这意味着 ROC_SHMEM 提供了完全免 CPU 的通信机制，可将多 GPU 执行的整个流程移到设备上。

第二，ROC_SHMEM 使用 GPU 共享内存（AMD 术语中的本地数据存储 LDS）存储网络状态以加快访问。据我们所知，NVSHMEM 未实现此优化。虽然这很可能有利于执行时间，但增加的共享内存占用可能限制占用率（occupancy），对性能产生负面影响。

第三，早期版本的 ROC_SHMEM 要求将对称缓冲区分配为不可缓存（uncacheable），以防止传输陈旧数据。然而，AMD 最近引入了内核内缓存冲刷指令，数据可以在发起网络事务前冲刷，从而允许数据被缓存。NVIDIA 未提供此类指令，意味着 NVSHMEM 缓冲区很可能被分配为不可缓存。

4.3.3 Intel SHMEM

Intel 最近推出了 Intel SHMEM——首个面向 Intel GPU 的 GPU 感知 OpenSHMEM 实现。它允许 SHMEM 例程直接操作 GPU 内存，并通过将 SHMEM 调用嵌入 SYCL 内核来支持 GPU 发起的通信。Intel SHMEM 与 SYCL 编程模型集成，为异构系统提供可移植的 C++ 接口。相比之下，NVSHMEM 和 ROC_SHMEM 等现有方案提供类似能力，但绑定于厂商特定生态。

为保持与 OpenSHMEM 1.5 规范的兼容，Intel SHMEM 同时支持设备侧和主机侧 API：点对点操作、通过 teams API 的集合操作，以及针对工作组（work-group）和子组（sub-group）通信的 SYCL 特定扩展。跨节点通信通过主机代理线程处理，该线程将 GPU 发起的操作转发给标准 OpenSHMEM 后端。换言之：GPU 在 SYCL 内核内发起 SHMEM 操作；GPU 将工作队列条目写入 GPU 与主机均可访问的内存区域；运行在 CPU 上的主机线程检测到 GPU 发起的 RMA 操作；主机使用标准 OpenSHMEM 库执行实际的 RDMA。当前实现依赖 Sandia OpenSHMEM（SOS），利用其对 OFI/libfabric 传输的稳健支持以及在 GPU 内存中维护对称堆的能力。这种分层设计使 Intel SHMEM 能够提供与 NVSHMEM 和 ROC_SHMEM 功能等价的能力，同时与 Intel 更广泛的 oneAPI 和 SYCL 生态干净地集成。

对节点内通信，Intel SHMEM 可以直接执行 GPU 到 GPU 数据搬运而不涉及主机 CPU。其实现方式是允许 GPU 从 SYCL 内核内发起 SHMEM 操作，使用 Intel 基于 Level Zero 的后端将这些操作转换为：（1）直接 GPU 加载/存储传输（当 GPU 共享统一内存结构时），或（2）GPU 拷贝引擎传输（绕过主机 CPU，使用 GPU 上的专用 DMA 引擎）。

4.4 用户级库的比较与讨论

虽然 GPU-aware MPI、GPUCCL 和 GPUSHMEM 都提供了多 GPU 系统编程机制，但它们在语义和性能特征上存在显著差异。关键区别包括流支持、API 位置、编程方式和性能。

4.4.1 流支持

GPU 基于流的概念运作——流是保证 GPU 操作顺序的命令队列。GPU 调度器确保在流上启动的内核和其他操作按入队顺序执行，并保持正确的数据依赖。由于内核启动是异步的、不阻塞主机，GPU 运行时可以将内核启动流水化，将启动延迟隐藏在内核执行之后。MPI 与 GPU 模型之间的语义不匹配在于：MPI 对 GPU 流没有感知。因此，无法将 MPI 调用排入某个 GPU 流，也无法让 GPU 流等待某个挂起的 MPI 例程完成。由此推论，将 MPI 调用与 GPU 内核交错执行需要主机阻塞式同步以保持数据正确性。例如，在发起 MPI 发送前，程序员必须阻塞主机以同步所有操作发送缓冲区的流；类似地，等待挂起的 MPI 通信完成也需要主机阻塞式同步。无论哪种情况，这些强制同步都会损害内核启动流水线化、妨碍重叠机会，迫使程序员在通信与计算的块状交替阶段之间切换。

我们看到两条互不排斥的可能路径来解决语义不匹配。第一条是使 MPI 运行时流感知：为 MPI 例程添加显式的流参数。这将解决内核启动流水线受损的问题，使 MPI 调用无缝集成到 GPU 运行时中。第二条是提供设备发起 MPI 调用的选项。这将减轻程序员在两种不同编程模型之间周旋的负担，并额外提供隐式的通信-计算重叠。两个方向都已在文献中被有限规模地探索。FLAT 编译器自动将设备侧 MPI 调用转换为其主机侧等价物。dCUDA 实现了具有 MPI 语义的设备侧操作，但实际通信使用 CPU 辅助线程；它们依靠 GPU 固有的内存延迟隐藏能力隐式地将通信与计算重叠，最终优于 GPU-aware MPI 基线。Namashivayam 等人探索了在 MPI 中引入 GPU 流感知的新通信方案：他们使用 HPE Slingshot 11 互连的触发操作特性，允许 CPU 将通信和同步操作排入网卡，再由 GPU 触发。这减少了 CPU 在关键路径中的参与，消除了昂贵的同步。虽然跨节点实验显示出一定性能优势，但所提方案在节点内配置下表现不佳，因为需要推进线程（progress thread）来模拟延迟执行语义。后续工作为节点内通信消除了推进线程，改用基于 P2P 直接加载/存储的 GPU 内核和基于 GPU IPC 的机制，评估显示相比流无感知 MPI 基线有性能提升。

更近的 MPIX streams 允许应用显式映射其 GPU 执行流上下文并传给 MPI 库，使 MPI 实现能够直接在 GPU 流上操作，从而消除不必要的同步开销、提升 GPU 到 GPU 通信性能。然而，这些方案均尚未被纳入 MPI 官方标准。

4.4.2 主机侧 vs 设备侧 API

GPU-aware MPI 和 GPUCCL 都使用主机侧 API，要求通信例程及其参数（如大小和数据类型）在主机 CPU 上用主机参数表达。GPUSHMEM 同时提供主机侧和设备侧 API，允许直接从 GPU 设备内核发起通信，输入参数定义为设备变量。这种以设备为中心的方式使 GPUSHMEM 能更好地处理低延迟和动态通信模式：数据传输在计算内核内发起并直接发送到网络，最小化延迟，非常适合细粒度通信。

然而，设备侧调用引入了编程和性能测量挑战。在单个内核内混合计算与通信使编程模型复杂化；而且由于当前剖析工具仅在内核粒度上工作，性能剖析并非易事。此外，尤其是与持久内核结合时，该技术会导致资源争用——流处理器和线程必须在两者之间共享。当采用划分线程块（TB）实现重叠等技术时，专用通信 TB 与计算 TB 之间的同步变得困难。虽然线程块簇（Thread Block Cluster，TBC）可能缓解此同步问题，但其编程会进一步增加整体编程复杂度。

仿效 NVSHMEM，NVIDIA NCCL 2.28 引入了新的设备侧 API，以实现通信与计算的融合。这与早期版本（所有 NCCL 操作均由主机发起）相比是重大转变。新 API 允许 GPU 内核直接发起数据搬运，其使用需要设置带有对称内存窗口的数据缓冲区，以实现直接的 GPU 到 GPU 通信。

4.4.3 编程示例

代码清单 1、2、3 分别展示了使用 GPU-aware MPI、NCCL 和设备侧 NVSHMEM 的简化单向带宽基准。这些示例突出了三种通信库之间的语义差异。三者都接收缓冲区指针和大小，但 MPI 和 NCCL 依赖具有同步发送/接收语义的双边通信模型；相比之下，NVSHMEM 采用单边模型：put/get 操作相对于远程 GPU 是异步的。这一区别在清单 1、2 与清单 3 的对比中可见——NVSHMEM 要求发送方直接指定接收方的缓冲区地址。

清单 1：GPU-aware MPI 简单单向带宽基准

if (rank == 0) {
    for (int j = 0; j < window_size; ++j) {
        MPI_Isend(send_buf, message_size, MPI_FLOAT, 1, 100, MPI_COMM_WORLD, send_request + j);
    }
    MPI_Waitall(window_size, send_request, reqstat);
    MPI_Recv(recv_buf, 1, MPI_FLOAT, 1, 101, MPI_COMM_WORLD, reqstat);
} else {
    for (int j = 0; j < window_size; ++j) {
        MPI_Irecv(recv_buf, message_size, MPI_FLOAT, 0, 100, MPI_COMM_WORLD, recv_request + j);
    }
    MPI_Waitall(window_size, recv_request, reqstat);
    MPI_Send(send_buf, 1, MPI_FLOAT, 0, 101, MPI_COMM_WORLD);
}

清单 2：NCCL 简单单向带宽基准

if (rank == 0) {
    ncclGroupStart();
    for (int j = 0; j < window_size; ++j) {
        ncclSend(send_buf, message_size, ncclFloat, 1, comm, stream);
    }
    ncclGroupEnd();
    ncclRecv(recv_buf, 1, ncclFloat, 1, comm, stream);
} else {
    ncclGroupStart();
    for (int j = 0; j < window_size; ++j) {
        ncclRecv(recv_buf, message_size, ncclFloat, 0, comm, stream);
    }
    ncclGroupEnd();
    ncclSend(send_buf, 1, ncclFloat, 0, comm, stream);
}

清单 3：设备侧 GPUSHMEM 简单单向带宽基准

__global__ void comm_kernel_send(...) {
    nvshmemx_putmem_signal_nbi_block(recv_buf, send_buf, nx,
        signal_buf + blockIdx.x, 1, NVSHMEM_SIGNAL_ADD, 1);
    grid.sync();
    if (blockIdx.x == 0 && threadIdx.x == 0)
        nvshmem_signal_wait_until(signal_buf, NVSHMEM_CMP_GE, i + 1);
}

__global__ void comm_kernel_recv(...) {
    if (threadIdx.x == 0)
        nvshmem_signal_wait_until(signal_buf + blockIdx.x, NVSHMEM_CMP_EQ, i + 1);
    grid.sync();
    if (blockIdx.x == 0)
        nvshmemx_putmem_signal_nbi_block(recv_buf, send_buf, 1 * sizeof(real),
            signal_buf, 1, NVSHMEM_SIGNAL_ADD, 0);
}

if (rank == 0) {
    nvshmemx_collective_launch(comm_kernel_send, dim3(window_size), dim3(1024), kernelArgs, 0, stream);
} else {
    nvshmemx_collective_launch(comm_kernel_recv, dim3(window_size), dim3(1), kernelArgs, 0, stream);
}

如 4.4.1 节所述，MPI 不暴露流参数，而 NCCL 和 NVSHMEM 操作显式绑定到 CUDA 流。NCCL 还支持在 groupStart/groupEnd 之间将多个操作分组，以摊销启动开销。由于 MPI 缺乏流语义，MPI 程序必须显式同步 CPU 与 GPU 的推进——例如使用 Waitall 确保完成。NVSHMEM 内核必须使用专门的集合启动例程（如 nvshmemx_collective_launch）启动，由于 GPU 不支持抢占，这也限制了线程块数量。

支持多种通信库常常迫使开发者重新实现通信后端，降低生产率。为解决此问题，Sağbili 等人提出了一个统一通信接口，将这些模型整合到单一 API 之下，并证明了与朴素实现相当的性能。

4.4.4 性能

多项工作通过真实应用或微基准比较了不同用户级库的性能。有研究在 NVIDIA 和 AMD GPU 上使用 MPI、NCCL/RCCL 和 NVSHMEM 实现了标准和流水线化共轭梯度（CG）。他们发现，通过流避免 CPU-GPU 同步可将 CG 性能显著提升 5–15%。对 NVIDIA 系统，他们建议小消息 AllReduce 使用 NCCL、点对点使用 MPI。在 AMD GPU 上，MPI 优于厂商方案，归因于厂商软件不够成熟。

NVSHMEM 的 GPU 发起通信的好处在使用 GROMACS 的分子动力学模拟中得到了证明，性能相比 MPI 提升达 2 倍。这些收益归因于 PGAS 模型的适用性——它将 CPU 移出关键路径，并支持内核融合等实现级优化。然而，随着 NCCL Device API 的近期推出，新研究表明其在点对点和 all-to-all 通信上的性能现已可与 NVSHMEM 相当。

在图处理应用——特别是广度优先搜索（BFS）——中也观察到类似优势。在这类工作负载中，固有的不规则、细粒度访问模式与 NVSHMEM 编程模型的契合度优于标准 MPI。其他不规则工作负载也报告了类似发现，如分布式布隆过滤器的更新和大规模图中连通分量的搜索。

其他工作在三台超级计算机上比较了 NCCL（和 RCCL）与 GPU-Aware MPI 在集合和点对点操作上的表现。结果显示，MPI 在点对点操作上倾向于表现更好，而 *CCL 在集合操作上更好。然而，这也取决于具体系统和这些库提供的优化，有时最佳方案同时取决于节点数量和向量大小。此外，论文报告了 RCCL 在高节点数下的不稳定性，Frontier 超级计算机上 Cray MPICH 与 RCCL 性能的比较也指出了这一点。

针对 Open MPI 和 Cray MPICH 报告的另一项限制，涉及基于归约的集合操作（如 MPI_Allreduce 和 MPI_Reduce_scatter）的执行。除非专门配置加速器卸载组件，这些 MPI 实现往往默认采用主机暂存协议：数据拷贝到主机、用 CPU 归约、再拷回设备。这种往返数据搬运和对 CPU 算术能力的依赖，相比直接在 GPU 上归约的原生 GPU 集合库，显著降低了性能。一项优化 MFDn 核组态相互作用代码的研究也解决了该问题：论文表明，用原生 CUDA 内核替换基线 MPI/OpenACC 实现并采用高性能通信协议，带来了可观的收益。对大型的、受网络带宽限制的问题规模，NCCL/CUDA 方法最为有效，通过消除归约的主机暂存并充分利用 GPU 原生通信与并行性，相比基线实现了最高 4.9 倍加速。

尽管在编程和性能上存在这些差异，三个库都已成熟，有多种实现可用；同时，如 4.2.2 节所述，新的 GPUCCL 变体正从工业界和研究界积极涌现，以应对不断演进的 AI 应用需求。

4.4.5 库之间的交互

原文图 4 展示了从应用向下到硬件的软件栈层级。应用通常直接对接 MPI、GPUSHMEM 或 GPUCCL 等高层库。UCX 或 libfabric 等较低层通信库通常作为这些高层框架使用的中间件，而非被应用直接访问。虽然传统 MPI 实现常直接运行在传输接口（如 libfabric 或 libibverbs）之上，现代 GPU 感知 MPI 栈更加模块化：通常在 UCX 之上处理点对点操作、在 UCC 之上处理集合操作。某些情况下（如 MVAPICH），MPI 也可能将集合操作直接卸载到 GPUCCL，以利用其拓扑感知优化。

类似地，GPUSHMEM 采用混合架构：使用 UCX 或直接传输接口（libfabric/libibverbs）处理低延迟点对点操作，同时在 GPUCCL 之上加速集合例程。最后，GPUCCL 通常直接对接传输层（libfabric 或 libibverbs），但也可通过特定插件配置为运行在 UCX 之上以增强可移植性。UCC 作为更高层抽象，通过将集合操作卸载到 GPUCCL 或使用自己在 UCX 和底层传输上的实现来执行集合操作。

5. 讨论、挑战与展望

随着多 GPU 执行成为许多真实工作负载的必需品，我们在本节进行讨论与展望，呈现我们认为当前和未来 GPU 中心通信研究的沃土。

5.1 摆脱 CPU

我们认为摆脱 CPU 控制的执行有前景，原因如下。第一，它解决了由内核启动和内存拷贝开销引起的 CPU 延迟壁垒问题。这些壁垒在强扩展场景中变得更加显著——随着 GPU 数量增加、每 GPU 计算量减少。在受延迟限制的场景中，传统 CPU 控制的实现无法将通信与计算重叠，因为发起操作的延迟比操作本身更长，实际上将通信与计算串行化。而免 CPU 执行即使在延迟占主导时也能实现足够程度的重叠。第二，内核内通信提供的并行性非常适合持久内核利用。借助高效的通信和同步 API，这种执行模型可以实现比 CPU 控制实现更高的带宽和更低的延迟。此外，免 CPU 执行通过将通信与计算内联，非常适合细粒度通信应用。第三，ROC_SHMEM 首次允许将应用执行完全迁移到设备上，对 CPU 辅助线程零依赖。采用 IBGDA 传输层的 NVSHMEM 也能将大量执行迁移到 GPU，但仍须依赖 CPU 保证功能正确性。

然而，这种执行模型面临若干挑战。一个主要挑战是持久内核可能导致占用率降低，潜在地制约计算。就目前而言，如果需要全局设备或多 GPU 屏障，持久内核必须以协作方式启动，这意味着只能启动可同时并发运行的线程数量，硬件超额订阅成为不可能。其结果是，以往由硬件调度器处理的工作负载分解和调度，现在需要程序员手动完成。这种手动方式不太可能像基于硬件的调度那样高效，计算密集型应用很可能受影响。此外，长时间运行的持久内核会消耗更多寄存器，也可能使用共享内存，进一步限制占用率。尽管如此，我们看到了该问题的若干解决方案。第一，应用执行期间有大量高带宽共享内存可用，可能抵消占用率降低带来的性能损失。第二，我们预测，随着 GPU 厂商越来越追求更高的 GPU 自主性，他们将引入允许持久内核与硬件超额订阅结合的 API。手动分解也可以由优化的编译器/运行时系统处理。另一种方案是将集合操作卸载到网络组件，释放 GPU 资源用于计算。

NVSHMEM 和 ROC_SHMEM 的另一个潜在问题是与现有运行时集成的难易程度。两个库都围绕对称堆构建，所有通信缓冲区必须由所有 GPU 在同一对称堆上以集合方式分配。这种对称分配需要库特定的分配器。现有运行时可能因对称内存分配要求而难以添加对 NVSHMEM 和 ROC_SHMEM 的支持。

5.2 UCX 作为 GPU 感知的潜在路径

统一通信框架 X（Unified Communication X，UCX）是一个开源通信框架，对多种网络 API、编程模型、协议和实现进行抽象。其理念是提供一组高层原语，同时将底层实现细节隐藏在 UCX 运行时之后。UCX 组织为三个主要组件：UCP（UC-Protocol），提供消息传递和远程内存访问等高层通信原语；UCT（UC Transport），提供对网络硬件的低层访问；UCS（UC Services），提供内存管理和线程的通用工具。

运行时，UCP 层根据指针类型和系统拓扑动态选择合适的 UCT 传输。例如，当 CUDA 指针传给跨节点 UCP 通信时，可能选择 rc_mlx5 传输；而节点内通信通常首选 cuda_ipc 传输。因此，UCP 可以接受设备指针作为数据负载，并使用最合适的机制执行传输，使 UCX 天然具备 GPU 感知。UCX 的 GPU 感知使 OpenMPI 和 MPICH 等高层通信库能够在其通信原语中直接处理设备指针。例如，在 CUDA-aware MPI 中，UCX 通过允许将 GPU 内存指针直接传给 MPI 函数，消除了显式的主机内存拷贝。

使用 UCX 实现 GPU 中心通信是文献中一个开始扎根的新方向。最相关的例子或许是 MPI 实现的 ROCm 感知。早期 GPU-aware MPI 工作大多是针对 NVIDIA GPU 使用原生 CUDA 库完成，然后直接集成到 MPI 运行时中。或许是不愿重复同样的工作来使其运行时具备 ROCm 感知，大多数 MPI 实现仅通过 UCX 提供 ROCm 感知。OpenMPI 和 MPICH 除原生集成外，还通过 UCX 提供 CUDA 支持。在非 MPI 工作中，Choi 等人扩展了 Charm++ 中的 UCX 层，为 Charm++ 生态中的若干编程模型提供 GPU 感知通信。

我们预测，更多运行时将趋向使用 UCX 来添加 GPU 中心通信支持。使用 UCX API 使程序员摆脱对厂商特定原生 API 的依赖，并允许同时为 ROCm 和 CUDA 添加 GPU 感知通信。对原生 API 实现的 GPU 感知与 UCX 提供的 GPU 感知进行性能比较将很有帮助。在这个方向上的一项工作中，Khorassani 等人为 MVAPICH2 提供了原生 ROCm 感知运行时，在 AMD GPU 集群上优于"OpenMPI + UCX"。然而，性能差异可能源于 MPI 实现本身的差异，而非 UCX。UCX 目前的一个限制是它始终在主机上：不使用任何通信内核，而是依赖传统的 cudaMemcpy 函数家族以及零拷贝 RDMA。

5.3 免 CPU 联网

随着 GPU 中心通信和更高 GPU 自主性的趋势持续加速，多项工作建议将大部分或全部联网栈迁移到内核。典型做法是启动单个长时间运行的持久内核，将数据路径和控制路径都移到 GPU。

在最早的工作中，GGAS 提议对网络设备做出修改，以实现统一的全局地址空间，使控制路径完全移到 GPU。这通过使用包含计算、通信和同步（全部在设备侧）的持久内核实现。虽然该工作首开先河并相比 CUDA-Aware MPI 基线展示了性能提升，但实验是在两个各有一个 GPU 的节点上进行的，所提硬件修改在 FPGA 上模拟。后续工作表明，GGAS 凭借消除控制路径中的 CPU 参与，相比 CPU 控制基线可以进一步提升性能并降低能耗。

随后出现了更多免 CPU 联网工作。Oden 等人使用 GPUDirect RDMA 允许 GPU 在无主机 CPU 参与的情况下直接对接 InfiniBand 网络设备。其做法是将整个 InfiniBand 上下文映射到设备侧，用 GPU 生成并向 HCA 发送工作请求。然而，由于 GPU 上单线程工作请求生成性能缓慢，所提修改相比 CPU 控制基线反而降低了性能。后续工作改善了这些性能限制，展示了更有前景的结果。另一项工作将所提 GPU 侧 InfiniBand Verbs 与 CUDA 动态并行结合，优化内核内同步瓶颈。GPUrdma 也在 GPU 上实现 InfiniBand，提出一个 GPU 侧库，用于从持久 GPU 内核内以零 CPU 参与直接通信；该设计在一系列微基准上优于 CPU 控制基线，但因持久内核与 GPU-NIC 交互而遇到正确性问题。

Silberstein 等人实现了 GPUNet，提供 GPU 侧套接字抽象和联网原语。GPUNet 允许在 GPU 上发起通信，但未将控制路径完全迁移到设备，而是依赖 CPU 辅助线程执行实际通信。dCUDA 采用类似方法，提供具有 MPI 语义的设备侧 API，但将其转换为由 CPU 辅助线程执行的标准 MPI 调用。LeBeane 等人对 GPU 联网方法进行了分类，深入讨论了其缺陷；作为回应，他们提出 GPU-TN——一种网卡硬件机制，允许 CPU 创建并向网卡注册消息，GPU 从运行中的持久内核触发它们。另一项工作 ComP-Net 使用嵌入式 GPU 微处理器，将辅助线程从 CPU 卸载到 GPU。虽然 GPU-TN 和 ComP-Net 都展示了有前景的性能，但它们需要对网卡和 GPU 的硬件修改，因此依赖模拟获得结果。

5.4 更广泛的 GPU 自主性

GPU 中心通信的近期普及代表了迈向更广泛 GPU 自主性的总体趋势。早期和近期的多项工作试图将传统上属于 CPU 职权范围的领域交给 GPU 掌管。早期工作中，Stuart 等人提出允许 GPU 向 CPU 发起回调的方法。Silberstein 等人实现了 GPUfs，允许 GPU 从 GPU 内核内直接请求主机 CPU 上的文件。Veselý 等人通过对 Linux 内核的修改，实现了从 GPU 内核内调用 POSIX 系统调用的支持。NVIDIA 的 GPUDirect Storage 在 GPU 与存储之间提供直接数据路径，但仍依赖 CPU 编排执行。

SmartIO 通过允许远程机器借用并直接访问 NVMe、GPU 和网卡等设备（如同本地设备一样），实现了跨 PCIe 连接主机的低成本、高效 I/O 解聚合。在此思路基础上，NVIDIA 的 Qureshi 等人提出 BaM，允许 GPU 在完全无 CPU 参与的情况下直接访问存储；实验结果显示 BaM 在若干工作负载上优于 GPUDirect Storage。Turimbetov 等人展示了多 GPU 上免 CPU 的设备侧任务图执行，消除了主机侧的内核启动和调度开销。最后，Baydamirli 等人提出了面向多 GPU 系统自主执行的编译器支持，使用 NVSHMEM 在 Python 代码中实现 GPU 发起的通信。

这些及其他工作展示了迈向通用 GPU 自主性的清晰趋势。与此一致，我们预期 GPU 中心通信会有进一步优化。如 Punniyamurthy 等人所指出的，若干近期机制很有前景：第一，近期的线程块簇（TBC）抽象可能有利于设备侧通信-计算重叠和线程块间同步；第二，AMD 近期的缓存冲刷指令允许在发起网络通信前冲刷缓存，意味着通信缓冲区不再需要分配为不可缓存；第三，"更胖的"GPU 节点和紧密的 GPU-NIC 集成等近期硬件趋势也很有前景。例如，最新一代 NVSwitch 直接连接 256 个 Grace Hopper 超级芯片，实现了空前规模的直接 P2P 全互连通信。

虽然这些进步推动系统走向更广泛的 GPU 自主性，但 CPU 仍扮演互补角色，承担系统级职责：调试、剖析、监控和软件开发仍从根本上依赖 CPU 对系统更广泛的可见性、成熟的工具生态，以及与操作系统服务的紧密集成。

5.5 集合算法设计

多 GPU 系统为集合操作算法设计带来了前所未有的挑战。第一，同一节点内的 GPU 常以复杂且非传统的拓扑排列，使现有集合算法效果降低。第二，GPU 对之间的带宽存在显著异构性，无论节点内还是节点间：同一节点上不同 GPU 对之间带宽差异可达 4 倍，节点内与跨节点带宽差距可达 10 倍。第三，这些拓扑和连接特性随 GPU 代际频繁变化。这些因素使集合算法设计复杂化，可能导致可用带宽利用不足和集合操作性能不佳。

一些方法使用线性规划公式，根据网络规格和集合操作规模寻找最优集合算法。然而，这涉及求解一个规模呈指数增长的 NP 难问题。例如，为 128 个节点求解可能需要长达 11 小时，且如果节点数或集合操作规模变化，可能需要重新求解。这种复杂性使得为大型系统生成集合算法即便不是不可能、也极具挑战。为简化多 GPU 集合操作的实现，MSCCL 提供了一种表达集合操作的高级语言，随后编译为 NCCL 代码。

5.6 调试、剖析与基准测试支持

高效的编程工具对高效的多 GPU 编程至关重要。然而，当涉及 GPU 原生通信时，现有工具严重匮乏。NVIDIA 的旗舰系统级剖析工具 NSight Systems 虽然提供了主机控制通信的详细视图，但在提供设备原生通信信息方面存在不足——包括直接加载/存储 P2P 通信，以及 NCCL 和 NVSHMEM 等库引起的通信。

Palwisha 等人提出 ComScribe——一种可监控 NCCL 集合和 P2P 通信的工具。然而，它限于单节点，且依赖已弃用的 nvprof 工具。Snoopie 在该领域做出了新的努力，聚焦于 GPU 中心通信的剖析和可视化。Snoopie 能够将通信归因到源代码行和参与通信的对象，提供不同粒度级别，既能给出系统的粗粒度概览，也能给出特定对象或设备在数据搬运方面的详细视图。

另一类有助于调试 GPU 通信的工具是竞态检测器。鉴于上述许多通信库使用分区全局地址空间模型（相比消息传递），引入竞态隐患的可能性很高。竞态检测工具对于驾驭多 GPU 编程中共享数据的复杂性至关重要。尽管重要，现有工具无一能检测多 GPU 编程中的竞态隐患。Compute Sanitizer 的 Racecheck 工具限于片上共享内存，仅支持在单个 GPU 上下文内检测竞态。HiRace 有所改进，支持全局内存，可检测 Compute Sanitizer Racecheck 无法检测的许多竞态类型，但也限于单个 GPU 上下文。

已有多项基准测试工具被提出，用于测量 GPU-GPU 通信性能。标准套件如针对 NCCL 和 RCCL 的 *ccl-tests 常用于评估集合和点对点通信性能。类似地，OSU 基准套件（OMB）为 GPU-aware MPI 环境提供广泛支持。NVSHMEM 也提供专门的微基准来评估设备发起的 put 和 get 延迟。然而，这些工具通常报告跨多次迭代聚合的性能测量，掩盖了关键的网络性能波动和瞬态噪声。此外，它们往往缺乏显式点对点拷贝的基线。为解决这些限制，Blink-GPU 基准被提出，能够捕获逐迭代计时，暴露聚合指标遗漏的性能异常。

我们认为，引入能够检测节点内和跨节点细粒度设备原生传输的调试和剖析工具，以及能够捕获多 GPU 上下文中竞态隐患的竞态检测器，对 GPU 中心通信的进一步发展至关重要。

5.7 压缩加速的 GPU 通信

跨网络搬运大数据量的高成本——尤其是跨节点 GPU 通信——推动了传输前压缩数据技术的发展。虽然这种方法引入了压缩和解压缩的计算开销，但这一成本常被通信时间的大幅减少所摊销。为最大化数据缩减（即实现高压缩比），通常采用有损压缩。具体而言，误差有界有损压缩允许高压缩比，同时保证数据失真保持在用户定义的精度阈值内。该技术已被集成到面向 CPU 和 GPU 架构的集合通信库中，相比 MPI 和 NCCL 等传统库展示了显著加速。

虽然能有效减少数据量，这些第一代方案从根本上受限于"解压-操作-压缩"（DOC）工作流：该工作流要求在任何计算（如 Allreduce 操作中的归约）执行前将数据完全解压，随后立即为后续步骤重新压缩，产生显著的处理开销。这一限制推动了同态压缩的发展——直接在压缩数据上执行计算。近期方法已在 CPU 和 GPU 架构上成功实现该范式。通过允许 GPU 完全在压缩域中计算和通信，该方法代表了最先进的方案，通过同时解决数据量瓶颈和内部 DOC 处理开销来最大化吞吐。

5.8 成熟度与用户级视角

澄清支撑现代加速器通信的技术（如 NVIDIA 的 GPUDirect、GPUDirect RDMA、AMD 的 ROCm 点对点能力和 GDRCopy）的角色和成熟度至关重要。这些组件不是高层的、面向用户的 API，也不是"黑科技"（hack）。相反，它们是基础性的、由厂商支持的硬件和驱动级功能，代表直接集成到加速器核心运行时环境和系统内核驱动中的官方底层架构。

这些方案的成熟度和可靠性，体现在它们被整个高性能开源通信生态普遍采纳为基本构件：高层库——包括所有主流 MPI 实现（如 Open MPI、MVAPICH、Intel MPI）、底层通信中间件（如 UCX）以及厂商自己的 *CCL——都构建在这些 API 之上。这些库显式使用这些底层功能来启用其 GPU 感知通信路径。因此，这些技术的好处（即消除 CPU 瓶颈，通过绕过主机内存实现真正的零拷贝、低延迟通信）是有充分文档记载的、稳定的、基础性的机制，支撑着所有现代多节点、多 GPU 超级计算。

6. 结论

传统上，多 GPU 通信由 CPU 管理，但 GPU 中心通信的近期进步已开始转移这一职责，使 GPU 能够对通信任务施加更多控制。本文深入探索了 GPU 中心通信，重点强调厂商提供的机制和用户级库支持，力求揭开该领域复杂性和选项多样性的神秘面纱，定义关键术语，并对单个节点内和跨多节点使用的各种方法进行分类。

讨论包括对厂商提供的多 GPU 执行通信与内存管理机制的分析，以及对主要通信库的综述——包括 CUDA-aware MPI、NCCL/RCCL/oneCCL、NVSHMEM、ROC_SHMEM 和 Intel SHMEM——突出其优势、挑战和性能考量。此外，本文介绍了免 CPU 联网、通信调试工具等重要研究范式，讨论了未来方向和未解问题。通过全面考察软硬件层上的 GPU 中心通信技术，本文旨在为研究人员、开发者、工程师和库设计者提供充分利用多 GPU 系统所需的知识。

致谢

科奇大学的作者得到欧洲研究理事会（ERC）在欧盟"地平线 2020"研究与创新计划下的资助（资助协议号 949587）。Flavio Vella 和 Daniele De Sensi 得到欧盟"地平线欧洲"计划资助（资助号 101175702，NET4EXA）。Daniele De Sensi 还得到罗马第一大学 ADAGIO 和 D2QNeT 项目资助（Bando per la ricerca di Ateneo 2023 和 2024）。

参考文献

注：参考文献保留原文（英文），此处仅列出主要条目；完整列表见原文。

Agostini et al. (2018) — GPUDirect Async: exploring GPU synchronous communication techniques for InfiniBand clusters. Journal of Parallel and Distributed Computing 114.
Agostini et al. (2017) — Offloading communication control logic in GPU accelerated applications. CCGrid '17.
Agostini (2023) — Inline GPU packet processing with NVIDIA DOCA GPUNetIO. NVIDIA Developer Blog.
Akhtar et al. (2021) — ComScribe: identifying intra-node GPU communication.
Alperen et al. (2025) — Optimizing nuclear configuration interaction calculations on GPUs. ISC High Performance 2025.
Baydamirli et al. (2024) — Autonomous execution for multi-GPU systems: compiler support. SC24-W.
Ben-Nun et al. (2020) — Groute: asynchronous multi-GPU programming model. ACM Trans. Parallel Comput. 7(3).
Brooks et al. (2024) — IntelSHMEM: GPU-initiated OpenSHMEM using SYCL. SC24-W.
Cai et al. (2021) — Synthesizing optimal collective algorithms. PPoPP '21.
Chen et al. (2023) — MPI-xCCL: a portable MPI library over collective communication libraries for various accelerators. SC-W '23.
Chen et al. (2022) — Scalable irregular parallelism with GPUs: getting CPUs out of the way. SC '22.
Choi et al. (2021) — CharminG: a scalable GPU-resident runtime system. HPDC '21.
Chu et al. (2019) — Designing high-performance in-memory key-value operations with persistent GPU kernels and OpenSHMEM.
Daoud et al. (2016) — GPUrdma: GPU-side library for high performance networking from GPU kernels. ROSS '16.
De Sensi et al. (2024) — Exploring GPU-to-GPU communication: insights into supercomputer interconnects. SC '24.
Di et al. (2025) — A survey on error-bounded lossy compression for scientific datasets. ACM Comput. Surv. 57(11).
Doijade et al. (2025) — Redesigning GROMACS halo exchange: improving strong scaling with GPU-initiated NVSHMEM. SC Workshops '25.
Dryden et al. (2018) — Aluminum: an asynchronous, GPU-aware communication library. MLHPC 2018.
Gysi et al. (2016) — dCUDA: hardware supported overlap of computation and communication. SC '16.
Hamidouche and LeBeane (2020) — GPU initiated OpenSHMEM: correct and efficient intra-kernel networking for dGPUs. PPoPP '20.
Hidayetoglu et al. (2024) — CommBench: micro-benchmarking hierarchical networks with multi-GPU, multi-NIC nodes. ICS '24.
Hidayetoglu et al. (2025) — HiCCL: a hierarchical collective communication library. IPDPS 2025.
Hu et al. (2025) — Demystifying NCCL: an in-depth analysis of GPU communication protocols and algorithms. arXiv:2507.04786.
Huang et al. (2025) — GhZCCL: advancing GPU-aware collective communications with homomorphic compression. ICS '25.
Ismayilov et al. (2023) — Multi-GPU communication schemes for iterative solvers: when CPUs are not in charge. ICS '23.
Issa et al. (2024) — Snoopie: a multi-GPU communication profiler and visualizer. ICS '24.
Jacobson et al. (2024) — HiRace: accurate and fast source-level race checking of GPU programs.
Jia et al. (2024) — PCCL: energy-efficient LLM training with power-aware collective communication. ICCD 2024.
Kwack et al. (2025) — AI and HPC applications on leadership computing platforms. IPDPS 2025.
LeBeane et al. (2017) — GPU triggered networking for intra-kernel communications. SC '17.
LeBeane et al. (2018) — ComP-Net: command processor networking for efficient intra-kernel communications on GPUs. PACT '18.
Li et al. (2020) — Evaluating modern GPU interconnect: PCIe, NVLink, NV-SLI, NVSwitch and GPUDirect. IEEE TPDS 31(1).
Muthukrishnan et al. (2021) — Efficient multi-GPU shared memory via automatic optimization of fine-grained transfers. ISCA 2021.
Namashivayam et al. (2022/2023) — Exploring GPU stream-aware message passing using triggered operations.
Oden et al. (2014) — InfiniBand-Verbs on GPU. IPDPSW 2014.
Oden and Fröning (2013) — GGAS: global GPU address spaces for efficient communication in heterogeneous clusters.
Potluri et al. (2018) — Efficient breadth first search on multi-GPU systems using GPU-centric OpenSHMEM.
Punniyamurthy et al. (2023) — GPU-initiated fine-grained overlap of collective communication with computation.
Qureshi et al. (2023) — GPU-initiated on-demand high-throughput storage access in the BaM system architecture. ASPLOS 2023.
Sağbili et al. (2025) — UniConn: a uniform high-level communication library for portable multi-GPU programming. CLUSTER 2025.
Shah et al. (2022) — TACCL: guiding collective algorithm synthesis using communication sketches.
Shah et al. (2025) — MSCCL++: rethinking GPU communication abstractions for cutting-edge AI applications.
Shamis et al. (2015) — UCX: an open source framework for HPC network APIs and beyond.
Si et al. (2025) — Collective communication for 100K+ GPUs.
Silberstein et al. (2014) — GPUfs: integrating a file system with GPUs. ACM TOCS 32(1).
Silberstein et al. (2016) — GPUNet: networking abstractions for GPU programs.
Singh et al. (2025) — The big send-off: high performance collectives on GPU-based supercomputers.
Spector et al. (2025) — We bought the whole GPU, so we’re damn well going to use the whole GPU.
Trotter et al. (2025) — CPU- and GPU-initiated communication strategies for conjugate gradient methods on large GPU clusters. SC '25.
Turimbetov et al. (2025) — A device-side execution model for multi-GPU task graphs. ICS '25.
Venkata et al. (2025) — Unified Collective Communication: a unified library for CPU, GPU, and DPU collectives. IEEE Micro 45(02).
Wang et al. (2019) — Blink: fast and generic collectives for distributed ML.
Wang et al. (2014) — GPU-aware MPI on RDMA-enabled clusters. IEEE TPDS 25(10).
Wang et al. (2023) — MGG: accelerating graph neural networks with fine-grained intra-kernel communication-computation pipelining. OSDI.
Weingram et al. (2023) — xCCL: a survey of industry-led collective communication libraries for deep learning. JCST 38(1).
Xie et al. (2021) — Fast and scalable sparse triangular solver for multi-GPU based HPC architectures. ICPP '21.
Xu et al. (2025) — AutoCCL: automated collective communication tuning. NSDI 25.
Zhang et al. (2021) — The PetscSF scalable communication layer.
Zhou et al. (2022) — MPIX stream: an explicit solution to hybrid MPI+X programming. EuroMPI/USA '22.