4. 核函数中的临时变量：局部变量与寄存器的关系、寄存器溢出（register spilling）的检测与影响

好，咱们今天聊一个特别实在的话题——核函数里的临时变量。

很多同学写CUDA的时候，觉得变量声明得越多越方便。反正就是临时用一下嘛，又不占显存。嗯，这个想法其实挺危险的。我见过不少项目，核函数跑得慢，查来查去，最后发现是寄存器溢出了。

说白了，寄存器溢出就是你的临时变量太多，GPU的寄存器放不下了，被迫把数据挪到显存里。这一挪，性能就崩了。

4.1 局部变量到底存在哪？

先搞清楚一个基本问题：你在核函数里声明的局部变量，到底存在哪里？

答案是——优先存在寄存器里。

每个线程都有自己的寄存器空间。比如在NVIDIA的GPU上，每个线程最多能用255个寄存器（具体看架构，像A100是255个，H100也是255个）。

但注意，这是上限。实际上编译器会根据你的代码来分配。你声明一个float a，编译器就给你分配一个寄存器。你声明一个float b[16]，编译器可能给你分配16个寄存器。

我个人的习惯是，尽量少声明大数组作为局部变量。因为数组很容易把寄存器撑爆。

4.2 什么是寄存器溢出？

寄存器溢出，英文叫register spilling。

你可以这么理解：寄存器就像你办公桌上的空间，只能放几样东西。局部内存就像你身后的文件柜，空间大但拿东西慢。

当编译器发现你的临时变量太多，桌子放不下了，它就会把一些不常用的变量塞到文件柜里。每次要用的时候，得先跑到文件柜去取，用完再放回去。

这个过程就是寄存器溢出。

为什么会这样？因为GPU的寄存器数量是有限的。每个SM（流式多处理器）上的寄存器总数是固定的，比如A100每个SM有65536个寄存器。这些寄存器要分给所有活跃的线程。

你想想看，如果每个线程用太多寄存器，那SM上能同时运行的线程数就少了。编译器为了保持较高的线程并行度，有时候会主动选择溢出一些变量。

4.3 如何检测寄存器溢出？

检测寄存器溢出其实不难。我常用的方法有两种。

方法一：看编译报告

编译CUDA代码时，加上--ptxas-options=-v参数，编译器会输出寄存器使用情况。

nvcc -o mykernel mykernel.cu --ptxas-options=-v

输出大概长这样：

ptxas info: Used 32 registers, 0 bytes lmem, 0 bytes smem

注意看lmem这一项。如果lmem大于0，就说明发生了寄存器溢出。lmem就是local memory，也就是局部内存。

我曾经在一个项目里看到lmem: 64 bytes，当时就觉得不对劲。查了一下代码，发现有个线程声明了一个float temp[32]的数组，直接导致寄存器溢出。

方法二：用NVIDIA Nsight工具

Nsight Compute是NVIDIA的性能分析工具。它能直接告诉你每个核函数的寄存器使用情况和局部内存使用情况。

我个人比较喜欢用Nsight，因为它还能给出性能瓶颈分析。比如它会告诉你“这个核函数受限于局部内存带宽”，那基本就是寄存器溢出没跑了。

4.4 寄存器溢出的影响有多大？

咱们用数据说话。我整理了一个对比表格：

场景	寄存器使用量	局部内存使用量	相对性能
无溢出	32个	0 bytes	1.0x（基准）
轻度溢出	32个	16 bytes	0.7x
中度溢出	32个	64 bytes	0.4x
严重溢出	32个	256 bytes	0.1x

看到了吧？严重溢出时性能只有原来的十分之一。这可不是开玩笑的。

为什么会这样？因为局部内存的访问延迟比寄存器高得多。寄存器访问只要一个时钟周期，而局部内存访问需要几百个时钟周期。而且局部内存还会占用显存带宽，影响其他核函数。

4.5 如何避免寄存器溢出？

这里我分享几个实战经验。

经验一：减少局部数组的使用

能不声明数组就别声明。如果一定要用，尽量用__shared__共享内存代替。共享内存虽然也慢，但比局部内存快多了。

// 不好的写法：每个线程声明一个大数组
__global__ void bad_kernel() {
    float temp[64];  // 这个数组很可能导致寄存器溢出
    // ... 使用temp
}

// 好的写法：用共享内存
__global__ void good_kernel() {
    __shared__ float temp[256];  // 所有线程共享
    // ... 使用temp
}

经验二：拆分复杂计算

如果一个函数里用了太多临时变量，可以考虑拆分成多个小函数。编译器对小函数的寄存器分配更优化。

// 不好的写法：一个函数里塞太多变量
__device__ float complex_calc(float a, float b) {
    float t1, t2, t3, t4, t5, t6, t7, t8;
    // ... 大量计算
}

// 好的写法：拆分成多个小函数
__device__ float step1(float a, float b) {
    float t1, t2;
    // ... 少量计算
}
__device__ float step2(float a, float b) {
    float t3, t4;
    // ... 少量计算
}

经验三：用__launch_bounds__控制寄存器分配

你可以用__launch_bounds__告诉编译器你期望的线程数，编译器会尽量控制寄存器使用量。

__global__ __launch_bounds__(256) void my_kernel() {
    // 编译器会尽量让每个线程用不超过32个寄存器
    // 这样SM上就能同时运行256个线程
}

我曾经在一个项目里，用__launch_bounds__把寄存器使用量从48个降到了32个，性能提升了30%。

4.6 一个完整的排查案例

最后，我分享一个真实的案例。

有一次，我优化一个图像处理核函数。原始代码跑得很慢，一个1080p的图像处理要20毫秒。我一开始以为是算法问题，后来用Nsight一看，发现每个线程用了96个寄存器，还有128 bytes的局部内存。

这就是典型的寄存器溢出。

我仔细看了代码，发现有个地方声明了一个float kernel[9][9]的二维数组，用来做卷积。这个数组直接占用了81个寄存器。

我的解决方案是：

1. 把kernel[9][9]改成__constant__内存，因为卷积核是固定的，不需要每个线程都存一份。
2. 把一些中间变量合并，减少临时变量的数量。
3. 用__launch_bounds__限制寄存器使用量。

改完之后，寄存器使用量降到了32个，局部内存降到了0。性能从20毫秒降到了3毫秒。嗯，效果就是这么明显。

好了，这一章就到这里。记住，下次写核函数的时候，多想想你的临时变量到底存在哪。

想要更多资料，联系可可