📚 课程目录第4章占用率定义

4. 占用率定义：占用率计算公式、理论占用率与实际占用率

4.1 核心概念：从“占用”到“占用率”

在CUDA编程模型中，占用率（Occupancy）衡量的是GPU计算资源（尤其是SM内的执行单元）被活跃线程覆盖的程度。它并不直接代表“GPU有多忙”，而是反映每个SM上驻留的线程块（Block）与线程束（Warp）数量相对于硬件上限的比例。高占用率意味着SM中有足够多的Warp可供调度器切换，从而用计算掩盖访存延迟，是隐藏延迟、提升吞吐量的关键指标。

占用率的本质是资源供需匹配：硬件提供固定容量的寄存器文件、共享内存、线程槽位，而每个Block会按需申请这些资源。占用率就是“当前资源约束下，实际能塞进SM的Warp数”与“硬件能容纳的最大Warp数”的比值。

关键认知：占用率不是“性能分数”，而是“延迟隐藏能力”的代理指标。100%占用率并不保证最优性能，但过低的占用率（如低于30%）通常意味着延迟无法被有效掩盖。

4.2 占用率计算公式与资源约束模型

占用率的精确计算需要同时考虑四类硬件资源上限，最终占用率取四者中的最小值。以NVIDIA Ampere架构（如A100）为例，每个SM的典型资源为：

最大线程数：2048个线程（即64个Warp）
最大线程块数：32个Block（部分架构为16或24）
寄存器文件：65536个32位寄存器（256KB）
共享内存：164KB（可配置为最多228KB，但需与L1缓存平衡）

设每个Block包含T个线程，每个线程使用R个寄存器，每个Block使用S字节共享内存，则各资源约束下的最大Block数为：

约束1（线程槽位）: floor(2048 / T)
约束2（Block槽位）: 32（或硬件上限）
约束3（寄存器）:   floor(65536 / (T * R))
约束4（共享内存）: floor(164KB / S)

实际驻留Block数 = min(约束1, 约束2, 约束3, 约束4)
实际占用率 = (实际驻留Block数 * T / 2048) * 100%

注意：寄存器是逐线程分配的，而共享内存是逐Block分配的。因此，当寄存器成为瓶颈时，减小R（如通过__launch_bounds__或maxrregcount）能直接提升占用率；当共享内存成为瓶颈时，则需减小Block的共享内存申请量或调整Block尺寸。

4.3 理论占用率与实际占用率

理论占用率是在编译期根据内核的静态资源声明（寄存器数、共享内存量、Block尺寸）计算出的理想值，它假设所有资源都能被完全利用，且不存在运行时波动。CUDA提供了cudaOccupancyMaxActiveBlocksPerMultiprocessor API来计算理论值：

int numBlocks;
cudaOccupancyMaxActiveBlocksPerMultiprocessor(
    &numBlocks,   // 输出：每个SM可驻留的Block数
    my_kernel,    // 内核函数指针
    blockSize,    // 线程块尺寸
    smemSize);    // 动态共享内存字节数
// 理论占用率 = numBlocks * blockSize / 2048.0f

实际占用率则是在运行时通过cudaOccupancyMaxActiveBlocksPerMultiprocessor结合设备属性（如cudaDevAttrMaxBlocksPerMultiprocessor）以及实际分配到的资源（可能因驱动版本、上下文开销、其他并发内核而减少）计算出的真实值。实际占用率往往低于理论值，原因包括：

驱动或运行时为系统预留了少量寄存器或共享内存
内核使用了动态共享内存，但分配粒度导致浪费
Block尺寸不是Warp大小（32）的整数倍，造成线程槽位碎片
同时运行多个内核时，SM资源被分割

实用建议：使用NVIDIA Nsight Compute的“Occupancy”分析页，可以同时查看理论占用率（绿色虚线）与实际占用率（蓝色柱状图），并直接定位是哪个资源（寄存器/共享内存/Block槽位）限制了占用率。

4.4 案例：矩阵乘法内核的占用率计算

考虑一个简单的矩阵乘法内核，每个Block包含256个线程，每个线程使用40个寄存器，每个Block申请8KB共享内存（用于分块缓存）。在A100（SM资源：2048线程/32Block/65536寄存器/164KB共享内存）上：

约束1（线程）: floor(2048 / 256) = 8 个Block
约束2（Block）: 32 个Block
约束3（寄存器）: floor(65536 / (256 * 40)) = floor(6.4) = 6 个Block
约束4（共享内存）: floor(164KB / 8KB) = 20 个Block

实际驻留Block数 = min(8, 32, 6, 20) = 6 个Block
实际占用率 = (6 * 256 / 2048) * 100% = 75%

这里寄存器成为瓶颈。若通过__launch_bounds__(256, 8)强制编译器将每线程寄存器数限制在32个以内，则约束3变为floor(65536 / (256*32)) = 8，占用率提升至100%。但需注意，寄存器减少可能导致局部变量溢出到本地内存，反而降低性能，因此需要权衡。

资源类型	每Block需求	SM总容量	可驻留Block数	瓶颈？
线程槽位	256线程	2048线程	8	否
Block槽位	1	32	32	否
寄存器	256×40=10240	65536	6	是
共享内存	8KB	164KB	20	否

4.5 误区与复盘：占用率调优的常见陷阱

误区1：盲目追求100%占用率。当内核本身是计算密集型（算术逻辑单元已饱和）时，提高占用率不会带来额外收益，反而可能因寄存器溢出或缓存抖动导致性能下降。正确做法是先用Profile确认瓶颈是延迟还是吞吐。

误区2：忽略Block尺寸对占用率的影响。例如，Block尺寸设为128时，线程约束为floor(2048/128)=16，但若寄存器约束为10，则占用率仍受限于寄存器。调整Block尺寸有时能打破资源约束的“整数除法”困境。

误区3：混淆“理论占用率”与“实际性能”。理论占用率只反映资源驻留能力，不反映访存局部性、指令级并行或分支发散。一个占用率90%但访存完全随机的内核，可能比占用率60%但完美合并访存的内核慢得多。

复盘建议：每次调优时，先记录理论占用率、实际占用率、SM活跃Warp数、内存吞吐四个指标。若实际占用率远低于理论值，优先检查是否有隐式同步、动态共享内存分配失败或驱动限制；若两者接近但性能不佳，则应将注意力转向内存访问模式和指令混合，而非继续堆占用率。

4.6 小结

占用率是CUDA性能调优的第一道门槛，它决定了SM能否用足够的Warp掩盖内存与指令延迟。掌握四类资源约束的最小值计算，区分理论占用率（编译期静态值）与实际占用率（运行时真实值），是避免“纸上谈兵”的关键。在实际调优中，应始终结合Profile数据，让占用率服务于真实瓶颈——当延迟主导时提高占用率，当吞吐饱和时转而优化访存与指令效率。