寄存器优化：寄存器压力、寄存器溢出、__launch_bounds__、maxrregcount
寄存器，在GPU里是最快的存储资源。比L1 cache快，比shared memory快，比global memory快出几个数量级。但问题是——寄存器数量有限。
我刚开始做CUDA优化那会儿，总觉得寄存器越多越好。把能放的变量全塞进寄存器，结果性能反而下降了。后来才明白，寄存器不是越多越好，而是够用就好。
寄存器压力是什么？
说白了，就是你的kernel需要的寄存器数量，超过了硬件能提供的数量。每个SM（流式多处理器）有固定数量的寄存器，比如32K或64K。这些寄存器要在所有并发线程间分配。
举个例子。假设一个SM有65536个寄存器，每个线程需要32个寄存器。那这个SM最多能跑65536/32=2048个线程。但如果每个线程需要64个寄存器，就只能跑1024个线程了。
你想想看，线程数减少意味着什么？意味着occupancy下降，意味着GPU的并行能力被浪费了。
核心观点：寄存器压力过大，会直接导致occupancy下降。occupancy下降，意味着更多的线程切换延迟无法被隐藏。
寄存器溢出——性能杀手
当编译器发现你的kernel需要的寄存器数量超过了硬件限制，它会怎么做？它会做一件很无奈的事——把一部分寄存器变量 spill 到local memory里。
local memory听起来名字里带个"local"，但实际上它是在显存里。访问local memory的延迟，比访问寄存器高出两个数量级。这就是寄存器溢出。
我在项目中遇到过这样一个案例：一个深度学习推理kernel，原本跑得还行。我加了一些中间变量的缓存，想减少重复计算。结果寄存器压力暴增，编译器开始spill。性能反而下降了30%。
怎么判断发生了寄存器溢出？看PTX代码或者用nvcc的--ptxas-options=-v参数编译，会输出寄存器使用情况和spill信息。
// 编译时查看寄存器使用情况
nvcc --ptxas-options=-v kernel.cu
// 输出示例
// 512 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
// 如果spill stores/spill loads不为0，说明发生了寄存器溢出避坑指南：我曾经在一个项目里，看到spill stores达到几千字节，性能惨不忍睹。后来通过减少每个线程的工作量、拆分kernel，才把spill降下来。记住：spill是性能的敌人。
__launch_bounds__——手动控制寄存器分配
__launch_bounds__是CUDA提供的一个利器。它告诉编译器：我这个kernel最多会启动多少个线程，最少需要多少个线程。
语法很简单：
__global__ __launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor)
void myKernel() {
// kernel代码
}第一个参数是每个block的最大线程数。第二个参数是每个SM至少需要多少个block才能达到满occupancy。
编译器会根据这些信息，决定分配多少寄存器。你告诉它线程数多，它就会少分配寄存器，让更多线程能同时运行。你告诉它线程数少，它就可以多分配寄存器，减少spill。
我个人的习惯是：先不加__launch_bounds__，让编译器自动优化。然后看性能，如果occupancy太低，就加上__launch_bounds__限制寄存器数量。
实战技巧：如果你发现kernel有spill，可以尝试用__launch_bounds__减少每个线程的寄存器数量。比如默认用了32个寄存器，你强制限制到24个。编译器会想办法用更多local memory或者重新安排计算顺序。
maxrregcount——全局控制寄存器数量
__launch_bounds__是per-kernel的控制方式。如果你想让整个编译单元都使用相同的寄存器限制，可以用maxrregcount。
// 编译时指定最大寄存器数
nvcc --maxrregcount=32 kernel.cu这个参数会覆盖所有kernel的默认寄存器分配。我一般只在调试时用这个参数，生产环境更推荐用__launch_bounds__做精细控制。
为什么？因为不同kernel对寄存器的需求不一样。有的kernel计算密集，需要更多寄存器来减少访存。有的kernel访存密集，寄存器少一点也无所谓。一刀切往往不是最优解。
寄存器优化的核心策略
说了这么多，到底怎么优化寄存器？我总结了几条经验：
- 减少每个线程的工作量：一个线程处理的数据越少，需要的寄存器就越少。可以考虑用更小的block size。
- 拆分kernel：如果一个kernel太复杂，寄存器需求太高，拆成两个简单的kernel。虽然增加了kernel启动开销，但往往能提升整体性能。
- 用shared memory换寄存器：有些中间结果可以放到shared memory里，而不是寄存器。访问shared memory比spill到local memory快得多。
- 循环展开要适度：循环展开会增加寄存器压力。我见过有人把循环展开8倍，结果寄存器爆了，性能反而下降。
- 使用__launch_bounds__做精细调优：先测出最优的occupancy，然后反推寄存器限制。
经验之谈：寄存器优化没有银弹。每个kernel都需要单独调优。我一般会先跑一个baseline，然后逐步调整寄存器限制，观察occupancy和性能的变化曲线。
寄存器与occupancy的权衡
这里有个经典的权衡问题：寄存器多，occupancy低，但每个线程计算快。寄存器少，occupancy高，但每个线程计算慢。
到底哪个好？取决于你的kernel是计算密集型还是访存密集型。
| kernel类型 | 推荐策略 | 原因 |
|---|---|---|
| 计算密集型 | 适当增加寄存器 | 减少spill，提升单线程计算效率 |
| 访存密集型 | 减少寄存器 | 提高occupancy，隐藏访存延迟 |
| 混合型 | 需要实测调优 | 找到occupancy和单线程效率的平衡点 |
我记得有一次优化一个矩阵乘法的kernel。默认用了48个寄存器，occupancy只有25%。我强制限制到32个寄存器，occupancy提升到50%，性能提升了40%。虽然每个线程计算慢了，但更多的线程同时运行，整体吞吐量上去了。
实战：如何诊断寄存器问题
如果你怀疑kernel有寄存器问题，可以按这个步骤排查：
- 用
nvcc --ptxas-options=-v查看寄存器使用量和spill信息 - 用NVIDIA Nsight Compute查看occupancy和寄存器压力
- 尝试不同的__launch_bounds__参数，观察性能变化
- 如果spill严重，考虑重构kernel，减少每个线程的变量数量
小技巧：Nsight Compute的"Register Pressure"分析器会直接告诉你哪些变量被spill了。点进去看，一目了然。
知识体系总览
下面这张图展示了寄存器优化的核心逻辑和决策路径：
嗯，这张图把寄存器优化的核心脉络理清楚了。从四个方向入手，最终目标都是找到occupancy和单线程效率的平衡点。
寄存器优化说难不难，说简单也不简单。关键是要理解硬件的工作原理，然后用工具去验证你的猜想。我建议你每次修改后都跑一下profiler，看看occupancy和spill的变化。数据不会骗人。
想要更多资料，联系可可