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

1 引言

近年来，GPU已成为高性能计算（HPC）和机器学习（ML）领域众多应用的首选加速器。这种快速的普及，主要是由于GPU具有强大的并行处理能力和高速内存带宽，使得大多数现代云计算和HPC能力如今都集中在GPU集群中。截至2024年6月，Top500超级计算机排名中的前10名中，有9台依赖GPU集群进行加速，这一趋势很可能会继续下去。唯一没有依赖GPU的超级计算机，依赖的是高度向量化的CPU架构与高带宽内存。虽然使用大量GPU已经被证明能显著加速计算，但GPU之间的通信很快就可能成为可扩展性的瓶颈。传统上，节点内外的多GPU通信一直是由CPU负责的。从GPU的初衷来看，它们被视为能够提供大量计算能力的设备，但本质上也依赖于CPU来处理辅助任务，如通信。在这种以CPU为中心的执行模型中，传递数据给GPU进行计算的例程并未考虑GPU的存在。在过去的十年里，出现了几项进展，统称为GPU中心的通信，它们试图挑战CPU在多GPU执行中的主导地位。总体来说，这些进展减少了CPU在执行关键路径中的参与，使GPU在启动和同步通信方面有了更多的自主性，并试图解决多GPU通信与计算之间的语义不匹配问题。本文将对GPU中心的通信进行全面回顾，重点讨论供应商机制和用户级库的支持。我们进行此项调查的目的是帮助消除当研究人员初次涉足该领域时所产生的困惑。我们希望帮助程序员、工程师、编程模型和库设计者理解可用选项的复杂性和多样性，因为GPU中心的通信涵盖了广泛的技术方法，包括硬件创新，如专有的GPU-to-GPU互连技术，以及软件机制。这些机制各有优缺点，难以判断在何种情况下应该优先选择哪种方法。由于文献中使用的术语不一致，而且供应商的产品各异，这个领域的情况更加复杂。

本文的结构如下：

1. 第2节：我们定义了相关术语，提供了GPU中心通信的定义，并对节点内和节点间现有的通信方法进行了分类，以消除可能的混淆。
2. 第3节：我们讨论了供应商提供的机制，这些机制用于启用通信和网络功能，管理多GPU执行中的设备内存。这些机制作为高层GPU中心软件库的构建模块。
3. 第4节：我们列出了主要的通信库，包括节点内和节点间的设置，讨论它们的优点和挑战，并依赖现有的基准测试文献提供性能洞见。
4. 第5节：我们讨论了GPU中心通信的主要研究范式，展望该领域的发展，并提出了开放的研究问题。

我们还承认，某些方法和技术依赖于专有的生态系统。在这项工作中，我们主要关注NVIDIA的解决方案。然而，AMD也提供了类似的解决方案，我们特别强调了只有其中一个供应商提供的技术。英特尔也提供高端HPC/ML GPU，但在撰写时，关于英特尔产品线的前景公开信息较少。英特尔的Ponte Vecchio不会部署在新的集群上，而下一代Falcon Shores GPU预计在2025年末才会发布。

2 术语与通信类型

我们可以宽泛地定义GPU中心通信为一种减少CPU在多GPU执行关键路径中参与的机制。这是一个非常广泛的定义，涵盖了多种解决方案，涉及供应商层面的改进（使GPU在通信中具有更多自主性）以及用户层面的实现（利用这些改进）。为了明确这一区别，我们将它们分别讨论。在第3节中，我们重点讨论NVIDIA CUDA和AMD ROCm运行时中原生提供的通信机制和原语。在第4节中，我们讨论这些机制如何催生更高级别的GPU中心通信库。我们还指出了节点内通信（intra-node）和节点间通信（inter-node）之间的区别。单个GPU加速节点包括一个共享内存的主机和多个附加的GPU卡。当进行节点内通信时，任何给定的GPU都可以由单个线程或进程控制，且它们共享内存和地址空间。多节点系统由多个这样的节点组成，每个GPU由不同的进程控制，且在不同节点上运行的进程之间内存不共享。通信的全景图会根据所使用的设置发生变化，因为节点间通信需要处理GPU与网卡（GPU-NIC）的交互以及跨进程通信。尽管将通信方法分类为GPU侧和CPU侧的方法常见，并且对于终端用户来说是足够的，但这种分类并不总是解释清晰且准确。为了避免定义的模糊性，我们将通信方法分为几种类型。这些类型是根据在通信过程中执行每个操作的执行者来划分的。我们定义了两种在节点内场景中进行通信所需的主要操作和四种在节点间场景中进行通信所需的操作。

在节点内的情况下，通信调用的两个组成部分是：

1. API：定义了程序员或库发起通信API调用的位置。
2. 数据路径：指明了谁参与数据的移动，并展示相应的数据路径。

表1和图1展示了节点内(Intra-node) 通信机制的分类示例，以及描绘数据路径的图示

以GPU为中心通信的全景

在图1中，通信方法1被称为“主机原生通信”，它发生在主机端，不涉及设备之间的直接P2P（点对点）访问。这些方法都是在主机端启动的，且禁用了P2P访问。否则，如通信方法2所示称为"主机控制通信”，主机控制的通信方法不会涉及额外的内存拷贝到主机内存，而是直接通过PCIe、NVLink或Infinity Fabric互连。NCCL、GPU感知MPI和*memcpy操作具有主机端API，因此它们可以归属于1和2两类。通过直接访问对等设备的内存，设备端API允许在节点内通信中将CPU从数据路径和控制路径中移除，如通信方法3所示，称为“设备原生通信”。NVSHMEM和ROCSHMEM也提供主机端API，但需要P2P访问，因此它们的主机端API属于2类，而设备端API属于3类。内核级P2P直接加载和存储提供类似的功能，属于3类，但即使在禁用P2P访问的情况下也能工作，这属于通信方法4，在这种情况下数据路径会回退到通过主机。

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

注册/构建消息：这一步涉及构建数据包并在网卡上注册它们。

触发通信：它定义了谁在网卡上触发“门铃”以发起数据传输。

我们在分类节点间通信时识别了五个主要类别。如表2和图2所示，随着通信调用的更多组件移到GPU端，数据传输路径中所需的数据拷贝次数也减少了。多年来，通信方法和相应的技术展示了数据传输和通信控制的优化，这些将在第3节中讨论。首先，完全由CPU端处理的通信方法(1)是可用的，后来通过移除CPU-GPU和CPU-NIC缓冲区之间的额外拷贝，改进为方法(2)。之后，GPU端的优化使得NIC能够直接访问GPU内存-方法(3），最小化了它们之间的数据路径。方法(4)代表了GPU触发的通信技术，如GPUDirect Async和GPU-TN，其中GPU能够发起通信，前提是CPU提前准备好了网卡上的数据包。方法(5)将数据包准备和与网卡的交互完全移交给GPU，使得设备原生通信成为可能。

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

3  厂商机制

在这一节中，我们讨论了供应商提供的机制，用于启用通信和网络功能，并在多GPU执行中管理设备间的内存。这些机制由GPU编程模型的运行时提供，或作为扩展API的一部分。

图3总结了NVIDIA提供的技术，详细列出了它们的时间线和可用性。这些技术分为四类：内存管理器、GPUDirect技术、硬件和库。接下来，我们将介绍内存管理机制和GPUDirect技术，随后是作为先驱的硬件支持，这些硬件最终使这些通信方法成为可能。这些技术构成了更高级GPU中心库的基础，这将在第4节中进行讨论。

3.1 内存管理机制

3.1.1 页面锁定/固定内存(Page-Locked Pinned Memory)

默认情况下，使用cudaMalloc()（这里应该为改为Malloc)在主机上分配的内存是可分页的，GPU无法直接访问。当在可分页的主机内存与设备内存之间进行数据传输时，GPU运行时必须先通过临时缓冲区将主机数据暂存到页面锁定内存中，然后再将数据从页面锁定内存复制到GPU内存。为了避免可分页内存到页面锁定内存的拷贝，cudaMallocHost()允许直接分配页面锁定内存，从而跳过中间的复制阶段。因此，页面锁定内存也被称为零拷贝内存或固定内存。

固定内存因其在主机与设备之间传输时具有较高带宽和低延迟而著称，通过启用系统范围内的直接访问，能够高效地协调CPU和GPU的执行。它还与GPUDirect RDMA结合使用，以改善节点间的通信。然而，其物理内存锁定可能导致显著的内存消耗，过度分配可能会影响系统性能。

3.1.2 统一虚拟地址（UVA）

UVA是CUDA 4.0中引入的一种内存管理技术，它允许节点内的所有GPU和CPU共享相同的统一虚拟地址空间。在引入UVA之前，主机与设备之间以及设备与设备之间的复制必须显式地指定传输方向。而通过UVA，物理内存位置可以从指针值推断出来，从而减少了管理独立内存空间的开销，并使库能够简化其接口。

3.1.3 CUDA IPC

在早期的CUDA版本中，指针不能跨进程边界访问，因此GPU缓冲区之间的内存复制必须通过主机进行，从而形成了瓶颈。为了克服这一限制，CUDA 4.1引入了CUDA进程间通信（IPC），它允许同一节点上的进程在不进行额外复制的情况下访问其他进程的设备缓冲区。通过CUDA IPC，内存句柄在进程之间使用标准的IPC机制创建和传递，从而减少了通过主机进行拷贝时产生的延迟。然而，创建内存句柄的开销可能是显著的，并且可能会抵消延迟带来的好处。

3.1.4 统一虚拟内存(UVM)

UVM是在CUDA 6.0中引入的，它允许通过cudaMallocManaged()调用分配受管理的内存，创建一个对单个节点内所有处理器可访问的单一地址空间。UVM通过将请求的内存划分为驻留在CPU上的页面来工作。程序员可以访问设备上的内存，而不需要显式的内存复制。如果访问的内存页面不在设备上，UVM驱动程序会触发页面错误，并自动将该页面迁移到请求设备上。当总分页内存大小超过设备内存时，UVM驱动程序还可以将页面从设备驱逐回主机内存。

UVM在可编程性方面提供了几个优势。首先，程序员可以访问一个统一的地址空间，就好像整个分配的内存块都驻留在单个GPU上一样。系统中发生的任何内存拷贝都是隐式的，并且对程序员是不可见的。此外，UVM允许内存过度订阅，这意味着可以分配的内存总量超过所有多GPU设备内存的总和。这是可能的，因为大部分内存可以保留在CPU上，并在设备请求时进行分页加载。

背景知识

cudaMalloc()
、cudaMallocHost()
 和cudaMallocManaged()
 是CUDA中用来分配内存的不同函数，它们各自的功能和用途有显著的区别。以下是它们的详细解释：

1.cudaMalloc()

• 功能：在GPU设备上分配内存。

• 使用场景：用于分配GPU设备上的内存，当你需要在GPU上执行计算并访问或存储数据时使用。

• 行为：分配的内存是设备内存，程序员需要显式地将数据从主机内存复制到设备内存，或者从设备内存复制回主机内存。

• 示例代码：

  float *d_array;
  cudaMalloc((void**)&d_array, size);
  

• 注意：分配的内存是设备内存，它无法直接由主机访问。

2.cudaMallocHost()

• 功能：在主机内存中分配页面锁定内存（pinned memory），并返回指向该内存的指针。

• 使用场景：用于主机内存中的固定内存，适用于需要高带宽和低延迟的数据传输场景，例如高效的CPU与GPU之间的数据交换。

• 行为：分配的内存是主机内存，但它是“固定”的，无法被操作系统分页（即锁定）。这允许更高效的数据传输，因为GPU可以直接访问这些内存，而无需先将数据从主机内存复制到页面锁定内存。

• 示例代码:    

  float *h_array;
  cudaMallocHost((void**)&h_array, size);
  

  
  

• 注意：页面锁定内存不能随便大量使用，因为它会占用系统的物理内存，可能影响系统的整体性能。

3.cudaMallocManaged()

• 功能：分配统一内存（managed memory），即CPU和GPU都可以访问的内存。

• 使用场景：用于分配可以在CPU和GPU之间共享的内存，并且能够自动处理内存迁移（当数据从CPU访问到GPU时，或者反之时，内存会自动迁移）。

• 行为：分配的内存是统一内存，允许CPU和GPU共享同一个内存地址空间。当程序访问这些内存时，CUDA运行时会自动将数据迁移到相应的设备（如CPU或GPU），从而简化了内存管理。

• 示例代码：

  float *umem_array;
  cudaMallocManaged((void**)&umem_array, size);
  

  注意：统一内存简化了内存管理，但由于内存迁移的开销，可能会影响性能。它适合于不需要频繁数据传输或复杂内存管理的应用场景。

总结：

选择哪个函数取决于你的应用场景及对内存管理的需求。例如，cudaMalloc()
适用于传统的设备计算，cudaMallocHost()
适用于需要高速数据交换的场景，而cudaMallocManaged()
适用于需要简化内存管理的应用。

未完待续...     

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

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

相关文档和资料统一存放在知识星球，加入获得更多相关资料

本文根据以下资料撰写，加入星球可获得更多1500+详细资料

互动群加入，目前已经满200，先加微信后再加入