GPU 架构深入
更新时间：2026-10-01。GPU 和 CPU 面对的是两类完全不同的问题：CPU 用深流水线、乱序执行和分支预测把单线程延迟压到最低；GPU 反其道而行，用成千上万个简单线程和零成本的 warp 切换把长延迟「藏」起来，换极限吞吐。本篇按标准篇深度，从 SIMT 执行模型一路讲到 warp scheduler 的就绪判定、寄存器堆、分支分歧与 occupancy 权衡，帮你把 CUDA 层面的每一条优化经验都对回硬件结构。线程层次的软件视角可先看 CUDA 线程层次：grid 与 block，SM 内部结构可对照 SM 与 CU 的内部结构。
本文要回答的问题
- grid、thread block、warp、thread 四级层次里，硬件真正的调度单位是哪一级？为什么是 32 而不是 1？
- warp scheduler 怎么判定一个 warp 能不能发射？为什么说 GPU 的上下文切换「零成本」，而 CPU 做不到？
- 双发射（dual issue）和超标量在 GPU 里长什么样，和 CPU 的多发射有何不同？
- warp 内部分支走不同路径（divergence）时硬件如何序列化、又如何在汇合点 reconvergence？
- occupancy 是不是越高越好？寄存器、shared memory 怎样反过来限制驻留 warp 数？
一、SIMT：单指令多线程，以及 grid/block/warp 三级层次
GPU 编程模型给用户看到的是海量线程：一次 kernel launch 产生一个 grid，grid 由若干 thread block（线程块）组成，block 内线程数通常取 128 或 256。但这些线程不是独立调度的。硬件把 block 内的线程按连续编号每 32 个打包成一个 warp（NVIDIA 的叫法；AMD 称 wavefront，宽度为 64），warp 才是 SM（Streaming Multiprocessor，流式多处理器）上的最小调度单位。同一 warp 的 32 个线程共享同一个程序计数器（PC），锁步执行同一条指令，只是各自的数据和 lane（通道）不同——这就是 SIMT（Single Instruction, Multiple Threads，单指令多线程）。
这张层次图里最需要记住的是两个边界。block 是软件调度边界：整个 block 被分配到一个 SM，执行期间不迁移，所以 block 内的线程才能用 shared memory 和 __syncthreads() 同步；不同 block 之间则不能直接同步。warp 是硬件执行边界：32 个 lane 锁步，这一数字带来后面所有的现象——分歧要序列化、访存要按 32 个 lane 的地址合并、occupancy 按 warp 计数。软件层面的分歧机制详见 SIMT 与 warp 分支分歧。
为什么 warp 宽度恰好取 32？这是面积、能效和编程抽象三者折中的结果。加宽到 64，SIMD 数据通路和寄存器端口更宽、摊薄了取指与控制的开销，但分支分歧时被掩码空转的 lane 更多，对非规整负载更不友好；缩到 8 或 16，控制逻辑和指令缓存的占比上升，合并访存也凑不满一个内存事务。32 同时是一条 128 字节缓存线按 4 字节元素切分的数量，使「一个 warp 的连续访存」与「一次缓存线事务」天然对齐——这个数字不是任意选的，而是和内存事务粒度一起协同设计的。理解这一点也能避免两个常见误用：block 内线程数一定要取 32 的整数倍（128、256 最常见），否则最后一个 warp 里不足 32 个有效 lane，硬件仍按整 warp 发射，尾部算力白白浪费；一个 warp 也永远不会跨 block 切分，block 边界处的浪费只能靠选择合适的 block 尺寸来消除。
二、warp scheduler：记分牌、就绪判定与双发射
一个现代数据中心 GPU 的 SM 上常驻数十个 warp（上限 64 个，即 2048 个线程），每个周期由一个或多个 warp scheduler 决定发射谁。调度器不靠复杂的分支预测和乱序窗口，而是维护一张记分牌（scoreboard）：为每个 warp 记录它下一条指令的操作数依赖——典型实现是一个按 warp 索引的位表，标记「某 warp 是否有尚未返回的长延迟操作」。一条指令发射后，若它是全局内存 load（数百周期）、特殊函数（SFU，若干到数十周期）等操作，对应 warp 被标记为未就绪；操作返回（数据写回寄存器堆）时清除标记。每拍调度器在所有就绪 warp 中按策略（通常是轮转加少量公平性约束，思路与 round-robin 仲裁器 同源）挑一个发射。
这张图要读出三个要点。其一，判定是粗粒度的：只区分「就绪/未就绪」加有限的操作数依赖，不像 CPU 那样在数百条指令窗口内做寄存器重命名和全乱序调度，因此调度器本身面积和功耗都很小。其二，现代 SM 有 2 到 4 个调度器分区，每个分区每周期最多向执行端口双发射两条相互独立的指令（典型为一条 FP/INT 运算加一条 load/store 或地址生成指令）；这和 CPU 超标量的区别在于：发射的两条指令仍来自同一 warp 的静态顺序、锁步执行于 32 个 lane，没有跨 warp 的动态重排序。其三，执行端口是异构分桶的——FP32、INT32、LDST、SFU、Tensor Core 各自独立，一个 warp 也可能因为目标端口忙而暂时不可发射。
三、延迟隐藏：用 warp 切换代替乱序窗口
GPU 的核心魔法是上下文切换成本为零。当被调度的 warp 发出一条全局 load（典型延迟 300 到 800 个周期），调度器下一拍直接切到另一个就绪 warp，连寄存器都不用搬——因为所有驻留 warp 的寄存器本来就同时躺在物理寄存器堆里，warp 的「上下文」只是一个 PC 和它在 scoreboard 里的状态位。于是只要就绪 warp 足够多，执行单元在等待期间始终有活干，长延迟对吞吐不可见。
经验法则是：藏住一次 400 拍左右的 DRAM 延迟，大约需要 10 到 14 个交替发射的 warp 在飞行中；具体数字随每指令吞吐和代际变化，但量级如此。把这套机制和 CPU 对比，差异就一目了然。
| 维度 | CPU（乱序核） | GPU（SIMT + warp 调度） |
|---|---|---|
| 应对长延迟 | 大 ROB、寄存器重命名、MSHR、推测执行 | 切到另一个就绪 warp |
| 硬件线程上下文 | 每核 1-2 个 SMT 线程，切换需刷流水线 | 每 SM 数十个 warp，切换零周期 |
| 同时在飞指令 | 单线程内 200-400 条 | 数十 warp × 少量指令 |
| 主频 / 每核功耗 | 3-4 GHz，重 | 约 1-2 GHz，核数多但每核简单 |
| 优化目标 | 单线程延迟 | 聚合吞吐 |
| 依赖的并行类型 | ILP（指令级并行） | TLP/WLP（线程/warp 级并行） |
CPU 的乱序机制可回看 CPU 乱序执行 与 重命名与提交。一句话：CPU 花钱深挖一个线程里的 ILP，GPU 假设程序员能提供海量 TLP，用简单的在途 warp 数替代复杂的乱序窗口。GPU 侧更细的调度与抢占机制见 GPU 调度与抢占。
用一个简化模型可以把这个差异量化。假设每次内存访问延迟 400 拍，CPU 乱序窗口能容纳 200 条独立指令、且程序里真有这么多独立指令，它可以让约 200 次访问在飞而线程不切换；但多数真实程序的 ILP 在 2 到 4 之间，窗口再大也填不满，于是 CPU 还要叠加预取、多级缓存和 SMT。GPU 不假设单线程内有并行，它要求一次 launch 就给每个 SM 送来上千个线程：只要平均每十几拍能找到一个就绪 warp，400 拍的空洞就被填满，算术上需要的在途 warp 数约等于「延迟除以每次切换获得的有效执行拍数」。这也解释了为什么同样的代码在 GPU 上「核数没少、占用率却低」时性能断崖——warp 不够，调度器面前就是一片就绪荒漠，再宽的 FP32 阵列也要空转。
四、寄存器堆、bank 与 occupancy：资源约束的算术
零成本切换有个硬前提：所有驻留 warp 的寄存器必须物理在场。以数据中心级 GPU 为例，每个 SM 的寄存器堆约有 65536 个 32 位寄存器（256 KB），被划成多个 bank 供每拍多 lane、多指令并行读写下；每个线程最多可使用 255 个寄存器。每个 warp 独占 32 × 每线程寄存器数 的物理空间。再加上每 SM 最多驻留 64 warp（2048 线程）、block 数和 shared memory 容量三条上限，实际驻留 warp 数取四个约束里的最小值，这就是 occupancy（占用率）。
算一笔账：kernel 每线程用 64 个寄存器，则一个 warp 占 2048 个，寄存器维度最多放 32 warp；若每线程用 128 个，只剩 16 warp，occupancy 直接腰斩。若编译器再挤出寄存器到 local memory（spilling），性能损失按倍数计。寄存器与 occupancy 的工程取舍详见 occupancy 与寄存器压力。
| 每线程寄存器 | 寄存器允许的 warp 数 | 典型 occupancy |
|---|---|---|
| 32 | 64 | 100% |
| 64 | 32 | 50% |
| 128 | 16 | 25% |
| 255 | 约 8 | 约 12% |
注意高 occupancy 不是目的本身：把寄存器压得过少可能逼出 spilling，而一个计算密集、无长延迟停顿的 kernel 在 50% occupancy 下同样能跑满。正确姿势是先保证足够藏住访存延迟的 warp 数，剩余资源向每个线程多留寄存器、向 shared memory tiling 倾斜。
五、分支分歧与 reconvergence stack
锁步执行在 if (threadIdx.x % 2 == 0) 这类条件面前立刻出问题：同一 warp 的 32 个 lane 条件结果不同，硬件无法在一条指令里走两条路。GPU 的做法是谓词化 + 序列化：把两条路径都取到，用活跃掩码（active mask，一个 32 bit 向量）标明当前哪些 lane 真正执行——走 then 分支时奇数 lane 被掩码掉、空转不写回，走 else 分支时反过来，最后在汇合点（post-dominator）重新合流。掩码的嵌套管理由硬件的 reconvergence stack（汇合栈） 完成：每进一层分叉，把分支地址和对应 mask 压栈，逐层弹栈直到所有 lane 重新汇合。
这张图解释了为什么 warp 内分歧的代价是「路径数 × 执行量」：两条路径序列化，理论吞吐减半；三路 switch 最坏只剩三分之一。还要区分两种代价完全不同的分歧：warp 内分歧（同一 32 线程里条件不一致）是真分歧，必然序列化；而 warp 间分歧（整个 warp 一致地走 then、另一个 warp 走 else）不是分歧，不同 warp 在不同 SM 或不同拍上各走各的，硬件层面没有任何惩罚。这就是经典优化「warp 对齐」的含义——把数据重排，让条件在连续 32 个线程内保持一致，把分支推到 warp 边界之外。工程上的对策都是让条件按 warp 对齐——按连续线程分块处理数据、用循环而不是发散的 switch、或直接走谓词执行让两边都算（当两边都很短时反而比分支便宜）。值得注意的是近年 NVIDIA 架构引入了独立线程调度（线程可在 warp 内各自推进），收敛点不再完全由编译器静态保证，但「锁步 + 活跃掩码」仍是理解性能行为的正确一阶模型。
六、访存合并与 shared memory bank：带宽侧的呼应
warp 执行一条 load 时，内存子系统看到的是 32 个 lane 的地址。若这些地址落在连续的 128 字节区间内，硬件把它们**合并（coalesce）**成一次内存事务；地址散开则退化成多次事务，有效带宽可掉到十几分之一。对矩阵、图像这类规整数据，合并访存是所有优化的地基，详见 内存合并访问。
on-chip 的 shared memory 则被切成 32 个 bank，每个 bank 每拍服务一个 4 字节访问；同 warp 中多个 lane 命中同一 bank 的不同地址时发生 bank conflict，访问被拆成多拍序列化，详见 shared memory 与 bank conflict。巧合的是，32 个 bank 正好对齐 32 个 lane：错开步长（padding）让各 lane 落不同 bank，就能每拍无冲突服务整个 warp。完整内存层次与 HBM 带宽见 GPU 内存层次与带宽 与 HBM 接口物理设计，整体性能用 GPU roofline 模型 判定是算力受限还是带宽受限；矩阵类负载还要叠加 Tensor Core 矩阵加速 这条专用数据通路，它和脉动阵列的关系可对照 AI 加速器与脉动阵列。
常见坑
- 把 occupancy 当唯一指标：盲目压寄存器导致 spilling，或为凑满 occupancy 牺牲 tiling，实际更慢。
- 非合并访存：warp 内地址跨步或随机，带宽利用率可掉一个数量级。
- 忽略 warp 级分歧：warp 内条件不统一使路径序列化，跨 warp 的不均匀负载还会造成 SM 间长尾。
- bank conflict 没算步长：shared memory 行宽与线程映射错配，每拍只服务一个 lane。
- 以为双发射等于乱序：双发射的两条指令仍按静态顺序、须相互独立，编译器排不好就发射不出去。
- 寄存器用量只看单个数：还要乘以 block 内 warp 数核对 SM 物理上限，launch 才不会静默降 occupancy。
这一篇在知识体系里的位置
- 软件视角承接：CUDA 线程层次：grid 与 block 定义 grid/block，SIMT 与 warp 分支分歧 展开分歧，本篇补上调度器与资源的硬件机理。
- 资源与带宽：occupancy 与寄存器压力、shared memory 与 bank conflict、内存合并访问 是本篇三个硬件机制的优化落点。
- 横向对比：CPU 乱序执行 给出另一条延迟处理路线，AI 加速器与脉动阵列 展示把控制硬件省到极致的专用路线。
- 向后衔接：估算与 RTL 性能建模 把这些结构参数变成早期性能/功耗模型输入。
总结
GPU SIMT 架构以 warp（32 个锁步线程）为最小调度单位，warp scheduler 靠一张轻量记分牌在就绪 warp 间轮转、在长延迟操作下零周期切换，用驻留 warp 数（occupancy）替代 CPU 的乱序窗口来隐藏访存延迟；而 occupancy 被寄存器、shared memory、block 与 warp 四条资源上限共同卡死，分支分歧则用活跃掩码与 reconvergence stack 把不一致路径序列化——合并访存、消除 bank conflict、管理寄存器压力，本质上都是在给这套「简单核 × 海量 warp」的机器喂够并行、排空停顿。