Profile基础：GPU架构概览、SM与Warp调度、内存层次结构、Profile的基本流程
做性能优化，不懂硬件架构，就像开车不看仪表盘。你踩油门，但不知道发动机转速多少，档位对不对，那基本就是瞎忙活。
这一节，我带你把GPU的底裤扒开看看。不扯太深，但核心概念必须吃透。我当年刚接触CUDA时，就是吃了不懂SM和Warp的亏，写出来的kernel跑得比CPU还慢，后来一分析，发现warp divergence严重到离谱。
GPU架构概览：从宏观到微观
先看一张图，这是我手绘的GPU核心架构逻辑图。别嫌丑，但逻辑是清晰的。
这张图里，最核心的就是 SM（Streaming Multiprocessor）。你可以把SM理解成GPU里的「计算小分队」。每个SM内部有多个CUDA核心、共享内存、寄存器文件，还有Warp调度器。
多个SM组成一个GPC（图形处理簇），多个GPC共享L2缓存，最后通过内存控制器访问显存。嗯，这个层次结构你得记在心里，因为Profile时，你看到的很多指标都跟这些层级有关。
SM与Warp调度：别让你的线程「内讧」
SM是真正的执行单元。你的kernel里写了成千上万个线程，但SM一次只能处理有限个。那怎么办？
答案是：Warp。
Warp是SM调度的基本单位，固定32个线程一组。SM把线程块（Block）拆成多个Warp，然后轮流调度。我习惯把Warp比作「一列火车」，32个线程就是32节车厢。调度器每次拉一列火车出去跑。
关键点：同一个Warp里的32个线程，执行的是同一条指令。如果它们走的分支不一样（比如if-else），那就糟了——这叫Warp Divergence。
举个例子：
__global__ void bad_kernel(float* data, int n) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
if (idx < n/2) {
data[idx] = data[idx] * 2.0f; // 分支A
} else {
data[idx] = data[idx] * 0.5f; // 分支B
}
}这段代码，同一个Warp里，前16个线程走分支A，后16个走分支B。SM会怎么做？它会先执行分支A，把走分支B的线程「屏蔽」掉，然后再反过来执行分支B。说白了，一个Warp本来一次能搞定的事，现在要两次。性能直接腰斩。
我的经验：我曾经优化过一个图像滤波的kernel，原本有4个分支，性能惨不忍睹。后来我把数据按条件重新排列，让同一个Warp里的线程走同一个分支，性能提升了3倍。所以，写kernel时，尽量让Warp内的线程「步调一致」。
SM里一般有多个Warp调度器（比如4个），它们可以同时调度不同的Warp。当一个Warp在等内存访问时，调度器会立刻切换到另一个Warp，这就是所谓的零开销上下文切换。你想想看，这多聪明——计算和访存可以重叠，隐藏延迟。
内存层次结构：速度与容量的博弈
GPU的内存层次，说白了就是「越靠近计算单元越快，但容量越小」。我画个表格，你一看就明白。
| 内存类型 | 位置 | 延迟 | 容量 | 作用域 |
|---|---|---|---|---|
| 寄存器 | SM内部 | ~1 cycle | 256KB / SM | 单个线程 |
| 共享内存 | SM内部 | ~5 cycles | 48~164KB / SM | 单个Block |
| L1缓存 | SM内部 | ~10 cycles | 与共享内存共用 | 单个SM |
| L2缓存 | 芯片级 | ~200 cycles | 几MB ~ 几十MB | 所有SM |
| 全局内存 | 显存 | ~400~800 cycles | 几GB ~ 几十GB | 所有线程 |
看到没？全局内存的延迟是寄存器的几百倍。所以，优化内存访问是Profile的重中之重。
我个人习惯，在写kernel之前，先想清楚数据流：哪些数据是线程私有的？放寄存器。哪些是Block内共享的？放共享内存。哪些是所有线程都要读的？考虑用常量内存或纹理内存。
避坑指南：我曾经犯过一个低级错误——把频繁访问的查找表放在了全局内存里，结果每次查表都要等几百个cycle。后来改成常量内存，延迟直接降到几十个cycle。所以，别偷懒，该用共享内存就用，该用常量内存就用。
Profile的基本流程：先看什么，后看什么
好了，架构和内存都讲完了，那Profile到底怎么上手？我总结了一个四步流程，你照着做就行。
- 第一步：跑一次基线Profile
什么都不改，直接跑你的应用。用Nsight Systems或Nsight Compute抓一次数据。看什么？看整体时间分布：kernel占多少？数据传输占多少？CPU端占多少？
- 第二步：定位热点Kernel
找到最耗时的那个kernel。通常80%的时间花在20%的kernel上。别去优化那些只占1%时间的kernel，那是浪费时间。
- 第三步：分析Kernel内部瓶颈
用Nsight Compute打开那个热点kernel。看几个关键指标：
- Occupancy：SM上活跃Warp的比例。太低说明资源受限（寄存器或共享内存用太多）。
- Memory Throughput：显存带宽利用率。如果远低于理论峰值，说明访存模式有问题。
- Compute Throughput：计算单元利用率。如果很低，说明计算不是瓶颈，访存才是。
- Warp Divergence：分支分歧率。太高的话，考虑重构代码。
- 第四步：针对性优化，再Profile
根据第三步的分析结果，改代码。比如：合并内存访问、减少分支、用共享内存缓存数据。改完后，再跑一次Profile，看指标有没有改善。反复迭代，直到满意。
我的小技巧：每次优化只改一个点。不要一次改多个地方，否则你根本不知道哪个改动起了作用。我见过有人一次改了10处，性能提升了，但完全不知道是哪个优化生效了——这种经验没法复用。
嗯，这一节的内容就到这。GPU架构、SM与Warp调度、内存层次、Profile流程，这四个东西是后续所有优化的基础。你把这些吃透了，后面讲具体案例时，你才能听懂我在说什么。
想要更多资料，联系可可