揭秘GPU：从核心架构到可扩展系统

目录

• 动机
• GPU的优化目标
• GPU 的关键概念——软件和硬件
• 深入探索扭曲空间
• 张量核心与 CUDA 核心有何不同？
• 将软件映射到硬件
• 为什么会发生空间扭曲？
• 为什么一个代码块只能在单个SM上运行？
• L1缓存和共享内存有什么区别？
• 如果单个积木无法放入单个空间模型中怎么办？
• 跨GPU扩展计算能力
• 符号 16x 和 4x 的解释

动机

想了解GPU的工作原理，却被各种专业术语搞得一头雾水？本文将为您解答：

• GPU 针对哪些方面进行了优化
• 理解单个GPU工作原理的关键概念
• 单个节点内的GPU如何相互连接
• GPU节点是如何连接的？

感谢斯坦福大学 CS 336 课程的讲授，我在 Google Gemini 的帮助下撰写了这篇文章，以帮助我填补知识空白，尤其是在我缺乏充分理解课程材料的先决知识时。

GPU的优化目标

GPU 针对并行处理和高吞吐量计算进行了优化，因此能够高效地同时处理海量数据。这与 CPU 形成鲜明对比，CPU 主要针对顺序处理和以较低延迟处理各种任务进行了优化。

在机器学习领域，我们主要使用GPU进行矩阵乘法运算。目标是使矩阵乘法运算受限于计算能力，这可以通过高FLOPs/bytes（参见下方的屋顶线模型）来实现。当GPU受限于计算能力时，意味着CUDA核心得到了最大程度的利用，即GPU性能得到了最大化发挥。

“屋顶线”代表给定硬件平台可达到的最大性能。倾斜的“屋顶”表示程序受限于内存，即处理速度受制于数据从内存传输到处理器的速度。平坦的“天花板”代表处理器算术逻辑单元（在Nvidia GPU中称为CUDA核心）的理论峰值性能。这表明程序受限于处理器的原始计算能力。

来源：Nvidia Nsight

• 运算强度：定义为执行的浮点运算总数 (FLOP) 与从内存中读取的数据字节数的比值。运算强度越高，意味着算法对读取的每个数据单元执行的计算次数越多，这种情况通常发生在密集矩阵乘法中。相反，对于包含大量零元素的稀疏矩阵，涉及这些零元素的计算通常会被跳过，从而减少了从内存读取的单位数据所需的计算次数。
• 吞吐量：衡量运行特定程序时硬件的性能。

GPU 的关键概念——软件和硬件

• 注：SIMT 指的是单指令多线程。它是一种并行计算中使用的执行模型，将单指令多数据流 (SIMD) 与零开销多线程相结合，即硬件能够逐周期地在线程之间切换的多线程技术。

来源：揭秘GPU计算架构

深入探索扭曲空间

下图是英伟达 Hopper H100 架构白皮书中的扭曲示意图。

来源：揭秘GPU计算架构

我们可以看到处理单元顶部有一个“Warp Scheduler (32 thread/clk)”模块。“Warp Scheduler (32 thread/clk)”表示该调度器每个时钟周期可以调度32个线程。这表明了调度器的吞吐量，使其能够在每个时钟周期内为一个完整的线程束（32个线程）发出指令。

“Clk”是时钟周期（或时钟）的缩写。在GPU等处理器中，时钟周期是同步内部组件运行的基本时间单位。时钟速度（以MHz或GHz为单位）决定了每秒发生的时钟周期数。

INT32、FP32 和 FP64 的数量表示线程束中可用于对这些数据类型执行操作的物理硬件单元数量。每个单元都是一个 CUDA 核心。如果用于特定操作的单元数量少于 32 个，则整个线程束的操作将在可用单元上“串行化”执行，耗时多个时钟周期。单个线程束不能在单个时钟周期内同时执行 FP32、FP64 和 INT32 操作。

张量核心与 CUDA 核心有何不同？

Tensor Core 和 CUDA Core 都是 NVIDIA GPU 中的处理单元，但它们是为不同类型的计算任务而设计的。

• CUDA核心：这些是通用并行处理器。它们可以处理各种计算任务，包括图形渲染、物理模拟、视频处理和一般计算操作。它们主要处理单精度 (FP32) 和双精度 (FP64) 浮点数，因此用途广泛，适用于各种计算任务。
• 张量核心：这些是专门优化的硬件单元，旨在加速张量运算，特别是矩阵乘法和融合乘加 (FMA) 运算，这些运算是深度学习和高性能计算的基础。它们专为混合精度计算而设计，支持半精度 (FP16) 和整数 (INT8) 计算，从而在不显著降低精度损失的情况下，加快 AI 模型的训练速度并减少内存消耗。

无需显式指示 PyTorch 使用张量核心进行矩阵乘法。PyTorch 使用诸如 CuBLAS 和 CuDNN 之类的库，只要满足以下条件，这些库就会自动利用张量核心：

1. GPU支持Tensor核心。
2. 数据使用适当的数据类型，例如混合精度。
3. 这些操作适合由张量核心处理，例如矩阵乘法（例如 torch.matmul、线性层、卷积）和融合乘加运算。逐元素运算、激活函数等通常由 CUDA 核心处理。

将软件映射到硬件

程序员定义了内存Grid、Block内存和Thread内存结构。GPU 硬件将此结构映射到其物理内存SMs和内存CUDA Cores地址上，并使用Warp内存调度机制来保持硬件繁忙并隐藏内存延迟。

程序员的逻辑视角：

Grid {
  Block_1 { Thread_1, Thread_2, ..., Thread_1024 }
  Block_2 { Thread_1, Thread_2, ..., Thread_1024 }
  ...
}

GPU的物理执行视图：

GPU {
  SM_1 {
    - Schedules Warp_A, Warp_B, ... (from Block_1)
    - Its [CUDA Cores] execute instructions for threads in those warps
    - Has a pool of Shared Memory for Block_1
  }
  SM_2 {
    - Schedules Warp_X, Warp_Y, ... (from Block_2)
    - Its [CUDA Cores] execute instructions for threads in those warps
    - Has a pool of Shared Memory for Block_2
  }
  ...
}

为什么会发生空间扭曲？

由于指令的获取和解码是复杂的任务，GPU 通过将每个控制逻辑应用于 32 个 CUDA 核心组成的集合，从而简化了控制逻辑，以适应更多核心。这降低了控制逻辑的复杂性和功耗。

因此，程序员需要精心设计代码，以最大限度地利用线程束中的所有 32 个核心。例如，在以下if-else语句中，线程束在任何时刻都只使用了 16 个核心，这导致利用率仅为 50%。

// All 32 threads in a warp execute this code at the same time
if (threadIdx.x < 16) {
    // Path A: Executed by the first 16 threads
    do_something(); 
} else {
    // Path B: Executed by the last 16 threads
    do_something_else();
}

下面这张简单的图表展示了如何使用 if-else 语句来实现 50% 的线程利用率。

为什么一个代码块只能在单个SM上运行？

• 共享内存访问：共享内存和 L1 缓存的速度比 L2 缓存和全局内存快 10 倍。由于共享内存和 L1 缓存都位于共享内存模块 (SM) 内部，因此数据块可以利用这些高速内存的优势。内存周期的差异使得实现方式能够充分利用更快的共享内存和 L1 缓存来加速计算，例如矩阵乘法中的分块操作。分块操作需要在将新数据加载到共享内存之前，尽可能多地将有用数据加载到共享内存中，从而最大限度地提高共享内存的计算效率。这减少了从速度较慢的全局内存读取数据的次数。

来源：优化您的 GPU 基础架构

• 线程同步：屏障同步_syncthreads()由 SM 在硬件级别管理，如果一个代码块跨越多个 SM，则会显著降低同步速度。

• 高效调度：通过将一个块限制在一个 SM 中，GPU 的主调度器更容易找到一个拥有足够资源的 SM 来启动整个块，从而避免了将块的资源和线程分配到多个硬件单元所需的工作。

L1缓存和共享内存有什么区别？

在现代英伟达GPU中，L1缓存和共享内存都位于SM（子系统）中单个统一的SRAM块上。L1缓存和共享内存会根据内核动态划分。程序员可以设置cudaFuncCachePreferShared内核是否更依赖线程间通信，以及cudaFuncCachePreferL1内核是否具有分散的全局内存访问模式。

来源：NVIDIA A100 40GB GPU 的基本内存层次结构

如果单个积木无法放入单个空间模型中怎么办？

CUDA 运行时将返回错误，内核将无法运行。代码块必须遵守与 GPU 硬件的以下约定：

• 每个块的最大线程数：例如，现代 Nvidia GPU 每个块最多允许 1024 个线程。
• 每个块共享内存：每个 SM 都有固定量的物理共享内存，例如 64 KB、96 KB。
• 每个块的寄存器数量：每个SM（状态机）拥有有限数量的物理寄存器，例如65,546个32位寄存器。块中的每个线程都需要一定数量的寄存器，该数量由编译器决定。一个块所需的寄存器总数为(threads per block) * (registers per thread)……

跨GPU扩展计算能力

为了支持密集型并行处理，GPU 节点按以下设置互连，以实现节点内（节点内）和节点间（节点间）的高速数据传输。

来源：CS 336，并行性讲座 1，9 分 14 秒

简而言之，高性能 NVIDIA GPU 集群中的节点利用两个不同的并行网络层。对于最密集的 GPU 间通信，它们通过 NVLink 架构连接，该架构由 NVSwitch 管理和扩展。这种高带宽架构既可用于节点内通信（连接单个服务器内的所有 GPU），也可用于高级系统中的节点间通信，使 NVSwitch 能够连接数十个节点，形成一个单一、无缝、数据中心规模的 GPU 计算域。

同时，这些节点还使用主机通道适配器 (HCA) 连接到更广泛的集群网络，例如 InfiniBand 或以太网。这第二层网络处理所有其他节点间通信，包括连接到存储系统以及将集群扩展到数百或数千个节点，远远超出单个 NVLink 架构所能承载的范围。

来源：Nvidia Hopper架构深度解析

符号 16x 和 4x 的解释

符号16x和4x分别指相应互连（PCIe、xGMI）中的通道数。

• 通道：通道是一组差分线对，其中一对用于发送数据，另一对用于接收数据。这实现了全双工通信，意味着数据可以同时发送和接收。
• 带宽扩展：互连的总带宽与其通道数成正比。16x 连接有 16 条通道，因此其带宽是 1x 连接的 16 倍。

来源：维基百科

结合图表来看：

• 16x：表示具有 16 条通道的连接。这是一种高带宽连接，通常用于 GPU 等高要求组件以及 CPU 和 PCIe 交换机之间的链路。例如，PCIe 4.0 x16 连接的理论带宽为 32 GB/s。
• 4x：表示 4 通道连接。虽然带宽低于 16x 连接，但它仍然是高速链路，适用于图中所示的主机通道适配器 (HCA) 等设备。PCIe 4.0 x4 连接的理论带宽为 8 GB/s。