4、寄存器与Occupancy：寄存器压力对Occupancy的影响、使用__launch_bounds__限制线程块、寄存器溢出(Register Spill)的检测与优化

各位同学，咱们今天聊一个非常关键的话题——寄存器与Occupancy。

说实话，我在做CUDA优化的头两年，一直把Occupancy当成一个“越高越好”的指标。后来被现实狠狠教育了一顿。有一次我写了一个kernel，Occupancy跑到了100%，结果性能反而比Occupancy只有50%的版本还慢。当时我就懵了。

后来一查，问题出在寄存器上。嗯，今天咱们就把这个坑彻底填平。

4.1 寄存器压力：Occupancy的隐形杀手

先问大家一个问题：一个SM（流式多处理器）能同时运行多少个线程块？

答案不是固定的。它取决于三个资源：寄存器、共享内存、线程块数量上限。其中，寄存器是最容易被忽视的瓶颈。

每个SM的寄存器总量是有限的。以V100为例，每个SM有65536个寄存器。假设你的kernel每个线程用了64个寄存器，那么一个SM最多能容纳：

65536 ÷ 64 = 1024 个线程

但V100每个SM最多支持2048个线程。你看，寄存器直接砍掉了一半的Occupancy。

我个人的习惯是，在写kernel之前，先估算一下每个线程的寄存器用量。怎么估？看代码里的局部变量数量。每个float、int、double都会占用寄存器。如果你写了一大堆临时变量，寄存器压力自然就上去了。

4.2 使用__launch_bounds__限制线程块

CUDA提供了一个非常实用的工具——__launch_bounds__。它可以告诉编译器，你的kernel最多会用多少个线程，以及每个线程块最少需要多少个线程块。

语法很简单：

__global__ void __launch_bounds__(maxThreadsPerBlock, minBlocksPerMultiprocessor)
myKernel(float* data) {
    // kernel代码
}

这里有两个参数：

• maxThreadsPerBlock：每个线程块的最大线程数。编译器会根据这个值来分配寄存器。
• minBlocksPerMultiprocessor：每个SM至少需要多少个线程块。这个参数告诉编译器，你要保证一定的Occupancy。

举个例子。假设你希望每个SM至少跑4个线程块，每个线程块256个线程：

__global__ void __launch_bounds__(256, 4)
myKernel(float* data) {
    // ...
}

编译器会怎么做？它会尽量压缩每个线程的寄存器用量，确保4个线程块能同时运行。如果寄存器不够，编译器会主动把一些变量“溢出”到局部内存。

4.3 寄存器溢出(Register Spill)的检测与优化

寄存器溢出，说白了就是寄存器不够用了，编译器把一些变量存到了局部内存里。局部内存虽然也在显存里，但访问速度比寄存器慢得多。

怎么检测寄存器溢出？

方法一：看编译器的输出。在nvcc编译时加上--ptxas-options=-v参数：

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

输出中会显示：

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

如果lmem（局部内存）不为0，说明发生了寄存器溢出。

方法二：用NVIDIA Nsight Compute分析。在Profiler的“Register”页面，会直接显示每个线程的寄存器用量和溢出情况。

我记得有一次，我写了一个复杂的矩阵乘法kernel，寄存器用量高达96个。用Nsight一看，lmem用了128字节，性能惨不忍睹。后来我做了三件事：

1. 减少局部变量：把一些中间结果直接复用，而不是定义新变量。
2. 使用__launch_bounds__：强制编译器降低寄存器用量。
3. 拆分kernel：把一个大kernel拆成两个小kernel，每个kernel的寄存器压力自然就小了。

优化后，寄存器降到了32个，lmem变成了0，性能提升了2.3倍。

4.4 知识体系图

下面这张图展示了寄存器、Occupancy、性能之间的关系：

4.5 实战建议

最后，给大家几个实战建议：

• 先测后调：不要凭感觉优化。先用Nsight Compute看寄存器用量和Occupancy。
• 合理使用__launch_bounds__：从保守值开始，逐步增加minBlocksPerMultiprocessor。
• 警惕寄存器溢出：lmem不为0时，优先考虑减少局部变量，而不是盲目增加__launch_bounds__的参数。
• 权衡Occupancy与指令级并行：有时候降低Occupancy，让每个线程做更多工作，反而更快。

嗯，今天就到这里。寄存器这块内容比较多，大家回去可以拿自己的kernel试试，看看寄存器用量和Occupancy的关系。有什么问题随时交流。

想要更多资料，联系可可