第 6 课：Shared Memory 布局与异步流水

参考答案版：包含完整代码实现与问答题深度参考解析。建议先独立完成练习题后再展开核对。

本课目标与完成标准

学完后你应能：安排 global→shared→register 多 stage 流水，并解释 fence/wait/barrier 的正确顺序。

通过必须同时满足：

独立补齐本课唯一的代码填空题，并通过给定检查；
三个问答题均说明因果链，而不是只报术语；
能指出至少一个正确性边界和一个性能取舍；
总分不低于 8/10，且没有一票否决级概念错误。

前置关系

课程前置：CUDA 线程模型、GEMM、C++ 模板基础
本课在路线中的作用：GEMM mainloop 用 shared memory 缓冲 A/B tile；多 stage 让下一 tile 的加载与当前 tile 的 MMA 重叠。

核心心智模型

1. 它是什么，解决什么问题

多阶段共享内存异步流水线（Multi-Stage Shared Memory Asynchronous Pipeline） 是现代 GPU（Ampere、Ada、Hopper）彻底打破“冯·诺依曼访存墙”、实现 Tensor Core 极限算力满载的核心生命线。

传统同步 GEMM 的停滞危机 vs. 异步环形流水（结合图 6-1 与图 6-2）： - 串行等待的算力暴跌：在早期 GPU 编程中，循环遵循“读取 Tile $\to$ 同步等待 $\to$ MMA 矩阵乘 $\to$ 同步等待”的串行模式。由于全局显存（HBM）的访问延迟高达 200~400 个时钟周期，计算单元在等待显存数据的漫长周期内完全处于冻结饥饿状态，硬件算力利用率极低； - 异步拷贝引擎（cp.async，图 6-1）：

图 6-1：cp.async 异步拷贝原理（图源自 NVIDIA Ampere 白皮书）

Ampere 引入了硬件级异步拷贝引擎。它允许线程仅发射拷贝指令，由显存控制器后台直接将 HBM 数据写入片上 Shared Memory，全程彻底绕过通用寄存器（Zero Register Overhead）； - 环形流水线（Circular Buffer Pipeline，图 6-2）：

sequenceDiagram
    autonumber
    participant Stage0 as SMem Stage 0
    participant Stage1 as SMem Stage 1
    participant MMA as Tensor Core 计算

    Note over Stage0,MMA: 异步多阶段双缓冲/多缓冲流水
    Stage0->>Stage0: 异步加载第 k 步数据 (cp.async)
    Stage0->>MMA: 读取第 k 步数据并执行矩阵乘
    Stage1->>Stage1: 同时异步预取第 k+1 步数据
    MMA-->>Stage0: 计算完毕释放 Stage 0
    Stage1->>MMA: 立即计算第 k+1 步，计算完全隐藏数据传输

图 6-2：Ampere/Hopper 多级异步环形流水线（Circular Pipeline）执行流

将共享内存划分为 $S$ 个环形阶段（Stages）。在当前时钟周期计算 Stage $k$ 的同时，硬件已在后台并行预取 Stage $k+1, \dots, k+S-1$ 的数据，使显存传输延迟被计算时间完全掩盖。

2. 它如何工作（结合图 6-1 与图 6-2）

CUTLASS 主循环采用标准的三阶段异步软件流水（Software Pipelining）：

序幕预取阶段（Prologue / Pipeline Fill）：
- 在进入计算主循环之前，线程连续向硬件下发 $S-1$ 个 Tile 的全局加载请求，分别填入 Shared Memory 的 Stage $0 \dots S-2$ ；
- 每轮预取通过 cp.async.commit_group 绑定为一个独立的硬件事务组（Group）；
稳态循环流水阶段（Steady-State Mainloop）：
- 等待就绪：调用 cp.async.wait_group<S-2>。硬件仅阻塞等待最老的一个 Stage 加载完成，其余后续 Stage 仍在后台异步传输；
- 可见性屏障：执行 __syncthreads() 或异步事务栅障（mbarrier.try_wait），确保该 Stage 物理数据对当前 Block 全体线程完全可见；
- 预取下一轮：调用 cp.async 发射第 $k + S - 1$ 步的数据搬运，写入环形缓冲区的下一个空闲槽位，调用 commit_group；
- 重叠计算：将当前已就绪的 Stage $k$ 数据读入寄存器，全速发射 Tensor Core MMA 计算指令。数据传输与矩阵乘加在硅片上完全物理并发！
尾声收口阶段（Epilogue / Pipeline Drain）：
- $K$ 维数据全部加载完毕后，依次等待并清空环形缓冲区中残留的最后几个 Stage，累加出最终完整的输出矩阵。

3. 正确性条件与常见误区

写后读（RAW）与读后写（WAR）双向冒险：
- RAW 冒险：未等待 cp.async 事务组完成就强行读取 Shared Memory，会读出未初始化的显存随机脏数据；
- WAR 冒险：在当前 Stage 的 MMA 计算完全结束之前，过早发起新的 cp.async 覆写该阶段内存，会导致计算中的线程读到未来时刻的新数据；
栅障作用域的聚合一致性：
- 栅障指令必须在整个 CTA 范围内严格对齐。绝不能在存在线程分流的分支语句内部调用同步，否则会导致硬件屏障计数器永久无法归零而死锁。

4. 性能与工程取舍（流水级数与 Occupancy 权衡）

最小流水级数（Minimum Stages）的物理推导：
- 设全局内存加载延迟为 $T_{\text{copy}}$ （通常约为 300 cycles），单个 Tile 的 Tensor Core 计算耗时为 $T_{\text{compute}}$ ；
- 若要完全隐藏显存加载延迟，在流水线中并发飞行的预取批次必须满足： $\text{In-flight Stages} \ge \left\lceil \frac{T_{\text{copy}}}{T_{\text{compute}}} \right\rceil$
- 加上当前正在被 MMA 计算占用的 1 个 Stage，实现零气泡流水所需的总阶段数为： $S_{\text{min}} = \left\lceil \frac{T_{\text{copy}}}{T_{\text{compute}}} \right\rceil + 1$
Stage 激增带来的“Occupancy 坍塌”陷阱：
- 共享内存总消耗量随 Stage 数呈线性严格增长： $\text{SMEM} = S \times (\text{Size}_A + \text{Size}_B) \times \text{sizeof}(T)$ ；
- 若从 2-stage 盲目增加到 5-stage，单 Block 的 SMEM 占用可能从 40 KB 暴涨至 100 KB。单个 SM 允许驻留的活跃 Block 数量可能从 4 个骤降至 1 个；
- 当 SM 仅剩 1 个活跃 Block 时，一旦遭遇任何未预期的 Cache Miss，Warp 调度器将失去任何备用线程束来掩盖延迟，算子整体耗时反而成倍恶化！

具体演示：流水线阶段配比与延迟掩盖演算

设某 GPU 体系下，一个 Tile 的全局显存拷贝延迟 $T_{\text{copy}} = 300\text{ 周期}$ ，单个 Tile 的 Tensor Core MMA 运算时间 $T_{\text{compute}} = 120\text{ 周期}$ ：

流水级数测算： $\left\lceil \frac{300}{120} \right\rceil = \lceil 2.5 \rceil = 3 \text{ 个预取阶段}$ 加上当前正在参与计算的 1 个阶段，最低所需总阶段数为： $S = 3 + 1 = 4 \text{ Stages}$
稳态执行微观时序图：
- 周期 0 ~ 120：MMA 计算 Stage 0；硬件在后台同时传输 Stage 1（进度 120/300）、Stage 2（进度 0/300）、Stage 3（新发射）；
- 周期 120 ~ 240：MMA 计算 Stage 1；硬件后台传输 Stage 2、Stage 3、Stage 4；
- 周期 240 ~ 360：MMA 计算 Stage 2；
- 因果结论：只要始终维持 4 个 Stage 环形流转，每次 MMA 计算完成时，下一个 Stage 的数据都恰好已经在片上 Shared Memory 准备就绪，计算引擎实现 100% 无气泡满载运行。

实践任务：唯一代码填空题

补齐覆盖 copy 延迟所需的最小 stage 数（包含正在计算的一 stage）。

规则：只能修改 TODO/______ 所在位置；不要删除断言或放宽误差。代码注释说明了每个边界条件。

Code

def minimum_stages(copy_cycles, compute_cycles):
    """两个参数必须为正整数。"""
    # TODO：只补齐下面这个表达式。
    return ______

assert minimum_stages(300,120) == 4
assert minimum_stages(120,120) == 2
assert minimum_stages(1,120) == 2

检查方法

运行本单元格；所有 assert 必须通过。另手工构造一个边界输入，解释预期结果。

提交时请给出：补齐后的代码、实际运行输出（环境不可用时注明“仅静态审查”）以及对失败用例的解释。

Q1

不要背定义：请从输入、状态变化和输出三个阶段解释“Shared Memory 布局与异步流水”的工作机制。

你的答案：

Q2

只增加 stage 数而不检查 shared-memory occupancy，为什么可能变慢？

你的答案：

Q3

在 profiler 中看到 memory dependency stalls，如何区分 stage 不足与访存不合并？

你的答案：

评分与通过规则

代码 4 分：正常输入 2 分，边界输入 1 分，解释实现 1 分；
Q1～Q3 各 2 分；
一票否决：结果碰巧正确但核心因果链错误、删除边界检查、把未运行结果说成实测。

需要提示时按四级机制请求：概念区域 → 具体方向 → 关键局部 → 完整答案。

💡 点击展开参考答案与详细代码实现

参考答案（仅 answer 分支）

先完成题目再核对。即使代码一致，也要能解释关键步骤，并尝试更换一个输入规模。

Code

def minimum_stages(copy_cycles, compute_cycles):
    """两个参数必须为正整数。"""
    # 参考实现：表达式直接对应上文不变量。
    return (copy_cycles + compute_cycles - 1) // compute_cycles + 1

assert minimum_stages(300,120) == 4
assert minimum_stages(120,120) == 2
assert minimum_stages(1,120) == 2

Q1 参考答案

输入阶段（Prologue Multi-Stage Pipeline Priming）：
- 在启动 GEMM 计算之前，算法在序幕阶段（Prologue）向硬件异步拷贝引擎发射连续 $S-1$ 个 Tile 的全局内存加载指令；
- 硬件直接发起 cp.async，将数据从 HBM 批量注入片上环形共享内存的 Stage $0 \sim S-2$ 槽位中，全程不占用通用寄存器；
状态变化阶段（Steady-State Overlapped Execution）：
- 异步消费等待：主循环每轮迭代调用 cp.async.wait_group<S-2>，仅同步等待最早发射的一个 Stage 加载完毕，并通过屏障保证该 Stage 数据在片上对所有线程可见；
- 物理并发推进：线程一方面向后台下发下一轮的异步预取请求（Stage $k + S - 1$ ），另一方面直接从已就绪的 Stage $k$ 中加载操作数切片，全速发射 Tensor Core MMA 计算指令。数据传输延迟在时钟周期级别被纯粹的计算时间完全吸收；
- 缓冲区槽位轮转：计算完成后，释放当前 Stage $k$ ，指针沿环形拓扑向前滚动；
输出阶段（Pipeline Drain & Epilogue Store）：
- 当全局收缩维度遍历完毕后，流水线进入尾声，清空并在寄存器中累加最后几个飞行中的 Stage 数据，最终将累加结果合并写回输出显存。

Q2 参考答案

若脱离硬件资源容量限制，盲目增加 Stage 级数（例如由 3-Stage 增加到 6-Stage），很可能导致整体性能严重劣化，原因在于触发了共享内存引发的 Occupancy 坍塌（Occupancy Collapse）：

共享内存容量的线性挤压：
- 单 Block 所需的共享内存容量随 Stage 数量呈严格线性比例膨胀： $\text{SMEM Total} = S \times (M_{\text{tile}} \times K_{\text{tile}} + K_{\text{tile}} \times N_{\text{tile}}) \times \text{sizeof}(T)$
- 以 NVIDIA Ampere A100 为例，单个 SM 的用户可用共享内存上限为 164 KB；
- 若 $S=3$ 时单 Block 占用 48 KB SMEM，则单个 SM 能够并发驻留 $\lfloor 164 / 48 \rfloor = 3$ 个活跃 Block，提供了充裕的并发 Warp 来掩盖各种偶发停顿；
- 一旦将 Stage 提高到 $S=6$ ，单 Block 占用暴涨至 96 KB，单个 SM 只能物理容纳 $\lfloor 164 / 96 \rfloor = 1$ 个活跃 Block！SM 理论 Occupancy 当场腰斩；
延迟隐藏能力的彻底丧失：
- 当单 SM 仅剩 1 个活跃 Block 时，一旦该 Block 遭遇指令 Cache Miss、全局同步栅障等待或分支发散，SM 的 Warp 调度器将没有任何其他就绪的 Warp 可供切换；
- 尽管异步拷贝的预取深度变深了，但硬件核心在遇到突发微架构停顿时陷入绝对空转，算子的端到端执行延迟反而急剧飙升。

Q3 参考答案

在 Nsight Compute (NCU) 性能分析器中，当看到突出的“Memory Dependency Stalls”（内存依赖停顿）时，区分究竟是“流水线 Stage 深度不足”还是“全局访存未合并”必须依靠显存传输扇区利用率与流水线飞行队列的交叉对比：

第一诊断维度：显存扇区利用率（DRAM Sector Utilization / Coalescing Efficiency）：
- 审查指标：dram__sectors_read.sum 与算法请求总字节数的比率，或者 l1tex__t_sectors_pipe_lsu_mem_global_op_ld 与 requests 的比值；
- 若为访存不合并（Uncoalesced Access）：该比值显著偏低（例如相邻线程访问跨度很大，导致每个 32 字节扇区只读取了 4 字节有效数据，扇区有效利用率仅 12.5%~25%）。此时显存总线充斥着碎片化的事务，实际搬运的物理字节数远超理论请求量，说明根本瓶颈在于地址映射不连续，增加 Stage 完全无济于事；
第二诊断维度：Warp 停顿状态与异步队列饱和度（Pipeline Queue Saturation）：
- 审查指标：Warp State Statistics 中的 Stall Long Scoreboard、Stall Wait 比例，以及 cp.async 飞行事务组队列深度；
- 若为 Stage 不足（Pipeline Starvation）：显存扇区利用率达到近乎 100%（即每次读入都是完美对齐的 128 字节合并传输），但 Warp 却长期停顿在等待内存加载结果的同步屏障上（Stall Wait 极高），且异步队列经常处于清空断流状态。这证明内存合并完美，但由于 Stage 级数太浅，数据往返延迟无法被计算完全掩盖，此时正是增加 Stage 级数的标准指征。

参考资料

资料用于建立事实基线；面试回答仍需用自己的语言组织。

← 上一课
第 5 课：MMA Atom、Tiled M… 下一课 →
第 7 课：混合精度、Accumulator…

--- title: "第 6 课：Shared Memory 布局与异步流水" jupyter: python3 --- ::: {.callout-note appearance="simple"} **参考答案版**：包含完整代码实现与问答题深度参考解析。建议先独立完成练习题后再展开核对。 ::: ## 本课目标与完成标准学完后你应能：安排 global→shared→register 多 stage 流水，并解释 fence/wait/barrier 的正确顺序。通过必须同时满足： - 独立补齐本课唯一的代码填空题，并通过给定检查； - 三个问答题均说明因果链，而不是只报术语； - 能指出至少一个正确性边界和一个性能取舍； - 总分不低于 8/10，且没有一票否决级概念错误。 ## 前置关系 - 课程前置：CUDA 线程模型、GEMM、C++ 模板基础 - 本课在路线中的作用：GEMM mainloop 用 shared memory 缓冲 A/B tile；多 stage 让下一 tile 的加载与当前 tile 的 MMA 重叠。 ## 核心心智模型 ### 1. 它是什么，解决什么问题 **多阶段共享内存异步流水线（Multi-Stage Shared Memory Asynchronous Pipeline）** 是现代 GPU（Ampere、Ada、Hopper）彻底打破“冯·诺依曼访存墙”、实现 Tensor Core 极限算力满载的核心生命线。 **传统同步 GEMM 的停滞危机 vs. 异步环形流水（结合图 6-1 与图 6-2）**： - **串行等待的算力暴跌**：在早期 GPU 编程中，循环遵循“读取 Tile $\to$ 同步等待 $\to$ MMA 矩阵乘 $\to$ 同步等待”的串行模式。由于全局显存（HBM）的访问延迟高达 200~400 个时钟周期，计算单元在等待显存数据的漫长周期内完全处于冻结饥饿状态，硬件算力利用率极低； - **异步拷贝引擎（`cp.async`，图 6-1）**： ::: {.img-card} ![异步拷贝流水图](assets/figs/fig_06_cp_async_拷贝原理_图源自_NVIDIA_Ampere_白皮书_.png){width="90%"} 图 6-1：cp.async 异步拷贝原理（图源自 NVIDIA Ampere 白皮书） ::: Ampere 引入了硬件级异步拷贝引擎。它允许线程仅发射拷贝指令，由显存控制器后台直接将 HBM 数据写入片上 Shared Memory，**全程彻底绕过通用寄存器（Zero Register Overhead）**； - **环形流水线（Circular Buffer Pipeline，图 6-2）**： ```{mermaid} sequenceDiagram autonumber participant Stage0 as SMem Stage 0 participant Stage1 as SMem Stage 1 participant MMA as Tensor Core 计算 Note over Stage0,MMA: 异步多阶段双缓冲/多缓冲流水 Stage0->>Stage0: 异步加载第 k 步数据 (cp.async) Stage0->>MMA: 读取第 k 步数据并执行矩阵乘 Stage1->>Stage1: 同时异步预取第 k+1 步数据 MMA-->>Stage0: 计算完毕释放 Stage 0 Stage1->>MMA: 立即计算第 k+1 步，计算完全隐藏数据传输 ``` 图 6-2：Ampere/Hopper 多级异步环形流水线（Circular Pipeline）执行流 将共享内存划分为 $S$ 个环形阶段（Stages）。在当前时钟周期计算 Stage $k$ 的同时，硬件已在后台并行预取 Stage $k+1, \dots, k+S-1$ 的数据，使显存传输延迟被计算时间完全掩盖。 ### 2. 它如何工作（结合图 6-1 与图 6-2） CUTLASS 主循环采用标准的三阶段异步软件流水（Software Pipelining）： 1. **序幕预取阶段（Prologue / Pipeline Fill）**： - 在进入计算主循环之前，线程连续向硬件下发 $S-1$ 个 Tile 的全局加载请求，分别填入 Shared Memory 的 Stage $0 \dots S-2$； - 每轮预取通过 `cp.async.commit_group` 绑定为一个独立的硬件事务组（Group）； 2. **稳态循环流水阶段（Steady-State Mainloop）**： - **等待就绪**：调用 `cp.async.wait_group<S-2>`。硬件仅阻塞等待最老的一个 Stage 加载完成，其余后续 Stage 仍在后台异步传输； - **可见性屏障**：执行 `__syncthreads()` 或异步事务栅障（`mbarrier.try_wait`），确保该 Stage 物理数据对当前 Block 全体线程完全可见； - **预取下一轮**：调用 `cp.async` 发射第 $k + S - 1$ 步的数据搬运，写入环形缓冲区的下一个空闲槽位，调用 `commit_group`； - **重叠计算**：将当前已就绪的 Stage $k$ 数据读入寄存器，全速发射 Tensor Core MMA 计算指令。数据传输与矩阵乘加在硅片上完全物理并发！ 3. **尾声收口阶段（Epilogue / Pipeline Drain）**： - $K$ 维数据全部加载完毕后，依次等待并清空环形缓冲区中残留的最后几个 Stage，累加出最终完整的输出矩阵。 ### 3. 正确性条件与常见误区 1. **写后读（RAW）与读后写（WAR）双向冒险**： - **RAW 冒险**：未等待 `cp.async` 事务组完成就强行读取 Shared Memory，会读出未初始化的显存随机脏数据； - **WAR 冒险**：在当前 Stage 的 MMA 计算完全结束之前，过早发起新的 `cp.async` 覆写该阶段内存，会导致计算中的线程读到未来时刻的新数据； 2. **栅障作用域的聚合一致性**： - 栅障指令必须在整个 CTA 范围内严格对齐。绝不能在存在线程分流的分支语句内部调用同步，否则会导致硬件屏障计数器永久无法归零而死锁。 ### 4. 性能与工程取舍（流水级数与 Occupancy 权衡） 1. **最小流水级数（Minimum Stages）的物理推导**： - 设全局内存加载延迟为 $T_{\text{copy}}$（通常约为 300 cycles），单个 Tile 的 Tensor Core 计算耗时为 $T_{\text{compute}}$； - 若要完全隐藏显存加载延迟，在流水线中并发飞行的预取批次必须满足： $$\text{In-flight Stages} \ge \left\lceil \frac{T_{\text{copy}}}{T_{\text{compute}}} \right\rceil$$ - 加上当前正在被 MMA 计算占用的 1 个 Stage，实现零气泡流水所需的总阶段数为： $$S_{\text{min}} = \left\lceil \frac{T_{\text{copy}}}{T_{\text{compute}}} \right\rceil + 1$$ 2. **Stage 激增带来的“Occupancy 坍塌”陷阱**： - 共享内存总消耗量随 Stage 数呈**线性严格增长**：$\text{SMEM} = S \times (\text{Size}_A + \text{Size}_B) \times \text{sizeof}(T)$； - 若从 2-stage 盲目增加到 5-stage，单 Block 的 SMEM 占用可能从 40 KB 暴涨至 100 KB。单个 SM 允许驻留的活跃 Block 数量可能从 4 个骤降至 1 个； - 当 SM 仅剩 1 个活跃 Block 时，一旦遭遇任何未预期的 Cache Miss，Warp 调度器将失去任何备用线程束来掩盖延迟，算子整体耗时反而成倍恶化！ ## 具体演示：流水线阶段配比与延迟掩盖演算设某 GPU 体系下，一个 Tile 的全局显存拷贝延迟 $T_{\text{copy}} = 300\text{ 周期}$，单个 Tile 的 Tensor Core MMA 运算时间 $T_{\text{compute}} = 120\text{ 周期}$： - **流水级数测算**： $$\left\lceil \frac{300}{120} \right\rceil = \lceil 2.5 \rceil = 3 \text{ 个预取阶段}$$ 加上当前正在参与计算的 1 个阶段，最低所需总阶段数为： $$S = 3 + 1 = 4 \text{ Stages}$$ - **稳态执行微观时序图**： - **周期 0 ~ 120**：MMA 计算 Stage 0；硬件在后台同时传输 Stage 1（进度 120/300）、Stage 2（进度 0/300）、Stage 3（新发射）； - **周期 120 ~ 240**：MMA 计算 Stage 1；硬件后台传输 Stage 2、Stage 3、Stage 4； - **周期 240 ~ 360**：MMA 计算 Stage 2； - **因果结论**：只要始终维持 4 个 Stage 环形流转，每次 MMA 计算完成时，下一个 Stage 的数据都恰好已经在片上 Shared Memory 准备就绪，计算引擎实现 100% 无气泡满载运行。 ## 实践任务：唯一代码填空题补齐覆盖 copy 延迟所需的最小 stage 数（包含正在计算的一 stage）。规则：只能修改 `TODO`/`______` 所在位置；不要删除断言或放宽误差。代码注释说明了每个边界条件。 ```{python} #| course_role: exercise def minimum_stages(copy_cycles, compute_cycles): """两个参数必须为正整数。""" # TODO：只补齐下面这个表达式。 return ______ assert minimum_stages(300,120) == 4 assert minimum_stages(120,120) == 2 assert minimum_stages(1,120) == 2 ``` ### 检查方法运行本单元格；所有 `assert` 必须通过。另手工构造一个边界输入，解释预期结果。提交时请给出：补齐后的代码、实际运行输出（环境不可用时注明“仅静态审查”）以及对失败用例的解释。 ### Q1 不要背定义：请从输入、状态变化和输出三个阶段解释“Shared Memory 布局与异步流水”的工作机制。 **你的答案：** ### Q2 只增加 stage 数而不检查 shared-memory occupancy，为什么可能变慢？ **你的答案：** ### Q3 在 profiler 中看到 memory dependency stalls，如何区分 stage 不足与访存不合并？ **你的答案：** ## 评分与通过规则 - 代码 4 分：正常输入 2 分，边界输入 1 分，解释实现 1 分； - Q1～Q3 各 2 分； - 一票否决：结果碰巧正确但核心因果链错误、删除边界检查、把未运行结果说成实测。需要提示时按四级机制请求：概念区域 → 具体方向 → 关键局部 → 完整答案。 ::: {.callout-tip collapse="true" title="💡 点击展开参考答案与详细代码实现" icon="false"} ## 参考答案（仅 answer 分支）先完成题目再核对。即使代码一致，也要能解释关键步骤，并尝试更换一个输入规模。 ```{python} #| course_role: solution def minimum_stages(copy_cycles, compute_cycles): """两个参数必须为正整数。""" # 参考实现：表达式直接对应上文不变量。 return (copy_cycles + compute_cycles - 1) // compute_cycles + 1 assert minimum_stages(300,120) == 4 assert minimum_stages(120,120) == 2 assert minimum_stages(1,120) == 2 ``` ### Q1 参考答案 1. **输入阶段（Prologue Multi-Stage Pipeline Priming）**： - 在启动 GEMM 计算之前，算法在序幕阶段（Prologue）向硬件异步拷贝引擎发射连续 $S-1$ 个 Tile 的全局内存加载指令； - 硬件直接发起 `cp.async`，将数据从 HBM 批量注入片上环形共享内存的 Stage $0 \sim S-2$ 槽位中，全程不占用通用寄存器； 2. **状态变化阶段（Steady-State Overlapped Execution）**： - **异步消费等待**：主循环每轮迭代调用 `cp.async.wait_group<S-2>`，仅同步等待最早发射的一个 Stage 加载完毕，并通过屏障保证该 Stage 数据在片上对所有线程可见； - **物理并发推进**：线程一方面向后台下发下一轮的异步预取请求（Stage $k + S - 1$），另一方面直接从已就绪的 Stage $k$ 中加载操作数切片，全速发射 Tensor Core MMA 计算指令。数据传输延迟在时钟周期级别被纯粹的计算时间完全吸收； - **缓冲区槽位轮转**：计算完成后，释放当前 Stage $k$，指针沿环形拓扑向前滚动； 3. **输出阶段（Pipeline Drain & Epilogue Store）**： - 当全局收缩维度遍历完毕后，流水线进入尾声，清空并在寄存器中累加最后几个飞行中的 Stage 数据，最终将累加结果合并写回输出显存。 ### Q2 参考答案若脱离硬件资源容量限制，盲目增加 Stage 级数（例如由 3-Stage 增加到 6-Stage），很可能导致整体性能严重劣化，原因在于触发了**共享内存引发的 Occupancy 坍塌（Occupancy Collapse）**： 1. **共享内存容量的线性挤压**： - 单 Block 所需的共享内存容量随 Stage 数量呈严格线性比例膨胀： $$\text{SMEM Total} = S \times (M_{\text{tile}} \times K_{\text{tile}} + K_{\text{tile}} \times N_{\text{tile}}) \times \text{sizeof}(T)$$ - 以 NVIDIA Ampere A100 为例，单个 SM 的用户可用共享内存上限为 164 KB； - 若 $S=3$ 时单 Block 占用 48 KB SMEM，则单个 SM 能够并发驻留 $\lfloor 164 / 48 \rfloor = 3$ 个活跃 Block，提供了充裕的并发 Warp 来掩盖各种偶发停顿； - 一旦将 Stage 提高到 $S=6$，单 Block 占用暴涨至 96 KB，单个 SM 只能物理容纳 $\lfloor 164 / 96 \rfloor = 1$ 个活跃 Block！SM 理论 Occupancy 当场腰斩； 2. **延迟隐藏能力的彻底丧失**： - 当单 SM 仅剩 1 个活跃 Block 时，一旦该 Block 遭遇指令 Cache Miss、全局同步栅障等待或分支发散，SM 的 Warp 调度器将**没有任何其他就绪的 Warp 可供切换**； - 尽管异步拷贝的预取深度变深了，但硬件核心在遇到突发微架构停顿时陷入绝对空转，算子的端到端执行延迟反而急剧飙升。 ### Q3 参考答案在 Nsight Compute (NCU) 性能分析器中，当看到突出的“Memory Dependency Stalls”（内存依赖停顿）时，区分究竟是“流水线 Stage 深度不足”还是“全局访存未合并”必须依靠**显存传输扇区利用率与流水线飞行队列的交叉对比**： 1. **第一诊断维度：显存扇区利用率（DRAM Sector Utilization / Coalescing Efficiency）**： - 审查指标：`dram__sectors_read.sum` 与算法请求总字节数的比率，或者 `l1tex__t_sectors_pipe_lsu_mem_global_op_ld` 与 `requests` 的比值； - **若为访存不合并（Uncoalesced Access）**：该比值显著偏低（例如相邻线程访问跨度很大，导致每个 32 字节扇区只读取了 4 字节有效数据，扇区有效利用率仅 12.5%~25%）。此时显存总线充斥着碎片化的事务，实际搬运的物理字节数远超理论请求量，说明**根本瓶颈在于地址映射不连续，增加 Stage 完全无济于事**； 2. **第二诊断维度：Warp 停顿状态与异步队列饱和度（Pipeline Queue Saturation）**： - 审查指标：`Warp State Statistics` 中的 `Stall Long Scoreboard`、`Stall Wait` 比例，以及 `cp.async` 飞行事务组队列深度； - **若为 Stage 不足（Pipeline Starvation）**：显存扇区利用率达到近乎 100%（即每次读入都是完美对齐的 128 字节合并传输），但 Warp 却长期停顿在等待内存加载结果的同步屏障上（`Stall Wait` 极高），且异步队列经常处于清空断流状态。这证明**内存合并完美，但由于 Stage 级数太浅，数据往返延迟无法被计算完全掩盖，此时正是增加 Stage 级数的标准指征**。 ## 参考资料 - [CuTe Layout Algebra](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/cute/01_layout.html) - [CuTe Tensors](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/cute/03_tensor.html) - [CuTe Algorithms](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/cute/04_algorithms.html) - [CUTLASS GEMM API](https://docs.nvidia.com/cutlass/latest/media/docs/cpp/gemm_api.html) - [CUTLASS repository](https://github.com/NVIDIA/cutlass) 资料用于建立事实基线；面试回答仍需用自己的语言组织。 ::: ::: {.course-pagination} [**← 上一课** 第 5 课：MMA Atom、Tiled M...](lesson05_tiled_mma.qmd){.pagination-link .pagination-prev} [**下一课 →** 第 7 课：混合精度、Accumulator...](lesson07_mixed_precision.qmd){.pagination-link .pagination-next} :::