H100 · Kernel Launch（kernel 启动约束与依赖启动协议）

目录 · ← l9 · l11 →

H100 · Kernel Launch（kernel 启动约束与依赖启动协议）

本文档基于课程讲义《8.2 Kernel Launch.pdf》（Lesson 8.2 / Constraints，12 页）整理， 讲解 kernel 启动相关的重要约束/属性，以及多 kernel 协作的依赖启动协议 （griddepcontrol）、L2 预取技巧与重叠窗口调优。 前置阅读：H100-Kernel-Design.md、H100-Stream-K.md、H100-异步与屏障.md。

目录

1. 重要的约束与函数
2. 为什么需要多 kernel 设置
3. Stream 串行化问题
4. 硬件信令：GDC 指令
5. 主机启动授权（Host Launch Authority）
6. 依赖方的墙：wait
7. 生产方的绿灯：launch_dependents
8. L2 预取技巧（Split DMA）
9. 调优重叠窗口
10. 总结与学习衔接

1. 重要的约束与函数

2. 为什么需要多 kernel 设置

• Epilogue 有硬上限：超出会毁性能——融合操作与 mainloop 的累加器 tile 竞争寄存器，造成寄存器压力。
• 跨 tile 依赖的操作（LayerNorm、Softmax）与 GEMM epilogue 的”独立 tile 处理”架构上不兼容： 数学上必须有一个 kernel 边界才能正确归约。
• 因此必须接受 kernel 边界与中间的 DRAM 写。工程挑战不是硬塞成一个巨型 kernel， 而是让 kernel 之间的交接几乎零成本。
• 手段：依赖启动（dependent-launch）协议、L2 预取策略、跨 grid 的 mbarrier。

3. Stream 串行化问题

• 通常同一 stream 上的下一个 grid 必须等上一个 grid 完全退役才能开始发工作。 这种严格完成顺序，在第一个 grid 进入低占用率的尾部时造成利用率缺口。
• 硬件调度器把同一 stream 的 grid 当成单一有序队列：在 Kernel A 报告全局完成之前，不会给 Kernel B 分配线程块槽位， 即使某些 SM 已经欠载。
• 当 Kernel A 接近完成时，只有越来越少的 SM 仍有活跃工作；其余 SM 因默认 stream 规则下没有下一 grid 的块可发而闲置。
• 交接是整 grid 退役事件，不是逐 tile 转移——除非启用显式依赖启动控制，否则所有依赖安全的、可重叠的机会都被忽略。
• 关键代价是”尾延迟放大”：硅片在场，却在 Kernel A 的最后阶段不发有用指令。

4. 硬件信令：GDC 指令

GDC（griddepcontrol）指令控制 GPU 执行时间线里的 grid 交接时机。 物理上它把两件事分开：“调度器可以启动依赖方” vs “依赖方可以读依赖敏感的内存”。

• griddepcontrol.launch_dependents：由正在运行的 warp 发出的、面向调度器的信号。 它告诉分发逻辑：依赖 grid 现在有资格启动，即使当前 grid 尚未全局退役。
• griddepcontrol.wait：依赖 grid 里的硬执行栅栏。到达该指令的 warp 被扣住， 直到启动依赖条件被满足，防止过早读取未解析的数据。

二者合起来是一个两段式协议：提前启动许可 + 内存安全门控。 调度器可以重叠启动工作，而栅栏为依赖敏感的 load 保持正确性。 它们是指令流里的物理控制点，放置位置直接改变重叠时机与硬件争用。

5. 主机启动授权（Host Launch Authority）

• 只有 CPU 的 launch packet 才能授权”依赖重叠”的宽松 stream 行为。物理目的是：除非主机软件显式选择加入，否则保持默认 stream 语义不变。
• 启动时，CPU 把属性写进GPU 命令处理器消费的命令描述符。其中一个属性授予”依赖 grid 在前驱完全退役前被调度”的许可。
• 若该许可位存在 → 调度器状态机把 launch_dependents / wait 当作依赖控制指令来执行； 若缺失 → 调度器强制普通串行化，这些指令无启用效果。
• 因此交接授权是”主机发起、硬件强制执行”的：设备端信令不能覆盖 stream 策略，除非启动元数据授权它。
• 这保护了混合工作负载的正确性（有的 stream 需严格顺序，有的需受控重叠）。

6. 依赖方的墙：wait

• 该阶段定义依赖 grid 里 producer warp 的”提前启动停顿点”。物理目的：让依赖 kernel 尽早预留执行上下文，同时阻塞不安全的内存流量。
• 依赖 grid 可被接纳、在可用 SM 上开始执行设置指令；producer warp 跑到 griddepcontrol.wait 时被硬件停驻（park）。
• 停驻期间，这些 warp 不能发依赖敏感的全局读（如前一个 kernel 产出的激活张量）， 防止因读”尚未最终确定”的数据造成缓存填充与内存序违规。
• 一旦依赖信号满足，栅栏释放，这些 warp 恢复发 load——无需在 producer 完成后再付完整冷启动延迟。
• 代价是预留资源：停驻的 warp 仍占调度槽，可能减少其它工作的即时余量。

7. 生产方的绿灯：launch_dependents

• 该阶段是活跃 kernel 发出”打开依赖调度”的精确信号。物理目的：在”最后一个正确性安全的时刻”触发重叠，仍能暴露启动延迟隐藏。
• 活跃 grid 的 producer warp 在剩余写已经足够推进、依赖方可安全启动时执行 griddepcontrol.launch_dependents。 该信号面向的是调度器资格，不是”立即可读内存”的许可。
• 在持久 kernel 里，正确放置点与最后一个被调度工作 tile 的生命周期绑定： 信号应对齐真实的流水线完成，而不是代码区域的词法结尾。
• 信号之后，依赖块可被分发，同时活跃 grid 排空剩余工作；依赖 grid 的 wait 栅栏仍守护依赖敏感读，直到条件满足。
• 发太晚 = 浪费重叠；发太早 = 增加并发压力、可能降低净吞吐。

8. L2 预取技巧（Split DMA）

• 该阶段通过在依赖 kernel 内把角色分到不同 warp 来实现重叠。物理目的：在依赖受缚的 producer 被栅栏挡住时，让内存结构忙于无依赖的传输。
• 一个 producer warp 到达依赖屏障、在激活读之前暂停； 另一个 prefetch warp 继续跑，对不依赖 producer 完成的静态权重发 DMA 式请求。
• 这些请求提前填充 L2，减少后续 miss 惩罚；内存系统因此在本来停顿的依赖时间里预热有用的缓存行。
• 交接是非对称但安全的：激活流量在栅栏后等待，权重预取不用。 栅栏释放时，producer 因部分工作集已在缓存而看到更好的有效延迟。
• 前提：预取的数据要有高复用且能在 L2 存活到被消费。

9. 调优重叠窗口

• 重叠比 = 依赖启动信号发射 与 真正依赖就绪 之间的时间差。更大的差 = 潜在隐藏更多，但前提是资源未拥塞。
• 信号太早 → 两个 kernel 争 SM 发射槽、寄存器文件容量、共享内存分配、内存端口，可能把每个 kernel 拖慢到总时间反而增加。
• 预取太激进 → L2 行在被用之前就被颠掉，DRAM 流量因回填而飙升，带宽花在把被逐出的数据搬回来，抵消预取重叠的收益。
• 实用调优靠硬件计数器：在最晚安全点启动、只预取稳定/复用的张量、 以 L2 命中率 + 内存队列压力 作为主要护栏。

10. 总结与学习衔接

10.1 核心脉络速记

10.2 与本课程其他内容的衔接

10.3 一句话记忆

Kernel Launch = 用约束/属性（cluster、maxnreg、grid_constant、restrict）给编译器交底， 再用”主机授权 + GDC 双指令（launch_dependents / wait）+ Split DMA 预取”把 kernel 之间的交接做成 几乎零成本——既让依赖 grid 提前启动隐藏尾延迟，又用 wait 栅栏守住内存正确性。

参考来源：8.2 Kernel Launch.pdf（Lesson 8.2 / Constraints，12 页）。