4. 寄存器与本地内存：寄存器溢出与本地内存的陷阱
好，咱们今天聊点硬核的。寄存器（Register）和本地内存（Local Memory），这两个东西在CUDA里就像是你的左膀右臂。用好了，性能起飞；用不好，嗯……我见过不少项目，明明算法没问题，就是慢，最后查下来，都是寄存器溢出（Register Spilling）在搞鬼。
说白了，寄存器是GPU里最快的内存，没有之一。每个线程都有自己的一套寄存器，访问它只需要一个时钟周期。但问题是，寄存器数量是有限的。你想想看，一个SM（流多处理器）就那么点寄存器，分给所有线程。一旦你的kernel里定义的变量太多，编译器就会说：“兄弟，寄存器不够用了，我把一些变量挪到本地内存去吧。”
这一挪，性能就崩了。本地内存虽然名字里带个“内存”，但它实际上用的是全局内存，还带缓存。延迟比寄存器高了一个数量级不止。这就是寄存器溢出。
核心观点：寄存器溢出是性能杀手。一个变量从寄存器被踢到本地内存，访问延迟从1个周期变成几百个周期。你想想看，这差距有多大。
4.1 如何检测寄存器溢出？
检测寄存器溢出，我个人的习惯是三步走。第一步，看编译器报告。第二步，看汇编代码。第三步，用profiler跑一下。
方法一：编译器报告
编译的时候加上 --ptxas-options=-v 参数，编译器会告诉你每个kernel用了多少寄存器，有没有发生溢出。
nvcc -arch=sm_80 --ptxas-options=-v mykernel.cu输出里你会看到类似这样的信息：
ptxas info: Used 64 registers, 0 bytes lmem, 0 bytes smem
ptxas info: Used 32 registers, 16 bytes lmem, 0 bytes smem注意那个 lmem。只要它大于0，就说明发生了寄存器溢出。16 bytes lmem意味着有4个float变量被挤到本地内存里了。
方法二：查看PTX或SASS代码
用 --keep 选项保留中间文件，然后打开 .ptx 文件。搜索 local 关键字，你会看到类似 st.local.f32 或 ld.local.f32 的指令。这些就是访问本地内存的指令。
方法三：NVIDIA Nsight Compute
这个工具最直观。打开profiler，找到你的kernel，查看“Register Usage”和“Local Memory”指标。如果Local Memory那一栏不是0，那就得警惕了。
我的经验：我曾经优化过一个图像处理kernel，一开始profiler显示有32 bytes的本地内存访问。我以为是小事，结果一算，每个线程多花了200多个周期。后来我把一些中间变量合并了，本地内存降到了0，性能直接提升了30%。
4.2 如何避免寄存器溢出？
避免寄存器溢出，说白了就是“省着用”。但也不能盲目省，得讲究方法。
策略一：减少每个线程的变量数量
这是最直接的办法。你想想看，一个kernel里定义了几十个float变量，每个变量都要占一个寄存器。能不溢出吗？
- 复用变量：能用同一个变量就别定义新的。
- 使用数组索引：有时候一个数组比一堆独立变量更省寄存器。
- 把计算拆分成多个kernel：如果逻辑太复杂，拆成两个kernel，每个kernel的寄存器压力就小了。
策略二：使用 __launch_bounds__ 限制线程数
这个技巧我特别喜欢。你可以告诉编译器：“我这个kernel最多用32个寄存器，你看着办。” 编译器会尽量把寄存器使用量压下来，哪怕多用点本地内存也在所不惜。但要注意，这可能会降低occupancy（占用率）。
__global__ void __launch_bounds__(256, 4) myKernel() {
// 你的代码
}第一个参数是每个block的最大线程数，第二个参数是每个SM的最小block数。通过调整这两个参数，你可以控制寄存器分配。
策略三：使用 -maxrregcount 编译选项
这个选项可以全局限制所有kernel的寄存器数量。比如 -maxrregcount=32 表示每个线程最多用32个寄存器。超出部分自动溢出到本地内存。
nvcc -arch=sm_80 -maxrregcount=32 mykernel.cu注意：强制限制寄存器数量不一定总是好事。有时候寄存器用得多，但occupancy低，性能反而更好。我曾经在一个矩阵乘法的kernel里试过，把寄存器从64降到32，occupancy上去了，但每个线程的计算能力下降了，最终性能反而变差了。所以，一定要实测。
4.3 本地内存的陷阱
本地内存，名字听着挺“本地”的，其实它就是个幌子。它用的是全局内存，只不过每个线程有自己的私有地址空间。访问本地内存，本质上就是访问全局内存，只不过多了一层L1/L2缓存。
陷阱一：误以为本地内存很快
很多新手以为“本地”就是“快”。其实不是。本地内存的延迟是寄存器的10倍以上。如果频繁访问，性能会急剧下降。
陷阱二：编译器自动溢出
编译器有时候会自作主张。比如你定义了一个大数组，编译器觉得放寄存器里太浪费，就自动把它放到本地内存里了。你甚至都不知道。
__global__ void myKernel() {
float data[1024]; // 这个数组很可能被放到本地内存
// 对data进行操作
}这个 data 数组，如果每个线程都定义，那每个线程都要在本地内存里占4KB。1024个线程就是4MB。你想想看，这得多大的内存带宽压力。
陷阱三：本地内存的bank冲突
本地内存虽然每个线程私有，但在硬件层面，它还是通过全局内存的缓存系统来访问的。如果多个线程同时访问本地内存，可能会发生bank冲突，进一步降低性能。
避坑指南：我曾经在一个深度学习推理的kernel里，用了一个局部数组来存储中间激活值。结果profiler显示本地内存访问量巨大。后来我把这个数组改成了寄存器变量，并用循环展开来减少变量数量。最终本地内存降为0，推理速度提升了15%。
4.4 实战建议
说了这么多，总结几条实战建议：
- 先检测，再优化。 不要凭感觉。用profiler看看本地内存是不是0。
- 能放寄存器就放寄存器。 寄存器是最快的，没有之一。
- 大数组要小心。 如果数组大小超过几十个元素，大概率会被放到本地内存。考虑用共享内存（Shared Memory）或者全局内存来替代。
- 权衡occupancy和寄存器使用。 不是occupancy越高越好。有时候降低occupancy，提高每个线程的寄存器数量，反而能提升性能。
- 实测为王。 所有优化都要以实际运行时间为准。不要迷信理论。
嗯，寄存器与本地内存这块，说白了就是“资源管理”的艺术。你管得好，GPU就听你的；管不好，它就用慢速内存来惩罚你。下一章我们会聊共享内存，那又是一个新的战场。