GPU硬件架构基础：流多处理器（SM）、CUDA核心、线程束（Warp）、共享内存与全局内存层次结构
好，咱们今天聊聊GPU的硬件骨架。很多人写CUDA代码，上来就是<<
所以这一节，咱们把GPU的硬件结构拆开来看。你理解了SM、Warp、内存层次这些东西，后面写代码时心里就有底了。
1. 流多处理器（SM）——GPU的“计算单元”
SM的全称是Streaming Multiprocessor，你可以把它理解成GPU里的一个“小CPU集群”。一个GPU芯片上，通常有几十到上百个SM。比如NVIDIA的A100有108个SM，H100有132个。
每个SM内部，包含多个CUDA核心、共享内存、寄存器文件、调度器、缓存等。说白了，SM才是真正干活的地方。你的线程块（block）会被分配到某个SM上执行，而且一个SM可以同时驻留多个block。
关键点：SM是硬件资源分配的基本单位。一个block里的所有线程，必须在同一个SM上执行。如果block需要的资源（比如共享内存、寄存器）超过了SM的容量，那这个block就启动不了。
我记得有一次，我写了一个很大的kernel，每个block用了很多共享内存。结果一跑，发现occupancy（占用率）特别低。查了半天，原来是共享内存用太多，导致一个SM只能放一两个block。嗯，这就是典型的“资源过度消耗”问题。
2. CUDA核心——SM里的“小工兵”
每个SM内部包含多个CUDA核心。注意，这里的“CUDA核心”不是CPU那种复杂的核心，它其实是一个简单的浮点/整数运算单元。你可以把它想象成一个流水线上的工人，只会做加减乘除这些基本操作。
一个SM里有多少个CUDA核心？不同架构不一样：
| 架构 | 每个SM的CUDA核心数 |
|---|---|
| Fermi (GTX 480) | 32 |
| Kepler (GTX 680) | 192 |
| Maxwell (GTX 980) | 128 |
| Pascal (GTX 1080) | 64 |
| Volta/Turing (RTX 2080) | 64 |
| Ampere (A100) | 64 |
你可能会问：为什么核心数不是一直增加的？其实，架构设计不是堆核心那么简单。每个核心的吞吐量、频率、配套的调度逻辑都在变。我个人习惯是，不纠结具体数字，而是关注“每个SM能同时处理多少线程”。
3. 线程束（Warp）——GPU调度的最小单位
这是GPU编程里最核心的概念之一。Warp是32个线程组成的一组，它们是“绑在一起”执行的。什么意思呢？就是这32个线程在同一个SM上，执行同一条指令，但处理不同的数据。这就是SIMT（单指令多线程）模型。
举个例子：
__global__ void add(float *a, float *b, float *c, int n) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx < n) {
c[idx] = a[idx] + b[idx];
}
}假设你启动了128个线程，那它们会被分成4个Warp（128 / 32 = 4）。每个Warp里的32个线程，同时执行 c[idx] = a[idx] + b[idx] 这条指令。但注意，如果线程的idx不同，它们访问的内存地址也不同。
避坑指南：我曾经写过一个kernel，里面有个if-else分支。结果发现，同一个Warp里的线程，有的走if，有的走else。这就导致Warp里的线程“分叉”了，GPU不得不串行执行两个分支。性能直接腰斩。所以，写代码时尽量保证同一个Warp里的线程走相同的路径。
为什么会这样？因为Warp是调度单位，SM一次只能给一个Warp发一条指令。如果线程走不同分支，那SM就得先执行if分支的线程，再执行else分支的线程。这就是所谓的“Warp divergence”（线程束分化）。
4. 内存层次结构——从快到慢，从贵到便宜
GPU的内存不是铁板一块，它有清晰的层次结构。你想想看，如果所有线程都去访问同一个全局内存，那带宽再大也扛不住。所以，GPU设计了多级缓存和共享内存来缓解压力。
下面这张图是我自己画的，展示了GPU内存的层次关系：
从上到下，速度越来越慢，但容量越来越大。咱们逐个说：
4.1 寄存器（Register）
每个线程私有的，速度最快，没有延迟。但数量有限——每个SM的寄存器总数是固定的，比如A100每个SM有65536个寄存器。如果每个线程用太多寄存器，那SM能同时容纳的线程数就少了。这就是为什么有时候减少寄存器使用能提升性能。
4.2 共享内存（Shared Memory）
同一个Block内的线程可以共享。它的速度比全局内存快很多（大约快20-30倍），但容量很小，通常每个SM只有几十到一百多KB。我经常用它来做数据复用——比如矩阵分块计算时，先把一块数据加载到共享内存，然后让所有线程反复读。
注意：共享内存是手动管理的。你得自己用 __shared__ 关键字声明，然后显式地加载数据。不像缓存是自动的。用得好是神器，用不好就是坑。
4.3 L1/L2缓存
L1缓存是SM内部的，自动缓存全局内存的访问。L2缓存是所有SM共享的。它们对程序员是透明的，但了解它们有助于优化——比如，如果你能保证数据访问的局部性，那缓存命中率就会高，性能自然就好。
4.4 全局内存（Global Memory）
就是显存，容量最大（几个GB到几十GB），但延迟也最高（几百个时钟周期）。所有线程都能访问，但访问速度最慢。所以，尽量减少全局内存的访问次数，或者让访问模式满足“合并访问”条件——说白了，就是让同一个Warp里的线程访问连续的内存地址。
5. 总结一下
嗯，这一节的内容有点多，但都是硬核知识。我建议你记住几个关键点：
- SM是硬件调度单位，Block是软件调度单位。一个Block只能在一个SM上跑。
- Warp是32个线程的集合，是真正的执行单位。避免Warp内部分支。
- 内存层次要心里有数：寄存器最快但最少，共享内存次之但需要手动管理，全局内存最慢但容量大。
我个人习惯是，写kernel之前先画个草图，标出数据流和内存访问模式。这样能避免很多后期调试的麻烦。你试试看，真的有用。
一个小技巧：用 cudaOccupancyMaxPotentialBlockSize 这个API可以帮你算出最优的block大小。我经常用它来快速验证我的配置是否合理。
好了，这一节就到这里。下一节咱们会深入线程调度和Warp执行模型，到时候你会更清楚GPU是怎么“同时”处理成千上万个线程的。