以GPU为中心通信的全景 The Landscape of GPU-Centric Communication-3

戏说数据那点事 2025-01-16

445

4 GPU-CENTRIC COMMUNICATION LIBRARIES

在这一部分，我们讨论了近年来为简化多GPU编程而出现的主要GPU中心通信库，即GPU感知MPI、GPU为中心集合操作和GPU为中心OpenSHMEM。

4.1 GPU-Aware MPI

鉴于MPI是高性能计算（HPC）的事实标准语言，许多努力已投入到使MPI通信与现有GPU编程模型兼容。这个工作 culminated in GPU-Aware MPI，通常定义为能够区分主机和设备缓冲区的MPI实现。在GPU-Aware MPI出现之前，所有多GPU通信必须通过主机进行中转，这样在源GPU上会有设备到主机的复制，在目标GPU上则会有主机到设备的复制。另一方面，使用GPU-Aware MPI实现时，程序员可以将设备缓冲区作为参数传递给MPI调用，从而允许通信使用由GPUDirect RDMA或ROCnRDMA建立的直接GPU到GPU数据通道。在此过程中，GPU感知消除了冗余的主机↔设备复制，并通过省略主机缓冲区的需求简化了通信代码。

MVAPICH2是第一个开始积极将GPU感知集成到其运行时的MPI实现。在引入GPUDirect RDMA之前，早期的工作增加了基本的GPU感知，允许用户使用位于GPU上的缓冲区发起MPI调用。尽管这样做的通信仍然需要通过主机中转，但这一过程是透明的，由库实现，并通过管道化方案优化主机↔设备和设备↔设备的传输。这些管道化方案得益于UVA，它使得库能够区分主机和设备指针，而无需依赖用户的提示。随之而来的GPU感知使得性能较GPU盲视版本有所提升。后续工作使用CUDA IPC优化了节点内传输，以前这些传输必须通过主机内存中的缓冲区进行。最终，添加了对GPUDirect RDMA的支持，通过会话协议允许传输绕过主机并消除冗余的主机↔设备复制。这减少了延迟；然而，由于现有架构的限制，带宽仍然受到限制。后续工作进一步添加了对GPUDirect RDMA的支持，采用急切协议解决了带宽限制并进一步降低了延迟。此外，还使用了一种新的环回机制和GDRCopy的早期版本，从而消除了昂贵的主机↔设备cudaMemcpys。GDRCopy允许GPU内存映射到用户地址空间，优化了小消息大小的传输，并且开销最小。另外，一项工作扩展了点对点MPI调用，支持GPUDirect Async，允许GPU处理CPU排队的通信，从而优化了控制路径。其他工作也日益关注为MVAPICH2-GDR添加UVM感知。

尽管MVAPICH2无疑为GPU感知MPI铺平了道路，但其他领先的MPI实现也已集成了GPU感知支持。OpenMPI、HPE Cray MPICH、MVAPICH2和IBM Spectrum MPI都原生支持CUDA感知。OpenMPI还通过UCX支持CUDA。鉴于NVIDIA GPU的广泛部署，关于ROCm感知MPI的文献较少。现有的MPI实现中，除了MVAPICH2外，其他实现都依赖UCX来实现ROCm感知。另一方面，Khorassani等人提供了MVAPICH2的原生ROCm感知运行时，该实现优于在AMD GPU集群上使用UCX的OpenMPI。

4.1.1 GPU感知MPI中的流支持

尽管GPU感知MPI有很多好处，但MPI与GPU的集成存在固有的挑战。最重要的是，MPI和GPU编程模型之间存在根本的语义不匹配。GPU的操作基于流的概念，流是命令队列，保证GPU操作的顺序。GPU调度器确保在流上启动的内核和其他操作按队列中的顺序执行，并且按照正确的数据依赖关系执行。由于内核启动是异步的，并且不会阻塞主机，GPU运行时可以对内核启动进行管道化，并将启动延迟与内核执行重叠。MPI和GPU模型之间的语义不匹配在于，MPI并不了解GPU流。因此，不可能将MPI调用排入某个特定GPU流，或者让GPU流等待一个待完成的MPI例程。由此产生的影响是，将MPI调用与GPU内核交替使用时，必须进行主机阻塞同步，以保持数据正确性。例如，在发起MPI发送之前，程序员必须阻塞主机以同步所有在发送缓冲区上操作的流。类似地，等待待完成的MPI通信也需要主机阻塞同步。在这两种情况下，这些强制同步会妨碍内核启动的管道化，阻止重叠的机会，并迫使程序员进入交替进行通信和计算的批量阶段。

MPI与GPU之间的语义不匹配最终源于试图将一个以CPU为中心的通信运行时与基于GPU的计算相结合。后者的设计依赖于流（streams）来实现并发、同步和重叠，而前者则没有流的概念。这个根本问题不太可能通过零散的解决方案来解决。我们看到了解决语义不匹配的两条可能的非互斥路径。第一条是使MPI运行时流感知，通过向MPI例程添加显式的流参数。这将解决内核启动流水线受阻的问题，并允许MPI调用无缝集成到GPU运行时中。第二条是提供设备发起的MPI调用选项。这将减少程序员在处理两种不同编程模型时的负担，并且提供隐式的通信-计算重叠。文献中在有限范围内已经探索了这两条路径。FLAT编译器自动将设备端的MPI调用转换为其主机端的等效调用。dCUDA实现了具有MPI语义的设备端操作，但使用CPU辅助线程来执行实际的通信。它们依赖GPU固有的内存延迟隐藏能力，将通信与计算隐式重叠，从而最终超过GPU感知MPI基线的性能。Namashivayam等人探索了引入GPU流感知的MPI新通信方案。他们利用HPE Slingshot 11互连的触发操作功能，允许CPU将通信和同步操作排队到NIC，然后由GPU触发。这减少了CPU在关键路径中的参与，消除了昂贵的同步操作。虽然跨节点实验显示了一些性能优势，但该方案在节点内设置中出现问题，因为进度线程需要用来模拟延迟执行语义。后续工作通过使用基于P2P直接加载存储的GPU内核和GPU IPC机制，消除了节点内通信的进度线程，从而提高了性能，相比流不可知的MPI基线

4.2 GPU中心的集合通信

随着深度学习模型越来越大，它们的计算需求需要跨多个GPU进行训练。许多工作通过GPU中心机制来优化多GPU训练。鉴于深度学习训练中集合通信的普遍性，NVIDIA和AMD都提供了高效的集合通信库。

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

涉及计算的GPU感知集体操作（如reduce/allreduce）在MPI中通过GPU内核实现本地归约，并通过CPU发起的拷贝在GPU间进行聚合。此方法会产生多个内核启动和通信调用延迟，并且还需要主机上的中间缓冲区。NCCL采用了不同的方法，通过在单个内核中实现集体通信和计算[68]。NCCL的第一个版本（NCCL 1.0）仅支持节点内通信，使用P2P直接加载/存储实现。NCCL 2.0增加了对节点间通信的支持，依赖GPUDirect RDMA和运行在CPU上的代理线程。

NCCL有几个优点。首先，已证明NCCL能够实现高带宽和并行效率。第二，NCCL是拓扑感知的，因此减轻了程序员优化拓扑的负担。程序员可以选择强制指定特定的通信路径。第三，NCCL与CUDA编程模型原生兼容，可以无缝集成。NCCL的API调用带有流参数，允许通过在单独的流上排队集体通信来实现通信-计算重叠。

尽管NCCL有其优势，但也有一些固有的挑战。首先，使用GPU线程进行通信和计算的一个根本问题是，两个例程现在会争夺相同的有限资源。在这种情况下，如果计算调度在NCCL集体操作之前，它可能会独占所有GPU资源，从而有效地将集体操作序列化在计算之后。一种解决方法是在更高优先级的流上启动NCCL集体操作，以确保它始终首先被调度。

4.2.1 集合通信的进一步研究

尽管P2P通信和集体通信在高性能计算中都扮演着重要角色，但深度学习的需求使得优化集体通信负载成为一个重要且有吸引力的研究领域。由于NCCL的早期版本不支持节点间通信，因此提出了支持节点扩展的分层集体机制。后续的工作提出了新的流水线设计，用于MPI_Bcast集体操作，挑战了NCCL在深度学习集体通信中的主导地位。所提出的设计在小消息大小时超越了NCCL2的广播，并且在较大消息时提供了可比的性能。Blink是一个集体通信库，旨在实现最佳的链路利用率。为此，Blink检测底层拓扑，将拓扑建模为图形，然后使用一种叫做“打包生成树”的技术动态生成通信原语。因此，Blink被证明在图像分类任务中，比NCCL2减少了40%的模型训练时间。

Dryden等人提出了Aluminum，这是一个GPU感知的库，用于大规模训练深度神经网络。Aluminum旨在解决NCCL和MPI的固有缺点。Aluminum通过扩展NCCL的树形算法来避免其默认环形实现的延迟瓶颈。此外，它增加了对非阻塞NCCL allreduce操作的支持。对于MPI，Aluminum通过将单个GPU流与MPI通信器关联，并只对该流进行同步，从而解决了MPI-GPU语义不匹配引起的强制同步问题。这些优化带来了比GPU感知MPI和基于NCCL的实现更快的性能。

Awan等人使用新发布的NCCL实现了高效的节点内MPI_Bcast，作为MVAPICH2-GDR的一部分[11]。Chen等人最近的工作使用CUDA IPC优化MPI节点内的全对全通信。他们使用CUDA IPC跨进程边界传递设备缓冲区，然后使用P2P直接加载/存储进行小消息通信，使用P2P DMA拷贝进行大消息通信。

像UCX一样，UCC是一个开源项目，提供了一个高性能和可扩展的集体操作API和库实现，利用了拓扑感知算法和网络计算等技术。它依赖于UCX的点对点通信，并支持NCCL/RCCL、SHARP等.

4.3 GPU-centric OpenSHMEM

GPU-centric OpenSHMEM 运行时包括 NVIDIA 的 NVSHMEM 和 AMD 的 ROC_SHMEM 库。两者在基本原理上相似，但也存在一些差异。鉴于 NVSHMEM 诞生较早，首先讨论 NVSHMEM，同时介绍这两个库的基本概念。在 ROC_SHMEM 部分，我们将讨论它与 NVSHMEM 之间的重要区别。

4.3.1 NVSHMEM

NVSHMEM 是 NVIDIA 针对 CUDA 设备实现的 OpenSHMEM 规范。NVSHMEM 是一个分区全局地址空间（PGAS）库，提供高效的单边 put/get API，允许进程访问远程数据对象。NVSHMEM 支持 GPU 之间的点对点通信和集体通信，既支持节点内，也支持跨节点通信。NVSHMEM 基于对称堆的概念工作。在 NVSHMEM 初始化过程中，映射到 GPU 的每个进程（称为处理单元，PE）都会通过 nvshmem_malloc() 保留一块 GPU 内存。

在 NVSHMEM 中，所有内存分配必须是集体执行的，意味着对称内存区域的大小必须一致，并且必须在同一时刻进行分配。为了访问不同 PE 上的远程内存，某个 PE 需要提供该对称内存的偏移量以及远程 PE 的排名。

此外，NVSHMEM 提供了一些 API 用于同步一组 PEs。这些 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 执行一致性强制 API 调用（例如：cudaDeviceFlushGPUDirectRDMAWrites()）。该解决方案的有效性尚不清楚，仍需进一步研究。由内核内部强制执行 GPU-NIC 内存一致性是 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 并未实现此优化。尽管这一优化可能对执行时间有利，但增加的共享内存消息可能会限制占用率，从而对性能产生负面影响。

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

4.3.3 应用

近年来，NVSHMEM 已被集成为多个运行时的通信后端。PETSc 实现了基于 NVSHMEM 的可扩展通信层 PetscSF，以补充其基于 MPI 的方法，后者在 CUDA 流语义中表现不佳，限制了内核启动的流水线。Kokkos Remote Spaces 通过使用 NVSHMEM 作为通信后端，为 Kokkos 编程模型添加了分布式内存支持。Kokkos 共轭梯度求解器的 NVSHMEM 实现，性能超过了 CUDA-aware MPI 实现，并显著减少了代码量。Choi 等人使用持久内核和 NVSHMEM 实现了 CharminG，这是一个完全 GPU 居住的运行时系统，灵感来自 Charm++。利弗莫尔大规模人工神经网络（LBANN）使用 NVSHMEM 实现了空间并行卷积，性能超过了 MPI 和 Aluminium 实现。QUDA，一个用于格点量子色动力学（QCD）计算的库，使用 NVSHMEM 和持久内核提高了狄拉克算子的强扩展性。

NVSHMEM 还被应用于基于运行时的方法之外，以实现性能提升。Chu 等人将 NVSHMEM 与持久内核结合，实现了一个先进的基于 GPU 的键值存储。Xie 等人使用 NVSHMEM 实现了单节点多 GPU 稀疏三角求解器（SpTRSV），相比基于 UVM 的设计，表现出了良好的性能扩展性。Ding 等人结合持久内核和 NVSHMEM，在单节点和多节点环境下实现了稀疏三角求解器（SpTRSV）的出色性能。Atos 使用 NVSHMEM 基于通信的持久内核和离散内核，在跨节点的多 GPU 广度优先搜索（BFS）中达到了业界领先的性能。Wang 等人提出了 MGG，一个加速图神经网络（GNN）的系统设计，通过 GPU 集中式的软件通信计算管道和 NVSHMEM，实现了细粒度通信。Ismayilov 等人使用持久内核和设备端 NVSHMEM，实现了完全 GPU 端的 Jacobi 2D/3D 和 CG 求解器，性能超过了 CPU 控制的基准。为了实现显式的设备端通信计算重叠，他们保留了部分线程块用于通信，其他线程块用于计算，这一技术称为线程块专门化。Punniyamurthy 等人使用 ROC_SHMEM 和持久内核，在深度学习推荐模型中重叠嵌入操作和集体通信。

5 讨论、挑战与展望

随着多 GPU 执行成为许多现实工作负载的必然需求，我们将在本节中进行讨论和展望，提出我们认为当前和未来在 GPU 集中式通信领域的研究潜力。

5.1 摆脱 CPU 控制

我们认为摆脱 CPU 控制的执行模式具有潜力，原因有多个。

首先，这解决了由内核启动和内存拷贝开销引起的 CPU 诱导延迟障碍。在强扩展场景中，随着 GPU 数量的增加和每个 GPU 的计算量减少，这些障碍变得更加显著。在延迟bond的环境下，传统的 CPU 控制实现无法将通信与计算重叠，因为启动操作的延迟比操作本身还要长，从而有效地将通信和计算串行化。而 CPU-free 执行模式，即使在延迟主导的情况下，也能实现适当的重叠。

其次，内核(kernel)内部通信所提供的并行性非常适合持久内核的优势。在使用高效的通信和同步 API 时，这种执行模型可以实现比 CPU 控制实现更高的带宽和更低的延迟。此外，由于通信与计算的内联，CPU-free 执行模式特别适用于具有细粒度通信的应用。

第三，ROC_SHMEM 首次实现了完全将应用执行迁移到设备上，且零依赖 CPU 辅助线程。NVSHMEM 结合其 IBGDA 传输也可以将大量执行迁移到 GPU，但仍需依赖 CPU 来保证功能正确性。

然而，这种执行模式面临几个挑战。一个主要挑战是，持久内核可能导致占用率降低，从而可能造成计算瓶颈。目前，如果需要全局设备或多 GPU 屏障，持久内核必须以协作方式启动。这意味着只能启动足够多的线程，且这些线程必须能够并发运行，无法进行硬件过度订阅。因此，之前由硬件调度器处理的工作负载分解和调度，现在需要由程序员手动完成。这种手动方法可能不如基于硬件的调度高效，计算密集型应用可能会受到影响。此外，长时间运行的持久内核会消耗更多的寄存器，且可能使用共享内存，这会进一步限制占用率。尽管如此，我们认为有一种双重解决方案来应对这个问题。

首先，整个应用执行过程中有大量高带宽共享内存，这可能抵消由降低占用率带来的性能损失。

其次，随着 GPU 厂商越来越追求更高的 GPU 自主性，预计他们将引入支持将持久内核与硬件过度订阅结合使用的 API。手动分解也可以通过优化的编译器/运行时系统来处理。

另一个潜在问题是 NVSHMEM 和 ROC_SHMEM 的集成到现有运行时系统中的难易程度。两个库都围绕一个对称堆展开，所有通信缓冲区必须在同一对称堆上由所有 GPU 集体分配。这种对称分配要求库特定的分配器。现有运行时系统可能难以支持 NVSHMEM 和 ROC_SHMEM，因为它们要求对称的内存分配。

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

统一通信 X（UCX）是一个开源通信框架，它在多个网络 API、编程模型、协议和实现之间进行抽象。其目的是提供一组高层次的原语，同时将低级实现细节隐藏在 UCX 运行时中。UCX 对 GPU 集中式通信的相关性在于，其标记和流 API 可用于为 ROCm 和 CUDA 实现 GPU 集中式通信层。实际的 GPU 集中式通信将使用相应的本地库。

使用 UCX 实现 GPU 集中式通信是近年来文献中开始受到关注的一个方向。最相关的例子可能是针对 MPI 实现的 ROCm 感知。早期，GPU 感知 MPI 的大部分工作都是基于 NVIDIA GPU 和原生 CUDA 库进行的，然后直接集成到 MPI 运行时中。或许是因为不愿重复做相同的工作以使得他们的运行时支持 ROCm，大多数 MPI 实现仅通过 UCX 提供 ROCm 感知。OpenMPI 除了原生集成外，还通过 UCX 提供了 CUDA 支持。在非 MPI 的工作中，Choi 等人扩展了 Charm++ 中的 UCX 层，为 Charm++ 生态系统中的多个编程模型提供 GPU 感知的通信。

我们预测，越来越多的运行时系统将倾向于使用 UCX 来增加对 GPU 集中式通信的支持。使用 UCX API 可以让程序员不必依赖于原生的厂商特定 API，同时可以为 ROCm 和 CUDA 添加 GPU 感知通信。然而，也存在一些需要注意的问题。其一是需要澄清 GPU 感知 UCX 是否能够将通信与计算重叠。第二个问题是，UCX 提供的更高通用性和便利性可能会影响性能。进行一个原生 API 提供的 GPU 感知与通过 UCX 提供的 GPU 感知之间的性能比较将会很有帮助。在这方面的工作中，Khorassani 等人为 MVAPICH2 提供了一个原生的 ROCm 感知运行时，该运行时在一组 AMD GPU 集群上表现超过了通过 UCX 的 OpenMPI[104]。然而，性能差异可能源自 MPI 实现的不同，而非 UCX 本身。另一个当前 UCX 的限制是，它总是在主机上运行，并且不使用任何通信内核，而是依赖于传统的 cudamemcpy 系列函数以及零拷贝 RDMA。

5.3 无 CPU 网络

随着 GPU 集中式通信和更大 GPU 自主性趋势的加速，几项研究提出将大部分或所有网络栈迁移到内核中。这通常通过启动一个单一的长时间运行的持久内核，并将数据路径和控制路径都迁移到 GPU 来实现。

在最早的工作中，GGAS提出了对网络设备的修改，以实现一个统一的全局地址空间，从而将控制路径完全迁移到 GPU。这是通过使用一个包含计算、通信和同步的持久内核实现的，所有操作都在设备端完成。尽管该工作是同类中的首次尝试，并且与 CUDA 感知的 MPI 基线相比显示出性能提升，但实验是在两个 GPU 节点上进行的，每个节点只有一个 GPU，而提出的硬件修改是在 FPGA 上仿真进行的。后续工作表明，GGAS 通过消除 CPU 在控制路径中的参与，可以比 CPU 控制的基线实现更大的性能提升并减少能耗。

随后的几项关于无 CPU 网络的工作也取得了进展。Oden 等人使用 GPUDirect RDMA 使得 GPU 可以直接与 Infiniband 网络设备交互，无需主机 CPU 的参与。他们通过将整个 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，这是一种 NIC 硬件机制，允许 CPU 创建和注册消息到 NIC，GPU 从正在运行的持久内核中触发这些消息。另一项工作 ComP-Net 使用嵌入式 GPU 微处理器将辅助线程从 CPU 卸载到 GPU。虽然 GPU-TN 和 ComP-Net 都显示出有前景的性能，但它们需要对 NIC 和 GPU 进行硬件更改，因此依赖于模拟来获得结果。

5.4 更广泛的 GPU 自主性

最近 GPU为中心的通信的普及代表了向更广泛 GPU 自主性的总体趋势。早期和近期的若干工作尝试将 GPU 掌控传统上由 CPU 负责的领域。在一项早期工作中，Stuart 等人提出了允许 GPU 向 CPU 发出回调的方法。Silberstein 等人实现了 GPUfs，允许 GPU 从 GPU 内核内部直接请求主机 CPU 上的文件。Veselý 等人通过更改 Linux 内核实现了从 GPU 内核内部调用 POSIX 系统调用的支持。NVIDIA 的 GPUDirect Storage 提供了 GPU 与存储之间的直接数据通道，但仍然依赖于 CPU 来协调执行。在 NVIDIA 的一项近期工作中，Qureshi 等人提出了 BaM，这是第一个允许 GPU 在没有 CPU 介入的情况下直接访问存储的方法。实验结果表明，BaM 在多个工作负载上优于 GPUDirect Storage。

这些以及其他工作显示了 GPU 自主性日益增强的明确趋势。与此一致的是，我们预计 GPU 为中心的通信将进一步优化。正如 Punniyamurthy 等人指出的那样，若干近期的机制非常有前景。

首先，近期的线程块（TB）集群抽象可能有助于设备端通信-计算重叠和跨 TB 同步。

其次，AMD 最近的缓存刷新指令允许在启动网络通信之前刷新缓存，这意味着通信不再需要分配为不可缓存。

第三，像更强大的 GPU 节点和紧密的 GPU-NIC 集成等硬件趋势也很有前景。例如，NVSwitch 的最新版本直接连接 256 个 Grace Hopper Superchips，实现在前所未有的规模下进行直接的 P2P 全到全通信。

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 集中通信至关重要。

6 结论

传统上，多 GPU 通信由 CPU 管理，但近期 GPU 集中通信的进展已开始将这一责任转移给 GPU，让 GPU 在通信任务中发挥更多作用。本文深入探讨了 GPU 集中通信，重点分析了供应商提供的机制和用户级库的支持，旨在揭开这一领域的复杂性和多样性，定义关键术语，并对单节点和跨节点的各种方法进行分类。

讨论包括对供应商提供的多 GPU 执行通信和内存管理机制的分析，以及对 CUDA 感知 MPI、NCCL/RCCL、NVSHMEM 和 ROC_SHMEM 等主要通信库的回顾，突出它们的优势、挑战和性能考虑。此外，本文还介绍了 CPU 无关网络、通信调试工具等重要研究范式，并探讨了未来的方向和未解问题。通过全面考察 GPU 集中通信技术，涵盖软件和硬件层面，本文旨在为研究人员、开发人员、工程师和库设计人员提供全面的知识，帮助他们充分利用多 GPU 系统。

------------------------------------------------------------------------

本文相关论文存放在知识星球