寄存器压力与Spilling：寄存器分配机制

寄存器这东西，说白了就是GPU里最快的那一小块存储空间。每个SM（流多处理器）上的寄存器文件是有限的，比如A100每个SM有65536个寄存器。你写的每个线程能分到多少，取决于你开了多少个线程。

我刚开始做CUDA优化时，总觉得寄存器越多越好。恨不得把所有中间变量都塞进寄存器。后来发现，这想法太天真了。寄存器用多了，反而会拖慢性能。

编译器在编译你的核函数时，会做寄存器分配。它会分析哪些变量是热点，哪些变量生命周期短。然后尽量把热点变量放在寄存器里。

但编译器不是万能的。它不知道你的数据访问模式，也不知道你的线程块大小。所以有时候它会做出不太聪明的决定。

关键点：每个线程能使用的寄存器数量上限，由编译器根据你的代码和编译选项决定。你可以用 __launch_bounds__ 或 -maxrregcount 来限制。

当你的代码需要的寄存器数量超过硬件限制时，会发生什么？

嗯，编译器会把一部分寄存器变量"溢出"到局部内存中。局部内存听起来像内存，但实际上它存在显存里，而且没有缓存！

我遇到过这样一个项目：一个简单的矩阵乘法核函数，性能死活上不去。用nvcc的--ptxas-options=-v一看，每个线程用了64个寄存器，但SM上只能给48个。结果大量spilling，性能直接腰斩。

spilling的代价有多大？

说白了，一次spilling操作，相当于你本来1纳秒能拿到数据，现在要等几百纳秒。如果spilling频繁发生，你的核函数性能会惨不忍睹。

注意：spilling不仅增加延迟，还会增加显存带宽压力。因为每个线程都在读写局部内存，而局部内存没有缓存，直接走显存带宽。我曾经见过一个核函数因为spilling，显存带宽利用率从80%掉到20%。

怎么控制寄存器压力？最直接的方法就是调整线程块大小。

每个SM的寄存器总数是固定的。你开的线程越多，每个线程能分到的寄存器就越少。反过来，线程越少，每个线程能用的寄存器就越多。

举个例子：

SM寄存器总数	线程块大小	每个线程可用寄存器	是否spilling
65536	1024	64	可能（取决于代码）
65536	512	128	通常不会
65536	256	256	几乎不会

但注意，线程块太小也有问题。比如256个线程的块，SM上可能只能同时跑2个块，总共512个线程。而1024线程的块，SM上能跑1个块，总共1024个线程。线程少了，隐藏延迟的能力就弱了。

我个人的习惯是：先写一个功能正确的核函数，然后用--ptxas-options=-v看寄存器使用量。如果发现spilling，就尝试减小线程块大小，或者用__launch_bounds__限制寄存器数量。

// 用 __launch_bounds__ 限制每个线程最多使用32个寄存器
__global__ void __launch_bounds__(256, 2) myKernel(float* data) {
    // 你的代码
}

这里256是线程块大小，2是每个SM上最多同时运行的块数。编译器会根据这个信息调整寄存器分配。

小技巧：用nvcc --ptxas-options=-v编译时，会输出类似这样的信息：

ptxas info: Used 32 registers, 0 bytes smem, 48 bytes cmem[0]

如果看到spill字样，说明有寄存器溢出。这时候就要调整了。

我曾经在一个图像处理项目里，核函数用了48个寄存器，但SM上每个线程只能分到32个。结果spilling严重，性能只有理论峰值的30%。我把线程块从512改成256，每个线程可用寄存器变成64个，spilling消失了，性能直接翻倍。

你想想看，只是改了一个数字，性能就翻倍了。这就是寄存器压力控制的威力。

这张图把寄存器压力与Spilling的核心逻辑串起来了。左边是寄存器分配的手段，中间是Spilling的代价，右边是实战控制方法。底部是最终目标：在避免Spilling的同时，保持足够的线程占用率来隐藏延迟。

记住，没有银弹。每个核函数都需要你亲自去调。用工具看寄存器使用量，用实验找最佳线程块大小。这是我做CUDA优化这么多年，最实在的经验。

想要更多资料，联系可可