GPU架构基础：CUDA核心与SM、内存层次结构、Warp调度与线程束
大家好，我是你们这门课的主讲。今天咱们聊聊GPU的“骨架”——也就是它的硬件架构。说实话，很多同学写CUDA代码，跑得慢，往往不是因为算法不对，而是根本不了解手里的“武器”长什么样。你想想看，连GPU怎么干活都不知道，怎么可能写出高效的算子？
这一章，我会带你从三个核心维度拆解GPU：计算单元（SM与CUDA核心）、存储体系（各级内存）、以及执行模型（Warp调度）。搞懂这些，你才算真正入了GPU优化的门。
核心观点：GPU的强项不是单核性能，而是“人多力量大”。它通过大量轻量级线程并行执行，用延迟隐藏来掩盖访存慢的问题。理解这一点，是后续所有优化的前提。
1. SM与CUDA核心：GPU的“车间”与“工人”
先说说SM（Streaming Multiprocessor，流式多处理器）。你可以把SM想象成一个独立的“生产车间”。每个GPU里有很多个SM，比如我的老卡RTX 3090有82个SM。每个SM内部，才是一排排的CUDA核心——也就是真正干活的“工人”。
我个人习惯把CUDA核心看作“轻量级ALU”。它不像CPU核心那样有复杂的乱序执行、分支预测逻辑。CUDA核心很纯粹，就是做算术运算的。一个SM里通常有几十到上百个CUDA核心。比如Volta架构的SM里有64个FP32核心，Turing是64个，Ampere是128个。
避坑指南：我曾经以为CUDA核心越多，单线程就跑得越快。后来发现完全不是这么回事。GPU的单线程性能其实很弱，它的优势在于同时跑成千上万个线程。所以优化时别盯着“怎么让一个线程跑更快”，而要思考“怎么让所有线程都别闲着”。
这里有个关键点：SM才是真正的调度单位。你的线程块（Block）会被分配到某个SM上执行。一个SM可以同时驻留多个线程块，但具体能驻留多少，取决于线程块使用的资源——共享内存、寄存器数量。嗯，这个后面会细讲。
2. 内存层次结构：速度与容量的博弈
GPU的内存体系，说白了就是一个金字塔。越往上速度越快，但容量越小；越往下容量越大，但速度越慢。咱们从下往上捋一遍。
| 内存类型 | 位置 | 容量 | 延迟 | 作用域 |
|---|---|---|---|---|
| 全局内存 | 显存芯片 | GB级别 | ~400-800 cycles | 所有线程 |
| L2缓存 | 芯片上，所有SM共享 | MB级别 | ~200 cycles | 所有SM |
| L1缓存/共享内存 | 每个SM内部 | 几十到上百KB | ~30 cycles | 单个SM内的线程 |
| 寄存器 | 每个SM内部 | 每个SM几万个 | ~1 cycle | 单个线程 |
全局内存就是显存，你调用cudaMalloc分配的内存都在这里。它的特点是：容量大（8GB、16GB、80GB都有），但访问延迟极高。我记得有一次优化一个矩阵乘法算子，发现90%的时间都花在等数据从全局内存搬过来。后来改用共享内存做分块，速度直接翻了5倍。
共享内存是我个人最喜欢的特性。它是SM内部的一块高速存储，同一个线程块里的所有线程都能访问。你可以把它看作“手动管理的缓存”。为什么说“手动”？因为L1缓存是硬件自动管理的，你没法控制；但共享内存需要你在代码里显式地读写。这给了你极大的优化空间。
注意：共享内存虽然快，但容量非常有限。一个SM通常只有48KB到164KB。你想想看，一个线程块如果用了太多共享内存，SM能同时驻留的线程块数量就会减少，反而可能降低并行度。这是个经典的“空间换时间”的权衡。
寄存器是速度最快的存储，每个线程独享。但寄存器数量也是有限的。如果一个线程用了太多寄存器（比如声明了很多局部变量），编译器会做“寄存器溢出”，把多余的数据挤到局部内存（其实还是全局内存）里，性能会急剧下降。我见过一个项目，就因为多声明了几个double数组，性能掉了30%。
3. Warp调度与线程束：GPU的“最小执行单元”
这是很多初学者最容易忽略的地方。你写代码时觉得是32个线程在并行执行，但GPU实际执行时，是以Warp为单位的。一个Warp包含32个线程，它们在同一时刻执行同一条指令。
为什么会这样？因为GPU的设计哲学是“单指令多线程”（SIMT）。SM里的Warp调度器会从多个Warp中选一个，然后发射一条指令，让这个Warp里的32个线程同时执行。如果某个线程因为分支条件跳走了，那它对应的CUDA核心就会“空转”——这就是所谓的线程束发散。
关键洞察：Warp内的线程如果走不同的分支，性能会严重下降。因为GPU会串行执行每个分支，然后屏蔽掉不相关的线程。所以优化时，尽量保证同一个Warp里的线程执行路径一致。
我举个例子。假设你写了个if-else语句：
if (threadIdx.x < 16) {
// 分支A
} else {
// 分支B
}如果这个Warp的线程ID是0到31，那前16个线程走分支A，后16个走分支B。GPU会先执行分支A（后16个线程被屏蔽），再执行分支B（前16个线程被屏蔽）。总共用了2个周期。但如果所有32个线程都走同一个分支，1个周期就搞定了。这就是为什么我常说：“让Warp内线程保持统一”。
另外，Warp调度还有个“延迟隐藏”的机制。当当前Warp在等数据（比如从全局内存读数据）时，Warp调度器会立刻切换到另一个就绪的Warp去执行。只要你有足够多的Warp，GPU就能一直保持忙碌，看起来就像没有延迟一样。嗯，这就是为什么GPU需要大量线程——不是为了算得快，而是为了“填满”等待时间。
个人经验：我曾经优化一个图像处理算子，一开始只启动了256个线程，结果发现SM利用率只有30%。后来我把线程数增加到1024，利用率直接飙到95%。原因就是线程多了，Warp调度器有更多Warp可以切换，延迟被完美隐藏了。
4. 总结：把这些串起来看
好了，咱们把今天的内容串一下。你写一个CUDA程序，启动了一个Grid，它被拆成多个Block，每个Block被分配到某个SM上。SM把Block里的线程分成32个一组的Warp，然后Warp调度器不停地切换Warp，让CUDA核心干活。数据从全局内存经过L2、L1/共享内存，最终到寄存器里被计算。
你想想看，如果某个环节出了问题——比如Warp发散、共享内存用太多、寄存器溢出——整个流水线就会卡住。所以，优化的本质就是让这个流水线尽可能顺畅。
下一章，我们会动手写第一个CUDA程序，然后分析它的性能瓶颈。到时候你会发现，今天讲的这些架构知识，全都能用上。
一句话记住：GPU不是更快，而是更“宽”。理解它的宽度（Warp、SM、内存层次），你就能驾驭它。