当前位置：首页 > news >正文

CUDA NCU Occupancy学习笔记

news 2025/6/25 9:39:34

占用率是每个多处理器的活跃 Warp 数量与最大可能活跃 Warp 数量的比率。另一种查看占用率的方式是，硬件处理 Warp 的能力中，实际使用 Warp 的百分比。较高的占用率并不一定能带来更高的性能，然而，较低的占用率总是会降低隐藏延迟的能力，从而导致整体性能下降。执行过程中，理论占用率与实际占用率之间的巨大差异通常表明工作负载高度不平衡。

https://www.bilibili.com/video/BV1WNrHY5EeP

https://blog.csdn.net/feng__shuai/article/details/125665305

occupancy：占用率

一个SM中理论最大能支持M个active warp，实际上受限于SM中的资源只能支持N个

occupancy = N / M

占用率指的是 某个 GPU Streaming Multiprocessor（SM） 上实际处于活动状态的 warp 数，占该 SM 理论上能同时容纳的最大 warp 数的百分比。
warp 就是一组 32 条线程的集合；在一个 SM 里同时挂起的 warp 越多，就越能在一次指令延迟（如访存）期间切换到其他 warp，从而隐藏延迟、提高并行度。

SM可以同时处理多个warp，受限于资源（执行单元、寄存器、共享内存等），并不是所有的warp都能同时活跃。

active warp是指正在执行指令的warp，可能处于不同的执行阶段：

（1）正在执行计算任务；

（2）正在等待内存访问结果

（3）正在执行控制流指令

上面是Ada架构（NVIDIA RTX4060、RTX4090等系列GPU）的SM内部结构，一个SM中有128个Core，所以至少能放下128个线程，对应4个warp。

为了隐藏延迟，需要在一个SM中放入更多的warp。

根据显卡架构的不同，可以计算出一个SM中最多承载的warp数量。

（1）设置的block的大小。

同一个block必须跑在同一个SM中，同一个block不能跑在不同的SM上。同一个SM中可以容纳多个block。

（2）SM中寄存器文件的大小

CUDA延迟隐藏机制为：当一个warp处于等待状态，会立刻切换到下一个warp，保证SM中的计算单元cuda core上始终存在warp在执行，不至于cuda core存在空闲的时刻。GPU有大量的寄存器，每个thread都有自己的寄存器，不存在线程切换时切换上下文的开销，CUDA中warp切换的开销约为0。

（3）SM中共享内存的大小

（1）（2）（3）可以分别计算出一个SM中最多承载的warp数量，取三者之间的最小值。

可以使用CUDA Occupancy Calculator（CUDA 占用率计算器）计算Occupancy

一颗 GPU 里有很多 SM（Streaming Multiprocessor），每个 SM 像一间教室，里面坐着学生（线程）分组（warp）上课。每项上限就是教室的“消防规章”，决定一次能塞多少人、发多少书（寄存器）、用多大黑板（共享内存）。你写 kernel 时的线程块(block) 就得在这些规章内排座位，否则编译器/驱动直接报 “out of resources”。

#	Property (截图原文)	物理含义	对 kernel/占用率(occupancy) 的影响
1	Threads per Warp = 32	CUDA 规定 32 条线程组成一个执行单元 warp，硬件线程调度都是以 warp 为粒度。	决定了许多数值必须是 32 的倍数（如寄存器分配、分支同步）。
2	Max Warps per Multiprocessor = 48	一个 SM 最多同时“挂起”48 个 warp。	48 warp × 32 线程 ≈ 1536 线程/SM 就是下条上限。若寄存器/共享内存用太多导致只能挂更少 warp，occupancy 会降。docs.nvidia.com
3	Max Thread Blocks per Multiprocessor = 24	每个 SM 最多可驻留 24 个线程块（block）。	如果你的 block 很小，理论上 24 个也封顶；如果很大，往往被寄存器或线程数先卡住。docs.nvidia.com
4	Max Threads per Multiprocessor = 1536	48 warp×32 线程 = 1536。硬件能同时保留这么多线程上下文。	这是占用率 100 % 时的“天花板”。线程多≠更快，超过隐藏延迟所需就没收益。
5	Maximum Thread Block Size = 1024	单个 block 里线程总数上限。	典型调参：128 - 512 之间试；要破 1024 只能改算法拆块。
6	Registers per Multiprocessor = 65536 (32-bit)	每个 SM 共有 64 K 个 32 位寄存器。	大内存还债最贵？不是，全局访存才贵。寄存器是最快的，但总量固定——线程太多或每线程寄存器太多都会抢这个池子。docs.nvidia.com
7	Max Registers per Thread Block = 65536	同一个 block 能分到的寄存器总量上限（也是全部 64 K）。	如果 block 本身就把 64 K 寄存器吃光，只能独占 SM，occupancy 直接锁死到 1 块。
8	Max Registers per Thread = 255	单线程能用的寄存器数上限（编译器会给 256 向上取整）。	>255 时编译报错。高寄存器/线程可能触发溢出(spill)，把寄存器内容写到慢得多的本地内存。docs.nvidia.com
9	Shared Memory per Multiprocessor = 102 400 B (≈100 KB)	每个 SM 可用的共享内存总量（Ada 架构把它和 L1 Cache 共享，可动态“分房间”）。	这 100 KB 由所有驻留块瓜分，算 occupancy 时必须够用。docs.nvidia.com
10	Max Shared Memory per Block = 102 400 B	单个 block 能申请的共享内存。Runtime 还会帮你占掉 1 KB（见条 14），所以有效是 100 KB-1 KB ≈ 99 KB。	如果你申请 >48 KB，需要 `cudaFuncSetAttribute` 显式“开闸”才能发射 kernel。docs.nvidia.com
11	Register Allocation Unit Size = 256	寄存器按 256 个为最小“页” 分配给一个 warp。不够整页会向上取整。	例：一个 warp 真正只需 33 ×寄存器，硬件仍按 256 给；浪费越多，可驻留 warp 就越少。forums.developer.nvidia.com
12	Register Allocation Granularity = warp	寄存器配额以 warp 为粒度，而不是线程或 block。	这就是为啥上一条要整页向上取整。forums.developer.nvidia.com
13	Shared Memory Allocation Unit Size = 128 B	共享内存也是按 128 字节对齐、向上取整。	所以即便你只用 1 字节，也会吃掉 128 B，可能让多个 block 无法并存。hanhaowen.github.io
14	Warp Allocation Granularity = 4	当计算寄存器 & 共享内存限制时，warp 数会向上取到 4 的倍数。	举例：21 warp → 硬件按 24 warp 计；多出来的“幽灵 warp”仍要算寄存器，所以占用率计算会更紧。forums.developer.nvidia.com
15	Shared Memory Per Block (CUDA runtime use) = 1024 B	CUDA 运行时保留的 1 KB“私房钱”，用于内部参数、动态调度数据等。	实际可用共享内存 = 条 10 上限 − 1024 B。保证 kernel 不会把自己撑爆。docs.nvidia.com