2、GPU架构核心：流式多处理器（SM）的内部世界
好，咱们今天聊聊GPU里最核心的东西——SM。
SM的全称是Streaming Multiprocessor，中文叫流式多处理器。你可以把它理解成GPU里的“计算小分队”。一个GPU芯片里，少则几十个，多则上百个SM。每个SM都是独立的计算单元，能并行执行成千上万个线程。
我刚开始接触CUDA时，总觉得SM就是个黑盒子。后来做性能调优，踩了不少坑，才慢慢摸清它的脾气。说白了，不理解SM，你就没法真正驾驭GPU。
SM长什么样？
每个SM内部，主要由这几部分组成：
- CUDA核心：负责执行算术运算，比如加减乘除、逻辑判断。一个SM里通常有几十到几百个CUDA核心。
- 共享内存：一块极快的片上存储，同一个SM里的线程可以互相访问。我习惯叫它“线程间的快递站”。
- 寄存器文件：每个线程私有的“小本本”，存临时变量。容量有限，用多了会限制并发线程数。
- 线程束调度器：负责把线程分组，然后派发给CUDA核心去执行。这是SM的“大脑”。
- 加载/存储单元：负责从显存里搬数据，或者把结果写回去。
嗯，这里要注意：不同代GPU的SM结构差异很大。比如NVIDIA的Volta架构引入了Tensor Core，专门做矩阵乘法。但核心思想是一样的——SM是并行执行的基本单元。
线程是怎么被调度的？
你写一个CUDA程序，启动成千上万个线程。这些线程不会一股脑全塞进SM里。GPU会先把线程分成一个个线程块，然后每个线程块分配给一个SM。
一个SM能同时处理多个线程块，但数量有限。我记得有一次做项目，线程块设得太大，结果一个SM只能塞下两个块，并发度上不去，性能直接腰斩。
那SM内部怎么调度线程呢？答案是线程束。
线程束是SM调度的最小单位。一个线程束包含32个线程。这32个线程会一起执行同一条指令，但处理不同的数据。这就是所谓的SIMT（单指令多线程）模型。
说白了，就是32个人同时做同一件事，但每个人拿到的数据不一样。
核心要点：SM以线程束为单位调度，而不是单个线程。所以你的代码里，如果32个线程走不同的分支（比如if-else），那性能就会下降。这叫“线程束发散”。
SM的并行执行能力
一个SM能同时执行多少个线程？这取决于它的资源。
我举个例子。假设一个SM有65536个寄存器，每个线程需要32个寄存器。那一个SM最多能同时容纳65536 / 32 = 2048个线程。但别忘了，还有共享内存的限制。
所以实际能跑的线程数，是寄存器、共享内存、线程块大小三者共同决定的。我曾经在调优时，为了省共享内存，把数据拆成小块分批处理，结果寄存器又不够用了……嗯，这就是GPU编程的日常。
下面这张图，是我画的SM内部结构简图，帮你理清关系：
共享内存：SM里的“高速缓存”
共享内存是SM里的一块特殊存储。它比全局内存快得多，延迟只有几十个时钟周期。同一个线程块里的所有线程，都能读写这块内存。
我个人的习惯是：能用共享内存，就不用全局内存。比如做矩阵乘法时，把子矩阵先搬到共享内存里，然后让线程反复读取，能省下大量访存时间。
但共享内存容量很小，通常只有几十KB到一百多KB。你得精打细算。我曾经为了省共享内存，把数据从32位浮点改成16位浮点，精度损失在可接受范围内，性能却翻了一倍。
小技巧：用__shared__关键字声明共享内存变量。记得加__syncthreads()同步，不然线程之间数据没写完就开始读，会出大问题。
寄存器：线程的“私人保险箱”
每个线程都有自己的寄存器。寄存器是芯片上最快的存储，没有之一。但数量有限。
如果一个线程用了太多寄存器，编译器会把多余的变量“溢出”到本地内存（其实还是全局内存），速度就慢了。你想想看，本来一纳秒能搞定的事，现在要几百纳秒，性能能不崩吗？
我建议你在编译时加上--ptxas-options=-v参数，看看每个线程用了多少寄存器。如果超过32个，就得考虑优化了。
线程束调度：SM的“时间管理大师”
SM里通常有多个线程束调度器。每个调度器负责管理一组线程束。当一个线程束在等待数据（比如从显存读数据），调度器会立刻切换到另一个就绪的线程束去执行。
这就是GPU的零开销上下文切换。CPU切换线程要保存一堆寄存器，开销很大。GPU不用，因为每个线程束的寄存器是物理独立的，切换就是换个指针的事。
所以，线程束越多，SM越容易“藏”延迟。我见过一个项目，把线程块从128线程改成256线程，占用率从50%升到75%，性能提升了30%。
避坑指南：我曾经以为线程束越多越好，结果把线程块设得太大，共享内存不够用，反而限制了并发。记住：占用率不是越高越好，要平衡资源使用。
SM的占用率计算
占用率就是SM里活跃的线程束数量，除以它最多能支持的线程束数量。比如一个SM最多支持64个线程束，你现在跑了48个，占用率就是75%。
占用率越高，SM越能隐藏延迟。但也不是越高越好。我做过实验，占用率从100%降到75%，有些程序性能反而提升了。为什么？因为每个线程分到的资源更多了，单个线程执行得更快。
所以，调优时别死盯着占用率。要结合具体算法，找到那个平衡点。
| 资源 | 典型容量（以某代GPU为例） | 对占用率的影响 |
|---|---|---|
| 寄存器 | 65536个/SM | 每个线程用越多，能跑的线程越少 |
| 共享内存 | 96KB/SM | 每个线程块用越多，能同时跑的块越少 |
| 线程块大小 | 最大1024线程/块 | 块太大或太小都会影响占用率 |
| 最大线程束数 | 64个/SM | 决定了占用率的上限 |
总结一下
SM是GPU的“心脏”。理解它的内部结构，你才能写出高效的CUDA代码。
- 线程以线程束为单位调度，注意避免分支发散。
- 共享内存是加速利器，但容量有限，要省着用。
- 寄存器是线程的私有财产，别浪费。
- 占用率不是越高越好，要找到资源使用的平衡点。
嗯，今天就聊到这儿。下次咱们聊聊全局内存的访问模式，那又是一个大坑。