第二章：GPU架构基础：SM与CUDA Core的关系、Warp调度机制、内存层次结构对负载的影响
好，咱们正式开始聊GPU的硬核内容。这一章是基础中的基础，但也是很多人容易搞混的地方。我见过不少开发者，写了很久的CUDA代码，却连SM和CUDA Core到底啥关系都说不清楚。嗯，这其实挺要命的。
2.1 SM和CUDA Core：别被名字骗了
先说说SM（Streaming Multiprocessor，流多处理器）。你可以把它理解成GPU里的“计算小组”。每个SM里包含了很多CUDA Core，但注意——CUDA Core并不是一个完整的CPU核心。
说白了，CUDA Core只是一个执行单元。它负责做整数运算、浮点运算这些基础操作。我个人的习惯是，把SM看作一个“车间”，CUDA Core就是车间里的“工人”。工人多了，车间干活就快。
关键点：一个SM里通常有几十到几百个CUDA Core。但别以为Core越多就一定越快。你得看它们怎么配合工作。
举个例子。我手头有块NVIDIA A100 GPU，每个SM里有64个CUDA Core。但实际跑起来，并不是所有Core都能同时干活。为什么？因为还有调度、内存访问这些限制。你想想看，如果工人都在等材料，那再多人也没用。
2.2 Warp调度：GPU的“小组作业”机制
接下来聊Warp。Warp是GPU执行任务的最小单位，通常包含32个线程。这32个线程会一起执行同一条指令。这就是所谓的SIMT（单指令多线程）模式。
我在项目中遇到过一个问题：某个核函数里有个if-else分支，结果导致Warp里的线程走了不同的路径。嗯，这就麻烦了。GPU会串行执行这些分支，先走if的线程，再走else的线程。性能直接砍半。
避坑指南：我曾经在一个图像处理项目里，因为没注意Warp内的分支发散，导致性能从预期的2ms掉到了5ms。后来我把分支条件改成了基于Warp ID的判断，才把性能拉回来。
Warp调度器是SM里的一个关键部件。它负责从多个Warp里选一个来执行。为什么要选？因为内存访问有延迟。当某个Warp在等数据时，调度器会切换到另一个Warp去干活。这样就能隐藏延迟。
说白了，GPU就是靠“人多”来掩盖“慢”的。一个Warp等数据，另一个Warp顶上。只要你有足够多的Warp，就能把内存延迟给填满。
2.3 内存层次结构：速度与容量的博弈
GPU的内存层次，说白了就是“越近越快，越远越慢”。但很多人只记住了这句话，没理解它对负载的影响。
| 内存类型 | 位置 | 速度 | 容量 | 作用 |
|---|---|---|---|---|
| 寄存器 | SM内部 | 最快（~1周期） | 很小（每个线程有限） | 存局部变量 |
| 共享内存 | SM内部 | 快（~5-30周期） | 小（几十KB） | 线程间通信 |
| L1缓存 | SM内部 | 较快 | 几十KB | 缓存数据 |
| L2缓存 | 芯片共享 | 中等（~200周期） | 几MB | 全局缓存 |
| 全局内存 | 显存 | 慢（~400-800周期） | 大（GB级） | 主存储 |
你看这个表，从寄存器到全局内存，速度差了上百倍。我刚开始做GPU优化时，总想着把所有数据都塞进共享内存。结果呢？共享内存太小，放不下，反而导致线程数受限。
注意：共享内存虽然快，但它是SM内的资源。每个SM的共享内存是有限的。如果你用太多，SM能同时运行的Warp数就会减少。这叫“占用率下降”。
我个人的经验是：先算清楚你的数据访问模式。如果是频繁访问的小数据，放共享内存。如果是偶尔访问的大数据，放全局内存。别盲目追求“快”，要平衡。
2.4 内存访问模式：合并访问是关键
说到全局内存，就不得不提“合并访问”。GPU从全局内存读数据时，是以128字节为单位的。如果你让一个Warp里的32个线程，连续访问相邻的地址，那一次就能读完。这就是合并访问。
反过来，如果线程访问的是随机地址，那GPU就得多次读取。性能直接崩掉。
// 好的例子：合并访问
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float val = data[idx]; // 连续访问
// 坏的例子：非合并访问
int idx = threadIdx.x * gridDim.x + blockIdx.x;
float val = data[idx]; // 跨步访问，性能差我在做矩阵转置时吃过这个亏。一开始没注意访问模式，性能只有理论峰值的10%。后来改成合并访问，直接提升了4倍。
2.5 负载均衡：SM之间的任务分配
最后聊聊负载均衡。GPU有多个SM，每个SM处理不同的Block。如果某个Block的计算量特别大，那它所在的SM就会忙死，其他SM却闲着。
为什么会这样？因为Block是静态分配到SM的。一旦分配完，就不能动态迁移。所以你得保证每个Block的计算量差不多。
技巧：我建议把大任务拆成多个小Block。比如一个1024x1024的矩阵，拆成32x32的Block。这样每个SM都能均匀分配到任务。
另外，Warp级别的负载均衡也很重要。如果一个Warp里有些线程提前退出了，那剩下的线程还得等。嗯，这其实就是前面说的分支发散问题。
总结一下这一章的核心：
- SM是车间，CUDA Core是工人
- Warp是32人小组，要避免组内分歧
- 内存层次要会用，别让工人等材料
- 访问模式要连续，别让GPU白干活
- 任务要拆均匀，别让SM闲死忙死
下一章，咱们会深入聊聊如何用这些知识来优化实际代码。到时候我会拿几个真实项目里的例子，手把手带你调优。