H100 · Kernel Design（kernel 设计的”招式”）

目录 · ← l7 · l9 →

H100 · Kernel Design（kernel 设计的”招式”）

本文档基于课程讲义《8. Kernel Design.pdf》（Lesson 8 / The Tricks，98 页）整理， 系统总结 Hopper 上高性能 kernel 的设计方法与优化技巧：warp specialization、流水线、 ping-pong/cooperative 流水线、调度（persistent/Stream-K）、epilogue 融合。 前置阅读：H100-异步与屏障.md、H100-cp-async.bulk.md、H100-WGMMA.md、H100-WGMMA-part2.md。

目录

1. CUDA kernel 的两类与拐点
2. Compute-bound kernel 的优化手段
3. Warp Specialization（warp 特化）
4. 流水线（Pipelining）与环形缓冲
5. 两屏障握手与 ABA 问题
6. Producer / Consumer 流程与三阶段
7. 两种流水线：Cooperative vs Ping-Pong
8. Tile 尺寸与 RS/SS 选择
9. 调度（Scheduling）
10. Static / Grouped / Stream-K 调度器
11. Memory-bound kernel 的三类
12. Epilogue（尾声融合）
13. 总结与学习衔接

1. CUDA kernel 的两类与拐点

CUDA kernel 分两类：

• Compute-bound：受算术运算速率限制。
• Memory-bound：受数据搬移速率限制。

分界点是算术强度（Arithmetic Intensity）：AI = FLOPs / Bytes moved。

拐点（Ridge Point）= 峰值 FLOPS / 峰值带宽：

• AI < 拐点 → memory-bound；
• AI > 拐点 → compute-bound。

1.1 Compute-bound kernel 的特征

1. 高算术强度：每字节读写的 FLOPs 多，复用强（数据留在寄存器/shared，多次使用）。
2. Producer 空闲：最显著特征是 Producer/Consumer 失衡——producer warpgroup 在屏障处空等 consumer 赶上。
3. Occupancy “必要但不充分”：许多 compute-bound kernel 在中等占用率下即可接近峰值，只要发射足够 WGMMA。 发射计算指令不是问题，寄存器利用才是。

2. Compute-bound kernel 的优化手段

• Warp Specialization（warp 特化）
• Persistent Kernels + Tile Scheduling（持久 kernel + tile 调度）
• Circular buffer（共享内存多级环形缓冲 + 显式同步）
• Cluster-Level Optimizations（簇级优化）
• Register Pressure Management（寄存器压力管理）
• Megakernels（巨型 kernel）
• Epilogue Fusion（尾声融合）

3. Warp Specialization（warp 特化）

Hopper 上一个 SM 能容纳很多活跃 warp，但每个 cycle 只有少数能发射指令。

• 若同一 warp 内线程走不同分支，warp 会串行化这些路径（最坏 ~32×）。
• 但若 warp 0 负责加载数据、warp 1 对数据做计算，它们是不同 warp、有独立执行上下文，就避免了 SIMT 惩罚。

Warp specialization = 刻意让同一线程块内不同 warp 承担不同工作：

• producer warp：搬/准备数据。
• consumer warp：对数据做计算。

3.1 为什么 warp specialization 几乎是必须的

1. 资源约束逼的：你无法在不溢出的情况下把整个活跃状态（寄存器/谓词等）塞进每线程/每 warp，所以把工作拆到不同 warp。
2. 变延迟操作难以静态调度：内存等变延迟操作让编译器/静态调度难以让所有单元都忙。
3. 阻塞同步会卡住发射：若某 warp 要等屏障，特化后其它 warp 能立刻跑，SM 不浪费发射槽。

3.2 资源约束与 setmaxnreg

• WGMMA kernel 的首要瓶颈是寄存器文件：要最大化吞吐，线程必须在寄存器里持一大块输出矩阵； 但每线程分配 ~200+ 寄存器会大幅减少 SM 能容纳的 warp 数。
• Warp specialization 把工作拆成两个角色，做非对称资源分配：
  • Producer：发 cp.async，需要最少寄存器。
  • Consumer：执行 WGMMA，需要最多寄存器。
• setmaxnreg 指令让 warp 在运行时动态改变自己拥有的寄存器数：
  • 以低寄存器数启动 kernel → SM 容纳更多活跃 warp → 最大化内存带宽利用。
  • 进入重计算段前，consumer warp 执行 setmaxnreg 请求更多寄存器；producer 用它减少寄存器。

3.3 Warp group 数量选择

• Producer：总是恰好 1 个 warp group（单线程就能发一条搬几 GB 数据的 cp.async）。 producer 几乎不做数学，只管理 mbarrier；可用 set_maxnreg 把它们限制在 ~32 寄存器。
• Consumer：根据寄存器压力选 1、2 或（罕见）3 个；数量由累加器 tile 大小决定。
  • 预算：~232–240 寄存器（留出 barrier 等空间）。两种用法：
    • Pingpong/Basic：每个 WG 独立算一个完整 tile，EffectiveThreads = 128。
    • Cooperative：两个 WG 拆分同一 tile，EffectiveThreads = 256。

4. 流水线（Pipelining）与环形缓冲

流水化 = 重叠：把工作拆成阶段，让不同”项”的不同阶段同时进行，而不是一个项做完再做下一个。

无流水（串行）：Load → Wait → Compute → Wait → Load → Wait → Compute → Wait ...
                每个阶段都等上一阶段完成，传输时硬件空闲

有流水（重叠）：
  Load:  L0      L1      L2      L3
  MMA:         M0      M1      M2      M3
  算 tile N 的同时加载 tile N+1 —— 硬件一直忙

4.1 为什么必须做对

• 延迟鸿沟：全局内存 ~400–800 周期，WGMMA ~30–60 周期/次。不流水时，计算要等 20 倍于它运行的时间。
• 不流水：load（400 周期计算空闲）→ compute（20 周期内存空闲）→ load（又 400 空闲）→ 计算硬件约 5% 利用率。
• 流水化：计算嚼 tile N 时，内存已在取 tile N+1；若流水够深，计算永不等数据，从 5% → 近 100%。

4.2 Circular Buffering（环形缓冲）

• 让 producer 与 consumer 重叠，需要一个二者之间的缓冲：一组固定数量的 shared-memory “槽位（stage）”， producer 填、consumer 排空。
• 环形缓冲 = 轮转复用这些槽位：producer 写 stage i、i+1……绕回 0；consumer 读 i、i+1……绕回 0。
• 它让”算 tile N 的同时加载 tile N+1”成为可能。

4.3 Stages 与为什么需要多个

• 单缓冲会强制串行（Load→Compute→Store 只有一个阶段在忙）。
• 多缓冲允许多个在飞数据块：stage 1 填 buffer A、stage 2 处理 buffer B、stage 3 排空 buffer C。
• 环形意味着 buffer 走完一圈后复用——无需无限内存，只要足够 buffer 让所有阶段都忙。
• 需要多个 buffer 的根因是”所有权 + 重叠”：一个阶段不能安全覆盖另一个阶段仍在读的数据。

5. 两屏障握手与 ABA 问题

环形缓冲里每个 stage 有两个信号，producer/consumer 永不竞争：

1. FULL 屏障：consumer 读之前在此等待。对 TMA 流水线它是事务屏障——producer 调 mbarrier.arrive， TMA 引擎在字节落到 shared memory 时发完成信号。consumer 无需轮询就知道 stage 确实填好了。
2. EMPTY 屏障：producer 覆盖前在此等待；consumer 用完该 stage 后发信号。

Producer 等 EMPTY → 声明 FULL 的 expect_tx → 发 TMA → TMA 完成 → FULL 触发
→ Consumer 读/用 → Consumer arrive EMPTY → stage 可复用

5.1 ABA 歧义

• 环形缓冲跨多轮复用同一 stage 索引。若同步只按 stage_id ∈ [0..Stages-1] 索引， consumer 看到”stage 0 满”时无法区分是 第 1 轮 stage 0（旧数据） 还是 第 2 轮 stage 0（绕回后的新数据）。
• 解法：跟踪一个 phase bit，每当环形索引从 Stages−1 绕回 0 时翻转。 每个参与者（producer/consumer）维护状态三元组：index（当前槽位）、phase（绕回翻转的 1 位纪元）、 count（单调递增迭代计数，用于簿记）。

6. Producer / Consumer 流程与三阶段

6.1 Producer 流程

每个 warpgroup 只有一个被选中的线程真正碰 mbarrier 和 TMA。每次迭代做三件事：

1. 获取 stage：等 EMPTY 屏障，然后在 FULL 屏障上调 expect_tx。
2. 发 TMA 拷贝：DMA 引擎写完 shared 字节后自动给 FULL 发信号（硬件级 producer-consumer 信号）。
3. 前进：递增流水线状态，绕回时翻转 phase。

所有工作结束时，mbarrier.try_wait 等 consumer 释放每个剩余 stage 后才退出（不能有人还在读你的数据就退出）。

6.2 Consumer 流程

consumer 维护两个流水线指针：正在消费的 stage 与 可释放回 producer 的 stage。

• 因 WGMMA 是异步的，不立即完成；只有确认读已完成后（warpgroup_wait）才能释放它读的 buffer。
• 每次迭代：等数据（FULL 屏障）→ 发 WGMMA（fence, arrive, gemm, commit）→ 等最老 WGMMA 完成（wgmma.wait_group<N>）→ 释放最老 buffer（EMPTY）→ 前进两个指针。
• 每次 WGMMA 前的跨 proxy fence 防止编译器把累加器读写重排到 WGMMA 边界之外。

6.3 三阶段

• Prologue（填充）：consumer 发 N 条 WGMMA 且不释放任何 buffer，让流水线填满， 稳态时总有在飞 WGMMA 与 TMA load 重叠。第一条 WGMMA 把累加器初始化为零。
• Steady-State（稳态）：吞吐最优循环——对每个 k-tile：等数据 → 发 WGMMA → 等最老 → 释放最老。 TMA 与 WGMMA 完全重叠，producer 领先 consumer 的释放指针 Stages 步。
• Drain（排空）：最后一个 k-tile 消费后，warpgroup_wait<0>() 冲刷所有在飞 WGMMA， 再释放剩余 N 个 buffer，让 producer 干净退出。

7. 两种流水线：Cooperative vs Ping-Pong

两种主流水线都用 warp specialization（producer 加载 + consumer 计算 + 环形缓冲）。 区别只在 epilogue 阶段：Ping-Pong 把 wgmma 与 epilogue 重叠，Cooperative 不重叠。

7.1 Cooperative Pipeline

• 384 线程，3 个 warp group：
  • WG0：Producer（发 TMA 加载）
  • WG1：Consumer（WGMMA + epilogue）
  • WG2：Consumer（WGMMA + epilogue，同一输出 tile）
• 输出 tile 128×128 时，两个 consumer 平分几何工作：consumer 0 算上半（M 0–63），consumer 1 算下半（M 64–127）。 因同时处理同一输出 tile，它们同一时刻消费同一份 A、B tile。
• 持久调度：每个 CTA 循环抓取一串 tile，每个 warpgroup 走同一串 tile。
• Consumer 循环：维护 read 指针 + release 指针，偏移 N（通常 1）。这个”滞后”是因 WGMMA 异步—— 先等 read 游标 → 发 WGMMA → 前进 read 游标；只有对应 wgmma 完成后（wgmma.wait_group）才能释放更老的 stage 并前进 release 游标。
• 关键洞察：Cooperative 不归约 WGMMA 结果——”合作”是共享 smem tile，不是合并累加器。 每个 WG 独立做所有 K tile 的 wgmma，累加进自己的寄存器 accum（warpgroup 寄存器私有，物理上不可能跨 WG 通信）。 每个 WG 写自己不相交的 M 区域。
• 收益：单个更大的 tile 可由两个 WG 一起映射，实现更大的有效 MMA tile（256×128 而非 128×128），同时每个 WG 仍满足寄存器预算。

7.2 Cooperative 的缺口（The Gap）

• consumer 累加完所有 k-tile 后要做 epilogue（缩放、bias、激活、写 global）。
• 整个 epilogue 期间 Tensor Core 完全空闲。
• 对大 epilogue 或小 K 维，这段空闲时间占总运行时间的相当比例——流水线能重叠 load 与 compute， 却无法重叠 compute 与 epilogue（因为只有一个 consumer 两件事都做）。

7.3 Ping-Pong（补上缺口）

• 加第二个 consumer：一个跑 epilogue 时另一个对下一个 tile 做 MMA，二者交替，WGMMA 永不停。
• 384 线程，3 个 WG：WG0 producer，WG1 consumer（C0），WG2 consumer（C1）。 producer WG 释放自己的寄存器，把寄存器文件空间让给两个 MMA WG。
• 交替：

  C0: [MMA T0][Epi T0][MMA T2][Epi T2]
  C1:          [MMA T1][Epi T1][MMA T3][Epi T3]
  

• 持久调度下，CTA 被分配一串 tile T0..T5：producer 全部处理（步长 1）， consumer 0 处理偶数（步长 2，起点 0），consumer 1 处理奇数（步长 2，起点 1）。

7.4 Ping-Pong 的 barrier（2×2 网格）

• 3 个 stage 的流水线只分配恰好 3 个物理 barrier，跨迭代复用。
• 用 2×2 的 mbarrier 网格保证 C0 的 epilogue 与 C1 的 MMA 同时进行而不破坏 shared memory：
  • 行（Stages/Depth）：0 = MMA 阶段，1 = Epilogue 阶段。
  • 列（Groups）：0 = Consumer 0，1 = Consumer 1。
• 每个 consumer 有 group_id；arrive() 给另一个 consumer 的 barrier（当前深度）发信号， wait() 在自己的 barrier（当前深度）上等。
• 深度循环：MMA → Epilogue → MMA → …。
• 必须保证：① 一次只有一个 consumer 做 MMA；② 一次只有一个做 epilogue；③ 同一 tile 先 MMA 后 epilogue。

7.5 一个 warp group 迭代的完整时间线

1. ordered_barrier.wait()       ← 等 C1 的上一个 epilogue
2. WGMMA mainloop               ← K 循环：从 smem 读、累加进寄存器
3. ordered_barrier.arrive()     ← "我的 MMA 完了，C1 可开始它的 MMA"
4. mma_tail()                   ← warpgroup_wait<0>，释放最后几个 smem stage
5. ordered_barrier.wait()       ← 等 C1 的 epilogue 完成
6. epilogue.store()             ← 融合 + R→S 拷贝 + TMA store
7. epilogue.store_tail()        ← 等所有 TMA store 落地
8. advance pipeline states      ← 前进 2 步（每个 warp group 一步）
9. ordered_barrier.arrive()     ← "我的 epilogue 完了，C1 可开始它的 epilogue"
10. fetch next tile, loop 回 1

7.6 warpgroup 内的角色分工

• Producer warp group（128 线程 = 4 warp），每个 warp 一个角色：
  • Mainloop DMA warp：TMA 加载 A、B tile 到 shared；
  • Epilogue DMA warp：TMA 加载 C tile 到 shared（用于残差加）；
  • Scheduler warp：取下一个 tile 坐标；
  • Auxiliary warp：可选的额外加载。
• Consumer 0 / 1（各 128 线程）：偶/奇 tile 的数学 + epilogue。

注意：数学 warpgroup 会切换角色——MMA 时是 mainloop 流水线的 consumer，epilogue 时变成 store 流水线的 producer。

8. Tile 尺寸与 RS/SS 选择

8.1 输出 tile 尺寸

• 输出 tile 决定单个 CTA 计算多大块的 C。
• 目标：让 tile_m、tile_n 越大越好（越大算术强度越高，因为每载一次 A、B tile 就做 2×tile_m×tile_n×tile_k 次运算）。
• 小/中 tile（128×128 或 64×128）：单个 consumer WG 的寄存器够存累加器 → 用 Base 或 Ping-Pong。
• 大 tile（256×128 或 128×256）：单个 consumer WG 在 FP32 下装不下 → 用 Cooperative。

8.2 内层 tile_K 与流水线 stage 数

• tile_k：内层循环每次沿 K 维走多深；大小要”喂饱 Tensor Core 又不撑爆 shared memory”。
• 要让 epilogue 重叠高效，K 维必须够大（太小则 C1 的 MMA 会在 C0 的 epilogue 之前做完，又卡住）。 计算时间必须 ≥ 内存写时间。
• WGMMA 原生一次消费 K=16（FP16/BF16）或 K=16（TF32）。
• 通常要 3 或 4 个 stage 隐藏 TMA 加载延迟。

8.3 RS vs SS（A 放寄存器 vs 共享内存）

9. 调度（Scheduling）

调度 = 决定”谁在何时以何顺序做什么”的决策逻辑，即把工作单元（Tile）映射到工作者（线程/warp/CTA/SM）的策略+机制。

好的调度区分”硬件都忙”与”部分 SM 空闲”、区分”良好局部性”与”抖动/停顿”，保证可预测的完成时间而非尾延迟悬崖。

9.1 为什么调度重要

• Occupancy/Utilization：给 SM 足够多独立 tile 隐藏延迟。
• Load Balance：避免”长尾”（少数 CTA 拿到大 tile，其它早完成闲置）。
• Locality/Bandwidth：最大化 L2 复用、最小化冗余 global load、促进 multicast。
• 解耦 tile 大小与调度粒度：持久调度引入第三分解轴：K 维。

9.2 非持久调度（问题）

• Grid Size = 总工作量 / Block Size，硬件调度器分配 block，完成即退役。
• 三个问题：① 多次启动 kernel 有开销；② 尾效应（133 block vs 132 SM → 先跑 132 再跑 1）； ③ 硬件调度器通常线性（Block 0 算 (0,0)，Block 1 算 (0,1)……等扫到下一行时，所需数据早已从 L2 被逐出）。

9.3 持久调度（Persistent）

• 启动固定数量的线程块（通常 = SM 数），这些 block 常驻 GPU，循环计算下一个 tile 的索引并处理，直到全部完成。

9.4 波量化（Wave Quantization）问题

• 数据并行基线：CTA 轮转取 tile。150 tile / 132 SM = 1.136 tile-units/SM → 实际 2 波 → 利用率 56.8%。
• 消除波量化：把”余下的 8 个 tile”沿 K 维平分给所有 132 个 SM（每个 SM 算约 0.06 个 tile 的 K 工作）。
• 本质洞察：波量化不是硬件限制，而是”tile 并行分解”的结果；持久调度让你换一种分解。

10. Static / Grouped / Stream-K 调度器

10.1 Static Persistent Scheduler

• 标准 GEMM 的默认高吞吐调度器。”Static”指工作到线程的映射是数学预计算的，不是原子计数动态认领。
• 把输出矩阵看成 tile 网格，用光栅化（Rasterization / Swizzling）曲线（常为 Z 曲线或 U 曲线）给持久 block 分配 tile。
• 持久 block 算完第一个 tile 后，按”启动的总 block 数”跳步（Grid Stride Loop）找下一个 tile。
• 用于标准 GEMM、compute-bound kernel（开销最低、无块间同步、swizzle 最大化缓存命中）。

Rasterization（光栅化顺序）：

• 决定 CTA 被分配到输出 tile 的顺序，直接影响 L2 局部性（让共享输入数据的相邻 CTA 在时间上靠近执行）。
• 顺序错了 → 每跳一个新 tile 都要从 DRAM 重载整块矩阵。

路径策略 AlongN vs AlongM：

• Column-major（外=N，内=M）：固定列 n 扫完一整列再右移 → B 的列保持热（B 最大复用），A 行在轮换。
• Row-major（外=M，内=N）：固定行 m 扫完一整行再下移 → A 的行保持热（A 最大复用）。
• 结论：想让哪个矩阵留在缓存，就垂直于它的复用维度遍历。

Swizzling（打散遍历）：

• 只沿一个方向复用（如 A 一行保持热）时，B 是零复用流式；扫下一行时又要重载 B。
• 需要两个维度同时有局部性。swizzle 把”细线”遍历改成”厚块”遍历。
• swizzle size 可调：size=1 是细光栅，size=2 → 4 个 tile 厚。CUTLASS 常用 1、2、4、8。
• 规则：内层（快轴）取更长的维度（更长内循环 = 更多次迭代后才”迈大步”，迈大步是昂贵的缓存上下文切换点）。

Cluster 与 swizzling：

• 纯 swizzle 的复用靠”相邻 CTA 时间上靠近”（希望），但物理执行顺序会漂移。
• Cluster 解决：簇内 CTA 在空间+时间上一起执行、可同步协作；共享操作数面板可跨 CTA 复用——复用从”偶然”变”有意”。

10.2 Grouped Persistent Scheduling

• 为 Grouped GEMM（一次 kernel 启动算多个不同 GEMM）设计。
• 把”Group 0 再 Group 1…”概念上拼接成一条长线性序列；维护当前 group ID、该 group 起始线性索引、该 group 的 tile 数。
• CTA 的 linear_idx 前进时可能跨 group 边界，朴素标量搜索太贵 → 用 warp-level speculative search。

warp-level speculative search：

• 不在当前 group 时，warp 以 32 个 group 为一批扫描：每个 lane 载一个 group 的形状、算它的（cluster 对齐）tile 数。
• 用 warp 原语（__ballot_sync、__ffs、__shfl_sync）选出”范围包含 linear_idx”的 lane 并广播 GroupInfo；没命中则跳 32 个 group 重试。
• 找到所属 group 后：算局部偏移 k = linear_idx - start，swizzle k → (cluster_major, cluster_minor)，按光栅顺序转 (M,N)，再加簇内 CTA 偏移。
• 该扫描很便宜（持久化摊薄，主要在 group 边界触发）。

10.3 Stream-K Scheduling

要解决的问题（尾效应）：标准调度把完整输出 tile 分给线程块；若 tile 总数不是 SM 数的整数倍，最后一”波”只填一部分——活跃 SM 处理最后 tile，其余 SM 全闲置。

• Stream-K 把整个矩阵乘看成一条连续的 1D “数学迭代带”，单元是”一定数量的 MMA 运算”（而非”一个完整 tile”）。
• 把总数学量严格平分给所有处理单元；一个 block 可能算完整 tile、可能继承半成品、也可能预算用完时停在一半。
• Hybrid 实现（只打尾）：把 tile 沿 K 维拆分引入跨块通信开销，对全问题用 Stream-K 常适得其反。 最优是混合：早期波保持纯数据并行（最大吞吐），只对最后”尾”波用 Stream-K 均衡负载。 50% 启发式：尾波已较满（>50%）时回退到标准数据并行，避免不必要的归约开销。

Fixup（点对点归约）：

• Cluster A 算 Tile X 前 50%、Cluster B 算后 50% → 写 global 前要相加（Fixup）。
• 用一块 global scratchpad：前半写完部分累加器 + 置 flag；后半完成时检查，看到 A 完成就载入 A 的部分结果、加到自己的寄存器、写最终和。

反向 tile 迭代（Backward Tile Iteration）：

• 两个 block 共享一个 tile 时，”算 K 尾”的 block 要等”算 K 头”的 block 完成才能写。
• Stream-K 让 worker 按 K 逆序迭代：越靠后的块越晚算它的共享部分，等它要合并时，前面的块已算完 → 大幅减少等待。

保持 L2 局部性：

• Stream-K 的 1D 带天然破坏空间局部性。解法：把 worker 逻辑分组，让同组单元处理不同输出 tile 的重叠 K 区间—— 它们同时迭代完全相同的 K 切片，L2 里的输入读完美重叠。

11. Memory-bound kernel 的三类

1. Bandwidth-bound：DRAM/L2 吞吐接近峰值，运行时间由”搬了多少字节 + 效率”决定。
2. Latency-bound：性能受”依赖解析多久”（访存、原子、同步、长指令链）限制，带宽未饱和。
3. Locality-bound：搬了很多字节但大量浪费（L2 命中低、随机/间接访问、cache/TLB 抖动）。

bias、激活、缩放、残差加、amax 跟踪——在朴素实现里都是独立 memory-bound kernel； 好 kernel 里它们被折叠进 GEMM 的 epilogue。你永远不写独立的 bias+GELU kernel， 而是在 epilogue 里组合；一次 global 往返而非三次。

12. Epilogue（尾声融合）

Epilogue = GEMM kernel 的最终处理阶段：发生在寄存器里的 wgmma 完成之后、写回 global 之前。 它把原始累加结果变换成最终输出格式，并顺便做其它 memory-bound 操作——目标是在做必要数学 （缩放、bias、激活）时隐藏延迟且不溢出寄存器。

12.1 Epilogue 操作清单

1. alpha * Acc（只缩放累加器）、alpha * Acc + beta * C（线性组合）
2. Bias 加（每行/每列 bias）
3. 激活（ReLU / GELU / SiLU）
4. 残差式加法路径
5. TopK + Softmax 融合（列向 softmax 变体）
6. Aux 张量操作（aux load / aux store）
7. 规约融合（行/列/标量规约）
8. Absmax/amax 跟踪（FP8 路径常用）
9. 缩放融合（A/B/C/D 的每行/列 scale、alpha-beta 变体）
10. 块缩放因子生成（块缩放工作流）

12.2 Epilogue 操作顺序

1. 载入 tile 所需全部输入：Acc、可选 C、bias、scales、可选 AuxIn
2. 先建基础表达式：Z = alpha*Acc + beta*C（无 C 则 Z = alpha*Acc）
3. 再加仿射项：行/列 bias、残差/额外线性项
4. 对 Z 施加非线性：ReLU / GELU / SiLU
5. 施加输出缩放/量化（需要时收集 amax）
6. 物化输出：主输出 D、可选 AuxOut（激活前/后）、可选行/列/标量规约
7. 最终类型转换 + store

12.3 模板

1. 常见前向融合：Z = alpha*Acc + beta*C → Z += bias → Y = activation(Z) → D = cast/scale(Y)
2. 训练式（存 aux）：Z = alpha*Acc + beta*C + bias → AuxOut = Z → Y = activation(Z) → D = Y
3. 反向式：dY = alpha*Acc + beta*C → dX = dActivation(dY, AuxIn) → dBias = reduction(dX) → D = dX
4. Softmax/TopK 路径：S = alpha*Acc + beta*C → rowmax/rowsum/exp 归一化 → 可选 TopK → D

12.4 融合生命周期（六阶段）

SM90 的 epilogue 不是单个函数，而是一个围绕”分离变化的部分与不变的部分“设计的六阶段流水线：

visit() 的主要风险是寄存器过度使用：融合深度超过寄存器预算时，local memory spill 会抵消融合收益。 postreduce() 是”交接边界”：数据 staged 后，CUDA core 从流水线释放，TMA 硬件独立管理异步写 gmem。

13. 总结与学习衔接

13.1 核心脉络速记

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

13.3 一句话记忆

Kernel Design = 把”算得慢”（compute-bound）与”搬得慢”（memory-bound）都变成”永不停”： warp specialization 分 producer/consumer、环形缓冲 + 双屏障把 load 与 WGMMA 重叠、 ping-pong 让 MMA 与 epilogue 重叠、persistent+Stream-K 调度消除尾效应、epilogue 融合省掉全部多余访存。

参考来源：8. Kernel Design.pdf（Lesson 8 / The Tricks，98 页）。