4、寄存器（Registers）：最快的存储，但数量有限，如何避免寄存器溢出

4.1 寄存器到底有多快？

寄存器，说白了就是GPU里离计算单元最近的那一小块存储。有多近？就在ALU旁边，一个时钟周期就能读写。我经常跟团队里的新人说：寄存器是零延迟的黄金存储，你用好了它，性能就赢了一半。

为什么这么快？因为寄存器是直接用触发器（Flip-Flop）实现的，没有缓存、没有寻址延迟、没有bank冲突。你想想看，L1缓存访问要几十个周期，全局内存要几百个周期，而寄存器——一个周期搞定。这差距，不是一星半点。

4.2 寄存器数量：一个硬约束

但寄存器有个致命问题——数量极其有限。每个SM（流多处理器）就那么几十KB的寄存器空间，分给每个线程就更少了。以NVIDIA的架构为例：

• Maxwell/Pascal：每个SM 64KB寄存器，最多1024个线程，每个线程最多255个寄存器
• Volta/Turing：每个SM 64KB寄存器，最多1024个线程，每个线程最多255个寄存器
• Ampere：每个SM 64KB寄存器，最多1024个线程，每个线程最多255个寄存器

嗯，这里要注意：255个寄存器是上限，但实际能用多少取决于你的kernel。编译器会根据你的代码自动分配，你写多了变量，它就会多用寄存器。

4.3 寄存器溢出：性能杀手

我在项目中遇到过最典型的性能问题，就是寄存器溢出。什么叫溢出？就是你的kernel需要的寄存器数量超过了硬件能提供的上限，编译器被迫把一部分变量存到本地内存（Local Memory）里。

本地内存听起来像内存，实际上它被映射到了全局内存，还经过L1/L2缓存。你想想看，本来一个周期能搞定的事，现在要几百个周期——这性能能好吗？

怎么判断有没有溢出？用nvcc编译时加--ptxas-options=-v参数，它会告诉你每个线程用了多少寄存器：

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

// 输出示例：
// ptxas info: Used 32 registers, 0 bytes lmem, 0 bytes smem
// 如果看到 lmem 不为0，说明有溢出

4.4 如何避免寄存器溢出？

避免溢出，说白了就是控制每个线程的寄存器使用量。我总结了几个实用方法：

4.4.1 限制寄存器数量

用__launch_bounds__或maxrregcount强制限制每个线程的寄存器数：

// 方法1：在kernel声明时限制
__global__ __launch_bounds__(256, 4) 
void myKernel() {
    // 编译器会尽量控制在32个寄存器以内
}

// 方法2：编译时全局限制
// nvcc --maxrregcount=32 mykernel.cu

我个人习惯用__launch_bounds__，因为它更精细，可以针对不同kernel单独设置。但要注意：限制太狠也不行。我曾经试过把寄存器限制到16个，结果编译器把大量变量溢出到本地内存，性能反而更差。

4.4.2 减少局部变量

这个道理很简单：变量越少，寄存器用得越少。但实际操作中，很多人不注意：

// ❌ 不好的写法：太多临时变量
float a = data[i];
float b = data[i+1];
float c = a * b;
float d = c + offset;
float e = d * scale;
result[i] = e;

// ✅ 好的写法：复用变量
float tmp = data[i] * data[i+1];
tmp = tmp + offset;
result[i] = tmp * scale;

4.4.3 使用数组索引代替多个变量

如果你有多个相似的变量，用数组代替。编译器对数组的寄存器分配更高效：

// ❌ 不好的写法：多个独立变量
float x0, x1, x2, x3, x4, x5, x6, x7;
x0 = ...; x1 = ...; // 每个变量占一个寄存器

// ✅ 好的写法：使用数组
float x[8];
for(int i = 0; i < 8; i++) {
    x[i] = ...; // 编译器可能复用寄存器
}

4.4.4 把数据移到共享内存

如果确实需要很多变量，可以考虑把一部分数据放到共享内存。虽然共享内存比寄存器慢，但比本地内存快得多：

__shared__ float shared_data[256];

// 把频繁使用的数据放到共享内存
shared_data[threadIdx.x] = input_data[threadIdx.x];
// 后续操作使用 shared_data 而不是局部变量

4.5 寄存器与Occupancy的权衡

这里有个经典问题：寄存器用得多，occupancy就低。为什么？因为SM的寄存器总量是固定的，每个线程用得多，能同时运行的线程就少。

举个例子：SM有64KB寄存器，每个线程用32个寄存器，那就能跑2048个线程（64KB/32B=2048）。但如果每个线程用64个寄存器，就只能跑1024个线程。

但别急着把寄存器压到最低。我记得有一次优化，把寄存器从64压到32，occupancy翻倍了，但每个线程的计算效率也下降了（因为溢出）。最后测下来，64个寄存器的版本反而更快。

所以，别盲目追求低寄存器数。正确的做法是：

1. 先写一个功能正确的版本
2. 用profiler看寄存器使用量和occupancy
3. 尝试不同的maxrregcount值
4. 选吞吐量最高的那个配置

4.6 实战：一个寄存器优化的例子

最后，分享一个我实际做过的优化案例。一个矩阵乘法的kernel，最初每个线程用了80个寄存器：

// 优化前：80个寄存器
__global__ void matmul_v1(float *A, float *B, float *C, int N) {
    float sum = 0.0f;
    float a_reg[4], b_reg[4]; // 每个线程处理4x4块
    // ... 大量临时变量
}

我做了两件事：

• 把部分中间结果存到共享内存
• 用__launch_bounds__限制到48个寄存器

// 优化后：48个寄存器
__global__ __launch_bounds__(256, 4)
void matmul_v2(float *A, float *B, float *C, int N) {
    __shared__ float tile_A[16][16];
    __shared__ float tile_B[16][16];
    float sum = 0.0f;
    // 利用共享内存减少寄存器使用
}

结果呢？occupancy从25%提升到50%，性能提升了1.8倍。嗯，这就是寄存器优化的魅力。

记住：寄存器优化不是越少越好，而是刚刚好。多跑几次profiler，找到那个最佳平衡点，你的kernel就能飞起来。