顺序执行模型与处理器内部并行性#

CPU ISA 暴露给程序员的是一台按程序顺序执行指令的机器：

1
x = x + 1;
2
y = y + 2;
3
z = x + y;

从语言和 ISA 语义看，上述语句可以映射成某种顺序指令流；但在硬件里，只要依赖关系允许，处理器就会尝试把多条指令同时推进：

1
fetch multiple instructions
2
decode multiple instructions
3
rename registers
4
check data dependence
5
issue independent instructions together
6
execute out of order
7
retire in order

这就是 Instruction-Level Parallelism。逻辑门天生并行，所以“一个周期内只有一条指令活动”本来就不是硬件的自然状态；真正难的是在并行执行之后仍然维持顺序语义。乱序执行、寄存器重命名、按序提交本质上都在做同一件事：

1
内部尽量并行，外部看起来仍然像顺序机。

因此第 13 讲里 sum++ 之所以会暴露并发问题，并不只是“你开了多个线程”；即使单核 CPU 也可能在流水线、cache、一致性层面并行推进大量事件。现代系统的默认背景噪音本来就是并行。

指令级并行的能力边界#

ILP 很强，但它依赖处理器自动在有限指令窗口里挖掘独立性。最典型的对比是：

1
// 强依赖链
2
for (...) x = x * 3 + 1;
3

4
// 独立指令
5
for (...) {
6
    x0 = x0 * 3 + 1;
7
    x1 = x1 * 3 + 1;
8
    x2 = x2 * 3 + 1;
9
    x3 = x3 * 3 + 1;
10
}

前者每一步都依赖上一步结果，后者给了 CPU 更多可同时发射的独立操作。讲义里 IPC 实测的核心结论正是：

1
依赖链限制吞吐，独立指令让 CPU 的多发射能力显现出来。

但 ILP 的问题是：

• 能发现的独立性受限于局部指令窗口。
• 前后依赖复杂时，乱序引擎无能为力。
• 它优化的是“同一线程内部几条指令并行”，不是“同一操作对很多数据并行”。

当程序形态变成：

1
for (int i = 0; i < n; i++) c[i] = a[i] + b[i];

瓶颈就不再是“几条标量指令如何乱序”，而是“完全同构的计算重复了 n 次”。这时更自然的硬件抽象就是数据级并行。

SIMD 的基本抽象与硬件结构#

SIMD = Single Instruction, Multiple Data。它表达的是：

1
同一条指令
2
对多个独立数据
3
做同一种操作

例如把 128-bit 寄存器看成 4 个 int32_t 槽位：

1
[a0 a1 a2 a3] + [b0 b1 b2 b3]
2
-> [a0+b0 a1+b1 a2+b2 a3+b3]

这不是编译器把一条向量指令偷偷翻译成 4 次普通加法循环执行，而是硬件真的提供了：

• 更宽的寄存器；
• 更宽的数据通路；
• 多个 lane 的并行算术单元；
• 共享的取指、译码、发射与提交逻辑。

因此 SIMD 的关键不在“它也是电路实现”，因为标量 ALU 当然也是电路实现；真正关键的是：

1
SIMD 通过共享一套控制流成本，驱动多条 lane 同时运算。

SIMD 指令的成本摊薄机制#

把一条指令的成本拆开来看：

1
总成本 = 前端控制成本 + 后端执行成本

前端控制成本包括：

• 取指；
• 译码；
• 寄存器重命名；
• 发射；
• 提交。

这些环节更接近“按指令条数收费”。因此：

• 标量加法：1 条指令处理 1 个元素；
• 向量加法：1 条指令处理 4/8/16 个元素；
• 前端仍然大致只处理 1 条指令。

后端执行当然不是零代价。SIMD 需要更宽的寄存器文件、更宽的旁路网络、更高功耗和更难的时序设计。但它不是每多一个元素就复制一整套 CPU 控制逻辑，而是：

1
共享控制，扩宽 datapath。

这就是为什么 SIMD 常常看起来“几乎白送”。它并不免费，但增加的主要是数据通路宽度，不是每个元素都重复一次完整的 fetch/decode/commit 流程。

SIMD 的性能收益来源与适用条件#

SIMD 的收益并不只是“每条指令算得更多”，还来自：

• 更少的动态指令数；
• 更好的 fetch/decode 摊薄；
• 更少的循环控制开销；
• 更高的单位访存有效载荷；
• 更容易把连续数据组织成顺序访问。

但这也意味着 SIMD 并非无条件有效。常见限制包括：

• 数据必须有足够规则的布局；
• lane 内最好做相同运算；
• 分支、shuffle、mask、gather/scatter 会抬高代价；
• 宽向量指令可能受限于频率、端口或 load/store 带宽。

所以真正的优化问题不是“有没有 SIMD 指令”，而是：

1
能不能把算法重排成硬件喜欢的、规则的、同构的、连续的向量工作流。

位并行技术与 Subword Parallelism#

讲义里特别重要的一点是：很多人第一次接触 SIMD，以为它是某个神秘的新概念；其实你很可能早就在用它的思想。

例如 bitset：

1
arr[x / 64] >> (x % 64)

它的底层并不是“天然就是 SIMD 指令”，而是：

1
先用整数压位；
2
再用一个机器字的各个 bit 表示一组 0/1 状态。

也就是说，一个 uint64_t 可以看成 64 个布尔位：

• bit = 1：元素在集合中；
• bit = 0：元素不在集合中。

于是集合运算变成位运算：

1
A | B   // 并集
2
A & B   // 交集
3
A ^ B   // 对称差

这本质上是：

1
一次整数指令
2
同时处理 64 个 0/1 状态

这叫 subword parallelism，也常被理解为 SWAR（SIMD Within A Register）风格的技巧。它和真正硬件 SIMD 的关系是：

• bitset 先利用机器字级并行；
• 编译器/库实现还可能进一步把多个机器字一起向量化；
• 但 bitset 本身不等于 SIMD 指令集。

因此，“bitset 的底层是不是 SIMD”这个问题，更精确的回答是：

1
bitset 的本质是按机器字压位并做位并行；
2
SIMD 是它可能进一步吃到的一层硬件加速。

也正因为如此，bitset 在算法题里经常非常快，即使编译器根本没有显式生成 AVX 指令：光是把 64 个布尔状态压进一个 word、减少内存流量和循环控制，就已经赚到了。

Popcount 的分层并行累加模型#

经典 popcount 写法：

1
x = (x & 0x55555555) + ((x >> 1) & 0x55555555);
2
x = (x & 0x33333333) + ((x >> 2) & 0x33333333);
3
x = (x + (x >> 4)) & 0x0f0f0f0f;

乍看像在做一堆位运算，实际上它在同一个整数里分层合并多个子字段的部分和。这个思路和 SIMD 的一致性很强：

1
把大寄存器切成很多小槽；
2
每一轮让这些槽同步做同构运算；
3
最后汇总结果。

因此讲义把 bitset、popcount、Bloom filter 放在一脉里非常准确：它们都是“在一个字里做很多份同构计算”的技术，只是历史上软件位技巧先出现，后来硬件把这种模式做成了专门的向量 ISA。

SIMD 的程序执行模型#

理解 SIMD 时要抓住一个不变量：

1
SIMD 仍然是 CPU 模型。

也就是说：

• 仍然是一个线程的控制流；
• 仍然是一个 PC；
• 仍然共享同一份地址空间、cache 和同步原语；
• 只是某些寄存器和 ALU 被加宽了。

所以 SIMD 的程序员心智模型更接近：

1
把“for 循环的一部分迭代”
2
压缩进一条向量指令

这和 GPU 的线程模型很不一样。GPU 不是“一个线程里有很多 lane”，而是“很多线程共享一份控制流硬件”。

GPU 的历史动机与通用化转折#

GPU 的历史背景确实来自图形学，但这里不需要记太多渲染细节，只要抓住因果链即可：

1
画面越来越复杂
2
-> 固定图形流水线吞吐不够
3
-> 厂商开始堆大量面向图形任务的专用并行硬件
4
-> shader 让这套硬件变得可编程
5
-> 最终走向 GPGPU 和 CUDA

早期游戏机的很多图形技巧依赖固定 tile、sprite 和特定硬件寄存器；后来的 3D 图形需求迫使 GPU 从“固定功能单元”走向“可编程 shader”。这一步的意义在于：

1
GPU 不再只是会画图的专用芯片，
2
而是变成会并行执行用户代码的大规模吞吐机器。

图形学部分在本讲里更像历史动机，而不是后面 CUDA 模型的核心难点，因此知道“shader 是 GPU 通用化的拐点”就够了。

CPU 与 GPU 的设计目标差异#

CPU 的目标更偏向：

1
尽量加速单条复杂控制流

因此 CPU 核心里堆了很多昂贵能力：

• 复杂分支预测；
• 乱序执行；
• 大 cache；
• 寄存器重命名；
• 投机执行；
• 低延迟单线程优化。

GPU 的目标更偏向：

1
在单位面积和单位功耗下
2
把吞吐量做大到极致

因此 GPU 更愿意把晶体管预算花在：

• 更多 ALU；
• 更多寄存器；
• 更多可驻留线程上下文；
• 更高带宽利用能力；
• 更大规模的同构并行。

换句话说：

1
CPU 强在“聪明地处理复杂控制流”；
2
GPU 强在“宽而便宜地推动大量统一控制流”。

SIMT 的基本抽象#

GPU 的关键抽象不是 SIMD，而是 SIMT：Single Instruction, Multiple Threads。

它看起来和 SIMD 只差一个字母，但程序模型差别很大：

• SIMD：程序员在一个线程里显式打包多个数据；
• SIMT：程序员写很多“普通线程”的标量代码，硬件把同时执行同一指令的线程绑成 warp。

例如 CUDA kernel：

1
__global__ void vec_add(float *a, float *b, float *c, int n) {
2
    int i = blockIdx.x * blockDim.x + threadIdx.x;
3
    if (i < n) c[i] = a[i] + b[i];
4
}

从源码上看，每个线程只做一件很普通的事：算自己的 i，然后处理 c[i]。但在硬件里，很多线程会被按 warp 分组，同时执行同一条指令。其关键不变量是：

1
warp 内线程各有自己的寄存器状态和数据，
2
但共享取指与控制流推进。

这就是“一个 PC 管多个线程”的含义。

SIMT 的面积效率来源#

如果做 32 个完全独立的小 CPU 核，就意味着 32 套：

• fetch；
• decode；
• branch prediction；
• issue；
• control logic。

这部分控制逻辑很贵，占 die area 很大。SIMT 的思路是：

1
让很多线程共享同一套控制逻辑；
2
每个线程主要保留自己的寄存器上下文和运算通路。

于是相同芯片面积可以容纳远多于 CPU 的算术吞吐能力。这也是为什么 GPU 的“核心数”能远高于 CPU，但这些核心不能按 CPU 的方式理解：它们不是很多个全功能乱序通用核，而是大量受共享控制流组织的吞吐单元。

SIMT 与 SIMD 的程序模型差异#

两者都共享一条指令流，但“共享对象”不同。

SIMD：

1
一个线程
2
一个 PC
3
多个 lane
4
程序员显式管理打包/拆包

SIMT：

1
多个线程
2
一个 warp 级控制流
3
每个线程像写普通标量程序
4
硬件负责把线程捆成并行执行组

所以如果用一句话区分：

1
SIMD 是“一个线程里的数据并行”；
2
SIMT 是“很多线程上的统一控制流并行”。

这也是为什么 CUDA 程序表面看起来像“写了一百万个普通线程”，而不是显式写 _mm256_add_ps 这类向量 intrinsics。

Warp Divergence 的形成机制与性能后果#

这一点和我们刚才的提问直接相关。GPU 更怕 divergence，不是因为“GPU 没有分支能力”，而是因为它的吞吐模型默认：

1
同一 warp 内线程尽量走相同控制流。

例如：

1
if (tid % 2 == 0) A();
2
else B();

如果同一个 warp 里一半线程走 A，另一半线程走 `B“，硬件通常不能像 CPU 那样为每个线程都独立维护一套高性能分支预测与乱序引擎。它更常见的做法是：

1
先执行 A 路径，屏蔽不用跑 A 的线程；
2
再执行 B 路径，屏蔽不用跑 B 的线程；
3
最后重新汇合。

这就是 warp divergence。它的性能问题不在于算错，而在于：

1
本来一个 warp 的很多线程可以一起推进；
2
分支分家后，部分线程会在某个路径上空等；
3
统一控制流的吞吐优势被浪费。

因此 GPU 与 CPU 的恐惧对象并不相同：

• CPU 更怕 branch misprediction；
• GPU 更怕 warp 内控制流不一致。

这是因为 CPU 的设计哲学是“猜对分支、继续冲”，GPU 的设计哲学是“大家最好别分家，一分家就分阶段串着跑”。

GPU 控制流开销的精确表述#

把它说成“GPU 控制电路少，所以怕 divergence”方向没错，但更精确的版本应该是：

1
GPU 把更多硬件预算投到大规模并行数据通路和线程驻留能力上，
2
而不是为每个执行流都配一套昂贵的复杂控制逻辑；
3
因此它要求同一 warp 内线程尽量共享控制流，
4
一旦控制流分化，就会显著损失利用率。

这比简单说“GPU 核更弱”更准确。GPU 的强弱不是抽象意义上的“弱”，而是：

1
它不为复杂、不规则、强分支的单线程工作负载做豪华配置，
2
而是为规则、大规模、同构的吞吐任务做极致优化。

Memory Coalescing 与访存规则性#

除了 divergence，GPU 还特别在意同一 warp 的访存是否规则。理想情况是：

1
thread 0 访问 addr + 0
2
thread 1 访问 addr + 4
3
thread 2 访问 addr + 8
4
...

这种相邻线程访问相邻地址的模式容易被合并成较少的内存事务，这就是 memory coalescing。它和 SIMD/bitset 那种“把数据排整齐再一起处理”的思想完全同构：

1
规则布局
2
-> 更少访存事务
3
-> 更高带宽利用率

因此 GPU 编程里经常要同时照顾两件事：

• warp 内线程尽量走同样分支；
• warp 内线程尽量访问连续地址。

前者避免控制流利用率塌陷，后者避免带宽利用率塌陷。

CUDA 的程序抽象与硬件映射#

CUDA 最成功的地方，在于它没有强迫程序员直接操作向量寄存器，而是暴露：

• grid；
• block；
• thread；
• shared memory；
• barrier；
• kernel launch。

这让程序员可以先用“线程索引 -> 数据索引”的方式表达并行计算，再让硬件把这些线程编组成 warp 执行。因此 CUDA 程序设计常见的第一原则是：

1
让相邻线程做相邻工作；
2
让 block 内协作局部数据；
3
让分支尽量按 warp 对齐。

这和第 18 讲的“把共享热点移出高频路径，把同步边压到计算图边界”其实是同一个思路在硬件层面的映射：共享状态少、局部性强、边界同步清晰，吞吐才高。

从 ILP 到 SIMD 再到 SIMT 的并行层次#

把这几种并行机制放到一张图里看，会更容易把握：

1
ILP:
2
  硬件在单线程指令流里自动找局部独立性
3

4
SIMD:
5
  程序/编译器显式说明“这些数据可以做同构并行”
6

7
SIMT/GPU:
8
  程序显式暴露海量线程，硬件按 warp 组织统一控制流并行

因此它们不是替代关系，而是层次关系：

• ILP 负责榨干单线程局部独立性；
• SIMD 负责榨干规则的数据级并行；
• GPU/SIMT 负责把规则的大规模任务映射到极端吞吐硬件上。

同一个应用完全可能同时使用这三层能力。

与课程前序内容的联系#

这一讲和前面课程并不是割裂的。

与第 13-17 讲的同步问题相连：

1
sum++ 为什么不对？
2
因为“顺序执行”从来不该被当成物理事实。

与第 18 讲的并行算法相连：

1
计算图切得再好，
2
最终仍然要落到具体硬件如何执行节点。
3
SIMD 适合规则密集向量计算，
4
GPU 适合海量独立但同构的任务。

与第 19 讲协程相连：

1
协程是在软件中复用执行流和调度点；
2
warp 是在硬件中复用控制流。

它们当然不是同一层抽象，但都体现了一个共同主题：

1
不要把“一个执行流 = 一整套昂贵控制状态”当成唯一实现方式。

面向 AI 基础设施的迁移理解#

这讲对 AI infrastructure 特别重要，因为今天的大模型训练与推理，本质上就是在吃这些抽象的红利。

矩阵乘、卷积、attention、layernorm、softmax 等 kernel，几乎都要求：

• 大量同构计算；
• 规则的数据布局；
• 高带宽利用；
• 尽量少的分支分化；
• 尽量强的局部性和 tile 复用。

所以 GPU 优化和第 18 讲并行算法的桥梁非常直接：

1
tile = 更大的本地计算块
2
shared memory = 局部复用缓存
3
warp-coherent control flow = 更少分支损失
4
coalesced memory = 更少通信事务

如果把视野再放大一点，batching、kernel fusion、operator scheduling、tensor layout、sharding，本质上都在回答同一个问题：

1
怎样把工作重排成硬件喜欢的统一、规则、局部的并行形式？

分析与优化检查清单#

遇到 SIMD / GPU 类优化时，可以按下面顺序问自己：

1
1. 这段计算是控制流密集，还是数据并行密集？
2
2. 依赖关系限制在单线程内部，还是很多独立元素之间？
3
3. 数据布局是否连续、规则、易于打包？
4
4. 这件事更适合 ILP、SIMD，还是 GPU/SIMT？
5
5. warp / lane 内是否会出现严重分支分化？
6
6. 访存是否连续，能否 coalescing / vector load？
7
7. 优化后的收益来自减少控制成本，还是减少访存成本，还是两者都有？
8
8. 是否只是算得更宽了，但实际上被 memory-bound 限住？

真正危险的误区，是把“并行”理解成抽象口号，而不问：

1
谁在共享控制流？
2
谁在承担分支成本？
3
谁在承担访存成本？

本节小结#

1
1. CPU 的顺序执行只是 ISA 层面的语义抽象，现代处理器内部长期依赖 ILP 并行推进多条指令。
2
2. ILP 是硬件自动发掘的隐式并行，但受限于局部依赖窗口，不能替代显式的数据并行模型。
3
3. SIMD 的本质是共享一套控制流成本，驱动多个 lane 对独立数据执行同构操作。
4
4. 一条 SIMD 指令并非“零代价”，但其前端成本按指令条数摊薄，因此不会随元素个数线性增长。
5
5. bitset、popcount、Bloom filter 等本质上是 subword parallelism；bitset 先是按机器字压位，并不天然等于 SIMD。
6
6. GPU 的起点来自图形学吞吐危机，而 shader 让 GPU 从固定功能图形硬件演化为可编程并行平台。
7
7. GPU 的设计目标是极致吞吐，不是豪华单线程控制流；它把更多预算投向算力和线程驻留，而不是复杂控制逻辑。
8
8. SIMT 与 SIMD 的区别在于：SIMD 是一个线程里的多 lane，SIMT 是很多线程共享统一控制流硬件。
9
9. warp divergence 的本质不是“不会算”，而是 warp 内控制流分化导致统一执行宽度被浪费。
10
10. GPU 除了怕 divergence，也怕不规则访存；memory coalescing 和统一控制流一样，都是吞吐成立的必要条件。
11
11. 从 ILP 到 SIMD 再到 GPU，是并行责任逐步从硬件隐式挖掘转向程序显式暴露的过程。
12
12. 对 AI infra 而言，tile、layout、batching、fusion、sharding 都是在把问题重排成硬件喜欢的规则并行形式。