9.1 GPU 架构基础

要理解大模型为什么能在合理时间内完成训练，必须先理解支撑它的硬件——GPU。GPU 并不是一个"更快的 CPU"，而是一种在设计哲学上与 CPU 截然不同的处理器。CPU 追求的是单个任务的最低延迟，GPU 追求的是单位时间内完成任务的最大数量（吞吐量）。这一根本性的差异决定了 GPU 内部的架构组织、执行模型和编程范式，也决定了什么样的算法能在 GPU 上高效运行。

本节将从 CPU 与 GPU 的核心差异出发，逐步深入到 GPU 内部的流式多处理器（SM）架构和线程/线程束/线程块/网格的层次化执行模型，为后续章节中的并行训练策略和性能优化奠定硬件基础。

9.1.1 CPU vs GPU：延迟导向 vs 吞吐量导向

理解 GPU 的第一步，是理解它与 CPU 在设计目标上的根本分歧。

CPU 的设计哲学：最小化延迟。 CPU 的核心目标是让单个线程尽可能快地完成任务。为此，现代 CPU 在每个核心内部堆叠了大量复杂的控制逻辑：深度流水线（pipeline）、分支预测器（branch predictor）、乱序执行引擎（out-of-order execution）、推测执行（speculative execution）等。这些机制的共同目的是减少单条指令的有效执行时间。与此同时，CPU 配备了多级大容量缓存（L1/L2/L3），总缓存容量可达数十甚至数百 MB，用于将频繁访问的数据保留在离计算单元最近的位置，从而隐藏主存的高延迟。然而，这些复杂控制逻辑和大缓存占据了大量芯片面积，导致 CPU 只能容纳少量核心——典型的服务器 CPU 有 8 到 96 个核心。

GPU 的设计哲学：最大化吞吐量。 GPU 走了一条完全不同的路线。它放弃了复杂的控制逻辑和大缓存，将节省出的芯片面积全部用于堆叠计算单元。一个现代 GPU（如 NVIDIA A100）拥有 6912 个 CUDA 核心和 432 个张量核心（Tensor Core），能够同时执行数以万计的线程。GPU 处理延迟的方式不是"消除"它，而是"隐藏"它——当一组线程因为等待内存数据而停顿时，硬件调度器会立即切换到另一组已经准备就绪的线程继续执行，从而保持计算单元始终处于忙碌状态。这种策略的代价是单个线程的执行速度较慢，但好处是在有足够多的并行任务时，整体吞吐量远超 CPU。

下表从多个维度对比了 CPU 和 GPU 的核心设计差异：

维度	CPU	GPU
设计目标	最小化单任务延迟	最大化并行吞吐量
核心数量	少量强核心（8-96 个）	大量简单核心（数千至上万个 CUDA 核心）
单核复杂度	高：乱序执行、分支预测、深流水线	低：顺序执行为主，控制逻辑简单
缓存策略	大容量多级缓存（L1/L2/L3 共计数十 MB）	小缓存 + 海量寄存器，依赖线程切换隐藏延迟
内存带宽	较低（DDR5 约 100-200 GB/s）	极高（HBM3 约 2-3 TB/s）
延迟隐藏方式	硬件预取 + 缓存 + 乱序执行	大量线程快速切换（零开销上下文切换）
适合的任务类型	串行逻辑、复杂控制流、低延迟响应	大规模数据并行、矩阵运算、批量处理
典型功耗	100-350W（服务器级）	250-700W（数据中心级）
指令集	复杂指令集（x86）或精简指令集（ARM）	SIMT 执行模型 + 专用张量指令

表 9-1：CPU 与 GPU 核心设计对比。

用一个类比来理解：CPU 像一支由少数顶尖专家组成的精英团队，每位专家都能独立解决复杂问题，但团队规模有限；GPU 像一条由成千上万名工人组成的大型流水线，每个工人只负责一个简单操作，但整条流水线的总产出远超精英团队。深度学习的核心计算——矩阵乘法——恰好是典型的"大规模简单操作"，因此天然适配 GPU 的吞吐量导向设计。

为什么深度学习选择了 GPU？ 神经网络的前向传播和反向传播本质上是一连串的矩阵乘法和逐元素操作。以一个形状为 $(B, M) \times (M, N)$ 的矩阵乘法为例，它包含 $B \times M \times N$ 次乘加运算，且这些运算之间高度独立。一个大模型的单次前向传播可能涉及数万亿次浮点运算，而这些运算可以被高效地分解为数以百万计的独立线程，恰好填满 GPU 的计算单元。这就是 GPU 成为深度学习核心算力引擎的根本原因。

9.1.2 GPU 内部架构：SM 与 SP

理解了 GPU 的设计哲学后，接下来深入其内部架构。现代 NVIDIA GPU 的架构是一个层次化的结构，其核心构建块是流式多处理器（Streaming Multiprocessor, SM）。

SM 是 GPU 的基本执行单元。 可以将 SM 理解为 GPU 内部的一个独立"微型处理器"。一个 GPU 由数十到上百个 SM 组成——例如 NVIDIA A100 拥有 108 个 SM，H100 拥有 132 个 SM。GPU 的总计算能力可以通过增加 SM 的数量来线性扩展，这也是不同档次 GPU 之间性能差异的主要来源之一。

SM 的内部结构。 每个 SM 内部集成了多种关键组件：

流式处理器（Streaming Processor, SP）/ CUDA 核心：这是最基本的标量计算单元，负责执行整数和浮点运算。每个 SM 内部包含大量 SP（例如 A100 的每个 SM 包含 64 个 FP32 CUDA 核心）。SP 是 SIMT 执行模型的物理基础。
张量核心（Tensor Core）：从 Volta 架构（V100）开始引入的专用矩阵计算单元。张量核心的核心操作是混合精度矩阵乘法累加（ $D = A \times B + C$ ）：输入矩阵 A 和 B 使用低精度格式（FP16/BF16/INT8），乘法结果在 FP32 精度下累加，兼顾了计算速度和数值稳定性。一个张量核心在单个时钟周期内就能完成一个小矩阵（如 $4 \times 4$ ）的乘法运算。张量核心的引入使得 GPU 在矩阵乘法上的吞吐量比纯 CUDA 核心高出一个数量级以上，是现代深度学习 GPU 性能飞跃的关键。
线程束调度器（Warp Scheduler）：SM 的调度引擎，负责管理其内部数以千计的线程。调度器的核心策略是：当一个线程束（Warp，32 个线程的组）因为等待内存数据而停顿时，零开销切换到另一个就绪的线程束继续执行。这种快速切换正是 GPU 隐藏内存延迟的核心机制。
共享内存（Shared Memory）与 L1 缓存：每个 SM 拥有一块私有的高速片上 SRAM（通常数十 KB），其访问延迟仅约 20-30 个时钟周期，远快于全局内存的数百个周期。同一线程块内的所有线程可以通过共享内存进行快速数据交换。在性能优化中，将数据从全局内存预加载到共享内存（即"分块"技术）是最核心的策略之一。
寄存器文件（Register File）：每个 SM 拥有一个大型寄存器文件（A100 约 256 KB/SM），供其内部所有线程使用。寄存器是 GPU 中最快的存储单元，访问延迟仅约 1 个时钟周期。

GPU 的整体结构层次。 从宏观到微观，GPU 的结构可以概括为：

GPU 芯片
 ├── GPC（图形处理集群）× 若干
 │    └── SM（流式多处理器）× 若干
 │         ├── CUDA 核心（SP）× 若干（如 A100: 64 FP32）
 │         ├── 张量核心 × 若干（如 A100: 4 个第三代）
 │         ├── 线程束调度器
 │         ├── 共享内存 / L1 缓存
 │         └── 寄存器文件
 ├── L2 缓存（所有 SM 共享，A100: 40 MB）
 └── HBM（全局内存，A100: 40/80 GB，带宽 ~2 TB/s）

GPU 的内存层级。 与 CPU 类似，GPU 也有多层内存结构，但设计侧重不同。从快到慢依次为：

寄存器：每个线程私有，延迟约 1 个周期，速度最快但容量最小。
共享内存 / L1 缓存：每个 SM 私有，线程块内共享，延迟约 20-30 个周期。程序员可以显式管理。
L2 缓存：所有 SM 共享，对程序员透明，延迟约 200-300 个周期。
全局内存（HBM/GDDR）：容量最大（数十 GB），但延迟最高（数百个周期）。所有线程可访问，也是 CPU-GPU 之间数据传输的通道。

GPU 的计算能力增长速度远超内存带宽增长速度，这导致了日益严峻的**"内存墙"问题**——计算单元经常因为等待数据而空闲。这也是 FlashAttention 等算法被设计出来的根本动机：它们通过重构计算顺序来最小化全局内存访问，从而将计算从"内存受限"推向"计算受限"。

9.1.3 执行模型：线程、线程束、线程块与网格

GPU 的硬件架构决定了它需要一套与 CPU 完全不同的执行模型。在 NVIDIA 的 CUDA 编程框架中，这套模型由四个层次构成：线程（Thread）、线程束（Warp）、线程块（Block） 和网格（Grid）。

理解这四个层次——尤其是它们如何映射到硬件——是理解 GPU 性能特性的关键。

线程（Thread）：最小的逻辑执行单元。 每个线程执行相同的核函数（kernel）代码，但通过唯一的线程标识符（threadIdx）处理不同的数据。例如，在向量加法 $C [i] = A [i] + B [i]$ 中，第 $i$ 个线程负责计算第 $i$ 个元素。每个线程拥有自己私有的寄存器来存储局部变量。

线程束（Warp）：硬件的基本调度单元。 这是理解 GPU 执行行为最关键的概念。一个 Warp 由 32 个连续线程组成。Warp 是 SM 调度器实际操作的最小粒度——在任意时刻，SM 都是以 Warp 为单位发射指令的。一个 Warp 内的 32 个线程在同一个时钟周期执行完全相同的指令，只是操作的数据不同。这就是 SIMT（Single Instruction, Multiple Threads，单指令多线程） 执行模型。

SIMT 可以理解为 SIMD（单指令多数据）的一种更灵活的变体。在传统 SIMD（如 CPU 的 AVX 指令集）中，程序员需要显式地将数据打包成向量并使用向量指令操作；而在 SIMT 中，程序员只需编写针对单个线程的标量代码，硬件自动将 32 个线程组成的 Warp 以 SIMD 方式执行。这大大降低了并行编程的门槛。

Warp 的存在带来两个重要的性能影响：

控制流分歧（Control Divergence）：如果一个 Warp 内的线程在 if-else 语句中走向了不同分支，硬件无法同时执行两条路径。它会串行化执行——先执行 if 分支（不满足条件的线程被屏蔽等待），再执行 else 分支（满足条件的线程被屏蔽等待），最后在分支结束处重新汇合。在最坏情况下，一个 Warp 的有效并行度降为 1/32。因此，编写高效 GPU 代码的重要原则之一是尽量避免 Warp 内部的分支分歧。
内存合并（Memory Coalescing）：当一个 Warp 中的 32 个线程访问全局内存时，如果它们的地址是连续且对齐的，硬件可以将这 32 次独立访问合并为一次或几次大的内存事务，极大提升有效带宽。反之，如果访问地址分散随机，每次访问都会独立执行，带宽利用率急剧下降。

线程块（Block）：协作的线程集合。 线程块是一组可以相互协作的线程。一个线程块被整体调度到单个 SM 上执行，其生命周期完全在一个 SM 内部——这意味着同一线程块内的线程可以通过共享内存进行高速数据交换，也可以通过同步原语（__syncthreads()）协调执行顺序。线程块是程序员控制数据局部性的主要手段：通过精心设计线程块的大小和数据访问模式，可以最大化共享内存的利用率。

一个线程块的大小（即包含的线程数）由程序员指定，通常是 Warp 大小（32）的整数倍，常见的选择有 128、256 或 512 个线程。线程块内部被硬件自动划分为若干个 Warp（例如，一个 256 线程的线程块包含 8 个 Warp）。

网格（Grid）：线程块的集合。 一次核函数调用会创建一个网格，网格由多个线程块组成。关键特性是：不同线程块之间是相互独立的——它们不能直接通信或同步，任何跨线程块的数据交换都必须通过全局内存间接进行。这种设计看似限制，实则赋予了 GPU 极高的调度灵活性：硬件可以以任意顺序、在任意可用的 SM 上执行网格中的线程块，从而自动适应不同规模的 GPU（SM 数量不同的 GPU 可以运行完全相同的代码）。

下面将这四个层次的映射关系总结如下：

层次	逻辑概念	硬件映射	规模	可否协作
线程（Thread）	最小执行单元	单个 CUDA 核心的一次操作	1 个	—
线程束（Warp）	硬件调度单元	SM 调度器的最小粒度	32 个线程	隐式（锁步执行）
线程块（Block）	协作线程组	映射到单个 SM	数百个线程（如 256）	可以（共享内存 + 同步）
网格（Grid）	线程块的集合	分布到所有 SM	数千至数百万个线程	不可以（独立执行）

表 9-2：GPU 执行模型的四个层次。

一个完整的执行流程示例。 假设我们要对两个长度为 $N = 100,000$ 的向量做逐元素相加。我们设定每个线程块包含 256 个线程，则需要 $⌈ 100,000 / 256 ⌉ = 391$ 个线程块，它们共同构成一个网格。启动核函数后：

GPU 的全局调度器将 391 个线程块分配到可用的 SM 上。以 A100（108 个 SM）为例，第一轮可以将 108 个线程块分配到 108 个 SM 上并行执行。
每个 SM 收到一个线程块后，将其中的 256 个线程划分为 $256 / 32 = 8$ 个 Warp。
SM 的 Warp 调度器从这 8 个 Warp 中选择就绪的 Warp 发射指令。当某个 Warp 因等待全局内存数据而停顿时，调度器立即切换到另一个就绪的 Warp，保持 SM 的计算单元持续工作。
第一批 108 个线程块执行完毕后，调度器继续分配下一批线程块，直到全部 391 个线程块完成。

波形量化（Wave Quantization）效应。 上述分批执行的机制引出了一个容易被忽视的性能陷阱。GPU 的线程块是以"波"（wave）为单位执行的，每一波最多并行执行与 SM 数量相等的线程块。如果线程块总数恰好略超过 SM 数量的整数倍，最后一波将只有很少的线程块在执行，而大量 SM 闲置。

例如，A100 有 108 个 SM。如果一个任务需要 108 个线程块，一波就能完成，利用率 100%；但如果需要 120 个线程块，就需要两波——第一波 108 个块（满载），第二波仅 12 个块（利用率仅 11%）。任务量只增加了 11%，但执行时间几乎翻倍。这就是波形量化效应，它要求在设计并行任务时，尽量使线程块总数为 SM 数量的整数倍，或远大于 SM 数量以摊薄尾波的影响。

9.1.4 SIMT 模型与 CUDA 编程范式

理解了硬件架构和执行层次之后，最后简要介绍将软件映射到硬件的编程接口——CUDA。

CUDA（Compute Unified Device Architecture） 是 NVIDIA 推出的并行计算平台和编程模型。它允许开发者使用 C++/Python 等高级语言编写在 GPU 上运行的程序，而不需要关心底层的图形渲染 API。CUDA 的编程模型直接映射到前面介绍的硬件层次：

程序员编写核函数（kernel）——一段针对单个线程的代码。
启动核函数时，程序员指定网格大小（线程块数量）和线程块大小（每块的线程数），语法为 kernel<<<grid_size, block_size>>>(args)。
硬件自动将线程组织为 Warp，将线程块调度到 SM，并管理所有的并行执行和内存访问。

CUDA 的意义在于它开启了 GPGPU（General-Purpose computing on GPU） 时代——GPU 从专用的图形渲染芯片转变为通用的并行计算加速器。围绕 CUDA，NVIDIA 构建了完整的软件生态系统：cuBLAS（线性代数）、cuDNN（深度学习原语）、NCCL（多 GPU 通信）等高度优化的库，使得 PyTorch、TensorFlow 等深度学习框架能够透明地利用 GPU 算力。

SIMT 与 SIMD 的本质区别。 在传统 SIMD 编程（如 CPU 的 AVX-512 指令集）中，程序员必须显式地将数据打包成固定宽度的向量（如 512 位 = 16 个 FP32），并调用专用的向量指令。SIMT 则提供了一种更自然的抽象：程序员只需编写标量代码（单线程逻辑），GPU 硬件在底层自动将 32 个线程的执行"向量化"。这种设计使得 GPU 编程的门槛大幅降低，同时保留了灵活性——每个线程可以独立进行内存寻址和条件判断（尽管分支分歧会损失效率）。

9.1.5 本节总结

本节介绍了 GPU 架构的基础知识，核心要点如下：

CPU 与 GPU 的根本差异在于设计目标不同——CPU 是延迟导向的（少量复杂核心 + 大缓存 + 复杂控制逻辑），GPU 是吞吐量导向的（大量简单核心 + 海量线程 + 线程切换隐藏延迟）。深度学习的核心计算（矩阵乘法）具有高度数据并行性，天然适配 GPU 的设计。
SM（流式多处理器）是 GPU 的核心构建块，内部集成了 CUDA 核心（通用计算）、张量核心（矩阵加速）、Warp 调度器、共享内存和寄存器文件。GPU 的性能通过增加 SM 数量来扩展。
执行模型由四个层次构成：线程（最小逻辑单元）→ 线程束/Warp（32 线程，硬件调度的最小粒度，SIMT 执行）→ 线程块（映射到单个 SM，可通过共享内存协作）→ 网格（所有线程块的集合，跨 SM 独立执行）。
Warp 是理解 GPU 性能的关键：控制流分歧会导致 Warp 内线程串行化执行；内存合并要求 Warp 内线程访问连续地址；波形量化效应要求线程块总数与 SM 数量合理匹配。
GPU 的内存层级从快到慢依次为寄存器、共享内存/L1、L2、全局内存（HBM），跨度达数百倍延迟差距。"计算增长快于带宽增长"导致的内存墙问题，是后续章节中算子优化和 FlashAttention 等技术的核心驱动力。

这些概念将在后续章节中被反复引用——无论是理解张量并行的通信开销、分析混合精度训练的数值行为，还是设计高效的自定义 CUDA 算子，都需要以本节的 GPU 架构知识作为基础。

9.1 GPU 架构基础 ​

9.1.1 CPU vs GPU：延迟导向 vs 吞吐量导向 ​

9.1.2 GPU 内部架构：SM 与 SP ​

9.1.3 执行模型：线程、线程束、线程块与网格 ​