1. 寄存器分配概述：什么是寄存器分配、为什么重要、寄存器分配与性能的关系

1.1 什么是寄存器分配？

寄存器分配，说白了就是编译器在生成机器码时，决定把哪些变量放到物理寄存器里，哪些放到内存里。

你想想看，一个GPU核函数里可能有几十个甚至上百个变量。但物理寄存器就那么几十个（比如NVIDIA的寄存器文件每个线程最多255个）。怎么把这些变量塞进有限的寄存器？这就是寄存器分配要解决的问题。

我刚开始做编译器后端时，觉得这步就是个「映射」工作。后来踩了坑才明白——寄存器分配是编译器优化中最复杂、最影响性能的环节之一。

1.2 为什么寄存器分配如此重要？

原因很简单：寄存器和内存的速度差距太大了。

在GPU上，寄存器访问延迟大约是1个时钟周期。而local memory（虽然叫「local」，但实际在DRAM里）的延迟是几百个时钟周期。global memory更慢，延迟能达到几百到上千个周期。

我曾经优化过一个矩阵乘法的核函数。最初版本用了大量local memory，性能惨不忍睹。后来通过寄存器分配优化，把关键变量都塞进寄存器，性能直接提升了3倍多。

为什么会这样？因为每次访问内存，GPU都要等待几百个周期。如果频繁溢出，你的核函数大部分时间都在「等数据」，而不是「算数据」。

存储层次	访问延迟	带宽	容量（每线程）
寄存器	~1 cycle	极高	255个（Maxwell+）
Shared Memory	~30 cycles	高	48KB~96KB（每块）
Local Memory	~500 cycles	低	512KB（每线程）
Global Memory	~800 cycles	低	几十GB

嗯，这里要注意：寄存器分配直接影响GPU的occupancy（占用率）。每个SM（流多处理器）的寄存器总数是固定的。如果每个线程用太多寄存器，能同时运行的线程块数量就会减少。

1.3 寄存器分配与性能的关系

寄存器分配对性能的影响，主要体现在三个方面：

1. 访存延迟：寄存器溢出导致额外的内存访问，增加延迟
2. Occupancy：寄存器使用量影响并行度
3. 指令调度：寄存器分配影响指令级并行（ILP）

我个人的习惯是，在优化核函数时，先看寄存器使用量。如果超过32个，就要警惕了。超过64个，基本可以断定性能会受影响。

举个例子，一个简单的向量加法核函数：

__global__ void vec_add(float *a, float *b, float *c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        float a_val = a[idx];
        float b_val = b[idx];
        float c_val = a_val + b_val;
        c[idx] = c_val;
    }
}

这个核函数只用了3个寄存器（idx, a_val, b_val）。编译器可能会把c_val复用a_val的寄存器。所以实际寄存器使用量是3个。occupancy可以达到100%。

但如果改成这样：

__global__ void vec_add_complex(float *a, float *b, float *c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        float a1 = a[idx];
        float b1 = b[idx];
        float c1 = a1 + b1;
        float a2 = a[idx+1];
        float b2 = b[idx+1];
        float c2 = a2 + b2;
        // ... 更多变量
        c[idx] = c1;
        c[idx+1] = c2;
    }
}

变量一多，寄存器就不够用了。编译器会开始把一些变量溢出到local memory。性能就会下降。

1.4 寄存器分配的基本策略

编译器在做寄存器分配时，主要用两种策略：

• 图着色（Graph Coloring）：把变量看作节点，如果两个变量同时活跃，就在它们之间连一条边。然后用图着色算法分配寄存器。这是最经典的方法。
• 线性扫描（Linear Scan）：按变量的活跃区间顺序分配。实现简单，速度快，适合JIT编译。

GPU编译器（如NVCC）通常使用图着色算法。因为GPU核函数通常比较小，图着色虽然慢一点，但分配质量更高。

我记得有一次调试一个复杂的核函数，发现性能瓶颈竟然是寄存器分配导致的。编译器把两个不相关的变量分配到了同一个寄存器，导致后续指令产生了不必要的依赖。这就是图着色算法没处理好的典型问题。

1.5 小结

寄存器分配不是简单的「变量→寄存器」映射。它直接影响：

• 访存延迟（溢出代价）
• 并行度（occupancy）
• 指令级并行（寄存器依赖）

在后续章节中，我会深入讲解如何通过寄存器分配优化核函数性能。包括如何分析寄存器使用量、如何手动干预寄存器分配、以及如何平衡寄存器和occupancy的关系。

记住一句话：寄存器是GPU上最宝贵的资源之一。用好它，你的核函数性能就能上一个台阶。