GPU Architecture — 从晶体管到线程
更新于 2026-04-03
在 AI Compute Stack 全景 文章中，我们从软件栈的角度理解了从推理框架到硬件 ISA 的七层结构。那篇文章的最底层是”Hardware ISA — GPU 能执行的唯一东西”。
本文往下钻一层：GPU 硬件本身长什么样。理解硬件架构是写高性能 GPU 代码的前提 — 你需要知道指令在什么样的物理结构上执行，才能做出正确的优化决策。
Section 1: GPU vs CPU — 两种设计哲学
CPU 和 GPU 面对的是两种截然不同的计算需求，因此走了完全不同的设计路线。
CPU 的设计目标：最小化单个任务的延迟（latency）。 为此，CPU 把大量晶体管预算投入到：
- 大容量多级缓存（L1/L2/L3，有时占芯片面积 50%+）— 让数据尽可能靠近计算单元
- 复杂的控制逻辑 — 分支预测、乱序执行（out-of-order execution）、寄存器重命名，让单个线程跑得尽可能快
- 少量高性能核心（4-8 个）— 每个核心都是”重装步兵”
GPU 的设计目标：最大化总吞吐量（throughput）。 完全不同的策略：
- 海量简单 ALU — 数千个小计算单元，每个都很弱，但合在一起吞吐量极高
- 最小化控制逻辑 — 没有复杂的分支预测和乱序执行，控制逻辑占比极小
- 小缓存，靠线程切换隐藏延迟 — 遇到内存等待不是靠缓存命中来加速，而是立即切换到另一组线程继续执行
延迟隐藏：两种截然不同的策略
CPU 遇到 cache miss（数据不在缓存中）时，线程会 stall（阻塞） — 什么都不做，干等数据从内存返回。这段等待时间完全浪费了。CPU 的应对策略是：用更大的缓存减少 miss 的概率。
GPU 的应对策略完全不同：不减少等待，而是在等待期间做别的事。GPU 上同时有数千个线程（分组为 warp），当一组线程等内存时，硬件 warp scheduler 立即切换到另一组已经准备好的线程。只要同时活跃的线程足够多，就能把等待时间完全填满。
这就是为什么 GPU 需要大量线程 — 不是因为有那么多独立的任务要做，而是需要足够多的线程来掩盖内存延迟。这个概念叫 latency hiding，是理解 GPU 架构最核心的思想。
Section 2: NVIDIA GPU 全局结构
以 NVIDIA H100 为例，从芯片级到最小计算单元的层级结构：
GPU 芯片 → GPC (Graphics Processing Cluster) → TPC (Texture Processing Cluster) → SM (Streaming Multiprocessor)
- GPC 是最大的逻辑分组。H100 有 8 个 GPC。
- 每个 GPC 包含若干 TPC（Texture Processing Cluster）。TPC 是历史遗留名称（来自图形渲染），在计算场景下主要是 SM 的容器。
- 每个 TPC 包含 2 个 SM — SM 是 GPU 的核心计算单元，也是我们最需要理解的层级。
- 所有 SM 共享一个大的 L2 Cache（H100 约 50MB）。
- L2 之外是 HBM（High Bandwidth Memory）— GPU 的主显存。
H100 SXM 全芯片： 8 GPC × 9 TPC × 2 SM = 144 SM（实际启用 132 个），50 MB L2，80 GB HBM3。
对比：RTX 4090 有 128 SM，L2 96 MB，24 GB GDDR6X。架构核心相同（都基于 Ada / Hopper SM），但规模和显存技术不同。
Section 3: SM 内部解剖
SM（Streaming Multiprocessor）是 GPU 中真正执行计算的单元。理解 SM 的内部结构，就理解了 GPU 的执行模型。
四个 Processing Block
每个 SM 被分为 4 个 Processing Block（也叫 Sub-partition），每个 Processing Block 有：
- 1 个 Warp Scheduler + 1 个 Dispatch Unit — 选择并发射指令
- FP32 CUDA Core — 做浮点乘加（Hopper 每 block 32 个，整个 SM 128 个）
- INT32 Core — 整数运算
- Tensor Core — 矩阵乘加速单元（下一篇文章详细讲）
- SFU (Special Function Unit) — 计算 sin/cos/exp 等超越函数
- Load/Store Unit — 从内存加载/写入数据
Processing Block 本地资源
每个 Processing Block 还有自己的 register file：
- Register File — 整个 SM 合计 256 KB，物理上分布在各 processing block 中（NVIDIA 架构白皮书的 SM 结构图显示每个 sub-partition 有独立的 register file）。这是 GPU 中最快的存储，每个线程最多使用 255 个 32-bit 寄存器。Warp 的寄存器存储在它所属 processing block 的本地 register file 中，warp 一经分配就不会在 processing block 之间迁移。Register file 很大是因为 GPU 需要同时维持数千个线程的寄存器状态（实现零开销 warp 切换的关键）
SM 级共享资源
4 个 Processing Block 共享：
- Shared Memory / L1 Cache — 256 KB（Hopper）的片上 SRAM，其中最多 228 KB 可配置为 shared memory，其余用作 L1 cache。物理上属于整个 SM，但逻辑上按 thread block 分区 — 同一 thread block 内的所有线程（跨所有 warp、跨所有 processing block）可以访问同一块 shared memory，不同 thread block 之间互不可见。可以配置 shared memory 和 L1 cache 的分配比例
代际演进
| 资源 | Ampere (A100) 2020 | Hopper (H100) 2022 | Blackwell (B200) 2024 |
|---|---|---|---|
| GPU 中 SM 数量 | 108 | 132↑ | 192↑ |
| FP32 Core / SM | 128 | 128 | 128 |
| INT32 Core / SM | 128 | 64↑ | 64 |
| Tensor Core / SM | 4 (3rd gen) | 4 (4th gen)↑ | 4 (5th gen)↑ |
| Tensor Core 精度 | FP16/BF16/TF32/INT8/INT4 | FP16/BF16/TF32/FP8/INT8↑ | FP16/BF16/TF32/FP8/FP4/INT8↑ |
| Register File / SM | 256 KB | 256 KB | 256 KB |
| Shared Memory / SM | 164 KB | 228 KB↑ | 228 KB |
| L1 Cache / SM | 192 KB 共享 | 256 KB 共享↑ | 256 KB 共享 |
| Max Warps / SM | 64 | 64 | 64 |
| Max Threads / SM | 2048 | 2048 | 2048 |
绿色箭头 ↑ 表示相比前一代有变化。Hover 高亮行。
Section 4: Warp — GPU 的最小执行单位
什么是 Warp
Warp = 32 个线程组成的一组，是 GPU 调度和执行的最小单位。Warp 中的 32 个线程在硬件层面锁步执行同一条指令（SIMT — Single Instruction, Multiple Threads）。
这就像一个班的士兵齐步走 — 32 人同时迈左脚、同时迈右脚。他们执行相同的”指令”（步伐），但各自踩在不同的”数据”（地面位置）上。
每个 Processing Block 的 Warp Scheduler 每个周期选择一个 ready warp（不在等待内存的 warp），为它发射一条指令。4 个 Processing Block 可以同时各选一个 warp，所以一个 SM 每周期最多发射 4 条指令。
Warp Divergence
32 个线程必须执行同一条指令 — 那如果代码里有 if/else 呢？
这就是 Warp Divergence：当 warp 中的线程走了不同的分支路径时，硬件无法同时执行两条路径。解决方式是串行化 — 先执行 if 路径（不走 if 的线程被 mask 掉），再执行 else 路径（不走 else 的线程被 mask 掉）。两条路径的时间加在一起。
一个 Warp 包含 32 个线程，执行同一条指令（SIMT）。所有线程同步执行 z[i] = x[i] + y[i]。
这就是为什么 GPU 代码要尽量避免分支 — 不是说不能写 if/else，而是同一个 warp 的 32 个线程最好都走同一条路径。
Warp Scheduler 与延迟隐藏
Warp scheduler 是 SM 的”调度大脑”。它的工作很简单：在 ready warp 中挑一个，发射它的下一条指令。
当一个 warp 执行了内存加载指令后，需要等待数百个周期才能拿到数据。Scheduler 不会空等 — 它立刻切换到另一个 ready warp 继续执行。这个切换是零开销的（zero-overhead context switch），因为所有 warp 的寄存器状态都常驻在 SM 的 register file 中，不需要保存/恢复。
Warp Scheduler 选择 Warp 0 执行计算指令。执行 2 个周期后，遇到全局内存访问 — 需要等待 ~数百个周期。
硬件概念 vs 软件抽象：对应关系
GPU 编程中有很多名字带 “block” 的概念，容易混淆。关键在于区分硬件实体和软件抽象：
| 概念 | 类型 | 说明 |
|---|---|---|
| SM | 硬件 | GPU 的基本计算单元 |
| Processing Block | 硬件 | SM 内部的子分区（各有 warp scheduler、CUDA cores、Tensor Core） |
| Grid | 软件 | 一次 kernel launch 的所有线程 |
| Thread Block | 软件 | 程序员定义的一组线程（如 256 个），共享 shared memory，可以用 __syncthreads() 同步 |
| Warp | 桥梁 | Thread block 内每 32 个连续线程自动编为一组，由硬件锁步调度 |
映射关系：
- Thread Block → SM：一个 thread block 整体分配到一个 SM，共享该 SM 的 shared memory。一个 SM 可以同时运行多个 thread block（取决于寄存器和 shared memory 资源是否充足，即 occupancy）
- Warp → Processing Block：thread block 内的线程按 32 个一组自动分成 warp，各 warp 被分配到 SM 内某个 processing block 的 warp scheduler。warp 一经分配就固定在该 processing block 上，不会迁移 — 因为 warp 的寄存器状态存储在该 processing block 的本地 register file 中，迁移意味着搬运大量寄存器数据（一个 warp 最多 32 × 255 × 4 = ~32 KB），开销巨大。而同一 processing block 内的 warp 切换是零开销的 — scheduler 只是切换到另一个 warp 的寄存器区域，无需保存/恢复。这与 CPU 不同，CPU 线程可以在核心间迁移（OS 做 load balancing），而 GPU 的 warp 靠静态分布和大量并行来保证负载均衡
- 寄存器：每个线程私有，同一 warp 内的线程不能直接访问彼此的寄存器（但可以通过
__shfl_sync等 warp shuffle 指令交换值）
注意 “Processing Block” 和 “Thread Block” 完全不同 — 前者是 SM 内部的硬件分区，后者是程序员定义的软件分组。名字相似但没有一一对应关系。
为什么需要 warp 和 thread block 两层抽象？
- Warp（32 线程）解决执行效率：SIMT 模型下，一个 warp 内 32 个线程共享一套取指/译码逻辑，硬件统一调度发射指令，但每个线程有独立的寄存器和 PC。这样只需 1 份控制逻辑就能驱动 32 个 ALU，大幅节省晶体管面积。一个 256 线程的 thread block 如果不分 warp，就需要 256 套独立的取指/译码单元 — 这是不现实的
- Thread Block（如 256 线程）解决协作需求：很多算法需要超过 32 个线程共享数据（例如 GEMM tiling 需要大量线程一起搬数据到 shared memory）。Thread block 提供了 shared memory 和
__syncthreads()同步机制，让多个 warp 可以协作
简言之：thread block 管合作，warp 管执行。
Intel GPU 的不同之处
以上是 NVIDIA 的 SIMT 模型。Intel Xe2 GPU 采用了根本不同的方式 — 编译器驱动的 SIMD。对应概念的对比：
| NVIDIA (SIMT) | Intel Xe2 (SIMD) | 关键差异 |
|---|---|---|
| Thread Block | Work-group | 类似：共享本地内存，可同步 |
| Warp (32 线程) | Sub-group (8/16/32 lanes) | NVIDIA 由硬件在运行时编组；Intel 由编译器在编译时向量化 |
| SM | Xe-core | 类似：包含多个执行单元和共享内存 |
| Processing Block | EU (Execution Unit) | NVIDIA 有 warp scheduler；Intel EU 直接执行 SIMD 指令 |
| Shared Memory | SLM (Shared Local Memory) | 类似：work-group 内共享 |
| 分支处理 | 分支处理 | NVIDIA：运行时硬件 masking；Intel：编译时 predication（编译器将 if/else 转为 predicated 指令，所有 lane 都执行但由 mask 控制写回） |
核心区别：NVIDIA 程序员写标量代码，硬件自动将 32 个线程编为 warp 锁步执行，分支时动态 mask。Intel 程序员同样写标量代码，但编译器负责将其向量化为 SIMD 指令，分支在编译时转为 predicated 执行。NVIDIA 的灵活性更高（硬件处理 divergence），Intel 的方式对编译器依赖更大但硬件更简单。
详细的 Intel Xe2 执行模型、Sub-group 操作和 SYCL 编程映射，参见 Xe2 执行模型与编程抽象。
Section 5: 内存层次
GPU 的内存层次是影响性能最大的因素。从快到慢：
| 层级 | 容量 (H100) | 带宽 | 延迟 | 作用域 |
|---|---|---|---|---|
| Register File | 256 KB / SM | 极高（片上） | 0 周期 | 线程私有 |
| Shared Memory / L1 | 256 KB / SM（最多 228 KB 可配置为 shared memory） | 极高（片上） | ~20-30 周期 | Thread Block 内共享 |
| L2 Cache | 50 MB | ~12 TB/s（理论计算值） | ~200 周期 | 全局共享 |
| HBM3 | 80 GB | 3.35 TB/s | ~400-600 周期 | 全局 |
关键数字：从 register 到 HBM 的延迟差距是 数百倍，带宽差距是 5-6 倍。这意味着：
- 让数据尽可能停留在 register 和 shared memory 是优化的核心
- 从 HBM 读一次数据很贵 — 读来之后应该尽可能多次复用
- 这就是 tiling（分块）策略的本质动机，我们将在 GEMM 优化文章中详细展开
Flash Attention 的核心创新正是基于这个内存层次 — 把 Attention 计算分块到 SRAM（shared memory）中完成，避免写回巨大的中间矩阵到 HBM。详见 Flash Attention 分块原理。
总结
GPU 架构的核心设计思想可以归纳为三点：
- 吞吐优先 — 数千个简单核心，牺牲单线程延迟换取总吞吐量
- 延迟隐藏 — 不减少等待时间，而是用大量线程（warp）填满等待。这要求 register file 常驻所有 warp 状态
- 内存层次 — 从 register 到 HBM 延迟差数百倍，优化 = 尽量让数据留在快速存储中
下一篇文章将深入 GPU 中最重要的专用加速单元 — Tensor Core 与 XMX，理解它们为什么能让矩阵乘法快一个数量级。