The Landscape of GPU-Centric Communication

发表时间: 2024-09 · arXiv:2409.09874 (Survey)

作者/机构: DIDEM UNAT, Koç University, Turkey; ILYAS TURIMBETOV, Koç University, Turkey; MOHAMMED KEFAH TAHA ISSA, Koç University, Turkey; DOĞAN SAĞBILI, Koç University, Turkey; FLAVIO VELLA, University of Trento, Italy; DANIELE DE SENSI, Sapienza University of Rome, Italy; ISMAYIL ISMAYILOV, Koç University, Turkey

A1 主要贡献

近年来，GPU因其大规模并行性和高内存带宽，已成为高性能计算（HPC）和机器学习（ML）应用的首选加速器。截至2025年11月，Top500超级计算机排名前10的系统中，有9个依赖GPU集群进行加速【79, TOP500 Supercomputer Sites, 2023, https://www.top500.org/】 。

核心问题: 尽管大量使用GPU能显著加速计算，但GPU之间的通信很快成为可扩展性的瓶颈【136, Peta-Scale Phase-Field Simulation for Dendritic Solidification on the TSUBAME 2.0 Supercomputer, 2011, SC ’11】【149, The Persistent Challenge of Data locality in Post-Exascale Era, 2025, Computing in Science and Engineering】。传统上，多GPU通信（无论是节点内还是节点间）一直由CPU负责，这种以CPU为中心的执行模型将GPU视为仅提供大量计算的设备，而在通信等辅助任务上依赖CPU。

研究目标: 过去十年中，一系列被称为“以GPU为中心的通信”的进步旨在挑战CPU在多GPU执行中的主导地位。这些进步从高层次上减少了CPU在关键执行路径上的参与，赋予GPU在发起和同步通信方面更多的自主权，并试图解决多GPU通信与计算之间的语义不匹配问题。

本文贡献: 本文对以GPU为中心的通信进行了全面的综述，重点关注供应商提供的机制和用户级库的支持。本文旨在：
1. 澄清复杂性：帮助研究人员、程序员、工程师以及库设计者理解该领域中各种可用选项的复杂性和多样性。以GPU为中心的通信涵盖了从专有硬件（如GPU间互连）到软件机制的广泛方法，这些方法各有优缺点，使得选择变得困难。
2. 定义术语和分类：定义“以GPU为中心的通信”，并对现有节点内和节点间的通信方法进行分类和抽象，以消除文献和供应商之间术语不一致造成的混淆。
3. 综述技术和库：介绍和讨论供应商为实现多GPU通信、网络和内存管理所提供的机制，这些机制是更高级别软件库的构建块。同时，回顾主要的通信库，讨论它们的优势、挑战和性能特点。
4. 展望未来：探讨以GPU为中心的通信背后的主要研究范式、未来前景和开放的研究问题。

本文通过广泛描述跨越软硬件堆栈的以GPU为中心的通信技术，为研究人员和开发者提供了如何最佳利用多GPU系统的见解。

A3 背景知识、关键观察与设计原则

我们可以将以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展示了节点内通信机制的分类示例及其数据路径。

Table 1. 节点内通信方法的类型

1中被称为主机原生（host native）的通信方法在主机侧进行，并且不涉及设备之间的直接P2P（点对点）访问。这包括所有可以在禁用P2P访问的主机侧启动的方法。否则，如2（主机控制，host-controlled）所示，通信不涉及额外复制到主机内存，而是直接通过PCIe、NVLink或Infinity Fabric互连。GPUCCL（即NCCL、RCCL、oneCCL等）、GPU-aware MPI和*memcpy操作具有主机侧API，因此它们可能属于1和2。通过启用对对等设备内存的直接访问，设备侧API消除了CPU在节点内通信中数据和控制路径的参与，如图1中的3（设备原生，device native）所示。NVSHMEM、ROCSHMEM、Intel SHMEM也提供主机侧API，但需要P2P访问，因此它们的主机侧API属于2，设备侧API属于3。内核内P2P直接加载和存储提供类似功能并属于类型3，但即使在P2P访问禁用的情况下也能工作，即类型4，此时数据路径回退到主机。

2.2 节点间通信

节点间场景更加多样化，因为需要与NIC（网络接口卡）进行交互。每种方法的实现细节可能涉及复杂的数据路径和决策。为了简化分类，我们区分了节点间通信的四个主要组成部分。除了节点内场景中使用的API和数据路径（到NIC），还有两个涉及与NIC交互的额外组件：

• 注册/构建消息：此步骤涉及构建数据包并在NIC上进行注册。
• 触发通信：它定义了谁在NIC上“敲门”（ring the doorbell）以发出数据传输。

我们在分类节点间通信时确定了五个主要类别。如表2和图2所示，从1到5，通信调用的更多组件被移至GPU侧，而数据传输路径到达NIC所需的数据复制次数减少。多年来，通信方法和相应技术描绘了数据传输和通信控制的优化，我们将在第3节中讨论。首先，1中完全在CPU侧的通信方法可用，然后在2中通过在GPU和NIC之间使用共享固定内存，消除了CPU-GPU和CPU-NIC缓冲区之间的额外复制，从而得到改进。这降低了GPU-NIC传输的延迟。之后，GPU RDMA（远程直接内存访问）3促进了NIC通过PCIe直接访问GPU内存，最小化了它们之间的数据路径。4代表了GPU触发的通信技术，如GPUDirect Async和GPU-TN【73, GPU Triggered Networking for Intra-Kernel Communications, 2017, SC ’17】，其中GPU变得能够发起通信，前提是CPU预先在NIC上准备好数据包。5将数据包准备和与NIC的交互也完全移至GPU，使设备原生通信成为可能。

表2和图2中给出的类型并未反映所有可能的组合，因为一些库根据配置和可用硬件可能会导致不同的数据路径和控制组合。例如，如果没有RDMA技术，即使使用GPU侧通信控制，数据路径也会涉及主机内存。

Table 2. 节点间通信方法的类型。粗体单元格表示发生变化或优化的位置。D/H表示设备侧和主机侧API调用都可能属于此类型。

例如，Intel SHMEM即使通信调用是在GPU上运行的设备内核内发出的，也会在主机上使用一个代理线程通过标准的OpenSHMEM库来执行实际的RDMA。

A2 方法细节

3 供应商机制

本节我们讨论供应商为实现多GPU执行中的通信、网络和设备间内存管理所提供的机制。这些机制由GPU编程模型运行时或作为扩展API的一部分提供。这些技术分为四类：内存管理器、GPUDirect技术、硬件和库。图3总结了NVIDIA提供的技术，详细说明了它们的时间线和可用性。

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

3.1 内存管理机制

• 3.1.1 页锁定/固定内存（Page-Locked / Pinned Memory）
  默认内存为可分页内存：默认情况下，使用设备malloc（例如cudaMalloc()、hipMalloc()）在主机上分配的内存是可分页的，GPU无法访问。当在可分页主机内存和设备内存之间进行传输时，GPU运行时必须首先将主机数据暂存到页锁定内存中的一个临时缓冲区，然后再将数据从页锁定内存复制到GPU。为了避免这种可分页内存到页锁定内存的复制，cudaMallocHost()允许直接分配页锁定内存，跳过中间的复制阶段。因此，页锁定内存也被称为零拷贝或固定内存【52, How to Optimize Data Transfers in CUDA C/C++, 2012, NVIDIA Developer Blog】【93, CUDA Programming Guide Release 12.2, 2023, NVIDIA】。

  固定内存的优势与劣势：固定内存以其在主机-设备传输中的高带宽和低延迟而闻名【115, Evaluating Characteristics of CUDA Communication Primitives on High-Bandwidth Interconnects, 2019, ICPE ’19】，通过实现直接的系统范围访问，有效协调CPU-GPU执行。它还与GPUDirect RDMA一起用于改善节点间通信【75, Evaluating Modern GPU Interconnect: PCIe, NVLink, NV-SLI, NVSwitch and GPUDirect, 2020, IEEE Trans. Parallel Distrib. Syst.】。然而，其物理内存锁定会导致高内存消耗，过多的分配可能会影响系统性能【52, How to Optimize Data Transfers in CUDA C/C++, 2012, NVIDIA Developer Blog】。

• 3.1.2 统一虚拟寻址（Unified Virtual Addressing, UVA）
  UVA实现共享虚拟地址空间：UVA是一种内存管理技术，它允许一个节点内的所有GPU和CPU共享同一个统一的虚拟地址空间【87, CUDA 4.0 Release Notes, 2011, NVIDIA】【93, CUDA Programming Guide Release 12.2, 2023, NVIDIA】。在UVA之前，主机↔设备和设备↔设备的复制必须明确指定传输方向。有了UVA，物理内存位置可以从指针值推断出来，从而减少了管理独立内存空间的开销，并使库能够简化其接口【128, Peer-to-Peer & Unified Virtual Addressing, 2011, NVIDIA Developer Webinar】。

• 3.1.3 进程间通信（IPC）
  IPC解决跨进程访问限制：在早期的GPU运行时版本中，指针无法跨进程边界访问，因此GPU缓冲区之间的内存复制必须通过主机进行，这造成了瓶颈。为了克服这一限制，进程间通信（IPC）使得同一节点上的进程能够无需额外复制即可访问其他进程的设备缓冲区【90, CUDA 4.1 Release Notes, 2017, NVIDIA】。通过IPC，内存句柄被创建并通过标准IPC机制在进程间传递，从而实现比通过主机暂存复制更低的延迟。

• 3.1.4 统一虚拟内存（Unified Virtual Memory, UVM）
  UVM提供单一地址空间：UVM允许通过cudaMallocManaged()调用分配受管内存，创建一个可由单个节点内所有处理器访问的单一地址空间。UVM的工作原理是将请求的内存划分为驻留在CPU上的页面。程序员可以在设备上访问内存而无需显式复制。如果内存访问属于不在设备上的页面，UVM会触发一个页错误，自动将该页面迁移到请求的设备。当总分页内存大小超过设备内存时，UVM驱动程序还可以将页面从给定设备驱逐回主机内存【93, CUDA Programming Guide Release 12.2, 2023, NVIDIA】。

  UVM的编程优势：UVM在可编程性方面提供了几个好处。首先，程序员面对的是一个单一的统一地址空间，他们可以像整个分配的内存块都驻留在单个GPU上一样访问它。系统中发生的任何复制都是隐式的，对程序员是透明的。此外，UVM允许内存超订（oversubscription），即可以分配比所有多GPU设备内存加起来还要多的内存。这是可能的，因为大部分内存可以保留在CPU上，并在给定设备请求时分页调入【92, Improving GPU Memory Oversubscription Performance, 2021, NVIDIA Developer Blog】【134, Oversubscribing GPU Unified Virtual Memory: Implications and Suggestions, 2022, ICPE ’22】。

3.2 GPUDirect技术

• 3.2.1 GPUDirect 1.0 (NIC)
  GPUDirect 1.0共享固定内存区域：GPUDirect 1.0允许GPU和NIC共享相同的固定内存区域。在此之前，用于GPU和NIC的系统内存中的固定内存区域是分开的。这意味着，要跨节点通信GPU数据，GPU首先将数据复制到其固定内存区域，然后CPU将其复制到NIC的内存区域，之后NIC才能访问并将其发送到网络，如图2中的类型1所示。GPU → NIC固定内存区域之间的中间CPU发起的复制增加了CPU开销并增加了GPU通信的延迟。GPUDirect 1.0引入了一个共享的GPU-NIC固定内存区域，从而避免了中间的CPU发起的复制【125, State of GPUDirect Technologies, 2016, GTC Presentation】【132, The development of Mellanox/NVIDIA GPUDirect over InfiniBand—a new model for GPU to GPU communications, 2011, Computer Science - Research and Development】。

• 3.2.2 GPUDirect 2.0 (Peer-to-Peer)
  GPUDirect P2P实现直接GPU间通信：随着UVA的引入，CUDA 4.0版本增加了对单个节点内共享相同PCIe根联合体（root complex）的GPU之间直接点对点通信的支持【87, CUDA 4.0 Release Notes, 2011, NVIDIA】。此功能被封装在一项称为GPUDirect 2.0或GPUDirect P2P的技术中。GPU现在可以直接通过PCIe访问彼此的内存，而不是通过主机暂存数据，从而首次建立了直接的GPU到GPU数据路径。这些变化带来了两种新的通信机制：P2P DMA复制，即cudaMemcpy调用会直接在源和目标GPU内存之间触发DMA传输；以及P2P直接加载/存储，GPU可以通过解引用指向远程GPU缓冲区的指针来直接访问数据。当NVLink技术（第3.4节）被引入时，GPUDirect P2P也增加了对其的支持【88, NVIDIA GPUDirect™ Technology, 2012, NVIDIA】【98, NVIDIA GPUDirect Family, 2023, NVIDIA】【125, State of GPUDirect Technologies, 2016, GTC Presentation】。

  GPUDirect P2P的优势：GPUDirect P2P提供了两个主要好处。它消除了冗余的GPU ↔ CPU复制和主机缓冲区，这些在通过CPU暂存传输时是必需的。此外，通过避免在主机上维护通信缓冲区并提供一种新的通信机制（P2P直接加载/存储），GPUDirect P2P提高了多GPU编程的便利性【88, NVIDIA GPUDirect™ Technology, 2012, NVIDIA】。

  UVA与P2P的关系：我们注意到，P2P DMA复制也可以在没有UVA支持的情况下工作。如果未启用UVA，可以通过使用cudaMemcpyPeer()变体并明确指定目标GPU来执行P2P DMA复制。然而，如果没有UVA，P2P直接加载/存储将无法工作，因为直接访问远程GPU的指针假定了一个统一的地址空间【93, CUDA Programming Guide Release 12.2, 2023, NVIDIA】。

• 3.2.3 GPUDirect RDMA
  GPUDirect RDMA实现跨节点直接通信：随着CUDA 5.0中GPUDirect RDMA的引入，NVIDIA GPU之间的跨节点直接通信变得可行。GPUDirect RDMA通过标准的PCIe功能促进了GPU与第三方设备之间的直接通信通道。该技术将GPU内存段暴露在PCIe内存资源上，称为基地址寄存器（BAR）区域。这使得NIC能够直接读/写GPU内存，而无需通过主机路由【96, GPUDirect RDMA, 2023, NVIDIA】。类似地，AMD提供ROCm RDMA（以前称为ROCnRDMA）【10, ROCK-Kernel-Driver, 2023, AMD】【12, ROCnRDMA, 2023, AMD】。GPUDirect RDMA为数据路径提供了几项优化，即消除了到主机内存的额外复制，减少了GPU-NIC交互固有的延迟，增加了带宽并减少了CPU开销。

• 3.2.4 GPUDirect Async
  GPUDirect Async优化控制路径：以前的GPUDirect技术专注于改善数据路径，而GPUDirect Async则优化GPU和NIC之间的控制路径。它在CUDA 8.0中引入，使GPU能够发起和同步网络传输，从而减少CPU在关键路径中的参与。GPUDirect Async的工作方式是让CPU预先注册消息，然后GPU内核可以通过在NIC上“敲门”来触发这些消息。因此，GPU可以在触发通信的同时继续执行，而不需要像以前那样停下来等待CPU发起通信【3, Offloading Communication Control Logic in GPU Accelerated Applications, 2017, CCGrid ’17】【4, GPUDirect Async: Exploring GPU synchronous communication techniques for InfiniBand clusters, 2018, J. Parallel and Distrib. Comput.】。

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

3.3 GPUNetIO

• GPUNetIO技术简介：GPUNetIO【1, Inline GPU Packet Processing with NVIDIA DOCA GPUNetIO, 2023, NVIDIA Developer Blog】是NVIDIA作为DOCA（数据中心单芯片架构）[34]一部分提出的技术解决方案。DOCA是一个全栈软件框架，旨在促进为NVIDIA BlueField数据处理单元（DPU）【22, Nvidia Data Center Processing Unit (DPU) Architecture, 2021, IEEE Hot Chips 33 Symposium (HCS)】开发应用程序。在非RDMA网络上，GPUNetIO允许GPU发送、接收和处理网络数据包【1, Inline GPU Packet Processing with NVIDIA DOCA GPUNetIO, 2023, NVIDIA Developer Blog】【2, Optimizing Inline Packet Processing Using DPDK and GPUDirect with GPUs, 2023, NVIDIA Developer Blog】。在RDMA网络（包括RoCE和InfiniBand）[34]上，从DOCA v2.7开始，GPUNetIO允许GPU不仅在内核边界执行RDMA发送和接收，而且可以在内核执行期间的任何时间点执行。简而言之，GPUNetIO允许GPU与NIC交互，无需任何CPU干预。

• GPUNetIO的同步机制：在RDMA网络上，GPU内核可以（以阻塞或非阻塞模式）等待RDMA接收操作的完成。在非RDMA网络上，GPUNetIO提供信号量，可以在内核内显式使用，用于在发送和接收数据包时与NIC同步。信号量也可以用于GPU内核与CPU的同步（如果数据包处理在CPU和GPU之间分配）或与其他CUDA内核的同步（如果数据包处理跨多个内核分配）。

3.4 现代以GPU为中心的互连技术

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

• NVLink：NVLink是NVIDIA GPU的专有互连技术。其设计旨在解决PCIe的带宽限制，PCIe在GPU加速应用中被认为是传输瓶颈【75, Evaluating Modern GPU Interconnect: PCIe, NVLink, NV-SLI, NVSwitch and GPUDirect, 2020, IEEE Trans. Parallel Distrib. Syst.】【91, NVIDIA DGX-1 With Tesla V100 System Architecture, 2017, NVIDIA】。表3展示了各代NVLink的规格【45, Ultra-Performance Pascal GPU and NVLink Interconnect, 2017, IEEE Micro】【75, Evaluating Modern GPU Interconnect: PCIe, NVLink, NV-SLI, NVSwitch and GPUDirect, 2020, IEEE Trans. Parallel Distrib. Syst.】【99, NVLink and NVSwitch, 2023, NVIDIA】。

  

• NVLink的演进：此外，NVLink也曾用于连接GPU与IBM Power8和Power9 CPU，但随着Grace Hopper超级芯片的推出，NVLink被用作芯片到芯片（C2C）互连，提供900 GB/s的双向带宽【75, Evaluating Modern GPU Interconnect: PCIe, NVLink, NV-SLI, NVSwitch and GPUDirect, 2020, IEEE Trans. Parallel Distrib. Syst.】【103, NVIDIA Grace Hopper Superchip, 2024, NVIDIA】。后来，随着Grace Blackwell超级芯片的引入，NVLink-C2C用于连接Grace CPU与2个Blackwell GPU，总共提供3.6 TB/s的双向带宽【102, NVIDIA GB200, 2024, NVIDIA】。NVIDIA Blackwell上的第五代NVLink为每个GPU提供1.8TB/s的双向吞吐量，为多达576个GPU之间提供高速通信。

• NVSwitch：NVLink的引入优化了NVIDIA GPU之间的带宽，使P2P通信成为节点内通信的可行机制，并将数据路径重度倾向于GPU。NVLink的一个缺点是它不是自路由的，这意味着如果任意两个GPU没有直接的NVLink连接，通信必须通过一个中间GPU进行路由【75, Evaluating Modern GPU Interconnect: PCIe, NVLink, NV-SLI, NVSwitch and GPUDirect, 2020, IEEE Trans. Parallel Distrib. Syst.】。这个限制被NVSwitch【99, NVLink and NVSwitch, 2023, NVIDIA】克服，NVSwitch是一种背板技术，可以实现所有GPU之间的全连接。例如，一个DGX-2节点包含16个V100 GPU，它们通过NVLink和NVSwitch全连接【95, DGX-2, 2023, NVIDIA】。从第三代开始，NVSwitch支持SHARP【104, NVIDIA NVLink SGXLS10 Switch Systems User Manual, 2024, NVIDIA】，它将allreduce操作卸载到NVSwitch，使allreduce能够以全线速运行【60, The NVLink-Network Switch: NVIDIA’s Switch Chip for High Communication-Bandwidth Superpods, 2023, HotChips Presentation】。

• AMD和Intel的方案：AMD的替代方案是Infinity Fabric/xGMI，用于现代AMD GPU加速器（例如AMD Instinct MI300X）。xGMI为节点内的GPU之间提供高带宽通信。该架构为一个8-GPU系统支持896 GB/s的总双向GPU间带宽。与NVSwitch不同，AMD当前的互连网格是扁平的，不依赖于交换结构。这在一定程度上限制了每个节点扩展到8个以上GPU。最近，成立了超加速器链接（UALink）联盟，旨在开发一个更开放的共享内存加速器互连，与多种技术和供应商兼容【81, Key Hyperscalers And Chip Makers Gang Up On Nvidia’s NVSwitch Interconnect, 2024, The Next Platform】。

• Intel Xe-Link：Intel的Xe-Link fabric是用于Aurora超级计算机等系统中的高带宽、全连接的节点内互连，将Intel数据中心GPU Max（Ponte Vecchio）设备连接成一个统一的拓扑。在典型的Aurora节点中，六个GPU以全连接配置连接，尽管Xe-Link也可以支持多达8路排列，其中每个GPU都与其他所有GPU有直接链接。这些链接支持加载/存储访问、复制引擎传输和远程原子操作，使GPU能够直接访问彼此的内存而无需主机参与。

3.5 供应商机制讨论

• 3.5.1 GPUDirect P2P和直接加载/存储对编程的影响
  范式转变：GPUDirect P2P的引入标志着多GPU执行范式的重大转变，使得在内核内使用加载和存储操作实现GPU之间的直接通信成为可能。基于直接加载/存储的通信提供了几个好处。首先，它允许程序员将通信与计算内联。程序员不再需要依赖于独立的通信和计算模型，而是可以在GPU内核内将它们结合起来【71, NVSHMEM: GPU-Integrated Communication for NVIDIA GPU Clusters, 2021, GTC Presentation】【117, Efficient Breadth First Search on Multi-GPU Systems Using GPU-Centric OpenSHMEM, 2018, OpenSHMEM Workshop】【120, Exploring OpenSHMEM Model to Program GPU-Based Extreme-Scale Systems, 2015, OpenSHMEM Workshop】。其次，直接加载/存储利用了GPU提供的高并行度，可以实现比DMA复制更高的带宽和更低的延迟【18, Groute: Asynchronous Multi-GPU Programming Model with Applications to Large-Scale Graph Processing, 2020, ACM Trans. Parallel Comput.】【114, Interconnect Bandwidth Heterogeneity on AMD MI250x and Infinity Fabric, 2023, arXiv】。第三，直接加载/存储可以通过GPU固有的延迟隐藏能力隐式地重叠通信与计算。鉴于GPU提供的高并行度以及现代互连技术提供的日益增长的带宽，GPU不仅有能力隐藏本地内存的延迟，也有能力隐藏远程内存的延迟【116, GPU-Centric Communication on NVIDIA GPU Clusters with InfiniBand: A Case Study with OpenSHMEM, 2017, HiPC】【117, Efficient Breadth First Search on Multi-GPU Systems Using GPU-Centric OpenSHMEM, 2018, OpenSHMEM Workshop】【120, Exploring OpenSHMEM Model to Program GPU-Based Extreme-Scale Systems, 2015, OpenSHMEM Workshop】。这对程序员来说是另一个福音，因为实现重叠的方法从程序员通过流和事件手动实现的基于软件的方法转变为自动的基于硬件的重叠。由于通信/计算重叠的责任从程序员转移到硬件，另一个启示是，随着硬件在隐藏内存延迟方面变得更好，重叠将得到改善。第四，直接加载/存储扩展了可以通过多个GPU加速的应用范围。传统上，具有细粒度通信模式的应用在多GPU系统上可扩展性较差，因为计算经常需要中断和同步，以便CPU发起通信。通过内核内的直接加载/存储，GPU可以很好地适应细粒度的通信模式。最后，直接加载/存储允许在不离开GPU的情况下在内核内触发通信。当与持久内核【61, Multi-GPU Communication Schemes for Iterative Solvers: When CPUs Are Not in Charge, 2023, ICS ’23】【165, PERKS: A Locality-Optimized Execution Model for Iterative Memory-Bound GPU Applications, 2023, ICS ’23】结合时，这个方向尤其有前景，其中单个内核被启动并在设备上通过内部循环维持其跨多个工作迭代的执行，从而最小化内核启动开销。事实上，Spector等人【141, We Bought the Whole GPU, So We’re Damn Well Going to Use the Whole GPU, 2025, Hazy Research Blog】在Llama-70B的张量并行推理中使用了直接加载/存储，并实现了一个巨型内核来执行异步通信，将其与计算和本地内存操作重叠。

  挑战：尽管直接加载/存储带来了改进，但仍存在一些固有的挑战。首先，一个根本的挑战是通信和计算争夺相同的有限资源，因为它们现在都需要大量的GPU线程来取得进展。当通信作为单独的内核实现时，这尤其成问题。如果计算内核先启动，它可能会垄断所有GPU资源，阻止通信内核启动，从而有效地消除了任何重叠的可能性。可以通过以更高优先级的流启动通信来缓解此问题，使其始终首先被调度。我们注意到，P2P DMA复制没有这个问题，因为它们使用GPU的DMA/复制引擎——一个物理上独立的资源——进行通信【19, Benchmarking multi-GPU applications on modern multi-GPU integrated systems, 2021, Concurrency and Computation: Practice and Experience】。其次，与单GPU内存访问类似，P2P直接加载/存储对内存合并高度敏感，随机非合并访问的性能远差于合并访问【18, Groute: Asynchronous Multi-GPU Programming Model with Applications to Large-Scale Graph Processing, 2020, ACM Trans. Parallel Comput.】。这种非合并的直接读取可能会暴露远程内存延迟，这超出了GPU调度器的隐藏能力，最终导致执行停滞。类似地，亚缓存行粒度的零星非合并直接写入可能会极大地未充分利用互连【83, Efficient Multi-GPU Shared Memory via Automatic Optimization of Fine-Grained Transfers, 2021, ISCA】。

• 3.5.2 GPUDirect RDMA的局限性
  内存一致性问题：GPUDirect RDMA的一个显著限制是，在内核运行时，GPU和NIC内存之间没有一致性保证。一致性只能通过将控制权返回给CPU，即拆除内核并启动一个新内核来保证，从而将通信限制在内核边界。这也意味着将持久内核与GPU发起的节点间通信结合将不可避免地导致数据正确性问题【96, GPUDirect RDMA, 2023, NVIDIA】。Chu等人【31, Designing High-Performance In-Memory Key-Value Operations with Persistent GPU Kernels and OpenSHMEM, 2019, OpenSHMEM Workshop】通过从NIC向GPU内存发出PCIe读取来绕过此限制，这会刷新之前NIC对GPU的写入并保证内存顺序。自11.3版本以来，CUDA也提供了cudaDeviceFlushGPUDirectRDMAWrites() API，可以类似地用于强制一致性【38, The Latest in GPUDirect, 2021, GTC Presentation】【94, CUDA Runtime - Device Management, 2023, NVIDIA】。虽然有用，但CUDA仍然依赖CPU来强制GPU-NIC的一致性。另一方面，AMD在设备侧通信的持久内核上下文中明确纠正了GPU-NIC的一致性问题，并将提议的修复集成到ROC_SHMEM中【51, GPU INitiated OPenSHMEM: Correct and Efficient Intra-Kernel Networking for DGPUs, 2020, PPoPP ’20】。我们将在第5.3节中进一步讨论这个问题。

• 3.5.3 在以GPU为中心的通信中启用触发能力
  CPU的角色与演进：在第2节介绍的类型3（GPUDirect/ROCn RDMA）中，CPU仍然负责系统的初始配置、数据传输准备和发起传输。第一阶段包括设置网络接口和加载GPU驱动程序。CPU向支持RDMA的NIC注册GPU内存。主机的这种注册允许NIC直接访问GPU内存，绕过了数据传输期间中间的CPU步骤。在第二阶段，CPU分配GPU内存缓冲区并确保它们正确对齐。这些缓冲区将用于与GPU之间的高效数据传输。然后，CPU设置GPU流和事件来管理和排序数据传输，并保证工作的编译。流用于对操作进行排队，确保它们按正确的顺序执行。然而，对于真正的低延迟应用，这种成本仍然可能成为瓶颈，因为该机制依赖于通过流的几个同步点【84, Exploring Fully Offloaded GPU Stream-Aware Message Passing, 2023, arXiv】【85, Exploring GPU Stream-Aware Message Passing using Triggered Operations, 2022, arXiv】【123, Optimizing Distributed ML Communication with Fused Computation-Collective Operations, 2023, arXiv】。

  GPU触发通信：第2节中定义的类型4和5中的GPU触发通信通过消除上述同步成本，促进了计算和通信控制路径向GPU的卸载。在这里，触发操作起着至关重要的作用，因为它们是仅在满足特定条件时才安排执行的特殊任务。在流触发（ST）策略中，这些操作通过GPU控制处理器管理数据传输和同步，从而减少CPU的参与。延迟执行是另一个重要方面，其中CPU创建具有延迟执行语义的命令描述符，并将它们附加到NIC命令队列。当GPU控制操作指定的条件满足时，这些描述符被执行。例如，HPE Slingshot 11 NIC支持这些延迟操作【84, Exploring Fully Offloaded GPU Stream-Aware Message Passing, 2023, arXiv】¹，包括当硬件计数器达到给定阈值时触发的发送和接收通信。这通过启用特定的命令队列（例如，Libfabric延迟工作队列）成为可能。GPU控制处理器和NIC之间的同步通过特殊机制确保通信操作的成功完成（例如，在NVIDIA DOCA中，这是通过GPUNetIO信号量实现的【1, Inline GPU Packet Processing with NVIDIA DOCA GPUNetIO, 2023, NVIDIA Developer Blog】）。

为了将这些供应商机制置于更广泛的通信生态系统中，图4总结了从应用程序和通信库到中间件、驱动程序和网络结构的软件堆栈。这个分层视图有助于将本节讨论的低级机制与接下来要审查的用户级通信库联系起来。

4 以GPU为中心的通信库

我们现在讨论近年来为简化多GPU编程而涌现的主要以GPU为中心的通信库，即GPU-aware MPI、GPU-centric Collectives和GPU-centric OpenSHMEM。

4.1 GPU-Aware MPI

• GPU-Aware MPI的演进：鉴于MPI是HPC事实上的通用语言，人们投入了大量精力使其与GPU编程模型互操作，最终形成了能够区分主机和设备缓冲区的GPU-Aware MPI实现。在GPU-Aware MPI之前，所有多GPU通信都必须通过主机进行暂存，在源GPU上产生设备→主机的复制，在目标GPU上产生主机→设备的复制。而使用GPU-Aware MPI实现，程序员可以将设备缓冲区作为参数提供给MPI调用，允许通信使用由GPUDirect RDMA或ROCnRDMA建立的直接GPU到GPU数据路径。在此过程中，GPU感知消除了冗余的主机↔设备复制，并通过省去主机缓冲区的需要简化了通信代码。MVAPICH2是第一个积极将GPU感知集成到其运行时的MPI实现。早期的MVAPICH2工作引入了基本的GPU感知，透明地通过主机暂存GPU驻留的缓冲区，并通过流水线方案优化主机↔设备和设备↔设备的传输。这些流水线方案之所以成为可能，是因为UVA允许库区分主机和设备指针，而无需依赖用户提示。随之而来的GPU感知带来了相对于GPU无感知版本的性能提升【156, GPU-Aware MPI on RDMA-Enabled Clusters: Design, Implementation and Evaluation, 2014, IEEE Transactions on Parallel and Distributed Systems】【157, MVAPICH2GPU: optimized GPU to GPU communication for InfiniBand clusters, 2011, Computer Science - Research and Development】【158, Optimized Noncontiguous MPI Datatype Communication for GPU Clusters: Design, Implementation and Evaluation with MVAPICH2, 2011, CLUSTER】。一项后续工作使用CUDA IPC来优化节点内传输，此前这些传输必须通过主机内存中的缓冲区进行暂存【121, Optimizing MPI Communication on Multi-GPU Systems Using CUDA Inter-Process Communication, 2012, IPDPSW】。最终，增加了对GPUDirect RDMA的rendezvous协议的支持，允许传输绕过主机并消除冗余的主机↔设备复制。这降低了延迟；然而，由于现有的架构限制，带宽受到限制【119, Efficient Inter-node MPI Communication Using GPUDirect RDMA for InfiniBand Clusters with NVIDIA GPUs, 2013, ICPP】。随后的工作增加了对GPUDirect RDMA的eager协议的支持，纠正了带宽限制并进一步降低了延迟。此外，一种新的环回机制和早期版本的GDRCopy【97, Magnum IO GDRCopy, 2023, NVIDIA】被用来消除昂贵的主机↔设备cudaMemcpys【135, Designing efficient small message transfer mechanism for inter-node MPI communication on InfiniBand GPU clusters, 2014, HiPC】。GDRCopy允许将GPU内存映射到用户地址空间，这对于小消息尺寸以最小的开销进行了优化【38, The Latest in GPUDirect, 2021, GTC Presentation】【97, Magnum IO GDRCopy, 2023, NVIDIA】【129, Designing a ROCm-Aware MPI Library for AMD GPUs: Early Experiences, 2021, ISC High Performance】。另一项工作扩展了点对点MPI调用以支持GPUDirect Async，允许GPU推进由CPU排队的通信，从而优化了控制路径【152, MPI-GDS: High Performance MPI Designs with GPUDirect-aSync for CPU-GPU Control Flow Decoupling, 2017, ICPP】。其他工作也越来越关注向MVAPICH2-GDR添加UVM感知【16, Designing High Performance Communication Runtime for GPU Managed Memory: Early Experiences, 2016, GPGPU ’16】【50, CUDA M3: Designing Efficient CUDA Managed Memory-Aware MPI by Exploiting GDR and IPC, 2016, HiPC】【77, Characterizing CUDA Unified Memory (UM)-Aware MPI Designs on Modern GPU Architectures, 2019, GPGPU ’19】。

• 其他MPI实现的GPU感知：其他主要的MPI实现也集成了GPU感知。Open MPI【111, Open MPI v5.0.x Documentation: CUDA, 2023, OpenMPI】【161, GPU-Aware Non-Contiguous Data Movement In Open MPI, 2016, HPDC ’16】在1.7.0版本中引入了CUDA感知支持。当使用其内部后端时，基于计算的集合操作（例如，MPI_Allreduce）可能仍会通过主机暂存数据【40, Exploring GPU-to-GPU Communication: Insights into Supercomputer Interconnects, 2024, SC’24】，这一限制可以通过UCX避免，UCX能够实现GPU驻留数据的移动和归约。Open MPI还利用GDRCopy【97, Magnum IO GDRCopy, 2023, NVIDIA】处理小消息，并利用CUDA IPC进行节点内通信。通过ROCm，GPU感知扩展到AMD设备，这是通过UCX实现的【9, GPU-aware MPI with ROCm, 2023, AMD】【11, ROCm Documentation: GPU-Enabled MPI, 2023, AMD】【112, Open MPI v5.0.x Documentation: ROCm, 2023, OpenMPI】。然而，Khorassani等人【129, Designing a ROCm-Aware MPI Library for AMD GPUs: Early Experiences, 2021, ISC High Performance】为MVAPICH2提供了一个原生的ROCm感知运行时，其在AMD GPU集群上的性能优于使用UCX的Open MPI。Open MPI集成了统一集合通信（UCC）框架【151, Unified Collective Communication: A Unified Library for CPU, GPU, and DPU Collectives, 2025, IEEE Micro】，这是UCX生态系统的一个组件，旨在为高性能集合操作提供统一接口。UCC建立在UCX的传输层之上，并利用其拓扑感知机制（如对NVLink、PCIe和共享内存层次结构的感知）来为GPU缓冲区选择高效的集合算法。这种集成使Open MPI能够使用最合适的GPU互连路径自动卸载或加速集合操作。

• MPICH及其他实现：其他主流实现也遵循了类似的发展。例如，MPICH在3.4版本中通过其CH4通信层引入了GPU支持，该层管理设备缓冲区并选择高效的通信路径。后来通过ROCm（自4.0版本起）将支持扩展到AMD GPU，Intel GPU支持正在开发中【48, MPICH for Exascale, 2022, ECPAM BoF】。在Cray系统上，用户必须通过设置MPICH_GPU_SUPPORT_ENABLED=1来显式启用GPU支持【55, Cray MPICH Documentation, 2021, HPE】。MPICH还集成了GDRCopy、GPUDirect RDMA和IPC以优化传输。MPICH已在其CH4设备层中采用UCX作为网络模块，并且正在进行的工作探索集成UCC以启用优化的GPU集合操作【33, GitHub - Pull Request for adding support for Unified Collective Communication (UCC), 2025, MPICH GitHub】。

4.2 以GPU为中心的集合通信 (GPUCCL)

• GPUCCL的动机与实现：随着深度学习模型越来越大，其计算需求需要在多个GPU上部署训练。鉴于集合通信在深度学习训练中的普遍性，NVIDIA、AMD和Intel都提供了针对其各自GPU架构优化的高效集合通信库。为简单明了起见，我们将这些特定于供应商的解决方案称为GPU集合通信库（GPUCCL）。它们已被集成为多个最先进的深度学习框架的通信后端，包括Pytorch、Tensorflow、MXNet、Caffe、CNTK和Horovod【160, xCCL: A Survey of Industry-Led Collective Communication Libraries for Deep Learning, 2023, Journal of Computer Science and Technology】。

• 与MPI的区别：在MPI中实现涉及计算的GPU感知集合操作（如reduce/allreduce）时，通常使用GPU内核进行局部归约，并由CPU发起的GPU间复制来执行聚合。这种方法会产生多个内核启动和通信调用的延迟，并且还需要主机上的中间缓冲区。GPUCCL采取了不同的方法，将集合的通信和计算一起在单个内核中实现【89, Fast Multi-GPU collectives with NCCL, 2016, NVIDIA Developer Blog】。接下来，我们重点介绍各供应商解决方案在支持功能方面的差异。

• 4.2.1 各供应商集合通信库的比较：表4列出了三个主要供应商支持的集合通信库的主要特性：NVIDIA的NCCL、AMD的RCCL和Intel的OneCCL。

  ![Table 4. 供应商集合通信库在集合原语、执行模型、加速器集成和传输语义方面的比较。缩写：B=广播, R=归约, AR=全局归约, RS=归约散射, AG/AGv=全局收集及其变体, AA=全局到全局, P2P=点对点。](images/1b2464f853548c3e99574d7fb4b15da184fbb416c