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

导语：这一课把调度层次串起来：从 CPU 提交 Kernel，到 Stream、Grid、Block，再到 SM 内部的 Warp 调度。

前面主要研究的是一个 Kernel 进入 GPU 后，Block 和 Warp 怎样执行。

但完整过程还包括：

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

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

这些调度发生在不同时间尺度，负责的问题也不同。

一、为什么不能只有一个调度器

GPU 同时面对几种不同粒度的任务：

一个 Kernel 可能包含几万个 Block；
一个 Block 可能包含多个 Warp；
一个 Warp 每隔几个周期就可能因为依赖而暂停；
GPU 还可能同时执行数据拷贝和多个 Kernel。

如果全部由一个中央调度器控制，会非常复杂，而且容易成为瓶颈。

因此 GPU 采用分层调度：

主机与运行时层：
决定提交哪些任务

GPU 前端层：
决定运行哪些 Kernel

Grid/Block 层：
决定 Block 放到哪个 SM

SM 内部：
决定下一周期发射哪个 Warp

二、CPU 和 GPU 的基本分工

在典型 CUDA 程序中，CPU 称为 Host，GPU 称为 Device。

CPU 负责：

分配和准备数据；
启动 Kernel；
提交内存拷贝；
管理多个计算阶段；
处理文件、网络和操作系统任务。

GPU 负责：

执行大规模并行 Kernel；
处理矩阵、向量和图形计算。

例如：

cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);

kernel<<<grid, block>>>(d_A, d_B);

cudaMemcpy(h_B, d_B, size, cudaMemcpyDeviceToHost);

可以理解为：

CPU：
把任务描述提交给 GPU

GPU：
按照任务描述自主执行大量线程

CPU 不会逐条指挥每个 Warp。

三、Kernel 启动通常是异步的

CPU 执行：

kernel<<<grid, block>>>();

通常表示把 Kernel 启动命令提交给 GPU。

CPU 不一定等待 Kernel 完成，而是可能继续执行后面的主机代码。

示意：

CPU 时间线：
提交 Kernel A
继续执行 CPU 工作
提交 Kernel B
等待结果

GPU 时间线：
        执行 Kernel A
                  执行 Kernel B

这种方式可以让 CPU 和 GPU 并行工作。

但如果后面的 CPU 操作必须使用 GPU 结果，就需要显式或隐式同步。

四、什么是命令队列

GPU 不能直接理解高级 C++ 代码。

CUDA Runtime 和驱动会把操作转换成 GPU 可执行的命令，例如：

内存复制命令
Kernel 启动命令
事件记录命令
同步命令

这些命令被写入队列，由 GPU 前端逐步获取并执行。

可以抽象为：

四、什么是命令队列

图示：四、什么是命令队列

五、CUDA Stream 是什么

CUDA Stream 可以理解为：

一条按顺序执行的 GPU 命令队列。

例如在 Stream 0 中提交：

Copy A
Kernel A
Copy B

这些操作在该 Stream 内保持顺序：

Copy A 完成
    ↓
Kernel A 开始
    ↓
Kernel A 完成
    ↓
Copy B 开始

同一个 Stream 中，后面的操作不会越过前面存在依赖的操作。

六、多个 Stream 有什么用

不同 Stream 中的任务，在资源和依赖允许时，可以重叠执行。

例如：

六、多个 Stream 有什么用

图示：六、多个 Stream 有什么用

可能形成流水线：

六、多个 Stream 有什么用

图示：六、多个 Stream 有什么用

理想情况下，可以让：

一批数据正在计算；
下一批数据同时传入；
上一批结果同时传出。

这称为计算与通信重叠。

七、Stream 不等于一个 SM

一个常见误区是：

Stream 0 → 固定使用 SM 0
Stream 1 → 固定使用 SM 1

这是错误的。

Stream 是命令执行顺序的逻辑概念。

一个 Stream 中的 Kernel 仍然可以使用整个 GPU 的多个 SM。

不同 Stream 的 Kernel 也可能共享同一组 SM，只要硬件资源允许。

因此：

Stream 是任务队列
SM 是计算硬件

二者不是一一对应关系。

八、什么是默认 Stream

如果程序不显式创建 Stream，Kernel 和内存操作通常提交到默认 Stream。

例如：

kernel_a<<<grid, block>>>();
kernel_b<<<grid, block>>>();

如果它们位于同一个有序 Stream 中，通常表现为：

Kernel A 完成
    ↓
Kernel B 开始

因为同一 Stream 中的命令保持顺序。

使用不同 Stream 后，A 和 B 才有可能并发，但并发并不一定真的发生。

九、事件 Event 的作用

CUDA Event 可以在 Stream 中记录一个执行位置。

例如：

Stream 0：
Kernel A
Record Event E
Kernel B

另一个 Stream 可以等待事件：

Stream 1：
Wait Event E
Kernel C

执行关系为：

Kernel A 完成
    ↓
Event E 生效
    ↓
Kernel C 可以开始

事件可以用于：

建立不同 Stream 之间的依赖；
测量 GPU 执行时间；
构建异步流水线；
避免 CPU 全局同步。

十、CPU 同步和 GPU 内部依赖的区别

假设 CPU 调用：

cudaDeviceSynchronize();

其含义通常是：

CPU 等待之前提交的相关 GPU 工作完成。

执行过程：

CPU 提交 Kernel
    ↓
CPU 调用同步
    ↓
CPU 暂停等待
    ↓
GPU 完成
    ↓
CPU 继续

而 Stream Event 依赖可以完全在 GPU 端解决：

Stream A 完成某阶段
    ↓
Stream B 自动继续

CPU 不需要中间介入。

因此高性能程序通常尽量减少不必要的全局同步。

十一、GPU 前端如何处理 Kernel

当 GPU 前端取得一个 Kernel 启动命令时，会获得：

Kernel 指令入口；
Grid 维度；
Block 维度；
Kernel 参数；
每线程寄存器需求；
每 Block Shared Memory 需求；
Stream 和依赖信息。

然后建立对应的 Grid 执行状态。

可以简化为：

十一、GPU 前端如何处理 Kernel

图示：十一、GPU 前端如何处理 Kernel

十二、Grid Scheduler 做什么

Grid Scheduler 可以理解为管理 Kernel 或 Grid 执行的上层调度逻辑。

它主要处理：

哪些 Kernel 已经满足依赖；
哪些 Kernel 可以开始；
Kernel 还有多少 Block 未调度；
是否允许多个 Kernel 并发；
是否还有可用 SM 资源。

不同 GPU 的具体实现和命名可能不同，但概念上存在这一层管理。

十三、Block Scheduler 做什么

当一个 Kernel 可以执行时，它的 Block 会逐步被分配到 SM。

Block Scheduler 需要寻找满足资源条件的 SM：

这个 SM 是否还有：
- 足够线程槽位
- 足够 Warp 槽位
- 足够寄存器
- 足够 Shared Memory
- 空闲 Block 槽位

只有所有条件都满足，Block 才能进入该 SM。

因此 Block 调度是资源感知的。

十四、Block 不会在 SM 之间迁移

一个 Block 被分配到某个 SM 后，通常一直在这个 SM 上执行到结束。

原因包括：

寄存器已经分配在该 SM；
Shared Memory 位于该 SM；
Warp 状态保存在该 SM；
Barrier 状态属于该 SM；
迁移成本很高。

所以：

Block 进入 SM
    ↓
在该 SM 上执行
    ↓
全部线程结束
    ↓
释放资源

十五、Warp Scheduler 是最细粒度的调度层

一个 SM 内可能同时驻留很多 Block。

这些 Block 被拆成大量 Warp：

Block A：Warp A0、A1、A2
Block B：Warp B0、B1、B2
Block C：Warp C0、C1、C2

Warp Scheduler 每个周期或若干周期从可执行 Warp 中选择：

周期 1：Warp A0
周期 2：Warp B1
周期 3：Warp C2

因此三个层次的关系是：

Grid Scheduler：
管理 Kernel/Grid

Block Scheduler：
把 Block 放进 SM

Warp Scheduler：
逐周期选择 Warp 指令

十六、为什么调度层次的时间尺度不同

不同层次的调度频率差别很大。

Kernel 层

Kernel 可能运行：

数微秒；
数毫秒；
甚至更久。

Kernel 调度相对低频。

Block 层

Block 通常运行较长的一段时间，完成后才释放资源。

Block 调度频率高于 Kernel，但远低于 Warp。

Warp 层

Warp Scheduler 可能几乎每个时钟周期都在选择指令。

这是高频、硬件实时调度。

所以不同层次需要不同复杂度和速度的控制逻辑。

十七、多个 Kernel 能否同时运行

可以。

这称为 Concurrent Kernel Execution，即 Kernel 并发执行。

假设 Kernel A 没有使用完所有 SM 资源：

Kernel A：
每个 SM 只占部分寄存器和 Shared Memory

那么 Kernel B 的 Block 可能同时进入剩余资源：

SM 0
├── Kernel A 的 Block
└── Kernel B 的 Block

不同 Kernel 的 Warp 可以被同一 SM 的调度器交替执行。

十八、Kernel 并发需要什么条件

通常需要满足：

Kernel 位于可并发的不同 Stream；
Kernel 之间没有未满足的数据依赖；
GPU 支持并发 Kernel；
SM 上仍有足够资源；
Grid 和调度限制允许；
没有全局同步阻止并发。

即使使用不同 Stream，也不保证并发。

例如 Kernel A 每个 SM 都占满：

Warp 槽位；
Shared Memory；
寄存器；
Block 槽位。

那么 Kernel B 仍然只能等待。

十九、一个大 Kernel 为什么可能阻塞其他 Kernel

假设 Kernel A 的每个 Block 使用大量 Shared Memory。

每个 SM 只能放一个 A Block：

SM 0：A Block
SM 1：A Block
SM 2：A Block
...

所有 SM 都没有剩余资源。

即使 Kernel B 很小，也无法进入。

这称为资源占满。

因此低优先级的大 Kernel 可能增加其他任务的等待时间。

二十、小 Kernel 为什么适合并发

假设 Kernel A 只使用 FP32 单元，并且资源占用较少。

Kernel B 主要使用 Load/Store 或其他资源。

如果它们的资源需求互补，可能提高总体利用率：

Kernel A：
计算较多，内存较少

Kernel B：
内存较多，计算较少

但实际能否有效重叠取决于：

执行单元是否竞争；
显存带宽是否竞争；
寄存器和 Shared Memory 是否足够；
调度开销；
Kernel 规模。

二十一、并发不等于性能一定提高

假设两个 Kernel 都已经把显存带宽跑满。

如果同时执行：

Kernel A 占用大量带宽
Kernel B 也占用大量带宽

它们只能争抢同一内存系统。

结果可能是：

每个 Kernel 都变慢；
总吞吐率没有提高；
延迟更加不可预测。

所以 Kernel 并发更适合互补资源，而不是完全相同的瓶颈。

二十二、Copy Engine 是什么

独立 GPU 通常有专门的数据拷贝引擎，用于：

CPU 到 GPU；
GPU 到 CPU；
GPU 内存之间；
某些点对点传输。

Copy Engine 可以在不大量占用 SM 算术单元的情况下完成数据搬运。

简化结构：

二十二、Copy Engine 是什么

图示：二十二、Copy Engine 是什么

二十三、计算和拷贝如何重叠

如果硬件支持，并且使用不同 Stream，可以形成：

二十三、计算和拷贝如何重叠

图示：二十三、计算和拷贝如何重叠

这样，数据传输时间可以被计算覆盖。

但需要满足：

使用异步拷贝接口；
使用合适的主机内存；
不同任务没有直接依赖；
Copy Engine 和计算资源均可用；
Stream 配置正确。

二十四、为什么需要页锁定主机内存

普通主机内存可以被操作系统换页或移动。

异步 DMA 传输更适合使用 Page-Locked Memory，也叫 Pinned Memory。

因为 Copy Engine 需要稳定的物理内存位置。

页锁定内存的优点：

支持更高效的异步传输；
更容易与计算重叠；
通常传输带宽更高。

缺点：

占用不可换出的物理内存；
分配成本较高；
使用过多会影响系统。

二十五、双缓冲数据流水线

假设不断处理输入批次。

可以准备两个缓冲区：

Buffer 0
Buffer 1

执行过程：

二十五、双缓冲数据流水线

图示：二十五、双缓冲数据流水线

这就是 Host-Device 层面的双缓冲。

前面矩阵乘法中的 Shared Memory 双缓冲是芯片内部流水线；这里是 CPU-GPU 之间的任务流水线。

思想完全一致：

用多个缓冲区把数据搬运和计算重叠起来。

二十六、三缓冲流水线

如果同时包含：

输入传输；
GPU 计算；
输出传输；

可以使用三个或更多缓冲区。

理想流水线：

阶段 n：

Copy Engine 0：
传入 Batch n+1

SM：
计算 Batch n

Copy Engine 1：
传出 Batch n-1

这样三个阶段可以同时工作。

总吞吐率最终由最慢阶段决定：

T_{\text{throughput}} \approx \max \left( T_{\text{H2D}}, T_{\text{compute}}, T_{\text{D2H}} \right)

而不是三者简单相加。

二十七、延迟和吞吐率的区别再次出现

不使用流水线时，一批任务耗时：

T_{\text{total}} = T_{\text{copy in}} + T_{\text{compute}} + T_{\text{copy out}}

使用多批流水线后，第一批的延迟仍包含多个阶段。

但稳定运行后的吞吐间隔可能接近最慢阶段。

例如：

输入：2 ms
计算：5 ms
输出：1 ms

单批延迟约：

2+5+1=8\text{ ms}

稳定流水线吞吐可能接近：

5\,\mathrm{ms/batch}

因为计算阶段最慢。

二十八、Kernel Launch Overhead

启动 Kernel 本身有固定开销，包括：

CPU Runtime 调用；
驱动处理；
命令提交；
GPU 前端解析；
建立 Grid 状态。

如果 Kernel 只执行非常少的计算：

Kernel 计算时间很短
启动开销占比很大

这时 GPU 可能不高效。

例如将一个大任务拆成几千个极小 Kernel，会产生大量启动开销和阶段间存储开销。

二十九、如何减少 Kernel 启动开销

常见方法包括：

Kernel Fusion

把多个小 Kernel 融合成一个大 Kernel。

CUDA Graph

预先记录一组固定任务和依赖，减少重复提交开销。

Persistent Kernel

启动长期运行的 Kernel，由其不断处理新任务。

增加每次处理批量

让一次 Kernel 处理更多数据。

这些方法本质上都是：

减少 CPU 与 GPU 前端反复提交小任务的次数。

三十、CUDA Graph 的基本思想

假设每次迭代都执行：

Copy Input
Kernel A
Kernel B
Kernel C
Copy Output

传统方式每次都逐项提交。

CUDA Graph 会把这些操作及依赖记录成图：

三十、CUDA Graph 的基本思想

图示：三十、CUDA Graph 的基本思想

之后重复执行整个图，减少运行时调度和提交开销。

适合：

深度学习训练迭代；
固定推理流程；
重复执行的多 Kernel 工作流。

三十一、优先级 Stream

某些 CUDA 环境允许创建不同优先级的 Stream。

高优先级任务可能比低优先级任务更早获得调度机会。

但需要注意：

优先级不一定能立即抢占正在运行的任意 Block。

如果一个低优先级 Kernel 的 Block 已经占满 SM，它可能需要等部分 Block 完成后，高优先级 Block 才能进入。

所以高优先级更接近“优先获得新资源”，而不一定是 CPU 式的任意时刻强制抢占。

三十二、GPU 抢占为什么困难

CPU 抢占一个线程时，需要保存相对有限的寄存器和状态。

GPU 一个 SM 可能同时驻留：

上千个线程；
大量寄存器；
多个 Block 的 Shared Memory；
Barrier 状态；
未完成内存请求。

完整保存和恢复这些状态成本很高。

因此 GPU 抢占通常比 CPU 复杂。

具体支持粒度取决于 GPU 架构和任务类型，可能包括：

Kernel 边界；
Block 边界；
指令级或线程级的更细粒度机制。

三十三、为什么长 Kernel 会影响响应延迟

假设一个 Kernel 的每个 Block 执行很长时间，并且占满 SM。

此时其他任务即使优先级较高，也可能需要等待资源释放。

所以实时或交互应用会关注：

Block 执行时间；
Kernel 时长；
是否支持抢占；
是否应拆分任务；
是否使用优先级 Stream。

数据中心多租户环境还需要更复杂的隔离和调度机制。

三十四、多进程如何共享 GPU

多个 CPU 进程可能同时使用同一块 GPU。

这时驱动需要管理多个 Context。

GPU 上可能采用：

时间片切换；
并发执行；
硬件分区；
软件服务层；

等方式共享资源。

不同 Context 之间切换可能需要切换：

地址空间；
页表；
Kernel 状态；
Cache 状态；
调度上下文。

这通常比同一进程内部多个 Stream 更重。

三十五、MPS 的基本概念

NVIDIA MPS 是 Multi-Process Service。

它的基本目的之一，是让多个进程更高效地共享 GPU，减少传统多 Context 调度开销，并提高小任务的并发利用率。

可以抽象理解为：

多个 CPU 进程
      ↓
共享的 GPU 服务层
      ↓
更统一地向 GPU 提交工作

适合多个 MPI 进程或多个小型 GPU 工作负载共享同一 GPU。

三十六、硬件分区与软件并发的区别

软件并发：

多个任务动态共享同一组 SM、Cache 和带宽

硬件分区：

把 GPU 的部分资源划分成相对独立的实例

硬件分区可以提供更强的：

资源隔离；
性能可预测性；
故障隔离；
多租户安全性。

但会降低资源动态共享的灵活性。

这类机制在数据中心 GPU 中更重要。

三十七、一次任务提交的完整流程

现在把全部过程串起来。

CPU 执行：

kernel<<<grid, block, shared_mem, stream>>>(args);

可能经历：

1. CPU 调用 CUDA Runtime
2. Runtime 准备 Kernel 参数
3. 驱动生成启动命令
4. 命令写入对应 Stream 队列
5. GPU 前端读取命令
6. 检查之前的 Stream 依赖是否完成
7. 建立 Kernel/Grid 状态
8. Grid Scheduler 标记 Kernel 可运行
9. Block Scheduler 查找有资源的 SM
10. Block 被分配到 SM
11. SM 分配寄存器和 Shared Memory
12. Block 被拆成 Warp
13. Warp Scheduler 选择 Eligible Warp
14. 指令被发射到执行单元
15. Warp 反复执行直到完成
16. Block 完成并释放资源
17. 新 Block 进入 SM
18. 所有 Block 完成
19. Kernel 完成
20. Stream 中下一条命令获得执行资格

三十八、多级调度的核心关系

可以用一张图概括：

三十八、多级调度的核心关系

图示：三十八、多级调度的核心关系

每一级负责不同粒度的问题：

层次	调度对象	主要目标
CPU/Runtime	Kernel、拷贝	构造任务和依赖
Stream	命令	保证顺序和并发
Grid Scheduler	Kernel/Grid	选择可运行 Grid
Block Scheduler	Block	分配 SM 资源
Warp Scheduler	Warp 指令	隐藏延迟、填充流水线

三十九、本课核心结论

第一，GPU 使用分层调度，因为 Kernel、Block 和 Warp 的粒度与时间尺度不同。

第二，CPU 主要负责提交任务，GPU 在提交后自主执行大量线程。

第三，CUDA Stream 是按顺序执行的命令队列，不是某个固定 SM。

第四，同一 Stream 中的操作保持顺序，不同 Stream 中的操作在条件允许时可以并发。

第五，Event 可以在不同 Stream 之间建立依赖，而不必让 CPU 全局等待。

第六，Grid Scheduler 管理 Kernel，Block Scheduler 分配 Block，Warp Scheduler 逐周期发射 Warp。

第七，多个 Kernel 可以同时执行，但必须有独立 Stream、无依赖并且 SM 资源足够。

第八，不同 Kernel 并发不一定提高性能，如果它们争用同一瓶颈资源，可能同时变慢。

第九，Copy Engine 可以使数据传输与 SM 计算重叠。

第十，页锁定内存通常是高效异步 CPU—GPU 传输的重要条件。

第十一，双缓冲和多缓冲可以构建输入、计算和输出流水线。

第十二，Kernel 启动有固定开销，极小 Kernel 可能被启动开销主导。

第十三，Kernel Fusion、CUDA Graph 和 Persistent Kernel 都可以减少任务提交或中间处理开销。

第十四，GPU 抢占比 CPU 困难，因为需要保存大量线程和片上资源状态。

第十五，多级调度的最终目标是：让数据搬运、Block 分配和 Warp 执行尽量持续进行，减少硬件空闲。

下一课适合进入多 GPU 系统：

第十四课：多 GPU 如何协同工作

主要包括：

PCIe、NVLink 和 NVSwitch
数据并行与模型并行
All-Reduce
张量并行与流水线并行
为什么通信会成为瓶颈
计算与通信重叠
GPU Direct
Blackwell GB200 NVL72 的基本组织方式

一、为什么不能只有一个调度器

二、CPU 和 GPU 的基本分工

三、Kernel 启动通常是异步的

四、什么是命令队列

五、CUDA Stream 是什么

六、多个 Stream 有什么用

七、Stream 不等于一个 SM

八、什么是默认 Stream

九、事件 Event 的作用

十、CPU 同步和 GPU 内部依赖的区别

十一、GPU 前端如何处理 Kernel

十二、Grid Scheduler 做什么

十三、Block Scheduler 做什么

十四、Block 不会在 SM 之间迁移

十五、Warp Scheduler 是最细粒度的调度层

十六、为什么调度层次的时间尺度不同

Kernel 层

Block 层

Warp 层

十七、多个 Kernel 能否同时运行

十八、Kernel 并发需要什么条件

十九、一个大 Kernel 为什么可能阻塞其他 Kernel

二十、小 Kernel 为什么适合并发

二十一、并发不等于性能一定提高

二十二、Copy Engine 是什么

二十三、计算和拷贝如何重叠

二十四、为什么需要页锁定主机内存

二十五、双缓冲数据流水线

二十六、三缓冲流水线

二十七、延迟和吞吐率的区别再次出现

二十八、Kernel Launch Overhead

二十九、如何减少 Kernel 启动开销

Kernel Fusion

CUDA Graph

Persistent Kernel

增加每次处理批量

三十、CUDA Graph 的基本思想

三十一、优先级 Stream

三十二、GPU 抢占为什么困难

三十三、为什么长 Kernel 会影响响应延迟

三十四、多进程如何共享 GPU

三十五、MPS 的基本概念

三十六、硬件分区与软件并发的区别

三十七、一次任务提交的完整流程

三十八、多级调度的核心关系

三十九、本课核心结论

第十四课：多 GPU 如何协同工作

相关文章

第二十七课：GPU 架构学习总结与统一分析框架

第二十六课：如何从零设计一个简化的 SIMT GPU

第二十五课：稀疏计算、不规则访存，以及 GPU 不擅长什么