第四章：寄存器优化——寄存器溢出处理、循环展开策略、指令级并行

大家好，欢迎来到第四章。这一章我们聊寄存器优化。

说实话，寄存器是GPU上最宝贵的资源，没有之一。我见过太多人，代码写得挺漂亮，一跑性能就是上不去。为什么？寄存器溢出了。说白了，就是你的变量太多，寄存器放不下，被迫往局部内存里塞。这一塞，延迟直接飙升几十倍。

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

先讲个我自己的经历。几年前优化一个期权定价的kernel，一开始每个线程用了大概80个寄存器。嗯，看起来还行？结果一测，occupancy只有25%。我当时就纳闷，明明计算量不大啊。后来用--ptxas-options=-v一看，好家伙，寄存器溢出严重，大量数据被spill到local memory里去了。

寄存器溢出（Register Spilling）是什么？简单说，就是编译器发现你需要的寄存器数量超过了硬件限制，它只好把一部分变量存到显存里。每次访问这些变量，都要走一遍内存通路。你想想看，寄存器访问延迟就几个周期，local memory可是几百个周期起步。

4.1.1 如何检测寄存器溢出

方法其实很简单。编译的时候加上这几个flag：

nvcc -arch=sm_80 -Xptxas=-v kernel.cu

输出里你会看到类似这样的信息：

ptxas info: Used 64 registers, 0 bytes spill stores, 0 bytes spill loads

如果spill stores和spill loads不是0，那就说明溢出了。数值越大，问题越严重。

我个人习惯，看到spill就一定要处理。哪怕只溢出一点点，也值得优化。

4.1.2 减少寄存器使用的技巧

• 拆分复杂表达式：把一个大表达式拆成多步计算，编译器有时反而能复用寄存器。
• 使用volatile要谨慎：volatile会阻止编译器优化，容易导致寄存器使用量上升。
• 手动限制寄存器数量：用__launch_bounds__或者编译选项-maxrregcount=32强制限制。

4.2 循环展开策略：用空间换时间

循环展开，说白了就是把循环体复制几份，减少循环控制的开销。但这里有个度的问题。

我记得有一次，一个同事把循环展开了32倍，结果寄存器直接爆了，性能反而下降了30%。这就是典型的「过度展开」。

4.2.1 什么时候该展开

• 循环次数少且固定（比如4、8、16次）
• 循环体内计算简单，控制开销占比大
• 展开后不会导致寄存器溢出

4.2.2 展开的几种方式

手动展开：

// 原始循环
for (int i = 0; i < 4; i++) {
    c[i] = a[i] + b[i];
}

// 手动展开
c[0] = a[0] + b[0];
c[1] = a[1] + b[1];
c[2] = a[2] + b[2];
c[3] = a[3] + b[3];

用#pragma unroll：

#pragma unroll 4
for (int i = 0; i < n; i++) {
    c[i] = a[i] + b[i];
}

我个人更推荐用#pragma unroll，编译器会帮你做权衡。但如果你对性能有极致要求，手动展开更可控。

4.3 指令级并行：让GPU的流水线跑满

指令级并行（ILP）是什么？说白了，就是让GPU在一个时钟周期内尽量多地执行独立指令。GPU的流水线很深，如果指令之间有数据依赖，流水线就会stall，性能就上不去。

我优化过一个矩阵乘法的kernel，原本的代码是：

float a = load(A[i]);
float b = load(B[i]);
float c = a * b;  // 这里必须等a和b都加载完

改成这样之后：

float a0 = load(A[i]);
float a1 = load(A[i+1]);
float b0 = load(B[i]);
float b1 = load(B[i+1]);
float c0 = a0 * b0;
float c1 = a1 * b1;  // 这里a0和a1的加载可以并行

性能提升了大概15%。为什么？因为加载和计算可以重叠了。

4.3.1 提高ILP的常用方法

• 减少数据依赖：把有依赖的指令尽量分开
• 使用独立变量：不要反复使用同一个临时变量
• 循环展开：展开后自然就有了更多的独立指令

4.4 知识体系总览

下面这张图是我自己总结的寄存器优化知识体系，你可以对照着看：

4.5 实战中的平衡艺术

说实话，寄存器优化没有标准答案。每个kernel的硬件环境、计算密度、访存模式都不一样。我一般会按这个顺序来调：

1. 先看occupancy，如果低于50%，优先压寄存器
2. 再看ILP，如果指令stall多，考虑展开
3. 最后看spill，如果展开后溢出，减少展开因子

嗯，这里要注意一点：不要盲目追求occupancy。有些计算密集型的kernel，occupancy低一点反而性能更好，因为每个线程有更多寄存器来存中间结果。

好了，这一章就到这里。下一章我们聊共享内存的bank conflict问题，那个坑更多，我到时候给你们讲讲我踩过的那些雷。

---