cp.async / ldmatrix 消融实验 —— 初步规划（SM80 专用）

目标设备：SM80（Ampere，如 A100）

目标：设计一套逐步叠加优化的消融实验（ablation），量化每一项优化（ldmatrix vs ld.shared、cp.async vs 同步 load）到底带来了多少性能提升。

本文档是 初步规划.md 的 SM80 裁剪版：已移除全部 TMA 相关内容（原 TMA 章节改为「SM80 能力边界」），并补充了 Ampere 特有的实验注意事项。 与原版的差异见文末 §6 与完整版的差异。

0. SM80 能力边界

0.1 各特性的可用性

结论：TMA 及一切依赖 TMA 的代码在 SM80 上不可用。

0.2 仓库中受影响的文件

⚠️ 注意：kernels/hgemm_sm120/readme.md 里的性能数字（cp.async 34.00 TF / TMA 34.31 TF 等）在 SM80 上无法复现，不能直接当作 baseline。所有 baseline 必须在你自己的 SM80 设备上重测。

0.3 可直接复用的 PTX 宏（SM80 安全）

来自 kernels/hgemm/hgemm.cu:22-40、kernels/sgemm/sgemm_tf32.cu:22-35：

#define CP_ASYNC_CG(dst, src)  asm volatile("cp.async.cg.shared.global.L2::128B [%0], [%1], 16;\n" ...)
#define CP_ASYNC_COMMIT_GROUP()      // cp.async.commit_group
#define CP_ASYNC_WAIT_GROUP_0()      // cp.async.wait_group 0
#define LDMATRIX_X4(R0,R1,R2,R3,PTR)   // ldmatrix.sync.aligned.m8n8.x4.shared.b16
#define LDMATRIX_X2(R0,R1,PTR)         // ldmatrix.sync.aligned.m8n8.x2.shared.b16
#define LDMATRIX_X2_TRANS(R0,R1,PTR)   // ldmatrix.sync.aligned.m8n8.x2.trans.shared.b16

多级流水版本需要的 wait_group<N>（N>0）在 hgemm_sm120 里有现成写法，但那个文件依赖 TMA 不可用 —— 需要把这段 PTX 抄到新文件里：

template <int N>
__device__ __forceinline__ void cp_async_wait_group() {
    if constexpr (N == 0)      asm volatile("cp.async.wait_group 0;\n" ::);
    else if constexpr (N == 1) asm volatile("cp.async.wait_group 1;\n" ::);
    else if constexpr (N == 2) asm volatile("cp.async.wait_group 2;\n" ::);
    else                       asm volatile("cp.async.wait_group 3;\n" ::);
}

0.4 已有对照数据（仅供参考，不能直接当 SM80 baseline）

kernels/hgemm/readme.md（4096³ bf16，来源机型未知）：

hgemm_cublas              4.210 ms  32.64 TF
hgemm_naive               5.191 ms  26.47 TF   ← ldmatrix + mma（无 swizzle / 无双缓冲）
hgemm_bcf                 4.337 ms  31.69 TF   ← + BCF swizzle
hgemm_bcf_dbf             4.096 ms  33.55 TF   ← + double buffer
hgemm_bcf_dbf_rw          4.076 ms  33.72 TF   ← + coalesced gmem r/w

kernels/sgemm/readme.md（4096³ TF32）：

sgemm_tf32_bt_swizzle_dbf      8.723 ms  15.76 TF   (非 cp.async：普通 load + st.shared)
sgemm_tf32_swizzle_bcf_dbf     8.276 ms  16.61 TF   (cp.async)

⚠️ 这组差异里 swizzle / dbf / grid-swizzle 三个因素是混在一起的，不能直接归因到 cp.async —— 这正是需要做受控消融的原因。

1. 消融实验设计

核心思路：把优化拆成两条独立的阶梯，每条阶梯逐档叠加，每档只改变一个变量，从而把增量收益归因清楚。

固定基线：GEMM 形状 M=N=K=4096、bf16、BM=BN=128、BK=32、256 线程 / 8 warp、warp tile 64×32、mma.m16n8k16。 （建议全部在 kernels/hgemm/hgemm.cu 的结构上做，保持其它变量一致。）

1.1 阶梯 A：smem → reg（ldmatrix vs ld.shared）

对应 ldmatrix 这一项。需要三档，且必须给 ld.shared 变体最有利布局，否则测到的是布局差异而非指令效率。

A3 即现有 kernels/hgemm/hgemm.cu:131/142 的写法。

关键：A3 相比 A1/A2 省的是 LSU 指令数和地址计算指令，不省带宽（数据搬运量相同）。 另外注意 ldmatrix.x4 是整个 warp 协作指令（32 个 lane 各提供一个 8×8 tile 的行地址，硬件自动 gather+转置+广播），而 ld.shared 是 per-lane 独立访存 —— 这正是 A3 的核心优势来源。

1.2 阶梯 B：gmem → smem（cp.async vs 同步 load）

SM80 上没有 TMA 作为「第三档」，B3 就是搬运路径的终点。若还想再往上走，下一档是 B4：cp.async + 寄存器级双缓冲（ldmatrix 结果在寄存器里预取下一 k_step，见 hgemm/hgemm.cu:355 的 bcf_dbf 思路）。

1.3 ⚠️ 三个会让结论作废的坑

坑 1：B1 单 buffer 下 cp.async 几乎无收益，甚至更慢。 若按 B0→B1→B2 顺序测，会得出"cp.async 没用"的错误结论。cp.async 的价值在于释放 LSU 和 issue slot，只有当有第二级 buffer 承接、让搬运算与 mma 重叠时才能兑现。

✅ 正确阶梯顺序：同步 load + 双缓冲 → cp.async + 双缓冲 → cp.async + 多级流水。 也就是说，"双缓冲"和"cp.async"必须拆成两步分别加，否则无法区分收益来自哪一项。

实现上：同步 load + 双缓冲 = 普通的 ld.global+st.shared 到 buffer[write]，__syncthreads() 后读 buffer[read]，再 __syncthreads() 换索引； cp.async + 双缓冲 = 把搬运算换成 cp.async.cg + commit_group/wait_group。

坑 2：在最终 kernel 上测 ldmatrix vs ld.shared，差异会被淹没。 若 kernel 被 mma 吞吐或 DRAM 带宽卡住，smem→reg 段的差异在 wall-clock 上看不见。 在 A100 上这个风险尤其高：A100 的 FP16 Tensor Core 峰值约 312 TFLOPS，而实测 GEMM 常在 150~250 TFLOPS，两者差距不像消费级卡那么大，LSU 段更容易被掩盖。

✅ 建议额外做一个 smem→reg 微基准（只 load 不 mma，循环 N 次），直接测 LSU 吞吐差异； 同时用 ncu 采集：

smsp__inst_executed_pipe_lsu_mem_shared_op_ld.sum
l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum
sm__inst_executed_pipe_lsu.avg.pct_of_peak_sustained_active

坑 3：只测一种 shape 会被误导。

• K 小（如 M=N=4096, K=256）→ latency-bound，cp.async / 多级流水收益最大
• K 大（K=8192）→ compute-bound，所有搬运优化收益都被 mma 吞吐盖住

✅ 建议扫 K ∈ {256, 512, 1024, 4096, 8192} × M/N ∈ {512, 4096}，画热力图。

另外两条硬性要求：

• 锁频：A100 上用 nvidia-smi -lgc <MHz> 锁核心频率、nvidia-smi -lmc <MHz> 锁显存频率（hgemm_sm120/readme.md 里有 A100 的做法），测完用 -rgc / -rmc 复位。否则重测噪声大于效应。
• 精度：cp.async 与同步 load 都不改数值；ldmatrix 与 ld.shared 也只重排不换算（ldmatrix.trans 仅做转置重排），因此正确性可直接 diff 对比，误差应为 0。

1.4 SM80 / A100 特有的注意事项

1. smem 容量：A100 每 SM 164 KB，单 block 最多 163 KB（需 cudaFuncSetAttribute(cudaFuncAttributeMaxDynamicSharedMemorySize, ...) 显式开启 >48KB）。 本实验 BM=BN=128, BK=32, bf16 → A tile = 128×32×2 = 8 KB，B tile = 32×128×2 = 8 KB，每 stage 16 KB。

   • STAGES=1 → 16 KB
   • STAGES=2 → 32 KB
   • STAGES=3 → 48 KB（正好卡在默认上限，必须显式 opt-in）
   • STAGES=4 → 64 KB（需 opt-in，且会压低 occupancy） → 建议 STAGES ∈ {2, 3, 4}，并同时记录 achieved occupancy 作为辅助解释变量。

2. cp.async.cg 的 16B 对齐：cp.async.cg 只支持 16B 粒度，且 gmem 地址需 16B 对齐。若 K 或 N 不是 8 的倍数（bf16 下 16B = 8 个元素），边界 tile 需要 fallback 路径。 实验时务必保证 M/N/K 都是 8 的倍数，把 fallback 从性能数据里排除掉。

3. cp.async 与 L2 行为：.cg 绕过 L1，命中 L2。A100 L2 为 40 MB（大），所以 K 较小的 shape 会大量命中 L2 —— 这也是为什么坑 3 里 K 的扫描很重要。

4. cp.async.wait_group 的语义是"按组计数"，而不是按字节。STAGES=3 时用 wait_group<1> 允许 1 组仍在飞行中，不能用 wait_group<0>（那会把流水退化成单 buffer）。这一点在写 B3 时最容易写错。

5. A100 上没有 wgmma，所有 mma 都是 warp 级 mma.sync，因此 ldmatrix 的收益不会被 Hopper 的 wgmma 掩盖 —— 这其实是 SM80 做这个实验的优势：在 SM90 上 wgmma 直接从 smem 取数，ldmatrix 的收益会大幅缩水。SM80 是观察 ldmatrix 价值更干净的平台。

6. TF32 的注意：如果同时在 sgemm_tf32.cu 上做对照，A100 上 TF32 Tensor Core 吞吐约为 FP16 的一半，别把两者的 TFLOPS 直接横向比较。

1.5 SIMT / Tensor Core 与访存指令的层级关系（重要澄清）

这一节回答一个容易踩坑的问题：SIMT 和 cp.async / ldmatrix 是同一层级的东西吗？ 答案：不是。 把 SIMT 误当作 ldmatrix / cp.async 的"对照组"会直接导致实验结论作废。

1.5.1 三个层级

SIMT = Single Instruction, Multiple Threads，是 NVIDIA 的硬件执行模型：一条指令由一个 warp 的 32 个线程同时执行。

它与 SIMD 的区别：SIMD 是"一条指令操作 N 个数据槽，程序员手动向量化"；SIMT 是"一条指令操作 1 个数据，硬件在 warp 内自动 SIMD 化"，程序员写标量代码。

关键：所有 CUDA kernel 都是 SIMT 的，包括 Tensor Core kernel —— 只是 mma.sync 指令本身的执行方式（一条指令被 warp 内多个线程分组协作完成）比普通 FMA 更"特殊"。

1.5.2 为什么 SIMT 不能当 ldmatrix / cp.async 的对照组

• ldmatrix 服务的对象是 Tensor Core 的 fragment 布局（m16n8k16 的 lane→(row,col) 映射），在 SIMT kernel 里根本用不了；
• cp.async 虽然技术上也能用于 SIMT kernel，但 SIMT kernel 通常不需要 smem staging（每个线程直接标量读 gmem 算），强行加上只会引入额外搬运成本。

如果把两者混在同一张消融表里，会犯一个经典错误：把 8 倍的架构收益算到 ldmatrix 头上。

1.5.3 SIMT 的正确用法

用途 1：性能上下界锚点（不是消融实验的一环）

// SIMT 版（kernels/sgemm/sgemm.cu）：每个线程算若干个输出元素
for (int k = 0; k < BK; ++k)
    sum[i][j] += As[i][k] * Bs[k][j];   // 标量 FMA，走 FP32 pipe

// Tensor Core 版（kernels/sgemm/sgemm_tf32.cu）
MMA_M16N8K8(...);   // mma.sync，走 Tensor Core

A100 峰值参考：

差 816 倍的量级会完全淹没 ldmatrix / cp.async 带来的 530% 提升 —— 这正是 §1.3「坑 2」的成因。

用途 2：架构级对照，回答"这套优化路线值不值得"

真正要回答的问题分两层：

1.5.4 仓库里的现成素材正好构成干净对照

kernels/sgemm/ 里是一对标准对照：

kernels/sgemm/sgemm.cu        ← SIMT 版（标量 FMA），已做到 tiling + swizzle + double buffer + coalesced r/w
kernels/sgemm/sgemm_tf32.cu   ← Tensor Core 版（mma.sync + ldmatrix + cp.async）

而且 SIMT 版本身也有一套完整的优化阶梯（见 sgemm/readme.md）：

sgemm_tiling                   ← 基础 tiling
sgemm_at_tiling                ← + A 转置入 smem
sgemm_at_bcf_swizzling         ← + swizzle 消 bank conflict
sgemm_at_bcf_swizzling_rw      ← + coalesced 写回
sgemm_at_bcf_swizzling_dbf_rw  ← + double buffer（反超 cuBLAS）

这是教科书级的消融阶梯写法 —— 每步只加一个东西，readme 里就有逐档 TFLOPS。 本实验的 cp.async / ldmatrix 阶梯应当照抄这个范式。

📌 公平性提醒：SIMT 版也可以做 shared memory tiling 和向量化解码（仓库里已经做了）， 所以 SIMT vs Tensor Core 的对比必须是"都优化到位"的两套实现之比，而不是"裸 SIMT" vs "花活 Tensor Core"。

1.5.5 小结

2. 代码组织建议

现在 kernels/hgemm/hgemm.cu 里有 4 个 kernel 几乎完全重复（naive / bcf / bcf_dbf / bcf_dbf_rw），继续加变体会爆炸。

建议新建 kernels/ablation/（基于 kernels/hgemm/hgemm.cu 改造，SM80 安全），用一个模板参数化的 kernel 覆盖全组合：

template <int LOAD_MODE,   // 0=ldg+st.shared, 1=cp.async
          int READ_MODE,   // 0=ld.shared, 1=ldmatrix
          int B_LAYOUT,    // 0=Bs[BK][BN], 1=Bs[BN][BK]
          int STAGES,      // 1,2,3,4
          int SWIZZLE>     // 0/1
__global__ void hgemm_ablation_kernel(...);

用 if constexpr 分支（编译器消掉未走分支，零运行时开销），变体命名 abl_L{lm}_R{rm}_bl{bl}_s{s}_sw{sw}。

实现要点：

• LOAD_MODE == 0 时用 float4/uint4 做 ld.global + st.shared，并保留 __syncthreads() 语义
• LOAD_MODE == 1 时用 CP_ASYNC_CG + CP_ASYNC_COMMIT_GROUP() + cp_async_wait_group<STAGES-2>()
• READ_MODE == 0 时按 mma fragment 的 lane→(row,col) 映射手工写 ld.shared（注意 A fragment 需 4 次 b32，B fragment 视 B_LAYOUT 决定是 4 次 b16 还是 2 次 b32）
• STAGES >= 2 时 smem 用 As[STAGES][BM][BK] / Bs[STAGES][BK][BN] 的 stage-major 布局，配合 read_idx / write_idx 索引

3. 预期结论骨架（待实测填数）

最终输出 3 张独立的表，不要把 SIMT 与 ldmatrix / cp.async 混排在同一张表里（理由见 §1.5）。

表 1：架构级对照（SIMT vs Tensor Core）

这一表回答"上 Tensor Core 值不值"。注意 SIMT 与 TC 两边必须都是优化到位的版本。

表 2：阶梯 A —— ldmatrix vs ld.shared（固定 cp.async + STAGES=3 + swizzle）

表 3：阶梯 B —— cp.async vs 同步 load（固定 ldmatrix）

关键读法："双缓冲"与"cp.async"必须拆成两步分别加，增量才可归因。 单 buffer 的 cp.async（B3 那一档）预期与同步 load 持平甚至略慢 —— 这是正常的，不是实验失败。

表 4：形状敏感性（每个变体跑一遍 shape 扫描）

K 小 → latency-bound，cp.async / 多级流水收益最大；K 大 → compute-bound，收益被 mma 吞吐盖住。

4. 待办清单

• 新建 kernels/ablation/，实现模板化 kernel（LOAD_MODE / READ_MODE / B_LAYOUT / STAGES / SWIZZLE）
• 实现 smem→reg 微基准（隔离出 ldmatrix 收益，避免被 mma 吞吐掩盖）
• 写 test.py：扫多个 shape（K ∈ {256,512,1024,4096,8192}）、锁频、输出 Markdown 表格（每档 TFLOPS + 相对上一档的增量百分比）
• 写 ncu 指标采集脚本（stall reason / LSU 指令数 / bank conflict / 带宽利用率）
• 补 readme.md 记录结论表格
• 复现一份 SM80 上的 cuBLAS baseline（kernels/hgemm/cublas.cu 已存在，直接用）
• 复现 sgemm.cu 的 SIMT 优化阶梯作为架构级对照（表 1），并与 sgemm_tf32.cu 配对比较

5. 附录：参考指标（ncu）

已有 profiling 脚手架：kernels/hgemm/run_ncu.sh、kernels/hgemm/test_for_ncu.py （用法：ncu -k regex:"gemm|Kernel" --set full -o profile_res -f python test_for_ncu.py）。

6. 与完整版的差异

7. 原版 TMA 章节（存档，SM90+ 才需要）

8. 更新记录