阶段 3 · 优化实战09 / 15约 22 分钟
Occupancy：为什么「越高越好」是个误解
占用率的三个限制因素、如何计算、以及什么时候该主动降低它
学完这一课你会
- 算清楚寄存器、共享内存、block 数三者如何共同限制占用率
- 会用 CUDA Occupancy API 自动选取 block 尺寸
- 理解 ILP 可以替代一部分 occupancy 的作用
- 知道 __launch_bounds__ 的用法和适用时机
Occupancy（占用率）= 每个 SM 上实际驻留的 warp 数 / 硬件支持的最大 warp 数。 上一阶段说过，GPU 靠超额并行来隐藏访存延迟，占用率衡量的就是「你给了调度器多少可切换的余量」。但它是手段而非目标——占用率 50% 却跑满带宽的 kernel，比占用率 100% 却在空转的 kernel 好得多。
▸打开 Occupancy 计算器选架构、调 block 尺寸、寄存器数和共享内存用量，直接看出到底是哪个资源在卡你的占用率。这是理解三者相互作用最快的方式。三个限制因素
每个 SM 的资源是固定的，能同时驻留多少个 block 由三者的最小值决定。以 Ampere（sm_80）为例，每 SM 最多 2048 个线程（64 warp）、65536 个 32 位寄存器、164 KB 共享内存、32 个 block。
| 限制因素 | 计算方式 | 举例（block=256） |
|---|---|---|
| 线程数上限 | 2048 / blockDim | 2048 / 256 = 8 个 block |
| 寄存器 | 65536 / (regs × blockDim) | 每线程 32 寄存器 → 65536 / 8192 = 8 个 block |
| 共享内存 | 164 KB / 每 block 用量 | 每 block 用 16 KB → 10 个 block |
| block 数上限 | 硬件常量 | 32 个 block |
| 实际驻留 | 取最小值 | min(8, 8, 10, 32) = 8 个 block = 64 warp = 100% |
让 CUDA 帮你算
CUDA C++
1// 1) 让运行时推荐一个能达到最高占用率的 block 尺寸2int minGridSize, blockSize;3CUDA_CHECK(cudaOccupancyMaxPotentialBlockSize(4 &minGridSize, &blockSize, myKernel,5 0, // 动态共享内存字节数（若与 blockSize 相关，改用带回调的重载）6 0)); // block 尺寸上限，0 表示无限制7printf("推荐 blockSize = %d\n", blockSize);8 9// 2) 查询某个配置实际能驻留多少 block10int numBlocks;11CUDA_CHECK(cudaOccupancyMaxActiveBlocksPerMultiprocessor(12 &numBlocks, myKernel, blockSize, 0));13 14cudaDeviceProp prop;15CUDA_CHECK(cudaGetDeviceProperties(&prop, 0));16float occupancy = (numBlocks * blockSize / (float)prop.warpSize)17 / (prop.maxThreadsPerMultiProcessor / (float)prop.warpSize);18printf("理论占用率 = %.1f%%\n", occupancy * 100.0f);用 __launch_bounds__ 约束编译器
编译器默认会尽量多用寄存器来减少指令数，但它不知道你打算用什么 block 尺寸。__launch_bounds__ 就是把这个信息告诉它：编译器会据此反推每线程的寄存器预算，必要时主动减少寄存器占用（代价可能是少量溢出，但换来更高占用率）。
CUDA C++
1__global__ void __launch_bounds__(256, 4) myKernel(float* data) {2 // 承诺：blockDim 不超过 256，且希望每个 SM 至少驻留 4 个 block。3 // 编译器于是把每线程寄存器预算压到 65536 / (256 * 4) = 64 个以内。4 ...5}6 7// 也可以直接卡死上限，不走 __launch_bounds__8// nvcc -maxrregcount=32 ... （全局生效，粒度太粗，一般不推荐）关键：高占用率不等于高性能
Vasily Volkov 在 2010 年那篇著名的 *Better Performance at Lower Occupancy* 里给出了反直觉的结论：通过增加每个线程的独立工作量（ILP，指令级并行），可以在更低的占用率下获得更高的性能。 因为延迟隐藏既可以靠「更多 warp」（TLP），也可以靠「一个 warp 内更多互不依赖的指令」（ILP）。
CUDA C++
1// 低 ILP：每线程一个元素。累加链完全串行，占用率必须很高才能填满流水线2__global__ void lowILP(const float* a, const float* b, float* c, int n) {3 int i = blockIdx.x * blockDim.x + threadIdx.x;4 if (i < n) c[i] = a[i] * b[i];5}6 7// 高 ILP：每线程 4 个元素，4 条乘法互不依赖，可以同时在流水线里飞8__global__ void highILP(const float4* a, const float4* b, float4* c, int n4) {9 int i = blockIdx.x * blockDim.x + threadIdx.x;10 if (i < n4) {11 float4 va = a[i], vb = b[i]; // 一条 128-bit 访存指令顶四条12 float4 vc;13 vc.x = va.x * vb.x; vc.y = va.y * vb.y;14 vc.z = va.z * vb.z; vc.w = va.w * vb.w;15 c[i] = vc;16 }17}- ▸访存密集型 kernel：占用率通常要 50% 以上才能有效隐藏延迟，优先保证它。
- ▸算力密集型 kernel：30%~50% 往往就够了，把资源让给寄存器分块（更大的 tile）反而更划算。
- ▸用了大量寄存器的 kernel（如 GEMM 的寄存器分块）：占用率可能只有 25%，但这是有意为之的正确设计。
- ▸判断标准只有一个：用 Nsight Compute 看
Achieved Occupancy和 warp stall 原因。如果 stall 主要是Long Scoreboard（等访存），那提升占用率有用；如果不是，提了也白提。
自测
先自己在心里回答一遍，再展开对照。答不上来的说明这一段值得重读。