第4章：寄存器优化——寄存器溢出、占用率与寄存器数量的权衡

寄存器，说白了就是GPU里最快的那块存储。我刚开始写CUDA时，总觉得寄存器越多越好，恨不得把所有中间变量都塞进去。结果呢？性能反而下降了。嗯，这里有个坑——寄存器溢出。

4.1 寄存器溢出：看不见的性能杀手

每个SM（流多处理器）的寄存器总量是固定的。比如在Turing架构上，每个SM有65536个32位寄存器。当你一个线程用太多寄存器，编译器就会把部分变量“挤”到局部内存里。这个过程就叫寄存器溢出。

我在项目中遇到过这样一个案例：一个卷积核函数，每个线程用了64个寄存器，占用率只有25%。我试着把一些中间结果重复计算而不是存起来，寄存器降到32个，占用率升到50%，最终性能反而提升了40%。

为什么会这样？因为寄存器溢出带来的访存开销，远大于重复计算的开销。你想想看，一个全局内存访问要400-800个周期，而多算几次乘加才几个周期。

4.2 占用率与寄存器数量的权衡

占用率，就是SM上活跃的线程束数量与最大线程束数量的比值。很多人以为占用率越高越好，其实不然。

我个人习惯是：先看你的kernel是计算密集还是访存密集。如果是访存密集，高占用率能隐藏延迟，我倾向于用32个寄存器以下。如果是计算密集，低占用率反而好，因为每个线程有更多寄存器，减少了溢出。

4.3 使用__launch_bounds__控制寄存器

CUDA提供了__launch_bounds__这个利器，可以告诉编译器你期望的线程块大小和每个SM的最大线程块数。编译器会根据这些信息来分配寄存器。

// 告诉编译器：每个线程块最多256个线程
// 每个SM最多4个线程块
__global__ __launch_bounds__(256, 4)
void myKernel(float* data, int N) {
    // 编译器会尽量让每个线程用 <= 32个寄存器
    // 因为 256 * 4 * 32 = 32768，没超过SM的65536
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < N) {
        float a = data[tid];
        float b = a * 2.0f;
        float c = b + 1.0f;
        data[tid] = c;
    }
}

我曾经接手过一个项目，同事写的kernel性能很差。我一看，每个线程用了80个寄存器，占用率只有12.5%。加上__launch_bounds__(128, 4)后，编译器自动把寄存器压到32个，占用率升到50%，性能翻了一倍。

4.4 寄存器优化的实战策略

这里我总结了几条实战经验，都是踩过坑才明白的：

• 减少局部变量：能用float就别用double，能用int就别用long long。每个寄存器32位，double要占两个。
• 重复计算代替存储：如果某个中间结果只用到两三次，而且计算量不大，直接重新算比存起来更划算。
• 使用#pragma unroll：循环展开有时能减少寄存器使用，因为编译器可以更好地调度。
• 检查编译器报告：用nvcc --ptxas-options=-v看寄存器使用量，如果超过32就要警惕了。

4.5 知识体系图

下面这张图展示了寄存器优化的核心逻辑：

这张图的核心逻辑很简单：先看寄存器使用量，超过32就查溢出，没超过就看占用率。然后根据情况选择优化策略，最后验证性能是否达标。

4.6 总结

寄存器优化说白了就是一场平衡游戏。寄存器多了，占用率低；寄存器少了，可能溢出。我个人的经验是：先用默认设置跑一遍，看寄存器使用量和占用率。如果寄存器超过32且性能不理想，就用__launch_bounds__压一下。如果寄存器在32以内但占用率很低，那可能是别的问题，比如线程块大小设置不合理。

记住一点：别迷信任何经验法则。每个kernel都有自己的特点，多测几组参数，找到最适合你场景的配置。