第十二课：Block 如何分配到 SM，以及 Occupancy 由什么决定

导语：这一课讲 Block 如何进入 SM、哪些硬件资源限制并发度，以及 Occupancy 为什么只是性能指标之一。

前面已经知道：

图示：第十二课：Block 如何分配到 SM，以及 Occupancy 由什么决定

现在需要回答几个关键问题：

1. 一个 SM 能同时运行多少个 Block？
2. 为什么有些 Kernel 并发线程很多，有些很少？
3. Occupancy 到底是什么意思？
4. Occupancy 越高，性能是否一定越好？

---

一、Block 是如何分配到 SM 的

Kernel 启动后，GPU 会把 Grid 中尚未执行的 Block 分配给有空闲资源的 SM。

假设 GPU 有4个 SM，Kernel 有12个 Block：

图示：一、Block 是如何分配到 SM 的

但是，一个 SM 并不一定只能驻留一个 Block。

实际可能是：

SM 0
├── Block 0
├── Block 4
└── Block 8

这些 Block 在逻辑上独立，但它们共同使用该 SM 的硬件资源。

“驻留”表示 Block 的状态和资源已经分配在 SM 上，不代表它的所有 Warp 每个周期都在执行。

---

二、一个 Block 一旦分配到 SM，会发生什么

假设一个 Block 有256个线程。

它会被拆成：

$$256\div32=8$$

个 Warp。

SM 需要为这个 Block 分配：

• 8个 Warp 的执行状态；
• 每个线程需要的寄存器；
• 该 Block 需要的 Shared Memory；
• 同步屏障状态；
• Block 和 Warp 管理项。

只有这些资源全部能够满足时，这个 Block 才能进入该 SM。

因此，Block 调度不是只看“SM 有没有空”，而是看：

SM 剩余的各类资源，能否完整容纳一个新的 Block。

---

三、限制驻留 Block 数量的主要资源

一个 SM 同时能驻留多少 Block，通常受到以下条件共同限制：

线程数量上限
Warp 数量上限
Block 数量上限
寄存器容量
Shared Memory 容量

最终可驻留 Block 数量是所有限制中的最小值。

可以写成：

$$N_{\text{blocks}} = \min \left( N_{\text{thread limit}}, N_{\text{warp limit}}, N_{\text{block limit}}, N_{\text{register limit}}, N_{\text{shared memory limit}} \right)$$

---

四、线程数量限制

假设某个 SM 最多支持2048个驻留线程。

如果每个 Block 有256个线程，则按线程数量计算：

$$2048\div256=8$$

所以最多容纳8个 Block。

如果每个 Block 有1024个线程：

$$2048\div1024=2$$

最多只能容纳2个 Block。

因此，Block 越大，一个 SM 通常能同时驻留的 Block 数量越少。

但一个 Block 太小，也可能导致：

• Warp 数量不足；
• SM 计算资源难以充分利用；
• Grid 中 Block 数量不足时无法覆盖所有 SM。

---

五、Warp 数量限制

SM 通常还有最大驻留 Warp 数量。

假设一个 SM 最多支持64个 Warp。

如果一个 Block 有256个线程，即8个 Warp：

$$64\div8=8$$

最多驻留8个 Block。

如果一个 Block 有512个线程，即16个 Warp：

$$64\div16=4$$

最多驻留4个 Block。

线程限制和 Warp 限制通常密切相关，但由于线程数需要向上取整到完整 Warp，它们并不完全等价。

例如，一个 Block 有100个线程：

$$\left\lceil\frac{100}{32}\right\rceil=4$$

硬件按4个 Warp，也就是128个线程槽位来组织。

其中最后一个 Warp 只有4个有效线程，其余28个位置闲置。

---

六、Block 数量本身也有上限

即使线程、Warp、寄存器和 Shared Memory 都足够，一个 SM 同时驻留的 Block 数量也有硬件上限。

例如假设某个 SM 最多驻留16个 Block。

如果每个 Block 只有32个线程：

每个 Block = 1 Warp

按 Warp 数量可能可以放64个 Block，但硬件 Block 上限只有16，所以实际只能放16个。

因此，过小的 Block 可能受 Block 数量上限限制，导致 Warp 数量仍然不足。

---

七、寄存器限制

每个 SM 有固定容量的寄存器文件。

假设一个 SM 有65536个32位寄存器。

如果每个线程使用32个寄存器，Block 有256个线程，则一个 Block 需要：

$$32\times256=8192$$

个寄存器。

按寄存器资源计算：

$$65536\div8192=8$$

最多驻留8个 Block。

如果每线程使用64个寄存器：

$$64\times256=16384$$

则：

$$65536\div16384=4$$

最多驻留4个 Block。

如果每线程使用128个寄存器：

$$128\times256=32768$$

最多只驻留2个 Block。

关系是：

每线程寄存器数量增加
          ↓
每个 Block 消耗的寄存器增加
          ↓
SM 可驻留 Block 和 Warp 数量减少

---

八、寄存器分配通常存在粒度

实际寄存器分配通常不是精确按单个寄存器进行，而可能按：

• Warp；
• Block；
• 固定分配粒度；

进行向上取整。

例如，一个 Block 理论上需要8200个寄存器，硬件可能按更大的粒度分配，例如8448或其他值。

因此手工计算只能得到近似结果。

实际开发中通常使用：

• CUDA Occupancy API；
• 编译器资源报告；
• Nsight Compute；

得到准确的架构相关结果。

---

九、Shared Memory 限制

每个 Block 可以申请静态或动态 Shared Memory。

例如：

__shared__ float buffer[4096];

4096个 FP32 占用：

$$4096\times4=16384\text{ Bytes}=16\text{ KB}$$

假设一个 SM 可用于 Shared Memory 的容量为64 KB。

则按 Shared Memory 计算：

$$64\text{ KB}\div16\text{ KB}=4$$

最多驻留4个这样的 Block。

如果每个 Block 使用32 KB：

$$64\div32=2$$

只能驻留2个。

如果一个 Block 使用超过 SM 可提供的单 Block 上限，Kernel 甚至无法启动。

---

十、综合计算示例

假设一个 SM 的限制为：

最大线程数：2048
最大 Warp 数：64
最大 Block 数：16
寄存器数：65536
Shared Memory：64 KB

Kernel 配置为：

每个 Block：256个线程
每个线程：64个寄存器
每个 Block：20 KB Shared Memory

逐项计算。

线程限制

$$2048\div256=8$$

Warp 限制

每个 Block：

$$256\div32=8\text{ Warp}$$

所以：

$$64\div8=8$$

Block 数量限制

$$16$$

寄存器限制

每个 Block 使用：

$$256\times64=16384$$

所以：

$$65536\div16384=4$$

Shared Memory 限制

$$64\div20=3.2$$

只能取整数，因此最多3个 Block。

最终：

$$\min(8,8,16,4,3)=3$$

所以每个 SM 最多驻留3个 Block。

驻留线程数：

$$3\times256=768$$

驻留 Warp 数：

$$3\times8=24$$

---

十一、Occupancy 是什么

Occupancy 通常定义为：

$$\text{Occupancy} = \frac{W_{\text{resident}}} {W_{\text{max}}}$$

继续使用前面的例子：

实际驻留 Warp：24
最大支持 Warp：64

则：

$$\text{Occupancy} = \frac{24}{64} = 37.5\%$$

Occupancy 衡量的是：

SM 上有多少 Warp 状态已经驻留，可以被调度器选择。

它不是：

• CUDA Core 利用率；
• 每周期实际执行线程比例；
• 显存带宽利用率；
• GPU 总体利用率。

这些概念必须区分。

---

十二、理论 Occupancy 和实际 Occupancy

理论 Occupancy

根据 Kernel 配置和硬件资源上限计算得到。

主要取决于：

• Block 大小；
• 每线程寄存器数；
• 每 Block Shared Memory；
• SM 架构限制。

实际 Occupancy

运行过程中实际达到的平均活跃 Warp 比例。

它可能受到：

• Grid 中 Block 数量不足；
• 不同 SM 执行时间差异；
• Kernel 尾部效应；
• Block 调度不均衡；

影响。

因此理论上可达到75%，不代表整个执行期间始终保持75%。

---

十三、为什么需要足够的 Occupancy

假设一个 Warp 发出 Global Memory 请求后停顿。

如果 SM 上还有很多其他 Warp：

Warp 0：等待内存
Warp 1：可执行
Warp 2：可执行
Warp 3：等待依赖
Warp 4：可执行

调度器可以继续发射 Warp 1、2、4。

如果 SM 只有少量 Warp：

Warp 0：等待内存
Warp 1：等待内存

就可能没有可执行 Warp，流水线空闲。

所以 Occupancy 的主要价值是：

提供足够多的候选 Warp，用来隐藏计算和内存延迟。

---

十四、为什么 100% Occupancy 不一定最快

高 Occupancy 只表示驻留 Warp 多，不表示这些 Warp 能持续有效执行。

例如，64个 Warp 全部在等待同一个长延迟内存请求：

Warp 0～63：Memory Dependency Stall

此时 Occupancy 为100%，但执行单元仍然空闲。

另外，为了追求高 Occupancy，可能强行减少每线程寄存器数量，导致寄存器溢出：

寄存器使用减少
       ↓
Local Memory 访问增加
       ↓
额外显存流量
       ↓
性能反而下降

因此：

$$100\%\text{ Occupancy} \neq 100\%\text{ performance}$$

实际中，50%或更低 Occupancy 的 Kernel，也可能已经能充分隐藏延迟。

---

十五、低 Occupancy 也可能有高性能

矩阵乘法 Kernel 往往每个线程使用很多寄存器，用来保存多个输出累加器。

例如：

每线程计算多个输出元素
→ 数据复用增加
→ 寄存器使用增加
→ Occupancy 下降

但由于：

• 每个 Warp 有大量独立计算；
• 数据复用率高；
• Tensor Core 持续工作；
• 内存访问被双缓冲隐藏；

即使 Occupancy 只有25%或50%，性能仍可能很高。

这说明 GPU 性能优化经常需要权衡：

更多驻留 Warp
与
每个 Warp 更高的工作效率

---

十六、寄存器与 Occupancy 的权衡

方案A：

每线程32个寄存器
Occupancy高
但每个线程只能保存较少中间结果

方案B：

每线程96个寄存器
Occupancy较低
但可以保存更多输出累加器
减少访存并增加指令级并行性

方案B 可能更快。

因此不能机械地要求编译器减少寄存器。

CUDA 编译器可以使用类似寄存器数量限制选项，但过度限制可能导致 Register Spilling。

---

十七、Shared Memory 与 Occupancy 的权衡

更大的 Shared Memory Tile 可以提高数据复用。

例如矩阵乘法：

较大 Tile
→ 每次从显存加载的数据被使用更多次
→ 计算强度提高

但同时：

每个 Block Shared Memory 增加
→ 每个 SM 驻留 Block 减少
→ Occupancy 下降

因此 Tile 大小需要在以下因素间权衡：

• 数据复用；
• Shared Memory 容量；
• Block 数量；
• Warp 数量；
• 同步开销；
• Bank Conflict。

---

十八、Block 大小如何影响性能

常见 Block 线程数包括：

128
256
512

它们都是32的整数倍。

Block 太小

例如每 Block 32个线程：

• 只有一个 Warp；
• Block 级协作能力有限；
• 容易受最大 Block 数量限制；
• 可能无法提供足够 Warp。

Block 太大

例如每 Block 1024个线程：

• 单个 Block 消耗大量寄存器和线程槽位；
• 一个 SM 可能只能驻留1～2个 Block；
• 如果 Block 在同步点等待，候选 Block 较少；
• 资源分配灵活性下降。

中等 Block

128或256通常是较常见的起点，但不是固定最优值。

最佳配置取决于：

• Kernel 计算结构；
• 寄存器使用；
• Shared Memory 使用；
• 数据布局；
• GPU 架构。

---

十九、为什么 Block 数量也要足够多

假设 GPU 有120个 SM，但 Kernel 只有30个 Block。

最多只有30个 SM 获得工作：

30个 SM 工作
90个 SM 空闲

即使每个 Block 内线程很多，也无法利用整个 GPU。

通常希望 Block 数量明显多于 SM 数量，使调度器有足够工作分配。

例如：

SM 数量：120
Block 数量：几百或几千

可以提高负载均衡能力。

---

二十、Tail Effect：尾部效应

假设一个 GPU 有4个 SM，一共10个 Block。

如果每个 SM 一次运行一个 Block：

第一轮：4个 Block
第二轮：4个 Block
第三轮：只剩2个 Block

最后一轮只有两个 SM 工作，另外两个 SM 空闲。

这就是尾部效应。

Block 数量越少，或者不同 Block 执行时间差异越大，尾部效应越明显。

因此提供更多较均匀的 Block，通常有助于负载平衡。

---

二十一、Block 执行时间不均衡

即使 Block 数量很多，如果执行时间差异很大，也可能出现负载不均衡。

例如：

Block 0：处理简单数据，执行1000周期
Block 1：处理复杂数据，执行10000周期

最后可能只剩少数复杂 Block 运行。

常见原因包括：

• 不同输入数据导致循环次数不同；
• 稀疏任务长度不同；
• 分支路径差异；
• 图计算中的节点度数差异；
• 动态工作量分布。

解决方法可能包括：

• 任务重排；
• 动态工作队列；
• 更细粒度 Block；
• Persistent Kernel。

---

二十二、什么是 Persistent Kernel

普通 Kernel 中：

一个 Block 完成
→ 退出 SM
→ 新 Block 被调度进来

Persistent Kernel 的思路是：

启动数量接近 SM 数量的长期驻留 Block，让它们不断从全局任务队列中获取工作。

示意：

SM 0 上的 Persistent Block
    ↓
取任务 0
取任务 7
取任务 15
...

SM 1 上的 Persistent Block
    ↓
取任务 1
取任务 8
...

这样可以动态平衡不规则任务。

---

二十三、Persistent Kernel 的优点

• 减少反复 Block 调度；
• 可以保留部分状态在寄存器或 Shared Memory；
• 支持动态工作分配；
• 改善不规则任务负载均衡；
• 适合持续服务型工作。

常见场景包括：

• 工作队列；
• 图计算；
• 稀疏计算；
• 某些融合算子；
• 长时间运行的推理服务。

---

二十四、Persistent Kernel 的缺点

• 实现复杂；
• 需要原子队列或任务分配机制；
• 可能长期占用 SM；
• 容易影响其他 Kernel 并发；
• 同步和退出条件复杂；
• 调度公平性需要考虑。

它不是普通 Kernel 的普遍替代方案。

---

二十五、Block 调度通常不能保证顺序

Grid 中：

Block 0
Block 1
Block 2
...

并不表示一定严格按编号顺序完成。

不同 Block 可能：

• 在不同 SM 上开始；
• 以不同速度执行；
• 以不同顺序完成。

因此普通 CUDA 程序不能依赖：

Block 0 一定先于 Block 1 完成

这也是为什么不同 Block 之间不能通过普通共享变量假定执行先后。

---

二十六、一个 SM 上多个 Block 如何共享执行单元

假设 SM 同时驻留 Block A 和 Block B。

它们分别包含多个 Warp：

Block A：
Warp A0、A1、A2、A3

Block B：
Warp B0、B1、B2、B3

Warp Scheduler 可以在这些 Warp 中选择：

周期 1：Warp A0
周期 2：Warp B2
周期 3：Warp A3
周期 4：Warp B0

执行单元并不永久属于某个 Block。

Block 的作用主要是：

• 资源分配；
• Shared Memory 隔离；
• 同步范围；
• 生命周期管理。

真正逐周期被调度的是 Warp。

---

二十七、同步如何影响多 Block 驻留

假设 Block A 执行到：

__syncthreads();

其中部分 Warp 已到达，其他 Warp 尚未到达。

已到达的 Warp 会等待。

此时 SM 可以执行：

• Block A 中尚未到达的 Warp；
• Block B 中可执行的 Warp；
• 其他驻留 Block 的 Warp。

如果 SM 只驻留一个 Block，而该 Block 大量 Warp 在同步点等待，隐藏同步延迟的能力可能较弱。

多个驻留 Block 可以提供更多独立工作。

---

二十八、Occupancy 与 Eligible Warp 的区别

Occupancy 高表示驻留 Warp 多。

但真正决定当前周期能否发射的是 Eligible Warp 数量。

例如：

驻留 Warp：64
Eligible Warp：1

Occupancy 很高，但只有一个 Warp 可以执行。

另一种情况：

驻留 Warp：32
Eligible Warp：16

Occupancy只有50%，但调度器有大量候选 Warp。

因此性能分析时，Eligible Warp 往往比单纯 Occupancy 更直接。

---

二十九、如何初步选择 Block 大小

可以采用以下经验起点：

1. 使用32的整数倍；
2. 首先尝试128或256线程；
3. 检查寄存器和 Shared Memory 使用量；
4. 计算理论 Occupancy；
5. 测试多个配置；
6. 使用性能分析工具确认瓶颈。

例如可以比较：

128 threads/block
256 threads/block
512 threads/block

最终选择执行时间最短的，而不是理论 Occupancy 最高的。

---

三十、为什么自动调优很常见

Block 大小、Tile 大小、每线程输出数量等参数互相影响：

Tile 增大
→ 数据复用增加
→ Shared Memory 增加
→ 寄存器增加
→ Occupancy 下降

很难仅靠公式找到全局最优点。

因此高性能库通常会：

• 为不同矩阵尺寸准备多个 Kernel；
• 根据 GPU 架构选择不同配置；
• 运行时选择合适实现；
• 使用自动调优搜索参数。

cuBLAS、cuDNN 和编译器框架都会采用类似思路。

---

三十一、从 RTL 角度看 Block 资源分配

一个简化的 SM Block 分配器需要维护：

可用线程槽位
可用 Warp 槽位
可用寄存器数量
可用 Shared Memory
可用 Block 槽位

当新 Block 到来时：

读取 Block 资源需求
        ↓
检查所有资源是否足够
        ↓
足够：
  分配 Block ID
  分配 Warp ID
  分配寄存器区域
  分配 Shared Memory 区域
  初始化线程 PC 和掩码

不足：
  暂时不接收该 Block

Block 完成后：

释放寄存器
释放 Shared Memory
释放 Warp 槽位
释放 Block 槽位

然后可以接收新的 Block。

---

三十二、资源碎片问题

假设 SM 还有：

10 KB Shared Memory

但新 Block 需要：

12 KB

即使仍有大量空闲线程和寄存器，这个 Block 也无法进入。

同样，可能还剩很多寄存器，但线程槽位不足。

GPU 资源分配要求所有必要条件同时满足，所以不同资源之间可能出现不平衡和碎片。

---

三十三、本课核心结论

第一，一个 Block 只能驻留在一个 SM 上，但一个 SM 可以同时驻留多个 Block。

第二，Block 进入 SM 时必须一次性分配线程、Warp、寄存器、Shared Memory 和管理资源。

第三，驻留 Block 数量由各类资源限制中的最小值决定。

第四，每线程寄存器越多，通常可驻留的线程和 Warp 越少。

第五，每 Block Shared Memory 越多，通常可驻留的 Block 越少。

第六，Block 过大会降低资源分配灵活性，Block 过小可能受 Block 数量上限限制。

第七，Occupancy 是实际驻留 Warp 数与最大 Warp 数的比值。

第八，Occupancy 反映可用于隐藏延迟的 Warp 数量，但不直接表示执行单元利用率。

第九，100% Occupancy 不保证最高性能，低 Occupancy 也可能通过高数据复用和高指令级并行获得高性能。

第十，Eligible Warp 数量通常比 Occupancy 更直接反映调度器是否有工作可发射。

第十一，Grid 中应有足够多的 Block 才能覆盖所有 SM 并降低尾部效应。

第十二，Persistent Kernel 让少量 Block 长期驻留，并动态获取任务，适合不规则工作负载。

下一课将讲 GPU 的多级并行与调度体系：

第十三课：一个 GPU 为什么会有多个调度层次

主要包括：

CPU 如何提交 Kernel
Command Queue 与 CUDA Stream
Grid Scheduler
Block Scheduler
Warp Scheduler
多个 Kernel 能否同时执行
Kernel 并发的资源条件
Copy Engine 与计算重叠
CUDA Stream 如何实现流水处理