读论文：The Landscape of GPU-Centric Communication

1. 术语与通信类型

节点内（intro-node）通信：

  • API：程序员或库在何处进行通信 API 调用。
  • Data Path：指示谁参与了数据移动，并显示相应的数据路径。

节点间（inter-node）通信：

  • Register /construct messages：这一步包括数据包的构造和在网卡（NIC）的注册。
  • Trigger communication：它定义了谁触发 NIC 发出数据传输。

2. 供应商机制

2.1 Memory Management Mechanisms

  1. Page-Locked / Pinned Memory.

    使用 cudaMalloc() 在 host 上分配的内存默认是分页的（pageable）且无法被 GPU 直接访问，当需要从分页的内存上传输数据到 GPU 时，需要经过两次拷贝：

    • host 分页内存 -> host 锁页（page-locked）内存
    • host 锁页内存 -> device 显存

    为了避免第一步的拷贝，cudaMallocHost() 运行直接分配 page-locked 内存，因此，page-locked 内存也叫：zero-copy or pinned memory。

    优缺点：

    • host-device 间数据传输时高带宽/低时延
    • 内存消耗大，可能因为大量内存占用影响系统性能

  2. Unified Virtual Addressing (UVA).

    CUDA 4.0 引入，允许节点内的所有 GPU 和 CPU 共享相同的统一虚拟地址空间。

  3. CUDA IPC.

    问题：在早期的 CUDA 版本中，指针不能跨进程访问，因此 GPU buff 之间的内存拷贝必须通过主机，从而产生瓶颈。

    Solution： CUDA 4.1 引入 CUDA IPC，允许同一节点上的进程访问其他进程的 GPU 显存，而无需额外的拷贝。

    优缺点：

    • 使用 CUDA IPC，内存句柄是使用标准的 IPC 机制在进程之间创建和传递的，这比通过主机进行分段复制的延迟要低。

    • 然而，创建内存句柄的开销可能是显著的，并且可能抵消延迟的好处。

  4. Unified Virtual Memory (UVM).

    CUDA 6.0 引入，通过 cudaMallocManaged() 创建 managed memory，创建一个可被单个节点内所有 processors 访问的单个地址空间。

    UVM 底层实现机制：

    UVM 的工作原理是将请求的内存划分为驻留在 CPU 上的页面。程序员可以在没有显式拷贝的情况下访问设备上的内存。如果内存访问是不在设备上的页的一部分，则 UVM 驱动触发一个页面错误，自动将页面迁移到请求设备。当总分页内存大小超过设备内存时，UVM 驱动程序还可以将给定设备中的页面驱逐回主机内存。

    优缺点：

    • 程序员暴露在一个统一的地址空间中，他们可以访问，就好像整个分配的内存块驻留在单个 GPU 上一样。系统周围发生的任何复制对程序员都是隐式的和隐藏的。
    • UVM 允许内存超额申请，因此可以分配比所有多 GPU 设备内存之和更多的内存。

2.2 GPUDirect Technologies

  1. GPUDirect 1.0 (NIC).

    CUDA 3.1 引入，允许 GPU 与网卡（NICs）共享同一块 pinned memory region，如上述图 2.① 所示，需要 GPU memory -> CPU pinned memory；CPU pinned memory -> NIC’s memory 两次拷贝。GPUDirect 1.0 引入了 GPU-网卡之间的共享内存区域，因此，避免中间 CPU 发起的拷贝。

  2. GPUDirect 2.0 (Peer-to-Peer).

    CUDA 4.0 引入 UVA，同时增加了同一节点内共享相同 PCIe 根的 GPU 之间直接点对点通信的支持，避免了 host 的参与，首次允许 GPU 间直接数据传输路径。由此引入两种新的通信机制：

    • P2P DMA Copies：cudaMemcpy 调用直接触发 GPU 之间的 DMA 数据传输。
    • P2P Direct Load / Stores：GPU 可以通过解引用指向远程 GPU buff 的指针来直接访问数据。

    优缺点：

    • 消除冗余的 GPU↔︎CPU 拷贝和 host buffer。
    • 通过避免在主机上维护通信 buff 并提供新的通信机制（P2P Direct Load / Stores），GPUDirect P2P 增加了多 GPU 编程的便利性。

  3. GPUDirect RDMA.

    CUDA 5.0 引入 GPUDirect RDMA，跨节点 GPU 间直接通信变得可行。

    优缺点：

    • 消除主机内存的额外拷贝，减少由 GPU-NIC 交互产生的固有延迟，增加带宽并减少 CPU 开销。

  4. GPUDirect Async.

    前文的的 GPUDirect 技术专注于改善数据路径，GPUDirect Async 优化 GPU 与网卡之间的控制路径。CUDA 8.0 引入，GPUDirect Async 的工作原理是让 CPU 预注册消息，然后 GPU 内核触发网卡上的这些消息。因此，在触发通信时，GPU 可以继续执行，而不是像以前那样需要停止让 CPU 发起通信。

    优缺点：

    • GPUDirect Async 使控制路径远离 CPU 的努力得到了改进。
    • 没有完全将控制路径转移到 GPU，仍需要 CPU 预先注册消息。

2.3 GPUNetIO

GPUNetIO 是 NVIDIA 提出的技术解决方案，作为 DOCA（Datacenter-On-a-Chip Architecture）的一部分。DOCA 是一个全栈软件框架，旨在促进 NVIDIA BlueField Data Processing Units (DPUs) 的应用程序开发。

从 DOCA v2.7 起，GPUNetIO 允许 GPU 不仅在内核边界上执行 RDMA 发送和接收，而且在内核执行期间的任何时候都可以执行。

On non-RDMA networks GPUNetIO allows the GPU to send, receive, and process network packets. On RDMA networks (both RoCE and InfiniBand), from DOCA v2.7, GPUNetIO allows the GPU to execute RDMA send and receive not only on kernel boundaries , but at any point during the kernel execution. In a nutshell, GPUNetIO allows the GPU to interact with the NIC without any CPU intervention on the critical path.

2.4 Modern GPU-centric Interconnects

NVLink，提出用于解决 PCIe 带宽有限的问题。各代 NVLink 的规格：

优缺点：

  • 高带宽。
  • NVLink 不是自路由的（-> NVSwitch），这意味着如果任何两个给定 GPU 没有直接的 NVLink 连接，通信将不得不通过中间 GPU 路由。

NVSwitch：一种背板（backboard）技术，可以实现所有 GPU 之间的 all-to-all 连接。

3. 以 GPU 为中心的通信库

3.1 GPU-Aware MPI

代表性库：

  • MVAPICH2
  • OpenMPI
  • HPE Cray MPICH
  • IBM Spectrum MPI

  1. Streaming Support in GPU-aware MPI.

    MPI 与 GPU Stream 语义不匹配。MPI 调用带来了强制同步，不利于 GPU kernel 流水线启动和 overlap。对这一问题的两种解决方案：

    • The first is making MPI runtimes stream-aware by adding an explicit stream parameter to MPI routines.
    • The second is providing the option of device-initiated MPI calls.

3.2 GPU-Centric Collectives

NCCL (NVIDIA Collective Communication Library) 是一个软件库，它为 GPU 间通信提供了拓扑感知的集体原语。此外，从 2.7 版本开始，NCCL 还提供了点对点通信 API。它已被集成为几个最先进的深度学习框架的通信后端，包括 Pytorch、Tensorflow、MXNet、Caffe、CNTK 和 Horovod。AMD 提供了一个类似的库，称为 RCCL。

3.3 GPU-centric OpenSHMEM

以 GPU 为中心的 OpenSHMEM 运行时是 NVIDIA 的 NVSHMEM 和 AMD 的 ROC_SHMEM 库。

  1. NVSHMEM.

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

    底层实现机制：

    NVSHMEM works on the concept of a symmetric heap . During NVSHMEM initialization, each process that is mapped to a GPU, referred to as a processing element (PE) reserves a block of GPU memory using nvshmem_malloc() . In NVSHMEM, all memory allocations must be performed collectively, meaning that all symmetric memory regions within the heap must have identical sizes and must be allocated at the same time. To access remote memory on a different PE, a given PE requires the offset for the symmetric memory as well as the rank of the remote PE.

  2. ROC_SHMEM.

    ROC_SHMEM 是 AMD 对 AMD GPU 的 OpenSHMEM 规范的实现。

4. 讨论、挑战与展望

  1. Moving Away from CPU
  2. UCX as a Potential Pathway for GPU-Awareness
  3. CPU-Free Networking
  4. Broader GPU Autonomy
  5. Collective Algorithms Design
  6. Debugging and Profiling Support
ML System
#GPU 通信
读论文：The Landscape of GPU-Centric Communication
https://arcsin2.cloud/2025/06/02/读论文：The-Landscape-of-GPU-Centric-Communication/
作者
arcsin2
发布于
2025年6月2日
许可协议