GPU架构基础回顾：SM、Warp、Memory Hierarchy、Occupancy概念
好，咱们正式开始。这一节我打算带大家快速过一遍GPU的核心架构概念。说实话，很多同学学了几个月CUDA，SM和Warp到底长什么样，心里还是模糊的。这不行。你想想看，连对手的底牌都不清楚，怎么优化？
我个人习惯，每次拿到一个新GPU，第一件事就是看它的SM配置。为什么？因为SM就是GPU的“心脏”。下面我们一个一个来拆。
1. SM：流多处理器
SM的全称是Streaming Multiprocessor。你可以把它理解成一个“迷你CPU核”，但它不是单核，它里面塞了很多小核心。
每个SM包含：
- CUDA Core：执行算术运算的单元。注意，不是所有Core都一样，有INT32、FP32、FP64之分。
- Shared Memory：一块可编程的片上缓存，速度极快。
- Register File：寄存器堆，每个线程私有，容量有限。
- Warp Scheduler：负责调度Warp的执行。
- Load/Store Unit：处理内存访问。
我在项目中遇到过一个问题：一个kernel在V100上跑得不错，换到A100上反而慢了。查了半天，发现是SM数量翻倍了，但每个SM的Shared Memory没变，导致 occupancy 反而下降了。嗯，这里要注意，SM数量多不一定总是好事。
核心要点：GPU的性能上限，往往取决于SM能同时处理多少线程块。不是线程越多越好，是SM能“喂饱”的线程才有效。
2. Warp：线程束
Warp是GPU执行的最小单位。一个Warp包含32个线程。这32个线程在同一个SM上，执行同一条指令，但处理不同的数据——这就是SIMT（单指令多线程）。
为什么会这样？因为GPU的设计哲学是“用面积换吞吐”。它不需要复杂的分支预测和乱序执行，它只需要让32个线程齐步走。
但问题来了：如果这32个线程走了不同的分支呢？比如：
if (threadIdx.x % 2 == 0) {
// 偶数线程走这里
} else {
// 奇数线程走这里
}这时候Warp会发生分支发散。SM会先执行偶数线程的分支，再执行奇数线程的分支。性能直接砍半。
避坑指南：我曾经在一个图像处理项目里，因为一个简单的if-else，导致性能掉了40%。后来我把数据重新排列，让同一个Warp内的线程访问连续的内存地址，分支自然就消失了。记住：Warp内尽量保持统一执行路径。
3. Memory Hierarchy：内存层次结构
GPU的内存层次，说白了就是一个“速度与容量”的权衡游戏。从上到下，速度越来越慢，容量越来越大。
| 层级 | 位置 | 延迟（大约） | 容量 | 作用域 |
|---|---|---|---|---|
| 寄存器 | SM内部 | 1 cycle | ~256KB/SM | 线程私有 |
| Shared Memory | SM内部 | ~20 cycles | ~48-164KB/SM | 线程块内共享 |
| L1 Cache | SM内部 | ~30 cycles | ~128KB/SM | SM内共享 |
| L2 Cache | 芯片全局 | ~200 cycles | ~6-40MB | 全局共享 |
| Global Memory | 显存 | ~400-800 cycles | ~16-80GB | 全局共享 |
你想想看，如果每次数据都从Global Memory读，那GPU大部分时间都在等数据。所以优化的核心思路就是：让数据尽量往上层走。
我个人习惯，写kernel之前先画一张数据流图：数据从哪里来，经过哪些计算，最终写到哪里去。然后问自己：能不能用Shared Memory做中间缓存？能不能让数据在寄存器里多待一会儿？
小技巧：如果你发现kernel的瓶颈在内存访问，试试用__ldg()指令（只读缓存）。我在一个矩阵乘法项目里，用这个指令白捡了15%的性能提升。
4. Occupancy：占用率
Occupancy，直译就是“占用率”。它指的是：每个SM上活跃的Warp数量，除以SM能支持的最大Warp数量。
举个例子：一个SM最多能同时放64个Warp（2048个线程）。如果你只放了32个Warp，那Occupancy就是50%。
为什么Occupancy重要？因为GPU靠“大量线程”来隐藏延迟。当一个Warp在等数据时，SM可以立刻切换到另一个Warp执行。如果Occupancy太低，SM就“无活可干”，只能干等。
但注意：100% Occupancy不一定最好。我见过很多同学拼命调参，把Occupancy从50%拉到100%，结果性能反而下降了。为什么？因为每个线程分到的资源（寄存器、Shared Memory）变少了，导致每个线程的执行效率下降。
我的经验：对于计算密集型kernel，Occupancy在50%-75%往往是最优的。对于内存密集型kernel，尽量拉高Occupancy到75%以上。具体多少，跑一下nvidia-smi和nvprof就知道了。
5. 几个关键公式
最后，我建议大家记住这几个公式。面试常考，调优也常用：
- 最大线程块数/SM = min(最大线程块数限制, SM资源限制)
- Occupancy = (活跃Warp数 / 最大Warp数) × 100%
- 理论带宽 = 显存频率 × 位宽 / 8
- 实际带宽利用率 = (实际吞吐 / 理论带宽) × 100%
嗯，这一节的内容就到这里。说白了，GPU架构就是一套“用空间换时间，用并行换延迟”的体系。理解SM、Warp、Memory Hierarchy和Occupancy，你就拿到了性能优化的钥匙。
下一节，我们会深入讲如何用NVIDIA Nsight工具定位瓶颈。到时候我会分享一些实战中的“翻车”案例，很有意思。