GPU的“心脏”：流式多处理器（SM）架构详解
各位同学，今天我们来聊聊GPU真正的核心——流式多处理器，也就是SM。我常说，不理解SM，就等于没入门GPU架构。它就像一颗心脏，负责把指令变成实实在在的计算结果。
你想想看，一块GPU里有成百上千个核心，但真正干活的组织单元，就是SM。每个SM内部，藏着大量的计算单元、调度器、缓存和寄存器。说白了，SM就是GPU里的“迷你CPU集群”。
核心概念：SM（Streaming Multiprocessor）是GPU内部最基本的硬件计算单元。一个GPU由多个SM组成，每个SM独立执行线程块（Thread Block）。
SM的内部结构长什么样？
我习惯把SM想象成一个“小工厂”。这个工厂里有几条生产线，每条线上有几个工人。工人就是CUDA核心，生产线就是Warp调度器。嗯，这里要注意，不同代的GPU，SM内部布局差异很大。
以NVIDIA的Fermi架构为例，一个SM包含：
- 32个CUDA核心：负责整数和浮点运算
- 16个LD/ST单元：处理内存加载和存储
- 4个SFU：专门处理特殊函数（如sin、cos）
- Warp调度器：每个时钟周期选择一条Warp执行
- 共享内存：线程块内快速数据交换
- 寄存器文件：每个线程私有，容量巨大
我在项目中遇到过一个问题：共享内存用多了，寄存器就少了。这俩是“抢地盘”的关系。你想想看，一个线程块如果占用了太多共享内存，能同时运行的线程块数量就会减少，最终影响整体吞吐量。
Warp：SM里的执行单位
SM不直接调度线程，它调度的是Warp。一个Warp包含32个线程，它们在同一时刻执行同一条指令。这就是SIMT（单指令多线程）的精髓。
为什么会这样？因为硬件设计上，32个线程共享一个指令解码器。这样省面积、省功耗。但代价是——如果32个线程走不同的分支，性能就惨了。
避坑指南：我曾经在写CUDA程序时，让一个Warp里的线程走了不同的if-else分支。结果性能直接腰斩。后来我学会了用“谓词执行”或者重新组织数据，让同一个Warp的线程尽量走同一条路。
Warp的调度过程是这样的：
- Warp调度器从就绪队列中选一个Warp
- 检查该Warp的指令是否准备好（数据依赖是否解决）
- 如果准备好，发射指令到执行单元
- 执行单元在4个时钟周期内完成32个线程的计算
说白了，SM通过快速切换Warp来隐藏延迟。当一个Warp在等内存数据时，调度器立刻切换到另一个Warp。这就是GPU高吞吐量的秘密。
共享内存与寄存器：SM的“高速缓存”
共享内存是SM内部的一块SRAM，所有线程都可以访问。它的延迟只有几十个时钟周期，远低于全局内存。我习惯用它来做线程间的数据交换，或者做数据重用。
寄存器呢？每个线程私有，速度最快。但数量有限。一个SM的寄存器总数是固定的，比如Kepler架构是65536个。如果每个线程用32个寄存器，那一个SM最多同时容纳2048个线程。
| 架构 | SM寄存器数 | 共享内存大小 | 最大线程数/SM |
|---|---|---|---|
| Fermi | 32768 | 48 KB | 1536 |
| Kepler | 65536 | 48 KB | 2048 |
| Maxwell | 65536 | 96 KB | 2048 |
| Turing | 65536 | 64 KB | 1024 |
你看这张表，寄存器数量从Fermi到Turing翻了一倍，但共享内存大小也在涨。我建议你在写kernel时，先用--ptxas-options=-v编译选项看看寄存器使用量。如果超过32个，就该优化了。
SM的指令流水线
SM内部有一条完整的指令流水线。我把它分成几个阶段：
- 取指：从指令缓存中读取Warp的指令
- 译码：解析指令类型和操作数
- 发射：将指令发送到对应的执行单元
- 执行：CUDA核心、SFU或LD/ST单元干活
- 写回：将结果写回寄存器
嗯，这里要注意，不同指令的延迟不一样。整数加法只要1个周期，但浮点乘法可能要4个周期。全局内存加载更是要几百个周期。SM通过Warp级并行来掩盖这些延迟。
警告：不要以为SM内部所有单元都能同时工作。CUDA核心和SFU是共享发射端口的。如果你在代码里大量使用sin/cos，SFU会成为瓶颈，CUDA核心反而闲着。我踩过这个坑，后来用查表法替代了部分三角函数。
SM的线程调度策略
SM怎么决定哪个Warp先执行？我见过两种策略：
- 轮询调度：每个就绪的Warp轮流执行一个周期
- 优先级调度：根据Warp的优先级或等待时间决定
实际硬件用的是混合策略。SM会维护一个就绪队列，每次选一个“最老”的Warp。这样能保证公平性，也能避免某些Warp饿死。
你想想看，如果一个Warp一直在等全局内存的数据，它就不会进入就绪队列。SM会一直调度其他Warp，直到数据回来。这就是为什么线程数越多，延迟隐藏效果越好。
SM架构的演进
从Fermi到Hopper，SM架构一直在变。我简单总结几个关键变化：
- Fermi：首次引入真正的SM概念，32个CUDA核心
- Kepler：SMX单元，192个核心，但调度器更复杂
- Maxwell：SMM单元，128个核心，分区调度提升效率
- Pascal：SM单元回归，但支持FP16和统一内存
- Turing：加入Tensor Core，专门做矩阵运算
- Ampere/Hopper：SM内部分区更细，支持异步指令
我个人觉得，Turing架构是个分水岭。Tensor Core的出现，让SM不再只是做通用计算，而是开始为AI推理和训练做优化。你想想看，一个SM里既有CUDA核心，又有Tensor Core，还有RT Core，调度复杂度可想而知。
SM性能调优的实战建议
说了这么多理论，来点实际的。我在调优SM性能时，会关注这几个指标：
- Occupancy：活跃Warp数 / 最大Warp数。低于50%说明资源利用率低
- 共享内存使用量：每个线程块用的共享内存，尽量控制在16KB以内
- 寄存器压力：每个线程寄存器数，超过32个会降低Occupancy
- Warp发散：同一个Warp里分支路径数，越少越好
小技巧：我常用NVIDIA的Nsight Compute工具看SM的stall原因。如果发现“long scoreboard”占比高，说明内存延迟没藏好。这时候可以增加线程数，或者用__ldg()指令启用只读缓存。
最后说一句，SM架构虽然复杂，但核心思想很简单：用大量线程并行执行，用快速切换隐藏延迟。你只要记住这个，后面学内存层次结构、学Tensor Core，都会轻松很多。
这张图展示了SM内部的核心组件。Warp调度器选好Warp后，指令分发单元把指令送到执行单元。CUDA核心、LD/ST单元和SFU各司其职。共享内存和寄存器文件为它们提供数据支持。
好了，SM架构的核心内容就这些。记住，SM是GPU的“心脏”，理解它，你就掌握了GPU硬件架构的钥匙。