MTT S4000 FP16 GEMM 优化实战:逼近 mcc 编译器的物理上限

我参加摩尔线程 GEMM 矩阵挑战赛——这篇文章是完整的 三段式工程复盘。

第一段:开发阶段(§2-§7)。 从官方两份示例(example_naive.mu 标量版 ~0.5 T、example_tensorcore_basic.mu WMMA 起点)出发,我自己主导方向 + AI 辅助快速生成代码,70+ 个版本(historical_versions/team3_gemm_v1 → v78)反复迭代:64×64×32 → 128×128×32 双缓冲 → 256×256×32 mega tile → 32×32×16 WMMA。整条路径覆盖 tile 拉大、协作加载、双缓冲流水线、MP22 原生 32×32×16 fragment 等所有 GEMM 优化套路。4 个 legacy 正确性用例每一版都过、8k 性能爬到 ~90 TFLOPS,在原评分脚本上拿到满分 100/100,1 月 22 日提交。本仓库 historical_versions/ 下 ship 了其中 6 个关键版本(v1 / v20 / v59 / v68 / v70 / v76),读者可以用 grader/test_cases_legacy.json 复现它们当年的状态——其中 v68 在 fresh build 上实测 93.12 TFLOPS,4 legacy 用例 max_err = 0,总分 100/100。

第二段:塌陷(§8)。 评审期脚本里新增了一条 8192×8192×16384 的正确性测试 correctness_8k,我那个 90T 提交版首次跑就 fail——75% 输出元素错,max_err ≈ 36。深挖发现:原 dispatch 把 M, N ≤ 1024 的 4 legacy 用例全部路由到一条独立的 16×16×16 单 warp 安全内核,而追性能的 256×256 mega tile 路径从未被正确性测试覆盖过。换句话说,90 TFLOPS 是建立在一个从未被校验的代码路径上的虚高数字——同一个 v68 用 4-用例配置看起来 100/100 满分,用 5-用例配置直接挂 correctness_8k。这一对矛盾结果你可以在仓库里 fresh clone 后用两条命令亲眼复现(README 末尾给出)。

第三段:重建(§9-§25)。 4 月 3 日下午到 4 月 8 日,6 天时间从头建。物理 warp = 128 修复(0 → 22 T)、stride-33 bank conflict 消除(22 → 48 T)、M-first 蛇形 CTA 排序(48 → 54 T)、256×256×32 + 1024 线程(54 → 57 T)、postRA + ILP 调度 flag(57 → 67 T)、FenceSetting=mixed(67 → 73 T)、4 个 load-store flag(73 → 73.8 T)。最终 participants/team03_gemm_work.mu 通过全部 5 个正确性用例(包括 correctness_8k max_err = 0)、绝对性能 73.8 TFLOPS / muBLAS 94% / 官方 100/100。

结论:73.8 TFLOPS 不是"代码还能更聪明",是触到了当前 mcc 3.1.0 + S4000 + load_matrix_sync 组合的编译器物理上限——VLIW 165/165 满打包(0 空 slot)、C++ 双缓冲 5 种方案全部低于单缓冲基线、手动 fragment 加载被 MMA 专用寄存器 bank 约束封死。过程中定位到两个可以直接立项的摩尔线程编译器/头文件 bug:machine-scheduler pass + internal-regs-ra 的 WMMA 操作数 bypass 冲突(73 → 79 T 的差距)、crt/mma.h:761 的 fragment 元素数定义可能多分配 3 倍寄存器。所有数字、代码、bug、复现命令都可以在本仓库里逐项交叉验证。

1. 起点：摩尔线程的题目和我手里的两份示例

先把题目摆清楚。比赛要求实现一个 FP16 GEMM，函数签名固定：

extern "C" void gemm_optimized(
    half const* d_A,   // M × K, column-major
    half const* d_B,   // K × N, column-major
    half* d_C,         // M × N, column-major
    int M, int N, int K);

A、B、C 全部列主序，输入输出 half，中间必须 FP32 累加，函数返回前必须调用 musaDeviceSynchronize()，不许调 muBLAS 或任何厂商 BLAS。

评分 30 + 70:

正确性 30 分考 4 个尺寸：small 256×256×512、medium 512×512×1024、large 1024×1024×2048、non_square 512×1024×2048，容差 0.01 × √(K/512)。
性能 70 分在 8192×8192×16384 上对 muBLAS 取相对比值，muBLAS 基线 86,012.6 GFLOPS，达到基线 30% 拿满 70 分。性能分的门槛很低，30% 就满——所以 100/100 不难，难的是绝对性能排前面。

官方在 participants/examples/ 下给了两份起点代码：

一份是 example_naive.mu——纯标量 1D 布局，每个线程算一个 C 元素，用 fmaf 做 FP32 融合乘加保证精度，连 shared memory 都没用。跑 8k 性能大概 0.5 TFLOPS 量级。它的作用是"保证能过正确性的兜底路径"，不是高性能参考。
另一份是 example_tensorcore_basic.mu——90 行、16×16×16 WMMA、每 block 32 个线程，主循环就是标准的 load_matrix_sync / mma_sync / store_matrix_sync 三件套。示例本身明说了"仅展示如何使用 MUSA WMMA API 进行最基本的矩阵乘法"，不是高性能参考。跑 8k 大概几个 TFLOPS。

我第一眼看就清楚要拿绝对性能靠前必须做这几件事：把 tile 拉大（example_tensorcore_basic 的 16×16 实在太小）提升算术强度 AI = BM·BN / (BM+BN)；加多 warp 让计算并行隐藏延迟；上双缓冲让 global→smem 搬运和 MMA 计算并行；最后再做 bank conflict、L2 利用、寄存器预算等微调。这些都是从 CUDA 上搬过来的通用套路，MP22 上的 MUSA 应该也适用。

问题是我了解CUDA，但 MUSA 和 MP22 对我是全新平台。MUSA WMMA API 我只翻过一遍，MP22 的寄存器文件大小、warp 粒度、shared memory bank 数、编译器 flag，这些细节我那时候没多少把握。从零写一个能到 80T 的 MUSA 内核我做不到。

我的策略是：快速迭代，每一版都编译跑 grader 看数字决定下一步。我负责判断下一步往哪改——拉 tile、上双缓冲、换 warp 布局——自己写代码、逐行核对、本地编译跑 grader，每一次改动都基于上一版的数字决定方向。

这套策略在 NVIDIA 上没问题。在 MP22 上它有一个盲区我没察觉到，后面会出事。

2. 第 1 版：64×64×32 + 4 warp（~5 TFLOPS）

当时的情况：刚跑通 example_tensorcore_basic，性能几个 TFLOPS。16×16×16 的 tile 太小，AI 只有 8，必须先把 tile 做大。

思路：最直接的改进是把 tile 从 16×16 拉到 64×64。4 倍的 tile 尺寸意味着 4 倍的 AI（理论上），搬运开销被摊薄到更多 MMA 上。每个 block 放 4 个 warp = 128 线程，按 2×2 布局，每个 warp 负责一个 32×32 子块，用 4 个 16×16×16 fragment 累加。这是 NVIDIA 上写 GEMM 的基本款。

关键技术点：

tile 尺寸 BM=64, BN=64, BK=32
block 尺寸 128 线程（4 个 32 线程 warp，按 NVIDIA 直觉分）
warp 布局 2×2 WARPS_M=2, WARPS_N=2
shared memory sA[32][64+PAD], sB[64][32+PAD]，PAD=8 按 NVIDIA 经验避 bank conflict
16×16×16 WMMA fragment
dispatch 分两路：gemm_baseline_safe（标量兜底）和 gemm_wmma_safe_kernel（WMMA 性能）

参考代码：

#define BM 64
#define BN 64
#define BK 32
#define WARPS_M 2
#define WARPS_N 2
#define PAD 8

__global__ void __launch_bounds__(128) gemm_wmma_safe_kernel(...) {

    const half* __restrict__ A,
    const half* __restrict__ B,
    half* __restrict__ C,
    int M, int N, int K)
{
    // ----------------------------------------------------------------
    // 1. 坐标计算
    // ----------------------------------------------------------------
    int bx = blockIdx.x;
    int by = blockIdx.y;
    int tid = threadIdx.x;
    int wid = tid / 32;// 按 32 线程分 warp   ← 致命假设
    int lane = tid % 32;

    // Warp 在 Block 中的位置
    int w_m = wid / WARPS_N;
    int w_n = wid % WARPS_N;
        // ... 4 个 warp 独立调用 WMMA ...
}

性能：~5 TFLOPS。相比起点的 ~0.5T 提了一个量级，离目标（muBLAS 86T）还有 17 倍的路。

踩过的坑：

第一次编译 link 报错。代码编译过了但 link 失败，报 undefined reference to musaDeviceSynchronize。漏了 -lmusart。加上之后才通过。官方 example 的 Makefile 是带这个 flag 的，但我起草新代码时没带。小事，但花了 10 分钟才反应过来。
dispatch 分两路的坑，后来要了命。最初的版本没分 dispatch，所有尺寸都走 WMMA 大 tile。结果 small 256×256 跑过，medium 512×512 过，但 non_square 512×1024 报错——边界处理没做对，最后一个 block 越界。我加一条 fallback：M 或 N 不能被 BM/BN 整除就走 scalar。这一条 fallback 后来被我进一步优化成"M, N ≤ 1024 就走 gemm_wmma_correct_16"——一个独立的 16×16×16 单 warp 安全内核。从此大尺寸和小尺寸走两条完全不同的代码路径。这个 dispatch 分流从第 1 版就长在代码里，70+ 版迭代一直没动它。它把所有 legacy 正确性测试都挡在了大 tile 路径之外——我当时完全没意识到这是一个盲区。
PAD=8 是拍脑袋的。我按 NVIDIA 经验加了 PAD=8 避 bank conflict，但那时候并不知道 S4000 是 16 banks（不是 32），也不知道 PAD_A=8 对应的 stride 对 smem 带宽会造成 1.67× 惩罚。PAD=8 看起来能工作就没细究，留了一个没排查的优化空间，直到后面做微基准才弄清楚。
__launch_bounds__(128) 的坑。我第一版加了 __launch_bounds__(128, 2) 希望 2 block/SM 提 occupancy。跑出来编译器报寄存器紧张，性能反而差。我改回 __launch_bounds__(128) 只指定线程数、让编译器自己决定 block/SM，性能才正常。回头看这是 MP22 和 NVIDIA 的一个差别——MP22 的寄存器预算分配模型和 NVIDIA 不完全一样，从 NVIDIA 经验外推 __launch_bounds__ 的第二参数不一定合适。这条经验建议摩尔线程未来在开发者文档里加一段"launch_bounds 最佳实践"，帮新用户避开这个外推陷阱。

怎么走过去的：第 1 版能跑、通过 small/medium 两个正确性用例,large/non_square 在 FP16 累加边缘(max_err=3.125e-02 对 tol=0.02),不同 build/GPU 状态下会偶发失败——当时我当成"浮点噪声"没深究。性能 ~5T,就算立住脚了。到 v2 把 tile 拉到 128×128、16 个"warp"分布更均匀之后,4 legacy 用例才全部稳定通过(下一节)。dispatch 路由"后面要验证"记了一笔但没真去验证(这一笔记错了地方)。继续往上推。

3. 第 2 版：128×128×32 + 512 线程（~25 TFLOPS）

当时的情况：第 1 版 64×64×32 跑到 ~5T，再往上推最直接的办法是继续拉 tile。

思路：AI = BM·BN/(BM+BN)，64×64 给 32，128×128 给 64——翻倍。tile 做大意味着每个 block 要负责更多输出，累加器数量增加，线程数也要跟着加。我决定 tile 拉到 128×128×32，block 线程数从 128 升到 512（16 个 32 线程 warp，按 4×4 布局）。每个"warp"负责一个 32×32 子块。

关键技术点：

tile 尺寸 128×128×32
block 512 线程 = 16 个"warp"，4×4 布局
每个"warp"负责 32×32 输出，用 4 个 16×16×16 fragment 累加
shared memory sA[32][128+8], sB[128][32+8] ≈ 9 KB，远远够用
协作加载：512 线程用 float4 向量化加载，一次 prologue 填满 sA/sB

性能：~25 TFLOPS，相比第 1 版加速 5 倍。

踩过的坑：

__launch_bounds__(512, N) 反复调整。第一次 __launch_bounds__(512, 1) 编过但性能只有 ~15T。我试着改成 __launch_bounds__(512, 2) 要两 block/SM 提 occupancy——结果寄存器溢出，private_memory 报 800+ 字节，性能暴跌到 8T。我反过来去掉第二个参数 __launch_bounds__(512) 让编译器自己决定——25T，回到正常。教训是 MP22 的 block/SM 约束和 NVIDIA 不一样，强制值经常把编译器逼进寄存器溢出的角落。
向量化加载越界。我写的 float4 协作加载逻辑里，(tid * 8) % BM 用来算每个线程在 sA 里的位置。跑小尺寸用例正常，跑 8k 性能测试报 illegal memory access。查了半天发现是边界处理：当 M % BM != 0 时最后一列 tile 的 A 指针会越过矩阵末尾。修复是在 gemm_optimized 里加 M % 128 == 0 的判断，不满足就走 fallback，同时协作加载逻辑本身不改。
PAD=8 和 PAD_A=8, PAD_B=8 的混乱。第 1 版用一个通用的 PAD=8，第 2 版我分成 PAD_A 和 PAD_B 分别控制。我没细想就接受了，但后来发现对 128×128 tile 两个值其实都应该是 8——分开写没带来任何好处，只是代码更啰嗦。这是一个没意义的改动。
epilogue 写回写错了第一次。第一版 epilogue 用 store_matrix_sync 把每个 warp 的 fragment 写回 smem 再转 half 写 global。第一次跑 medium 512×512 结果完全错。查代码发现 smem 中转区 float smem_out[128*128] 吃了 64 KB——一个 block 就占满整个 SM 的 smem 预算。我改成每个 warp 用自己的 32×32 中转区、串行 4 次 store 写回，性能和正确性都恢复正常。

怎么走过去的：25T 是一个稳定的新起点。这一版的代码结构（512 线程 + 16 "warp" + 4×4 布局）被后续所有版本继承，直到 4 月 3 日我发现根本错在哪里才被推翻。

4. 第 3 版：手写双缓冲流水线（~50 TFLOPS）

当时的情况：第 2 版 25T。反编译看 binary，主循环里计算（MMA）和搬运（DMA from global）是串行的——一个 K tile 的 DMA 完成后才开始 MMA，MMA 完才开始下一个 tile 的 DMA。这是明显的浪费：DMA engine 和 Tensor Core 应该并行。

思路：NVIDIA 上标准的双缓冲流水线——两份 smem 缓冲区交替，当前 tile 在 bufferA 上做 MMA，下一 tile 的数据预取到 bufferB。5 步循环体：

Loop body:
Issue Loads for Tile K+1 (Global -> Register)  [DMA engine]
Execute MMA for Tile K   (Shared -> Register)  [Tensor Core]
  → 1 和 2 并行
Barrier (wait MMA to finish reading Tile K)
Commit Tile K+1 (Register -> Shared)
Barrier (wait Tile K+1 ready)

这是 CUTLASS 里 GEMM kernel 的经典结构，在 NVIDIA 上对 memory-bound 和 compute-bound 都有效。我按这个流水线重写了 v2 的主循环。

关键技术点：

两份 smem 缓冲区 sA[2][...], sB[2][...]
4 个 float4 寄存器做 prefetch rA[2], rB[2]
手写 5 步流水线
两个 __syncthreads() 分别守 MMA 读完和 commit 写完
保持 512 线程 / 16 "warp" / 4×4 布局（第 2 版的结构）

关键主循环代码：

for (int k_step = 0; k_step < K - BK; k_step += BK) {
    // 1. Prefetch next tile to registers
    for (int i = 0; i < 2; ++i) rA[i] = *(const float4*)(pA + i * 16 * M);
    for (int i = 0; i < 2; ++i) rB[i] = *(const float4*)(pB + i * 128 * K);

    // 2. Compute current tile (from smem)
    for (int ki = 0; ki < 32; ki += 16) {
        // ... load_matrix_sync + mma_sync ...
    }

    pA += stride_A; pB += stride_B;

    // 3. Sync
    __syncthreads();

    // 4. Commit prefetch to smem
    for (int i = 0; i < 2; ++i) *(float4*)&sA[...] = rA[i];
    for (int i = 0; i < 2; ++i) *(float4*)&sB[...] = rB[i];

    // 5. Sync
    __syncthreads();
}

性能：~50 TFLOPS，相比第 2 版翻了一倍。

踩过的坑：

smem 双缓冲 vs 单缓冲的抉择。我一开始想做真正的双缓冲（两份完整 sA/sB 交替），但 128×128×32 tile 下双缓冲 smem 预算是 2 × (sA 9KB + sB 9KB) = 36 KB，占掉一半 smem 预算，而且每次访问都要多算一次缓冲区偏移。测试下来双缓冲版本 42T，比单缓冲 + 寄存器预取的 50T 还低 —— 动态索引的寄存器分配压力超过了双缓冲能省的 barrier。这是一个很反直觉的结果：NVIDIA 上双缓冲几乎总是赢，MP22 上不一定。我花了一天才接受"这个平台上单 smem 缓冲 + 寄存器预取就够了"这个结论。
两个 __syncthreads() 能不能省一个。我试过把 __syncthreads() 数量从 2 减到 1——放在 commit 之前等 MMA 读完，commit 之后不加 barrier 直接进下一轮。结果偶发错误：下一轮的 MMA 读到没写完的 smem。两个 barrier 必须都留。
prefetch 到寄存器 vs 直接 DMA 到 smem。最初的版本是直接从 global DMA 到 smem（不走寄存器中转）。测试性能 42T。我改成"先 DMA 到 float4 寄存器，然后寄存器写到 smem"——这样 DMA 指令和 MMA 指令可以在编译器层面真正交错。改完性能到 50T。关键是让编译器看到"搬运"和"计算"是两条独立的数据依赖链。
浮点寄存器越界。一次改动我把 prefetch 寄存器数从 float4 rA[2], rB[2] 扩成 rA[4], rB[4] 想让 prefetch 覆盖更大的 K 段。编译器立刻报 private_memory = 432 bytes 寄存器溢出，性能暴跌到 12T。我查 compiler_report.sh 看 temp register 数从 98 跳到 140，超出单 block 预算。改回 2 个立刻恢复。寄存器溢出的惩罚非常大，MP22 上任何一次增加寄存器变量的改动都要立刻跑 compiler_report 校验。
#pragma unroll 的 I-cache 坑。双缓冲主循环里 #pragma unroll 默认全展开，但 BK=32 下每次 MMA 展开后主循环体超过 4KB 的 L0 i-cache，性能比不展开还低。改成 #pragma unroll 1（显式禁用展开）或者 #pragma unroll 2 部分展开，性能才正常。

怎么走过去的：50T 之后我对"双缓冲流水线"这个方向已经没有新思路了，开始把精力放到"继续拉 tile 尺寸"上。

5. 第 4 版：256×256×32 Mega Tile（~70 TFLOPS）

当时的情况：第 3 版 50T，双缓冲已经吃完。下一步只有继续做大 tile。

思路：128×128 的 AI = 64，256×256 的 AI = 128——再翻倍。AI = 128 意味着每加载 1 字节数据可以驱动 128 次 FLOP，roofline 模型下彻底进入 compute-bound。要做到 256×256，smem 预算必须严格规划：

smem 布局:
  sA[BK=32][BM=256 + PAD]:  32 × 264 × 2 B ≈ 17 KB
  sB[BN=256][BK=32 + PAD]: 256 × 40  × 2 B ≈ 20 KB
  单缓冲合计: ~37 KB
  如果双缓冲: ~74 KB > 72 KB 上限 → 装不下

所以必须在"双缓冲 + 更小 tile"和"单缓冲 + 256×256 tile"之间二选一。第 3 版的经验告诉我单缓冲其实不差，最终选单缓冲 + 寄存器 prefetch 的方案。这就是 v59 "Mega Tile" 的起点。

关键技术点：

tile 256×256×32, BK=32
单缓冲 + 寄存器 prefetch
512 线程（和第 3 版一样，16 个"warp" 4×4 布局）
每个"warp"负责 64×64 输出 = 4 个 32×32 子块
shared memory sA[32][256+8] + sB[256][32+8] ≈ 37 KB
__launch_bounds__(512, 1) 1 block/SM
协作加载：512 线程 × float4 = 4 KB/iter，和 tile 数据量匹配

性能：~70 TFLOPS。v59 的代码注释里自信地写了 "Theoretical Peak: ~98 TFLOPS"——实际只摸到 70T，离 muBLAS 86T 还差 16T。

踩过的坑：

v58 的 smem 溢出灾难。v58 开头就写了 "CORRECTION FROM v58: v58 failed due to SMEM overflow (92KB > 72KB)"。v58 是我第一次尝试：想做 256×256×32 双缓冲，算下来 2 × (16 + 16) = 64 KB 加 padding 约 92 KB，超 72 KB 硬上限。编译器直接报错：shared memory size exceeds hardware limit。v59 才改成单缓冲。这次失败浪费了小半天。
累加器寄存器预算。256×256 每 block 的累加器总量是 256 × 256 = 65,536 个 float。512 线程平均每线程 128 个 float = 512 字节。看起来刚好在 S4000 单线程 256 寄存器（= 1024 字节）的预算内，但这是"平均"——实际编译器分配寄存器时会有大量中间变量和临时寄存器，真实需求远超平均。第一次编译 v59 编过了但 private_memory = 300 字节——有溢出。性能也只有 55T，比 v20 的 50T 好得不多。我反复调整，发现必须放弃几条中间变量、用宏把主循环展开写紧才能把 spill 消掉。最后 private_memory = 0，性能到 70T。
epilogue 的 smem 中转区从哪来。256×256 的 epilogue 需要把 65K 个 float 写回 global，如果每个 warp 用自己的 32×32 smem 中转区需要 16 × 4 KB = 64 KB 的 smem——和主循环的 sA/sB 合起来又爆了。我想到的方案是复用 sA/sB 作为 epilogue 的中转区——主循环结束后 sA/sB 数据已经没用了，拿来做 float 中转正好。这个 trick 是对的，但第一次写的时候漏了 __syncthreads()——主循环的最后一次 mma_sync 还在跑时 epilogue 就开始往 sA 写，c00 还没算完就被覆盖了。加一个 barrier 修好。
WARPS_M × WARPS_N = 4 × 4 的布局其实不对称。256×256 被 4×4 分成 16 个 64×64 子块，理论上每个 warp 覆盖 1 个子块。但 warp 之间 A 和 B 的加载比例是不对称的：warp(0,0) 和 warp(0,1) 共享同一行 A，warp(0,0) 和 warp(1,0) 共享同一列 B。4×4 布局下每行 4 个 warp 共享 A 但用 4 份不同 B；反过来也一样。MMA/load ratio 大概 1.0。第 3 版 128×128 的 4×4 布局 MMA/load 差不多也是 1.0。这个比例没有本质改善。
L2 缓存还没调。70T 之后我用 mthreads-gmi 看 L2 命中率——30% 左右。256×256 的 tile 虽然 AI 高，但 CTA 遍历顺序默认是 N-first，相邻 block 共享 A 列但 B 列不共享，8k 规模下 B 的总量远超 L2。这个问题我留到了后面，没在 v59 阶段解决。

怎么走过去的：70T 离 90T 还差不少，但 v59 的单缓冲 + 256×256 骨架是对的。我继续迭代，v60-v67 尝试了各种 128×256 / 256×128 / 非方 tile 的组合，性能都在 60-70T 之间波动。这些尝试基本没有实质进展，几天时间过去了，我卡在 70T 上下不去。

6. 第 5 版：MP22 原生 32×32×16 WMMA，提交 ~90 TFLOPS

当时的情况：在 70T 卡了很久，开始怀疑是不是 16×16×16 WMMA 的单条指令吞吐不够。

思路：我翻 crt/mma.h 时发现 MP22 原生支持 fragment<matrix_a, 32, 32, 16, half, col_major>——32×32×16 的 WMMA fragment，比 NVIDIA 的最大 16×16×16 大 4 倍（2×2 翻倍）。如果我把单条 WMMA 从 16×16×16 升级到 32×32×16，单条 mma_sync 的计算量就是原来的 4 倍，但 fragment 加载开销不是简单的 4 倍——因为 32×32 的 fragment 可能用同样数量的 load_matrix_sync 调用就能装下。MMA/load ratio 可以被显著提升。

32×32×16 也是 muBLAS 内部用的 MMA 尺寸，用它才能摸到 muBLAS 的性能水平。

关键技术点：

主循环用 fragment<..., 32, 32, 16, ...> 替代 16, 16, 16
每个"warp"（32 线程）负责 64×64 输出，由 2×2 = 4 个 32×32 累加器 c00/c01/c10/c11 组成
每次 kk 步（K=16）做 load + mma 序列：2 个 A fragment a0, a1 + 2 个 B fragment b0, b1 + 4 次 mma
其他结构沿用 v59：单缓冲 256×256 tile、512 线程、4×4 warp 布局

核心主循环代码：

for (int kk = 0; kk < 32; kk += 16) {
    fragment<matrix_a, 32, 32, 16, half, col_major> a0, a1;
    fragment<matrix_b, 32, 32, 16, half, col_major> b0, b1;

    int m_base = w_m * 64;    // w_m = wid % 4, wid = tid / 32
    int n_base = w_n * 64;

    load_matrix_sync(a0, &sA[kk][m_base + 0],  BM + PAD);
    load_matrix_sync(a1, &sA[kk][m_base + 32], BM + PAD);
    load_matrix_sync(b0, &sB[n_base + 0 ][kk], BK + PAD);
    load_matrix_sync(b1, &sB[n_base + 32][kk], BK + PAD);

    mma_sync(c00, a0, b0, c00); mma_sync(c01, a0, b1, c01);
    mma_sync(c10, a1, b0, c10); mma_sync(c11, a1, b1, c11);
}

4 次 load_matrix_sync 驱动 4 次 mma_sync——相比 16×16×16 的每次 load+mma = 1:1，32×32×16 是 load:mma = 1:1 但每条 MMA 的有效计算量是 4 倍。等效的 MMA/load ratio 翻了 4 倍。

性能:~88-90 TFLOPS 是我开发期间反复测到的典型值。本仓库 ship 时用 fresh build + legacy 4-用例配置重新跑了一次,得到 93.12 TFLOPS(完整数据见 reports/historical_versions_legacy_run.log)——稍高于当时的口头印象,但这是正常的 DVFS 波动范围,两个数字描述的是同一个版本。一步从 70T 跳到 88-93T,增幅 25-30%,是整个 v1-v78 迭代路径上单次改动最大的一次提升。

关键的诡异点(本节重点):这个 93 TFLOPS 是用 4-legacy 配置跑出来的。如果用包含 correctness_8k 的 5-用例配置(也就是 ship 在 grader/test_cases.json 里那份)跑同一个 v68,8k 正确性 FAIL,75% 元素错——评测会停在正确性阶段,根本进不了性能测试。也就是说:v68 的 93T 是一个真实的性能数字,但它建立在一个从未被正确性覆盖过的代码路径上。这就是文章 §8 的故事的物理证据——你可以 fresh clone 后亲自跑两次对比看到。

踩过的坑：

头文件的开关。crt/mma.h 里 32×32×16 的 fragment 类是被 #ifdef __MUSA_INCLUDE_COMPILER_INTERNAL_HEADERS__ 包起来的——必须在 #include <crt/mma.h> 前 #define 这个宏，否则看不到 32×32×16 的 overload。第一次编译时没加宏，fragment<matrix_a, 32, 32, 16, ...> 类直接没找到，编译报 incomplete type。这个坑 crt/mma.h 文档里没写，我是在 example_tensorcore_basic.mu:9-11 里看到示例代码用了这个宏才反应过来：

#define __MUSA_INCLUDE_COMPILER_INTERNAL_HEADERS__
#include <crt/mma.h>
#undef __MUSA_INCLUDE_COMPILER_INTERNAL_HEADERS__

--offload-arch=mp_22 是必需的。只加头文件宏还不够——编译时必须传 --offload-arch=mp_22 才能启用 FP16 WMMA 重载。不加这个 flag 编译器只暴露 fp8 和 int8 的 WMMA 接口。我第一次编译时把 --offload-arch=mp_22 漏了，mcc 给了一堆 no matching function for call to mma_sync，根本找不到 FP16 版本。加上 flag 后立刻过。官方 Makefile 是带这个的。
epilogue 大改重写。32×32×16 的 fragment 布局和 16×16×16 不一样——一个 32×32 累加器里 32 个 lane 各自持有的元素分布是 MP22 特有的。v68 的 epilogue 改成 store_matrix_sync(..., mem_col_major)，把 32×32 按列写入 smem，然后每个 lane 读一整列 float 转 half 再写 global。这个 epilogue 有一个非常隐蔽的问题：stride=32 的 smem 读回造成 bank conflict。我那时候没注意到，因为最终跑出来性能达标（~90T），正确性过（4 个 legacy 用例），就没深挖。这个 bank conflict 在我后面重建过程中被 compute-only 微基准暴露出来，是 stride-33 修复的源头。
双缓冲再次尝试再次失败。v68 跑通 90T 后我又想加双缓冲冲一下 95T。和第 3 版一样，动态 smem 索引加上去性能反而下降。这次我更确定 "在 MP22 的 C++ WMMA API 上双缓冲拿不到收益"——但还没拿到具体证据，要等后面专门做 5 种方案的对比实验。
v68 的很多变种都没收益。v68_swizzle4、v68_swizzle16、v68_twostage_pad4、v68_epilogue_pingpong、v68_mublas_inspired 这些都是我围绕 v68 基础结构做的微调尝试。每一个都花了 2-4 小时，结果都是 ±1-2 TFLOPS 的波动，没有本质改进。

怎么走过去的：v68 的 ~90T 已经超过 muBLAS 基线（86T）约 5%，达到竞赛满分所需的 30% 基线的 3 倍多。我把 v68 做最后清理——删掉历史注释、保留最干净的主循环、加上完整的 dispatch 分流——作为最终提交版 src/team03_gemm.mu（600 行，1 月 22 日提交）。

7. v69-v78 的微调实验（性能在 88-90T 之间波动）

当时的情况：v68 已经提交了 90T 的版本，我想继续优化冲到 95T 甚至 100T。接下来的一周我在 v68 基础上做了十几个微调版本 v69 到 v78。这一段没有实质突破，但踩了很多坑，值得记下来。

尝试过的方向（全部失败或边际收益）：

v70 instruction interleaving。historical_versions/team3_gemm_v70.mu:10-20 注释写 "Instead of separating Load and Compute phases (v59), we mix them. Inside the massive 32-step compute loop: Step 0..3 MMA, Step 4: Issue Global Load A[0] → Register...". 思路是手动把 MMA 和 Load 交错,让 Tensor Core 和 LSU 并行。结果灾难性失败——本仓库 fresh build 实测这个 v70 跑 8k GEMM 需要 27 秒,即 0.08 TFLOPS(完整数据见 reports/historical_versions_legacy_run.log)。手动交错严重破坏了编译器的 VLIW 打包决策,导致流水线停顿到几乎不可用。我当时口头印象是"~85T,比 v68 略慢",事实是这个 build 有 1000× 倒退——本应该立刻丢弃,却因为我"看起来还在 80T 区间"的错误印象保留了下来。这是项目里我对自己代码性能判断最离谱的一次失误,放在这里作为反面教材。
v72 all-loads-first 调度。反过来：所有 load_matrix_sync 先做完再做 MMA。思路是让 fragment 加载先占满 DMA 发射槽。实测 ~86T，略低。
v73 交错 load + MMA。手动一条一条交错 load → mma → load → mma，希望把每条 MMA 的延迟藏在 load 后面。实测 39T，大幅倒退。代码膨胀把 i-cache 撑爆了。
v74 BK=64 扩大 K 方向。思路：BK=64 每次主循环迭代处理 64 个 K 元素，屏障数量减半。实测性能暴跌到 22T。检查 compiler_report.sh：private_memory = 420 bytes——寄存器溢出灾难。BK=64 需要 4 个 prefetch float4 而不是 2 个，寄存器预算直接爆。
v76 降 tile 到 256×128 + 2 block/SM。historical_versions/team3_gemm_v76.mu:5-16 注释:"v68/v75 (256x256 Tile) Required 1KB registers/thread for Accumulators. Result: MASSIVE REGISTER SPILL. Capped performance at ~50T. Reduce Tile to 256x128x32. Accumulators per thread: 128 floats = 512 Bytes. Fits in registers." 思路是降 tile 减少累加器、换取 2 block/SM 提 occupancy。本仓库 fresh build 实测 40.6 TFLOPS(我开发期间口头记忆是 65-70T,这次重新跑发现实际是 40T——可能是不同 DVFS 状态下的差异,以本次实测为准)——比 v68 的 93T 差一半多。高 occupancy 比不过大 tile 的高 AI,这个结论在后面重建阶段又被验证了一次。
v77-v79 warp 数量扫描。分别测 1/2/4 warp per SM 的配置。1 warp 延迟掩盖不够（~20T），2 warp 稍好（~50T），4 warp 是 v68 的默认（90T）。再往上加 warp 数就超线程预算。
v80 persistent kernel。让 block 在 kernel 内部循环处理多个 tile，避免 block 启动开销。实测 tile 循环的额外开销抵消了调度收益，性能还低于 v68。
v81 双缓冲 256×128 + PAD_B=0。想绕开 "256×256 双缓冲装不下 72KB" 的限制，降 tile 到 256×128 再双缓冲。PAD_B=0 时 32 stride 触发 bank conflict，性能只有 ~55T。
v82 async double buffer BK=16。异步 copy + BK=16 双缓冲。实测 __syncthreads() 次数翻倍（BK=16 vs BK=32 每个 tile 多一次 sync），同步开销吃掉双缓冲收益。~70T。
v83 super-tile swizzle。在 N 方向做 N_BLOCKS_PER_STRIP=8 的"超 tile"遍历顺序，希望比默认的 N-first 更好。实测 ~90T，和 N-first 持平。没改进。
v84 交错 loads on v76 base。在 v76 的 256×128 结构上试 v73 风格的交错，实测 ~60T。代码膨胀 + 低 tile AI 双重惩罚。
v85 zero-smem 直接从 global 做 MMA。试着完全绕过 smem，每次 MMA 前直接 load_matrix_sync 从 global 读——load_matrix_sync 的第二个参数可以是 global 指针。实测 ~15T，因为 DMA 延迟远大于 smem 延迟。这条路根本不通。
v86 muBLAS-style 256 线程。想模仿 muBLAS 的 512 线程 / 4 物理 warp 架构，用 256 线程 / 8 "warp"（当时还没意识到 warp 是 128）。实测 ~75T，线程太少延迟掩盖不够。

这十几个 v69-v86 版本占了我大量时间，没有一个比 v68 的 ~90T 有实质改进。最后我还是用 v68 的清理版作为最终提交。这段时间最大的教训是"在一个错误的骨架上做微调不可能变对"——但我那时候还不知道骨架本身是错的。

8. 加测 8k 正确性那天：崩塌

当时的情况：1 月 22 日提交 v1（v68 清理版），官方评分 100/100，8k 性能 ~89-90 TFLOPS。我以为结束了。

然后是评审期。摩尔线程在原评分脚本里加了一条用例 correctness_8k——尺寸 8192 × 8192 × 16384（就是性能测试用的那个尺寸），容差 0.057（0.01 × √(K/512) = 0.01 × √32 ≈ 0.057）。

我拉新脚本本地跑，结果返回 FAIL：

correctness_8k: FAIL
  mismatched elements: ~49,000,000 / 67,108,864 (~73%)
  max_err: 36.2  > tolerance 0.057

第一反应是"容差算错了"。拉公式核对一遍：tolerance = base × √(K / reference_K) = 0.01 × √(16384 / 512) = 0.01 × √32 ≈ 0.0566。新脚本里写的 0.057 是对的。

第二反应是"浮点累加误差"。K=16384 是 4 个 legacy 用例里最大 K（2048）的 8 倍，累加次数多当然误差大。但 max_err = 36 远超 FP32 累加误差能造成的量级。即使把 16384 次累加全算上，FP32 rounding 误差量级也在 0.1 以下。36 不是"精度不够"，是"算错了"。

第三反应是去看 dispatch 代码。打开 src/team03_gemm.mu:577-595：

if ((M <= 1024) && (N <= 1024) && (K % 16 == 0)) {
    dim3 block(32, 1, 1);                                  // 32 threads = 1 "warp"
    gemm_wmma_correct_16<<<grid, block, 0, stream>>>(...); // 16x16x16 单 warp 内核
} else if ((M % 256 == 0) && (N % 256 == 0) && ...) {
    dim3 block(512, 1, 1);
    gemm_mp22_mega_tile_256x256x32<<<grid, block, 0, stream>>>(...); // v68 mega tile
}

4 个 legacy 正确性用例的 M, N 都 ≤ 1024——全部命中 gemm_wmma_correct_16。那是第 1 版写 dispatch 时就加进去的"安全小内核"，32 线程一个 block、16×16×16 单 warp、不会越界、不会触发 mega tile 的任何 bug。

correctness_8k 的尺寸 8192 × 8192 × 16384 走的是 else if 那条分支——第一次真正执行 gemm_mp22_mega_tile_256x256x32。

这个责任在我，不在评测工具。原评分脚本的 4 个正确性用例覆盖 256-1024 的小、中、大、非方阵尺寸——对评测工具来说是合理的尺寸覆盖。问题出在我自己的 dispatch 代码设计：我把"安全兜底 kernel"和"追性能 kernel"放在了两条完全不相交的代码路径上，而我自己从来没有在每一条 dispatch 分支上都放正确性用例去校验。评审期加测 correctness_8k 立刻暴露了这个 dispatch 分流的盲区——这恰恰说明评审团队对"尺寸覆盖之外还要看代码路径覆盖"有更细致的要求，是评测设计的一次升级。

换句话说，我做了 70+ 版本迭代，每一版 official grader 都跑了正确性测试、每一次都显示 4/4 pass——但我没注意到那 4 个用例始终走 fallback kernel，我自己的高性能代码从来没有被我自己放在正确性校验链路里。我以为自己在持续验证一份算对的代码，实际上是在持续"误以为验证了"一份从未被校验过的代码路径。

AI 辅助的过程也没有帮我挑出这个盲区——但这也不应该怪 AI。AI 的评估逻辑是"基于上一版的性能数字做评估"，而提出"每条 dispatch 分支都要有对应的正确性用例"这种元层面的测试覆盖要求，本来就不是 AI 的职责，是我作为项目负责人的职责。我该自己去 review dispatch 的分支覆盖矩阵，而不是默认 official grader 的尺寸覆盖已经足够。

那时候的情绪比纯粹的技术挫折更复杂。70多个版本的迭代、100/100 的满分、~90T 的绝对性能——全部建立在一个"正确性和性能走两条路"的自我设计失误上。说实话当时最难受的不是"要重写代码"，是意识到过去每一次看到"pass"的时候，那个"pass"对应的代码路径根本不是我真正在优化的那条。

2026 年 4 月 3 日下午。我在自己的实验笔记开头写了一句话：

Apr 3: 90T 是一个从未被正确性测过的数字。6 天之内要重建一份真正算对的代码。前 70+ 个版本全部作废。

然后把 src/team03_gemm_work.mu 重置——那时候它还和 src/team03_gemm.mu 一样——从 examples/example_tensorcore_basic.mu 重新开始。交出一份 correctness_8k 上 max_err = 0 的提交，而且性能不能掉太多。规则变了：这一次正确性是硬门槛，性能是在硬门槛之内能抢多少是多少。

9. 重建之前先立规矩

动手之前我给自己写死了三条项目硬规则：

正确性优先于性能：8k 正确性不通过就立即拒绝，不管性能多好看
每一步优化都要能解释"为什么这样改"：只记"做了什么 → 拿到什么"的路径走不远，必须写清"观察 → 假设 → 验证"
每个数字都能在代码里验证：不确定的东西先标"待验证"

之所以把这三条写死，是因为 70+ 版快速迭代的最大教训就是——我失去了对每一步改动"为什么"的控制。当每一版都"比上一版快一点"的时候，很容易默认"这个方向是对的"然后把验证外包给"下一版跑得更快"这个信号。但那个信号本身建立在从未被校验的代码路径上。

这一次我给自己的约束是：任何一步优化，动手之前必须能用一两句话说清"我为什么相信这会更快"。说不清就先不要改。

还有一条默认的 bug 怀疑顺序——是我过去排 CUDA bug 的经验，从最常见到最少见：

epilogue 写回竞争或重叠
smem 复用生命周期 bug
CTA 映射 / swizzle 在满 grid 下的碰撞
向量化 half 打包 / 写回顺序
最后才考虑主循环数学

这个顺序是拿 CUDA 经验外推的。后面我会看到：8k 的 bug 其实不在这 5 条里的任何一条。它属于 NVIDIA 的 checklist 上根本不存在的第六类——"硬件执行粒度和软件抽象不对齐"。但那天我还不知道。

10. 排查：五条怀疑全部不命中

4 月 4 日早上。我把两份代码放在编辑器里对照：一个是旧的 src/team03_gemm.mu（fail 8k），一个是历史版本里最早的 historical_versions/team3_gemm_v1.mu（64×64×32，128 线程，结构最简单）。按 bug 怀疑顺序逐条实验排查。

条 1（epilogue 写回竞争）：把 epilogue 全删，只留主循环加一个 musaMemset 清零。跑 correctness_8k——错误率不变，还是 ~75%。epilogue 不是原因。

条 2（smem 复用生命周期）：加额外 __syncthreads() 在 smem commit 前后、改单缓冲、延迟 prefetch——每次都跑一遍 8k 正确性，错误率都不变。也不是 smem 的问题。

条 3（CTA 映射 / swizzle 碰撞）：怀疑是不是大 grid 下 CTA 之间有空间碰撞。把 grid 缩到 1×1 单 block 跑一个 256×256 tile——错误率还是 75%。和 grid 大小无关。

条 4（向量化 half 打包）：把向量化的 float4 写回拆回逐元素 __float2half，跑——结果一模一样错。

条 5（主循环数学）：循环体是标准的 load_matrix_sync + mma_sync 序列，看不出数学错。在非常小的尺寸上（比如 32×32）单独跑，4 个 32×32 累加器自己的数值对比标量参考是正确的。算得对，是"怎么把它们拼回 C 矩阵"这一步出了问题。

5 条都否掉之后剩下的可能性只有一个：硬件调用约定和软件抽象对不上。load_matrix_sync / mma_sync 的硬件粒度是什么？我代码里用的软件粒度（32 线程 warp）和它是不是一致？

我回去翻 MUSA 文档，这次盯着"warp"相关的章节看。之前读的时候我默认跳过了——以为和 NVIDIA 一样。

11. 物理 warp = 128 的假设和最小复现

MUSA 文档里 MP22 的"物理 warp"是 128 线程——不是 32。

这个数字让我愣了一下。v68 的 block(512, 1, 1) 按 NVIDIA 模型是"16 个 32 线程的 warp"，按 MP22 模型是"4 个 128 线程的 warp，每个内部含 4 个 32 线程子 warp"。两个模型完全是两回事。

我提出一个假设：load_matrix_sync / mma_sync 在 MP22 上要求同一物理 warp（128 线程）内的所有 128 个线程传入一致的地址参数；如果 4 个子 warp 传入了不同地址，属于未定义行为。

这个假设里有两层——一层是硬件约束（物理 warp 粒度），这是我有把握的；另一层是"分歧时硬件具体会怎么做"——我最初的推测是"硬件采纳其中一个子 warp 的地址、其余子 warp 静默丢弃"，但这是实验观察到的外部行为，不是官方规格。真实的规格可能是"未定义行为"，也可能是"实现相关"。我能断言的是：统一地址的结果是对的，分歧地址的结果是错的——任何依赖分歧地址的代码都是错的，不管硬件具体怎么实现这个未定义行为。

如果假设成立：v68 mega tile 里每个物理 warp 的 4 个子 warp 分别算 4 个不同的 32×32 子块，硬件实际执行的结果只覆盖其中一个子块，其余 3 个子 warp 拿到了"不是自己请求的那份数据"的结果，再按各自的 (w_m, w_n) 写到 C 的 4 个不同位置——1/4 正确、3/4 是从另一个位置复制过来的数据。错误率 75%。

写了一个最小复现。两个 kernel 都是 <<<..., 128>>>（单 block 128 线程 = 一个完整物理 warp），都跑 32×32×16 WMMA，唯一差异是子 warp 之间是否使用分歧地址。文件 reproduction/test_subwarp_diverge.mu:15-37：

__global__ void __launch_bounds__(128) test_divergent(...) {
    int tid = threadIdx.x;
    int sub = tid / 32;                                    // 0..3 (sub-warp)
    int tile_m = (int)blockIdx.y * 64 + (sub % 2) * 32;    // divergent
    int tile_n = (int)blockIdx.x * 64 + (sub / 2) * 32;    // divergent

    fragment<accumulator, 32, 32, 16, float> acc;
    fill_fragment(acc, 0.0f);
    for (int k = 0; k < K; k += 16) {
        fragment<matrix_a, 32, 32, 16, half, col_major> fa;
        fragment<matrix_b, 32, 32, 16, half, col_major> fb;
        load_matrix_sync(fa, A + (long)k * M + tile_m, M);
        load_matrix_sync(fb, B + (long)tile_n * K + k, K);
        mma_sync(acc, fa, fb, acc);
    }
    // ... write acc to C ...
}

test_uniform 把 tile_m / tile_n 改成不依赖 sub 的固定值，其他完全一样。编译跑（K = 32，随机数据）：

=== UNIFORM: all sub-warps same addr (32x32 blocks) ===
  0/65536 errors, max=0.000977           ← 只有 FP16 rounding
=== DIVERGENT: sub-warps diff addr (64x64 blocks) ===
  64477/65536 errors (98.4%), max=2.553  ← 几乎全错

max_err = 2.553 在 K=32 下已经完全脱离 FP16 rounding 量级。放大到 K=16384 同样的错误结构会让 max_err 膨胀到 ~36——正好对上 correctness_8k FAIL 时看到的那个数字。假设的核心部分被证实：子 warp 间地址分歧的 WMMA 调用会产生错误结果，必须强制同一物理 warp 内 128 线程地址一致。

📌 一个诚实的逻辑 gap 说明:严格来说,这个最小复现只证明了"分歧地址的 WMMA 会出错",并没有直接证明 v68 的 8k fail 就是因为这个原因——它是通过代码模式匹配得到的间接证据:v68 的代码里 wid = tid / 32 + m_base = w_m * 64 这个结构和最小复现的 tile_m = blockIdx.y * 64 + (sub%2)*32 是同一种"子 warp 间地址按 sub 编号分歧"的模式;最小复现证明这种模式会错;所以 v68 这个具体的 bug 大概率就是这个原因。

如果你想要直接证据,严谨的做法是:把 v68 的 mega tile kernel 整段拷出来,只改一行(int wid = tid / 32 → int wid = tid / 128,然后相应重算 w_m / w_n 的 mask),其他不动,跑 correctness_8k,看是否从 fail 变 pass。这是我当时实际做的下一步——4 月 4 日早上写出来的"重建版 0"(下一节 §12)就是这个直接修复,5/5 用例 pass + 22T 性能,完全证实了根因。所以最小复现 + 直接修复两个证据合起来才闭合了整条逻辑链——单看最小复现是不够的,这一点我应该提前讲清楚。

关于 "硬件具体怎么处理分歧地址" 这一点:从我观察到的错误结构(75% 错,模式是 4 个位置被同一份数据填充)可以推测硬件等同于 "只接受了 4 个子 warp 中某一个的地址",但这个行为是未文档化的,我不应该依赖它;正确的工程实践是主动保证地址一致,而不是猜测 "反正硬件会选某一个"。

一个容易被误读的细节:本文里反复说的 "128 线程原子",指的是同一物理 warp 内全部活跃 lane 必须一致,而不是"block 必须至少有 128 线程才能合法调用 WMMA"。两者是很不一样的:

v68 的情况是 block(512, 1, 1) → 4 个物理 warp,每个 warp 全部 128 lane 都活跃。此时如果代码按 32 线程切"逻辑 warp",会让同一物理 warp 的 4 个 32 线程子 warp 活跃且分歧——非法,触发本文的 bug。
相反,如果 block(32, 1, 1) → 只有 1 个 32 线程"partial warp",硬件把不活跃的 96 lane mask off,活跃的 32 lane 之间天然一致(因为只有 32 个)——合法,WMMA 可以正确工作。

第二种情况是我在 dispatch 里给小尺寸写的兜底 kernel gemm_wmma_16(32 线程 / block)能正常跑的原因。读者看到 dim3 block(32, 1, 1) 不要以为它违反了本文的 "128 原子" 声明——它没有,partial warp 合法,divergent full warp 非法,这两者是本文故事的关键区别。

那"90 TFLOPS"到底是什么？ 算一下：v68 每个物理 warp 的 4 次 mma_sync 被硬件退化为 1 次，独立的浮点乘加只有理论 2·M·N·K 的 1/4。性能公式 TFLOPS = 2·M·N·K / runtime_seconds 的分子按理论尺寸算，分母按实际时间算——而"实际做 1/4 工作"的时间大约是"做完整工作"时间的 1/4，所以报告的 TFLOPS 被放大约 4 倍。真实的"做 1/4 工作"吞吐大约是 22 TFLOPS（这正好是后面第一版正确内核的性能），乘以 4 倍得到 ~90 TFLOPS。

90T 不是我跑得快，90T 是我算了 1/4 的东西然后除以 1/4 的时间。

那天晚上笔记上又多了一行：

Apr 3 晚: 找到了。MP22 物理 warp = 128，WMMA 是物理 warp 原子的，v68 的 "90T" = (1/4 工作量) × (4× 公式放大)。明天把 WARP_SIZE 改掉。

12. 重建版 0：物理 warp = 128 修复（0 → 22 TFLOPS）

当时的情况：4 月 4 日早上。假设已经被最小复现证实，但代码还没动。

思路：把 WARP_SIZE = 32 的假设从代码里彻底清掉。wid 按 128 线程分组，所有 WMMA 参数在同一物理 warp 内强制一致。这是一次性修复整个 bug 的唯一办法。

关键技术点：

改动分三处：

第一处，warp 编号的计算：

// Before (v1 broken):
int wid = tid / 32;        // 16 "warps" in 512-thread block
int w_m = wid % 4;         // 4x4 layout → divergent across sub-warps
int w_n = wid / 4;

// After (first correct version):
int wid = tid / 128;       // 4 physical warps
int w_m = wid & 1;         // 2x2 layout → uniform within each 128-thread warp
int w_n = wid >> 1;

第二处，tile 尺寸和 warp 布局重新算。v1 mega tile 是 256×256×32 按 4×4 布局（16 个逻辑 warp）。改成物理 warp 粒度后 512 线程只剩 4 个物理 warp——4 个 warp 一口气覆盖 256×256 撑不住，每个 warp 要扛 128×128 输出，累加器寄存器会爆。所以第一版我先退回一个更稳的配置：128×128×64 tile、512 线程、4 物理 warp 的 2×2 布局，每个 warp 负责 64×64，由 2×2 的 32×32×16 累加器组成。先把正确性立住再想性能。

第三处，主循环 WMMA 地址计算：

int m_base = w_m * 64;     // same for all 128 threads in a warp
int n_base = w_n * 64;
#pragma unroll
for (int kk = 0; kk < BK; kk += 16) {
    fragment<matrix_a, 32, 32, 16, half, col_major> a0, a1;
    fragment<matrix_b, 32, 32, 16, half, col_major> b0, b1;
    load_matrix_sync(a0, &sA[kk][m_base     ], BM + PAD);
    load_matrix_sync(a1, &sA[kk][m_base + 32], BM + PAD);
    load_matrix_sync(b0, &sB[n_base     ][kk], BK + PAD);
    load_matrix_sync(b1, &sB[n_base + 32][kk], BK + PAD);
    mma_sync(acc[0][0], a0, b0, acc[0][0]); mma_sync(acc[0][1], a0, b1, acc[0][1]);
    mma_sync(acc[1][0], a1, b0, acc[1][0]); mma_sync(acc[1][1], a1, b1, acc[1][1]);
}

性能：跑 5 个正确性用例：全过，包括 correctness_8k，max_err = 0。8k 性能 22 TFLOPS。

22T 看起来比之前的 90T 慢了 4 倍，但我清楚那 90T 本来就是幻觉——除以 4 正好是 22T，精准对上 §11 里"1/4 工作量 × 4× 公式放大"的数学。22T 是这个项目里第一个不带水分的数字。

踩过的坑：

一开始想直接保留 256×256 tile。我最初的第一个方向是"只改 warp 布局不动 tile 尺寸"——256×256 + 4 物理 warp 的 2×2 布局，每个 warp 覆盖 128×128 输出 = 16 个 32×32×16 累加器。改完编译 compiler_report.sh 报 private_memory = 1600 bytes——寄存器溢出灾难。性能 ~5T。我太贪心，想一步到位。退回 128×128 单 block 才稳住。
PAD_A = 4 尝试。改 warp 布局的时候我顺手把 PAD 从 8 减到 4 想省点 smem——编译过了，性能 ~18T，比 PAD=8 的 22T 低。反编译看发现 stride=4 对 load_matrix_sync 的 bank 访问造成了新冲突。改回 PAD=8 立刻回到 22T。没必要的改动在重建期千万别做，每次只改一个变量。
epilogue 没跟着改。v1 的 epilogue 是为 4×4 warp 布局 + 16 个小累加器写的。改到 2×2 warp 布局后每个 warp 的累加器结构完全变了，epilogue 的 store_matrix_sync 序列必须重写。第一次跑 max_err ≠ 0——不是 WMMA 错，是 epilogue 把结果写到了错的位置。重新核对每个 warp 的输出坐标才修好。
22T 让我怀疑自己。看到 22T 的那一刻说实话心里有点崩——原来的 90T 是幻觉，现在的 22T 离目标还差 4 倍。我纠结了一下是不是应该接受 22T 直接交差，至少 100/100 的分拿回来了。但我给自己立的硬规则里写了"正确性优先 + 性能是硬门槛内能抢多少是多少"——既然还有时间，就继续推。

怎么走过去的：22T 立住脚了，接下来每一步必须"先解释为什么会快，再改"。

13. 重建版 1：stride-33 bank conflict 消除（22 → 48 TFLOPS，+118%）

当时的情况：22T 离合理值还远。

思路：在瞎猜"哪里慢"之前先做瓶颈隔离。我把当前 kernel 复制一份，改成 compute-only 版本——保留主循环的 load_matrix_sync + mma_sync，把 epilogue 的所有写回代码全部注释掉、只留一个 __syncthreads() 防 dead-code-elimination。跑一次：52.6 TFLOPS。

compute-only kernel:   52.6 TFLOPS   ≈ 42 ms
full kernel (w/ epi):  22.0 TFLOPS   ≈ 100 ms
                                       ────
epilogue 独吞的时间:                     58 ms

epilogue 吃掉的时间比整个 MMA 计算阶段还多。这个数字让我意外——直觉上 epilogue 应该只占总时间的一小部分，毕竟主循环要跑 K/BK = 16384/32 = 512 次迭代、每次 4 个 MMA、总 2048 次 MMA 调用。怎么可能 epilogue 比主循环还慢？

关键技术点：

去看 epilogue 代码，原版是：

// Old epilogue (slow):
float my_smem[32 * 32];
store_matrix_sync(my_smem, c_frag, 32, mem_col_major);    // stride = 32
__syncthreads();
int col = lane;                                             // 0..31
for (int r = 0; r < 32; r++) {
    half val = __float2half(my_smem[col * 32 + r]);
    C[(long)(global_col_base + col) * M + global_row_base + r] = val;
}

读回时每个 lane 固定读一列 32 个 float。32 个 lane 并行访问起点地址 \{0*4, 32*4, 64*4, ..., 31*32*4} 字节——每个 lane 的起点相隔 128 字节。

S4000 的 shared memory 是 16 banks，每 bank 32-bit 宽。这条是官方文档 Ch02 明确写的，之前读 WMMA 章节时我没留意这条。把 bank 数学写清楚：

smem bank 索引 = (byte_address / 4) % 16（4 bytes / bank width）
我的访问模式：lane k 读 my_smem[k * 32 + r]（k = 0..31 是 lane，r = 行偏移）
lane k 第一个访问的地址：k * 32 * 4 + r * 4 = (k * 32 + r) * 4 字节
对应 bank：((k * 32 + r) * 4 / 4) % 16 = (k * 32 + r) % 16 = (0 + r) % 16

因为 k * 32 % 16 = 0（32 是 16 的倍数）——所有 32 个 lane 无论 k 取什么值都落到同一个 bank。这不是"间隔大所以冲突"，是"stride 是 bank 数的倍数所以冲突"——关键不是绝对字节距离，是 stride 相对 bank 数量的取模结果。32-way bank conflict，完全串行化。

（NVIDIA 上 smem 是 32 banks，stride=32 恰好让每个 lane 映射到不同 bank 是"恰好对齐"的特殊情况。我凭 NVIDIA 经验写 stride=32 在 MP22 的 16 banks 下就从"恰好对齐"翻转成了"32-way 冲突"——这是个典型的"pattern 平移时常数没跟着调"的 bug。）

修复一个常量：stride 从 32 改成 33（和 16 互质）：

// New epilogue: stride 33 breaks the bank collision
float my_smem[32 * 33];
store_matrix_sync(my_smem, c_frag, 33, mem_col_major);
// ... read-back loop uses stride 33 ...

改成 stride=33 之后，lane k 的 bank = (k * 33 + r) % 16 = (k + r) % 16（因为 33 % 16 = 1）——32 个 lane 被 "分摊" 到 16 个 bank 上，每个 bank 有 2 个 lane 同时访问，冲突从 32-way 降到 2-way。2-way 冲突硬件可以在 2 个 cycle 内处理，和完全不冲突（1-way）的性能差别只有 2 倍——但比 32-way 的 32 倍好 16 倍。实测性能从 22T 跳到 48T，接近理论的 2× 加速。

性能：22 TFLOPS → 48 TFLOPS，加速 2.2 倍。一个常量的修改，两倍多性能。这是整个项目里单次改动收益最大的一次。

踩过的坑：

我一开始没想到是 epilogue 慢。看到 22T 第一反应是"主循环肯定有问题"——可能 MMA 调度不好、可能 fragment 加载 bank conflict、可能 prefetch 没对上。如果没做 compute-only 隔离，我估计会在主循环里瞎折腾半天。compute-only 是 10 分钟写出来的，但它指向的瓶颈位置完全出乎意料。微基准隔离是整个项目里 ROI 最高的调试工具。
stride-33 vs stride-17 vs stride-65 的选择。改成 33 之前我也考虑过 17 或者 65——只要和 16 banks 互质都能避开冲突。测了一下三个值性能差不多（都到 47-48T），最终选 33 是因为 33 和 MMA 的 32 size"最接近"，epilogue 循环边界最自然写。没有本质区别。
之前 70+ 版里 epilogue 都用的 stride=32。这件事让我反思了很久。v68 的 90T 里其实也有 stride-32 bank conflict 的影子——但因为 v68 实际上只做 1/4 工作量、epilogue 的计算量也只有 1/4，bank conflict 的串行化惩罚也按比例减小。幻觉的 90T 里其实掩盖了 stride-32 的 bank conflict。正确性塌陷修好、真实吞吐恢复之后，这个 bank conflict 才成了可见的瓶颈。之前的迭代里从来没有注意过这个问题——因为我当时不知道 S4000 是 16 banks，也没想到要跑 compute-only 微基准来隔离瓶颈。这是我自己该注意而没注意的事。
store_matrix_sync 的 stride 参数是"逻辑列数"不是"字节数"。第一次写改动时我搞错了——把 stride = 32 直接改成 stride = 33 × 4 = 132（以为是字节），结果越界写坏 smem。看文档才明白这个参数是"每行多少个元素"（对 float 就是逻辑列数）。改回 33 才对。这条和 CUDA WMMA 的语义一致，我是凭惯性假设没仔细读签名导致的——和 NVIDIA WMMA 有用过经验的开发者不会踩这个坑。

怎么走过去的：一行改动、2.2 倍加速。48T 之后开始考虑下一个方向。

14. 重建版 2：M-first 蛇形 CTA 排序（48 → 54 TFLOPS）

当时的情况：48T，epilogue 已经不是瓶颈。下一个怀疑对象是 L2 cache 利用率。

思路：在 8192×8192×16384 规模下先把 B 矩阵的数据量算清楚：

B 的一条 N 条带（BN=128, K=16384, half）:     128 × 16384 × 2 B =  4 MB
B 的所有 N 条带（8192/128 = 64 条）:                  64 × 4 MB = 256 MB
L2 总容量（`musaGetDeviceProperties.l2CacheSize` 实测）:              24.0 MB

默认的 N-first 遍历（blockIdx.x 在 N 方向先动）让相邻 block 共享 A 列——但 A 列在整个 grid 上是不断滑过的（每个 M tile 完就换），没多少 L2 复用的可能；而 B 的 64 条 N 条带总共 256 MB，L2 装不下——每个 block 都是 L2 miss。

改成 M-first 遍历（先在 M 方向动、N 方向外层）让相邻 block 共享 B 的同一条 N 条带——一条 N 条带 4 MB 轻松塞进 24 MB L2，后续所有处理这条 N 条带的 block 全部 L2 hit。再加蛇形遍历（奇数列反转 M 方向）让列切换处的最后一个 block 和下一列的第一个 block 空间相邻，A 也能在列边界处复用一部分 L2。

关键技术点：

int bx = (int)blockIdx.x;           // N direction tile
int by = (int)blockIdx.y;           // M direction tile
if (bx & 1) by = (int)gridDim.y - 1 - by;   // serpentine: reverse odd columns

零算法改动，只改遍历顺序。

性能：48T → 54T，+12%。这次改动的推理全程不需要 profiler，只需要三个数字：L2 容量、一条 B 条带大小、所有 B 条带总大小。算完就知道该让哪一维留在 L2 里。

踩过的坑：

"super tile swizzle" 失败。看 mutlass 代码发现 NVIDIA 的高级技巧叫 "super tile swizzle"——不是简单的 M-first，而是把 grid 切成 SWIZZLE × SWIZZLE 的超大块，每个超大块内 M-first 遍历。我试了 SWIZZLE=2/4/8，全部没有比简单蛇形 M-first 更好。v83 就是这个尝试，性能和蛇形持平。过度设计没有收益。
L2 persistence API 尝试。MUSA 有 musaStreamAttributeAccessPolicyWindow 可以把一段内存标记为 "persisting"，希望 L2 缓存优先保留。我把 B 矩阵标记成 persisting 试了一下——完全没效果。musaStreamSetAttribute 对 S4000 来说只是软提示，实际 L2 替换策略由硬件决定，不受用户控制。这是纯浪费时间的尝试。
CTA 遍历模式的测试需要大 grid。第一次改完 CTA 排序我用 4096×4096×8192 小一点的 grid 测了一下——没看到明显差异（4096×4096 在这个规模下 L2 能装下全部 B，蛇形和非蛇形都 hit）。我以为改动没生效。换到 8192×8192×16384 才看到 +12% 的收益。CTA tile ordering 的收益只在超过 L2 容量的 grid 下才显现，小 grid 测不出来。
边界 tile 的处理。蛇形遍历下奇数列的 by 是倒序的——如果 gridDim.y 不是偶数会漏掉最后一行或者越界。幸好 8192/BM 永远是 16 或 32 的整数，没出问题。但我在代码里加了一条 assert assert(gridDim.x > 0 && gridDim.y > 0) 防御。

怎么走过去的：54T 的下一个方向是继续拉 tile。

15. 重建版 3：256×256×32 大 tile + 1024 线程（54 → 57 TFLOPS）

当时的情况：54T，tile 尺寸回到了 128×128×64，warp 布局 2×2。要继续推性能必须把 tile 做大回 256×256——但这次要在"正确的 warp 粒度"前提下重新设计。

思路：算术强度公式 AI = BM·BN / (BM+BN)：

Tile	AI	说明
`128×128`	64	当前
`256×256`	128	翻倍

AI = 128 意味着每从 smem 加载 1 字节数据可以驱动 128 次 FLOP，roofline 模型下彻底进入 compute-bound 区域。

代价是 256×256 tile 需要更多累加器和更多线程。512 线程 + 4 物理 warp 撑不住（参见 §12 踩坑 1 的 16 累加器寄存器爆炸）。必须升到 1024 线程 + 8 物理 warp——block 线程数翻倍，每个 warp 仍然覆盖 64×128 输出（8 个 32×32×16 累加器）。

关键技术点：

tile 256×256×32，BK=32
1024 线程 = 8 物理 warp
warp 布局 4×2（M 方向 4、N 方向 2），每个 warp 负责 64×128 输出
每个 warp 持有 8 个累加器 fragment acc[2][4]（2 行 × 4 列 32×32 子块）
shared memory sA[32][256+8] + sB[256][32+8] = 16896 + 20480 = 37376 bytes
单缓冲 + float4 prefetch 寄存器
__launch_bounds__(1024, 1) 1 block/SM
每次 kk 步（K=16）：2 个 A fragment + 4 个 B fragment → 8 次 mma_sync

每次 kk 迭代的 MMA/load ratio：

loads:   2 A + 4 B = 6
MMAs:    2 × 4 = 8
ratio:   8/6 ≈ 1.33

和 v68 的 1.0 比更好——意味着每个 fragment 加载能驱动更多 MMA，搬运开销被摊薄。

性能：54 → 57 TFLOPS，+6%。AI 翻倍带来的收益比预期小——因为 54T 阶段 kernel 已经不是 memory-bound 了，再提 AI 对 compute-bound 的改进是次线性的。

踩过的坑：

__launch_bounds__(1024, 2) 寄存器暴跌 3T 灾难。第一次编译我加了 __launch_bounds__(1024, 2) 想要 2 block/SM 提 occupancy——结果 compiler_report.sh 报 private_memory = 1120 bytes，性能从预期的 57T 暴跌到 3 TFLOPS。查了半天发现是 1024 线程 × 2 block/SM × 每线程约 130 寄存器 = 266K 寄存器需求，超过单 SM 的 262K 寄存器预算，编译器被迫溢出到 DRAM。改回 __launch_bounds__(1024, 1) 立刻回到 57T。MP22 上强制 2 block/SM 几乎总是得不偿失。
warp 布局 2×4 vs 4×2 vs 8×1 对比。1024 线程 = 8 warp 有三种布局方式：

布局	每 warp 覆盖	A loads	B loads	MMAs	ratio	实测
8×1（全在 M 方向）	32×256	1	8	8	0.89	~44 T
2×4（M 2, N 4）	128×64	4	2	8	1.33	~37 T
4×2（M 4, N 2）	64×128	2	4	8	1.33	57 T

表面上 2×4 和 4×2 的 MMA/load ratio 都是 1.33，按 roofline 推它们应该差不多——但实测 2×4 是 37T、4×2 是 57T，差了 20 TFLOPS。这个数字让我很意外，我怀疑过是不是 2×4 实现里有 bug，单独跑了一遍正确性测试（5 个用例全过），确认不是代码错。

反编译看汇编，2×4 的主循环多出一批和 A fragment 加载相关的 LDMA 指令——A 的 stride 比 B 大得多（sA[BK=32][BM+PAD=264] 的 M 方向 stride 是 264 half，sB[BN=256][BK+PAD=40] 的 N 方向 stride 是 40 half），2×4 布局下每个 warp 需要 4 次 A fragment 加载、每次加载的有效延迟因为大 stride 比 B 加载高。4×2 布局反过来——每个 warp 只需要 2 次 A 加载、4 次 B 加载，主要延迟由 B 加载控制，B 的 stride 小所以总延迟低。

我的解释是 "A/B stride 不对称让 A 加载成为瓶颈"，但这是反编译 + 直觉得到的推测，不是通过严格 profiling 证实的。也可能是编译器在两种布局下做了完全不同的寄存器分配，或者 2×4 触发了某个我没注意到的 smem bank conflict。准确的机制需要更细的 profiler 分析才能确定——我选了"实测最好的 4×2 直接上"这个务实路径，没有继续挖"为什么 2×4 差"。如果你复现这个实验得到不同的数字，很可能是和我实现上的细节差异导致的，不要当成 2×4 架构本身比 4×2 差 20T 的结论。

BK=64 寄存器溢出灾难。我试过把 BK 从 32 扩到 64 减半 barrier 数量——实测 22T，寄存器溢出 private_memory = 420 bytes。BK=64 需要 4 个 float4 prefetch 寄存器（而不是 2 个），加上累加器压力就爆了。接受 BK=32。
16% occupancy 的心理障碍。1 block/SM 意味着 1024 / (max 6144 threads/SM) ≈ 16% occupancy。这个数字在 NVIDIA 经验里是"很差"的——通常希望 50% 以上。我一开始怀疑 16% 是不是瓶颈。试了 __launch_bounds__(512, 2) 压回 512 线程 + 2 block/SM——性能 48T，不如 1024 + 1 block/SM 的 57T。高 occupancy 在 MP22 上不是目标，大 tile + 足够 warp 才是。8 个物理 warp 提供的延迟掩盖已经够用。

一个值得说明的 occupancy 计算 nuance:scripts/compiler_report.sh 实际上给出了两个 occupancy 估算:
- 乐观模型(只算 Temp R 寄存器):126 regs × 1024 threads ≈ 129K regs/block,占 262K 预算的一半 → 1 block/SM, 16% occupancy
- 保守模型(算上 Temp + Attribute + Internal + Coefficient 全部 per-thread regs):总寄存器需求超过 SM 预算 → 0 block/SM, 0% occupancy (理论上"装不下")
实际运行是 1 block/SM——也就是乐观模型对、保守模型过于保守。这意味着 MP22 的 AR/I/Coefficient 寄存器不完全占用主寄存器文件(可能有独立硬件 bank,或者按需分配)。这是 MP22 寄存器预算模型一个未公开文档化的细节,如果摩尔线程能补充 ISA 手册说明这点,后续开发者计算 occupancy 会准确得多。我在这里只能凭实测确认"乐观模型对",没法给出第一性原理的解释。
协作加载模式的改动。1024 线程加载 sA/sB 的模式和 512 线程完全不同。原来每个线程加载 float4 (8 half)，sA 需要 (32 × 256) / (8 × 512) = 2 次加载、sB 需要 (256 × 32) / (8 × 512) = 2 次加载。1024 线程只需要 1 次加载就够了，但模式要重新算。第一次写错了——用的是旧的 512 线程映射，导致 sA 的前一半线程加载了正确数据、后一半线程越界。正确性立刻挂。修改的时候我把线程→数据的映射单独拎出来用了两行注释写清楚：

// sA layout: sA[BK=32][BM+PAD=264], 1024 threads load 16896 bytes
// each thread handles float4 = 8 half; total elems = 32*256 = 8192
// 1024 threads * 8 half = 8192 → exact cover, 1 load per thread
int ta_k = tid >> 5;                // 0..31 (BK direction)
int ta_m = (tid & 31) << 3;         // 0, 8, 16, ..., 248 (BM direction)

这次的 debug 时间 30 分钟。

怎么走过去的：57T 之后代码层面的明显优化空间都吃完了，要继续往上推必须去找别的杠杆。

16. 重建版 4：postRA + ILP 调度 flag（57 → 67 TFLOPS，+18%）

当时的情况：57T，tile 尺寸 256×256、1024 线程、4×2 warp 布局都已经到位。代码层面没有明显空间了。

思路：我想知道代码的"物理极限"在哪，于是做了一件之前一直没做的事——把 team03_gemm_work.mu 编译产物反编译到 binary 级，逐条看主循环的 VLIW word 是怎么组织的。

MP22 的每条 VLIW word 是 12 字节 = 3 个 4 字节 slot，每个 slot 可以装一条来自 MMA / LDMA.RD / LDMA.WR / DMA.LD / INT / FOP / BIT / CTRL / AP 八种 pipe 之一的指令，同一 word 里的 3 条指令在同一 cycle 并行执行。编译器的核心工作就是把主循环的几百条指令塞进这些 3-slot word。

用 mcc --musa-device-only -S 拿汇编，然后把 .text 段切成 12 字节一条 VLIW word，对照 MP22 的 opcode 表逐条解码 pipe 类型。主循环一共 55 条 VLIW word，pipe 分布：

MMA:         16
LDMA.RD:     24  (fragment 读取)
LDMA.WR:     2   (smem 写入)
DMA.LD:      2   (global 读取)
INT + BIT:   7   (地址计算)
AP.BARRIER:  2
CTRL:        2
─────────────
Total:       55 VLIW words × 3 slots = 165 slots

0 个空 slot。编译器在"把指令塞进 VLIW word"这一层已经到顶——每一个 cycle 都在做有效工作，没有任何一个空槽能塞新指令进去。要再往上推，唯一的杠杆不是"塞更多指令"，而是"让现有指令以更好的顺序执行"——也就是调度。

MCC 是基于 LLVM 的，和 NVIDIA 的 nvcc 不同——它把大量调度相关的 LLVM 开关通过 -mllvm -mtgpu-* 暴露给用户。我翻 MCC User Manual 的 -mllvm 章节，挑出几个看起来相关的：

-mtgpu-enable-postra-sched                    # 寄存器分配后再跑一次调度器
-mtgpu-enable-max-ilp-scheduling-strategy     # ILP 最大化策略

然后做一次组合搜索：每个 flag 单独开/关测一次性能，再做组合。

关键技术点：

加两个 flag 就行，改 Makefile / 编译命令：

mcc --offload-arch=mp_22 -O3 \
    -mllvm --mtgpu-enable-postra-sched \
    -mllvm --mtgpu-enable-max-ilp-scheduling-strategy \
    src/team03_gemm_work.mu -o test

两个 flag 的作用：

postRA-sched：LLVM 默认在寄存器分配（RA）之前做一次指令调度。但 RA 阶段可能会因为寄存器冲突插入 spill / 重排指令，改变调度结果。postRA-sched 让编译器在 RA 之后再跑一次调度器，基于真实的寄存器布局进一步优化。对寄存器密集的 GEMM 内核收益明显。
max-ilp-scheduling-strategy：激进寻找指令级并行（ILP）机会。默认调度器偏保守（可能考虑功耗或关键路径长度），这个策略把 ILP 最大化作为主要目标。

性能：

组合	性能
默认（无 flag）	57 TFLOPS
`postRA` 单独	~63 TFLOPS
`ILP` 单独	~62 TFLOPS
`postRA + ILP`	67 TFLOPS

57 → 67T，+18%。两个 flag 的收益几乎独立相加。

踩过的坑：

flag 是否会影响正确性。启用新 flag 之后我先跑全部 5 个正确性用例——必须的习惯。correctness_8k max_err = 0 才敢确认性能数字。有一个 flag（-mtgpu-enable-internal-regs-ra）跳到 79T 但 max_err = 350——那个要留到 §19 单独讲。
单独启用每个 flag 收益不明显。postRA 单独 +10%，ILP 单独 +9%，组合起来 +18%。差不多是独立相加。但如果只看单独数字会觉得"不值"——完整收益要组合后才显现。这是为什么必须做组合搜索而不是单个测试。
flag 搜索的顺序。我一开始按字母表顺序测 flag——低效。后来改成按"可能最有用"的优先级：先测调度相关的（postRA、ILP、misched），再测 mem/fence 相关的（FenceSetting、load-store-opt），最后测 IR 级的（internal-opt）。优先级排序让搜索从 4 小时降到 2 小时。
反编译工具链的坑。mcc --musa-device-only -S 只给汇编不给 binary；要拿 binary 需要 mcc --musa-device-only -c 产 .o，然后 llvm-objdump -d <file.o>。第一次尝试的时候 objdump 报 "unknown architecture"——MTGPU 架构不在标准 LLVM 表里。摩尔线程提供了一个特殊的 llvm-objdump-mtgpu，必须用这个版本才能反编译 MTGPU 代码。

怎么走过去的：67T 之后我继续翻 MCC User Manual 找还有什么 flag 可以试。

17. 重建版 5：FenceSetting=mixed（67 → 73 TFLOPS，+9%）

当时的情况：67T。调度器已经尽力，但 barrier 相关的 VLIW slot 还有潜在空间。

思路：MCC 的 FenceSetting 控制内存 fence 的语义强度。默认是保守（强 fence，任何 memory 操作前都插入全量同步），这保证正确性但浪费指令。MCC 支持几个级别：strict（强）、mixed（混合，按需选择最弱的满足语义的 fence）、relaxed（弱）。

对 WMMA 内核来说，__syncthreads() 是主要的 fence 点，load_matrix_sync / mma_sync 内部的隐式 fence 其实很多都可以弱化——比如同一 warp 内的 MMA 之间不需要跨 warp 同步。换成 mixed 让编译器按需选择，理论上能省一些 fence 相关的指令 slot。

关键技术点：

-mllvm --FenceSetting=mixed

性能：67 → 73 TFLOPS，+9%。

踩过的坑：

FenceSetting=relaxed 正确性挂。relaxed 模式最激进，想一次到位试了一下——correctness_8k max_err = 12，挂。回滚到 mixed。mixed 是"按需选最弱合法 fence"，relaxed 是"直接用最弱 fence 不管合法性"，后者不安全。
FenceSetting=graph_color（图染色）更慢。MCC 文档里还提了一个 graph_color 模式——理论上应该更聪明，基于依赖图染色选择 fence。实测 72 TFLOPS，比 mixed 还慢 1T。我没再深究，直接用 mixed。
这个 flag 的收益区间需要经验指南。MCC User Manual 对 FenceSetting 的签名描述清晰——列出了 strict / mixed / relaxed / graph_color 四个级别。但对每个级别的典型收益区间、适用 workload、组合效果没有给出经验数据——这对新用户来说需要自己做一次完整扫描才敢选。我的 4 级别扫描结果（strict 67T / mixed 73T / relaxed 正确性挂 / graph_color 72T）可以作为未来官方"flag 最佳实践文档"的一个数据点。建议摩尔线程在 MCC 后续版本的文档里补一份 "GEMM/Conv 典型 workload 的 flag 推荐组合表"，帮助新用户绕过扫描过程直接到达最优配置。

怎么走过去的：73T 之后继续找下一批 flag。

18. 重建版 6：load-store-opt 等 4 个新 flag（73 → 73.8 TFLOPS，+1%）

当时的情况：73T。明显 flag 已经吃完，但我想穷尽 MCC User Manual 里所有可能的选项。

思路：MCC 文档里还有几个我没测过的 flag：

-mtgpu-load-store-opt                    # smem load/store 指令合并
-mtgpu-load-cluster-mutation             # 分配后指令重排
-mtgpu-memory-sched-mutation             # 内存调度优化
-mtgpu-tiny-offset-hint                  # 偏移 < 2^32 的编译器暗示

全部一起开试一下。

关键技术点：四个 flag 组合加到编译命令里。

性能：73 → 73.8 TFLOPS，+1%。

边际收益很小，但是值得——总共 7 个 flag 组合后性能从 57 推到 73.8，+29.5%，没改一行 C++。

把所有 flag 装进 mcc wrapper：

#!/bin/bash
REAL_MCC=/usr/local/musa/bin/mcc
exec "$REAL_MCC" \
    -mllvm --mtgpu-enable-postra-sched \
    -mllvm --mtgpu-enable-max-ilp-scheduling-strategy \
    -mllvm --FenceSetting=mixed \
    -mllvm -mtgpu-load-store-opt \
    -mllvm -mtgpu-load-cluster-mutation \
    -mllvm -mtgpu-memory-sched-mutation \
    -mllvm -mtgpu-tiny-offset-hint \
    "$@"

装到 /usr/local/bin/mcc，让所有调用 mcc 的地方（grader、Makefile、harness、test_cycle）都透明获得 flag 注入。这样单独的测试脚本不用每个地方都改命令行。

踩过的坑：

mtgpu_unroll_threshold 属性无效。MCC 文档里提到一个 kernel attribute __attribute__((mtgpu_unroll_threshold(N))) 可以控制循环展开阈值。我试了 N=4/8/16/32——性能完全不变。可能这个属性只对某些循环生效，对我用 #pragma unroll 1 显式禁展开的主循环没影响。
mtgpu_tiny_offset 属性无效。同样 __attribute__((mtgpu_tiny_offset)) 标在 kernel 上，性能不变。和上面一起进了"文档写了但实测无效"的清单。
-mtgpu-if-convert 风险。文档里 -mtgpu-if-convert 描述是"条件分支转 predicated execution"。启用后性能 73.2T，略低，可能是主循环里没什么条件分支，predicated 反而开销。不用它。
-mtgpu-maxregcnt=N 的脆弱性。想强制给每个线程更多寄存器（-mtgpu-maxregcnt=160）——编译器报 "cannot satisfy register constraint"，直接拒绝。改 maxregcnt=140 能编但触发溢出。这个 flag 只有在压寄存器下限时有用。

怎么走过去的：73.8T 是 flag 搜索的终点。接下来只剩两个可能方向：一个是跨不过去的 Pass #346 bug（§19），一个是头文件 bug 阻挡的 512 线程架构（§20）。

19. 没走通的路 1：`internal-regs-ra` 与 `machine-scheduler` 的 WMMA 正确性冲突

当时的情况：73.8T，flag 搜索到底。但 flag 搜索过程里我碰到一个诡异的开关，值得单独讲。

思路：搜 flag 的时候我试过 -mllvm --mtgpu-enable-internal-regs-ra。这个 flag 的描述是"启用 internal register allocator，使用 Internal bypass 寄存器提升性能"。

启用后 team03_gemm_work.mu 性能从 73T 跳到 79 TFLOPS——一步跨过 muBLAS 78.5T 的水平线。我看到 79T 的时候有点不敢相信，连着跑了 5 次确认不是波动。

然后 correctness 挂了：max_err ≈ 350。这不是 rounding 误差的量级，是 WMMA 累加器被彻底写错。

我纠结了几分钟要不要放弃——既然挂 correctness 直接用不了。后来决定花时间定位 bug，因为 +8% 的诱惑太强，如果能找到 workaround 就能一步跨过 muBLAS。

关键技术点：

LLVM 有个不太常用的调试工具叫 opt-bisect-limit：-mllvm -opt-bisect-limit=N 让编译器执行到第 N 个 pass 之后就停止，后面的 pass 全部跳过。通过二分 N 的值，可以精确定位到"从第几个 pass 开始引入错误"。

思路：N=∞（默认）会跑错，N=0（不跑任何 pass）应该是对的（因为根本没做任何优化）——那么中间必然有一个临界点 N*，从 N* 开始结果变错。

我跑了十来次二分：

N = 1000 (default):                       ❌ max_err=27.2734
N = 500:                                   ❌ max_err=27.2734
N = 250:                                   ✅ max_err=0.0000
N = 400:                                   ❌ max_err=27.2734
N = 340:                                   ✅ max_err=0.0000
N = 360:                                   ❌ max_err=27.2734
N = 350:                                   ❌ max_err=27.2734
N = 345:                                   ✅ max_err=0.0000
N = 346:                                   ❌ max_err=27.2734

临界点定位在 LLVM 的 machine-scheduler pass——pre-RA 阶段的 Machine Instruction Scheduler。在我本地的 mcc 环境下它的运行序号是 #346，但 pass 编号会随 LLVM 版本、启用的 MTGPU 定制 pass 集合而浮动，核心定位是 pass 名 machine-scheduler，不是数字 346。摩尔线程编译器团队看到这段时请用 pass 名定位，不要用编号。

机制分析（以下是我根据 opt-bisect 二分 + MCC 文档字面描述 + 外部可观察行为推测的机制，真实机制需要摩尔线程编译器团队确认）：

--mtgpu-enable-internal-regs-ra 启用的是一种特殊的寄存器替换策略——根据 MCC 文档字面描述，它会使用 "internal bypass register"（文档里称作 I[] 寄存器）。按我的推测，这类寄存器是 MTGPU 微架构里某种可以让数据在 pipe 之间短路传递的专用通道，延迟比走通用 Temp 寄存器短。这是对 MTGPU 微架构的外部推测，不是我通过查源码确认的事实。
pre-RA 的 machine-scheduler 重排指令时，会把某些寄存器的引用替换为 I[] 寄存器以利用这个短路通道。
根据实测，重排后 WMMA 的操作数被 route 通过 I[] 寄存器——这一步是实验观察到的行为（通过反编译 binary 看 WMMA 前后的寄存器分配差异）。
但 WMMA 硬件无法从 I[] 寄存器读取正确数据。这是最关键的一条——我通过"启用 IRA + 默认调度器 → max_err=350"这个组合观察到的外部行为，推测的硬件原因。更精确的机制可能是：MMA 从一个专用 fragment 寄存器 bank 读取操作数（参见 §21），I[] bypass 通道没有连接到这个 bank；或者是 I[] 和 fragment bank 在硬件实现上有 race condition。具体哪一种需要摩尔线程硬件团队确认。
结果：MMA 读到未定义数据，累加器数值损坏。

所有 workaround 尝试：

组合	正确性	性能
IRA + `-enable-misched=false`（禁调度器）	✅	70.4 T（小于基线）
调度器 + 无 IRA	✅	73 T（基线）
IRA + `-misched=converge`（收敛模式调度器）	✅	71.7 T（小于基线）
IRA + `--internal-opt-user-fop=false`（禁用 FP bypass）	❌	79 T（仍错）
IRA + `--internal-opt-user-int=false`（禁用 INT bypass）	❌	79 T（仍错）
IRA + `--internal-opt-user-bit=false`（禁用 BIT bypass）	❌	79 T（仍错）
所有 bypass 子类型全禁用	❌	79 T（仍错）
IRA + 默认调度器	❌	79 T ← 想要但不对

禁用单个 bypass 子类型全部无效——这说明 bug 不在"哪类 bypass 被启用"的层面，而是在"哪些寄存器被选为 I[] 替换候选"的决策阶段。我猜测是 RA 阶段决定替换候选时没有把 WMMA 操作数排除出候选集合。

给摩尔线程编译器团队的反馈（请核实再 action）：

这是一个可以精确复现的 bug：启用 -mllvm --mtgpu-enable-internal-regs-ra 后 team03_gemm_work.mu 的性能从 73T 跳到 79T、但 WMMA 累加器产生 max_err ≈ 350 的数值错误。opt-bisect 定位到 machine-scheduler pass。相对的正确/错误组合我在上表列全了。

我的建议方向（不是结论，因为机制是推测的）：在 machine-scheduler 或 RA 的候选选择逻辑里，对 opcode 为 MMA.* / WMMA.* 的指令的操作数寄存器加一条"不可被 I[] bypass 替换"的约束。具体实现位置和是否这样修对需要摩尔线程编译器团队根据真实的微架构语义判断——可能我推测的机制是错的，实际 bug 在另一个地方。

如果这个方向对、bug 修了，直接收益 73T → 79T（+8%），一步跨过 muBLAS 78.5T。这是我能给的最精确的一条可复现 bug report——但最终定性和修复方案需要摩尔线程内部工程师确认。

踩过的坑：

二分搜索找错临界点 2 次。opt-bisect-limit 的反馈有时候不是立即的——Pass 执行会缓存一些状态，有时候 N=346 报错但 N=345 不报错仅仅是因为 stale cache。我用 rm -rf /tmp/mcc_cache 清缓存才得到稳定的二分结果。
-opt-bisect-limit 和 -opt-print-after-all 组合输出巨大。第一次开 -opt-print-after-all 想看 machine-scheduler 前后的 IR diff——输出超过 1 GB，我的终端直接卡死。后来用 -opt-print-after=machine-scheduler 只打印那一个 pass 前后的 IR 才看清楚。
"IRA" 和 "internal-regs-ra" 的命名需要仔细区分。我查 LLVM 资料时一开始把 MCC 的 internal-regs-ra 和 LLVM upstream 的 IRA（Iterative Register Allocator）搞混了——前者是 MTGPU 特有的 internal bypass 机制，后者是独立于 MTGPU 的通用组件。这是一个理解细节，看明白之后后续工作顺畅很多。

怎么走过去的：没走通。留作给摩尔线程的 bug report。

20. 没走通的路 2：mma.h:761 头文件 bug

当时的情况：在 §19 的 Pass #346 路径卡住之后，我想另一个方向——muBLAS 的 512 线程 + 16 累加器架构。

思路：通过 §21 的 muBLAS 逆向工程（后面会讲）我发现 muBLAS 用的是 512 线程 / 4 物理 warp / 每 warp 16 个累加器 的配置。相比我当前的 1024 线程 / 8 warp / 每 warp 8 累加器，muBLAS 的架构：

每 warp 的 MMA/load ratio 是 2.0（比我的 1.33 高 50%）——意味着每个 fragment 加载驱动更多 MMA
更少线程意味着更少同步开销
250 寄存器/线程（vs 我的 126）——有空间做更激进的 prefetch 和 fragment 管理

我想在我的代码里重现这个架构。

关键技术点：

tile 256×256×32，block 512 线程 = 4 物理 warp，warp 布局 2×2，每 warp 覆盖 128×128 输出 = 16 个 32×32×16 累加器（4×4 布局）。

性能：第一次编译立刻出问题。compiler_report.sh 报 private_memory = 2080 bytes——寄存器溢出到 DRAM。性能暴跌到 58 TFLOPS，比 1024 线程版差。

踩过的坑：

试图绕过 2080 字节溢出。我先怀疑是自己的代码有问题——有没有意外声明局部数组？没有，代码干干净净。我试过：
- 减少累加器数量：16 → 8，溢出降到 1040 bytes，但 MMA/load ratio 也掉回 1.0
- __launch_bounds__(512, 1) 强制 1 block/SM：private_memory 没变
- -mtgpu-maxregcnt=256 提寄存器预算：没变
- register 关键字标注所有局部变量：没变
- #pragma clang loop unroll_count(2) 减少展开：性能 40T，更差

全部无效。private_memory = 2080 这个数字是编译器不管我怎么改代码都坚持给出的。

反推出 2080 的来源。2080 bytes = 16 × 32 × 4 + 32 bytes align。16 是累加器数量，4 是 float 字节数——所以每个累加器占 32 个 float。问题是累加器应该只占 8 个 float（32×32 / 128 = 8，MP22 warpSize = 128）。怎么变成 32？
追到 fragment 类定义。用 grep -rn "__frag_base<float" /usr/local/musa/include/——找到 /usr/local/musa/include/crt/mma.h:761：

// /usr/local/musa/include/crt/mma.h:761 (当前)
class fragment<accumulator, 32, 32, 16, float> :
    public __frag_base<float, 32> {};

__frag_base<float, 32> 意思是当前实现给每个 fragment 分配 32 个 float 元素的空间。如果按 NVIDIA 的 fragment 元素数公式 M × N / warpSize 推算：

NVIDIA（warpSize = 32）：32 × 32 / 32 = 32
MP22（warpSize = 128）：32 × 32 / 128 = 8

按这个推算 MP22 只需要 8 个 float / fragment，而当前头文件分配的是 32 个——多分配了 3 倍寄存器空间。这是一个可能的优化空间，但我不确定 32 这个数字背后有没有设计意图。可能的解释有几个：

(a) 从 NVIDIA 代码移植时没有按 warpSize 调整——按公式应该是 8
(b) 为了和 NVIDIA 源码的某种兼容模式保持结构对齐
(c) 为未来更大 warp size 的架构（比如 MP31+）预留 slot
(d) 硬件对齐要求，需要 32 float 的粒度

哪个解释对需要摩尔线程 compiler/header 团队内部确认。

实测改小之后的效果。作为一个验证实验，我在本地环境把这一行改成 __frag_base<float, 8> 后重新编译：private_memory 从 2080 降到 8 字节，512 线程版性能从 58T 提升到 69T，正确性 max_err = 0（5 个用例全过）。这说明至少在当前 MUSA SDK 3.1.0 的代码生成路径上，改成 8 不会破坏功能。但这不能作为"32 这个值是错的"的充分证据——可能存在我没观察到的其他使用场景依赖 32 这个 slot 数。

69T 仍低于 1024 线程版的 73T——因为 512 线程只有 4 物理 warp，延迟掩盖不如 8 warp 充分。所以这个改动的直接收益并不是"多 11T"。真正的价值是如果这 3 倍的寄存器预算确实可以安全收回，就能打开 512 线程 + 16 累加器 + 双缓冲这条更激进的架构路径——muBLAS 就是在这条路径上靠手写汇编达到 78.5T 的。当前实现下 512 线程架构被阻塞在寄存器预算之外。

不能把修改发到评测环境。评测用的是摩尔线程预装的 MUSA SDK 3.1.0，头文件是只读的。我本地改了验证原理，但最终提交的代码不能依赖这个改动——必须用 1024 线程版的 73T 作为最终数字。

给摩尔线程 compiler/header 团队的反馈：

请确认一下 /usr/local/musa/include/crt/mma.h:761 的 __frag_base<float, 32> 是否必须是 32。如果 32 只是一个历史保留值、可以安全改到 M × N / warpSize = 8，那么这是一个一行代码的优化，能为后续 512 线程 + 16 累加器 + 汇编双缓冲路径扫清寄存器预算障碍。如果 32 背后有我不知道的设计意图（比如对齐、兼容、未来硬件预留），请告知，方便我在文章里更新定性。

我能提供的证据是实测数据：改到 8 之后 private_memory 从 2080 → 8 字节、性能从 58T → 69T、正确性 max_err = 0。这证明至少在我的 workload 上功能不破坏。最终判断需要你们内部评估。

怎么走过去的：这条路在我的环境下不能走通（头文件只读）。结论留作给摩尔线程的一条可能优化 inquiry。

21. 其他踩过的坑汇总（踩坑清单）

重建过程中除了上面讲过的几个关键版本踩坑之外，还有一批"试了没用"的尝试。下面列一个汇总给后来者省时间。

C++ 双缓冲的 5 种方案（全部失败）：

方案	smem	性能	失败原因
动态 `buf` 变量 `sA[buf][k][m]`	66 KB	30 T	编译器无法常量折叠动态索引
静态展开 ×2（手动 buf=0/1）	66 KB	17 T	代码膨胀 → I-cache + 寄存器压力
XOR 指针切换 `base ^= BUF_SIZE`	66 KB	51 T	每次 smem 访问多一次 INT.ADD
模板常量 `template<int BUF>`	66 KB	49 T	模板内联代码膨胀 + 正确性 bug
3D 数组 + 优化 flags	66 KB	56 T	编译器仍生成动态偏移

全部低于单缓冲 73T 基线。原因每个略有不同，共同特点是动态索引带来的额外 INT.ADD 指令、代码膨胀、寄存器分配恶化，间接效果超过了双缓冲能省的那 1 个 barrier。muBLAS 之所以能做双缓冲是因为它用单条 BIT.LUT.XOR 指令实现基地址切换——零 per-access 开销，这在 C++/WMMA API 层面做不到。

Binary patch 的尝试。我甚至尝试过不改源码、直接 patch 编译后的 fatbin：

删除主循环的第 2 个 __syncthreads() barrier：成功 patch，正确性保持，性能完全不变。WMMA API 的 barrier 延迟（152 cyc）在 8 warp 的延迟掩盖下已经被消化掉了，删掉也没收益。
把 barrier 改成 NOP fence：同上，正确性保持，无性能提升。
直接替换整条 VLIW word：破坏了同 word 内其他 pipe 的指令，kernel 启动崩溃。放弃。

这条路整整用了 1 天时间，得到的结论是 "VLIW binary 已经是编译器优化后的结果，手动 patch 没空间"。具体工具链细节（fatbin 结构、AP.BARRIER 编码、VLIW word 格式）可以结合 docs/MUSA_WMMA_API_REFERENCE.md 和摩尔线程官方 MCC 手册里描述的二进制格式自行摸索。

Split-K 灾难。想把 K 维度拆成多个 stage、每个 CTA 只处理一段 K、最后 atomicAdd 累加。预期收益是更多 CTA → 更好的 SM 利用率。实测：

Split-K = 2：性能 33 T（慢了一半）
Split-K = 4：6.8 TFLOPS——atomicAdd(float*, float) 在 MP22 上是软件模拟的，每次 atomic 要几百 cycle

完全不可行。这条路 1 小时就放弃了。

77.5T 是零数据的假象。一次偶然我跑出 77.5T 的数字——超过 muBLAS 78.5T。兴奋了 10 分钟，然后开始怀疑：为什么这个数字只在我用 musaMemset(dA, 0, ...) 初始化的时候出现？换成随机数据立刻掉回 73T。

查了一下——零数据让 L2 命中率接近 100%，因为所有 A 和 B 的 tile 都是同一块零内存，L2 永远 hit。真实 workload（随机数据）下 L2 miss rate 正常，性能回到真实值。77.5T 是测量方法 artifact，不是真实性能。这个乌龙让我给自己立了一条规则："所有 > 73T 的历史测量都要怀疑是零数据或 cache warmup"。

PAD 值扫描。想系统扫描 PAD_A 和 PAD_B：

PAD_A	PAD_B	性能
0	8	71.5 T
4	8	72.0 T
8	8	73.0 T（基线）
16	8	72.8 T
8	0	36.9 T（灾难）
8	4	66 T
8	16	72.5 T

PAD_A=8, PAD_B=8 是最优的。PAD_B=0 尤其灾难（36.9T）——B 矩阵 stride=32 遇上 16 banks 的 bank conflict，连 warp scheduler 都掩盖不了。PAD_A=0 损失较小（71.5T）因为 A 的加载模式不那么敏感。

register 关键字反效果。C++ 的 register 关键字在现代编译器基本是 no-op——我试了一下，结果 mcc 居然对它敏感，给所有局部变量加 register 之后性能掉到 40.3T。可能 mcc 把它解释成了"强制放寄存器但不考虑预算"，导致溢出。不要用 register 关键字。

异步内存 copy。MUSA 有 __pipeline_memcpy_async 类似 NVIDIA 的 cp.async——理论上可以异步从 global 搬到 smem。实测 MP22 上这个是软件 fallback，不是硬件异步，性能和同步 DMA 一样甚至更差。

持久 kernel。让 kernel 内部跑一个 while 循环处理多个 tile，避免 block 启动开销。尝试了 cooperative launch——MUSA 3.1.0 的 musaLaunchCooperativeKernel 报 "FIXME in driver" 错误，直接不支持。退而求其次用 atomicAdd 实现任务队列——同样慢，atomicAdd 延迟太大。

手动 fragment 加载绕过 load_matrix_sync。最激进的一条路。在 §20 的 load_matrix_sync 约束下我想：如果我能自己把数据放进 fragment 寄存器，就不受 stride 约束，可以做任意 swizzle。代码写出来后测试数据值：手动加载的 fragment 和 load_matrix_sync 产生的 fragment，逐元素 bit-exact 对比——256/256 完全一致。然后测试 MMA 结果：相同数据值分别喂给 mma_sync，累加器输出——1024/1024 全部不匹配，max_diff = 311。

结论：MP22 的 MMA 硬件不是从普通寄存器读 A/B 操作数，而是从一个专用的 fragment 寄存器 bank 读。load_matrix_sync 的工作不只是"把数据加载到寄存器"，它还把数据放到 MMA 硬件能识别的那个特殊 bank 里。手动填的通用寄存器即使字节内容一致，MMA 也读不到。这一条一次性封死了所有"绕开 load_matrix_sync"的 C++ 路径——swizzle、双缓冲 PAD=0、手工 prefetch，全部不可行。

这个实验是整个项目里最关键的一次"证伪"——它证明了 MP22 的某些硬件约束是 C++ API 绕不过去的。

22. 73.8T 是我这个工具链路径的经验上限

73.8T 走到这里之后,我想确认一件事:这就是当前工具链和硬件的极限吗?能再往上推多少? 仅靠前面几节积累的实验笔记不够——那些笔记记录的是"我试了什么、得到什么数字",但回答不了"这个数字是不是硬件物理上限"。我做了三组面向硬件的直接实证:

22.1 VLIW 165/165 满打包

反编译 team03_gemm_work.mu 主循环的 55 条 VLIW word，每条 3 个 slot 全部填满，0 空 slot（详见 §16）。编译器在"往 VLIW word 里塞指令"这一层完全没有空间。

⚠️ 必须区分两个不同的 100% 指标——这是常被混淆的地方:

指标含义我的实测与硬件峰值的关系
VLIW slot 打包率 编译器把指令塞进 VLIW word 的密度 165/165 = 100% 不直接对应 MMA 单元利用率
MMA 执行单元利用率 Tensor Core 单元在多少 cycle 中实际在做 MMA ~73% (= 73 / 100 TFLOPS) = 性能 / 硬件峰值

VLIW 100% 打包说的是"编译器没有让任何一个 slot 空着",但每个 slot 里装的是什么才决定 MMA 利用率。在我的主循环里:

16 条 MMA 指令

24 条 LDMA.RD (fragment 加载)

7 条 INT/BIT (地址计算)

4 条 LDMA.WR + DMA.LD

2 条 AP.BARRIER + 2 条 CTRL/NOP

总 55 word × 3 slot = 165 slot,其中只有 16 个装了 MMA。所以即使 VLIW 打包率 100%,MMA 单元只在 16/55 ≈ 29% 的 word 时间里被驱动。剩下 71% 的 word 时间在做地址计算 / 数据搬运 / 同步——这些是支撑 MMA 必需的开销,但本身不创造 FLOPs。

真正的性能上限是 MMA 利用率,不是 VLIW 打包率。我的 73 TFLOPS / 硬件峰值 100 TFLOPS = 73% MMA 利用率;muBLAS 的 78.5 / 100 = 78.5% 利用率;两者相差 5.5%——和"VLIW word 差距 5/55 ≈ 9%"在数量级上一致(但不完全相等,因为 muBLAS 的 50 word 里 MMA 比例更高,等价于"用更少的非 MMA 开销完成同样的 MMA")。

正确的解读是:本文的 "VLIW 165/165 打包" 证明的是**"编译器调度已尽力,没有空 slot 可以给我塞新指令";它不证明** "MMA 单元已达硬件极限"。要进一步逼近 100T 硬件峰值,需要的不是"塞更多指令"(没空间了),而是**"让更高比例的 slot 装 MMA"**(降低非 MMA 指令的占比)——这正是 muBLAS 用手写汇编 + BIT.LUT.XOR 单条切换 buffer 做到的事情。

指标	含义	我的实测	与硬件峰值的关系
VLIW slot 打包率	编译器把指令塞进 VLIW word 的密度	165/165 = 100%	不直接对应 MMA 单元利用率
MMA 执行单元利用率	Tensor Core 单元在多少 cycle 中实际在做 MMA	~73% (= 73 / 100 TFLOPS)	= 性能 / 硬件峰值

22.2 muBLAS 的完整逆向数据

通过 LD_PRELOAD 拦截 musaLaunchKernel 和 muLaunchKernel，加上 ELF .note 段的 MessagePack 元数据解码，我拿到了 muBLAS hgemm 内核的完整画像。

⚠️ 方法论免责声明:下面这些数字是我在开发期间通过逆向工程得到的。完整方法和示例代码框架在 reproduction/mublas_inspect/ 目录里(包含可编译的 LD_PRELOAD wrapper、fatbin 解析步骤、ELF .note 提取脚本)。但出于摩尔线程 muBLAS 二进制的版权考虑,我没有 ship 一个"一键拆解 libmublas.so"的完整工具——只 ship 方法、框架代码和我得到的样例数据,读者需要自己按方法在本地复现。如果你复现时得到的数字和下面的样例不一致(可能是 SDK 版本变化、muBLAS 内部更新、或我的方法有偏差),请以你自己的实测为准,并在 issue 里告诉我。本文 "5 个 VLIW word 差距" 这个核心结论依赖下面表里 muBLAS 的 ~50 VLIW word 数字,任何对它的独立验证或反驳都对整篇文章的可信度有意义。

kernel name:          hgemm_nn_512_256x256_tensor_qy2
block:                (512, 1, 1) = 4 物理 warp
grid (8k 规模):       (32, 32, 1) = 1024 CTAs
temp_reg_count:       250
shared_reg_count:     74
shared_memory_size:   65,664 bytes   (= 2 × 32,768 + 128 padding = 双缓冲)
private_memory_size:  0              (零溢出)
main loop:            ~800 bytes ≈ 100 指令 ≈ 50 VLIW words
has_barrier:          true
code type:            手写汇编

和我的版本对齐（我的数据来自本次 session 实测，compiler_report.sh 输出可验证）：

指标	我（1024 线程）	muBLAS（512 线程）
线程 / block	1024	512
物理 warp / block	8	4
Temp regs / thread	126	250
累加器 / warp	8	16
MMA / load ratio	1.33	2.0
smem (static)	37,396 bytes 单缓冲	65,664 bytes 双缓冲（PAD=0）
主循环 VLIW word	55	~50
Private memory（溢出）	8 bytes（metadata 级，非 user data）	0
代码类型	编译器生成	手写汇编

差距 = 5 个 VLIW word（约 10%）。这 5 个 word 是地址计算冗余——muBLAS 用单条 BIT.LUT.XOR 指令实现双缓冲基地址切换，所有 smem 地址都相对于切换后的基地址计算，零 per-access 开销；C++ 编译器面对动态缓冲区索引时每次 LDMA 加载都要多一次 INT.ADD 计算偏移。这 5 个 word 对应 73T → 78.5T 的 6% 差距。

22.3 Fragment 寄存器约束（§21 讲过）

MP22 的 MMA 硬件从专用 fragment 寄存器 bank 读取 A/B 操作数，手动加载的通用寄存器即使 bit-exact 匹配 MMA 也读不到。这一条一次性封死了所有绕开 load_matrix_sync 的 swizzle / 双缓冲 PAD=0 路径。

综合三条证据：

编译器在 VLIW 打包层面已经到顶（0 空 slot）
进一步优化需要手写汇编（muBLAS 的 50 VLIW word 是手写汇编做的，C++ 做不到）
手写汇编路径需要 MTGPU inline asm 寄存器约束支持——当前 mcc 3.1.0 不支持（这是我给摩尔线程的另一条 API 建议）

结论：73.8 TFLOPS 是我这个优化路径（mcc 3.1.0 + C++ WMMA API + load_matrix_sync + 编译器生成代码）的经验上限。这里要把几个"上限"的层次讲清楚，避免被误读成"S4000 的硬件极限是 73.8T"：

层次	性能	来源
S4000 硬件 FP16/BF16 Tensor 理论峰值	100 TFLOPS	摩尔线程 `MTT S4000 产品规格书` 官方规格表，FP16/BF16 张量算力条目
官方比赛 muBLAS 基线（DVFS 禁用）	86,012.6 GFLOPS	`grader/test_cases.json` 的 `mublas_baseline_gflops` 字段
muBLAS 含 sync 的 8k 测量（我的 Docker 容器环境，DVFS 未禁用）	~78.5 TFLOPS	我用 `bench_mublas` 20 次中位数实测，和 73.8T 苹果对苹果比较
我的 C++ 编译路径上限	73.8 TFLOPS	本文讨论的路径，grader 报告记录 `absolute_gflops=73808.5`、`official_perf_gflops=73855.5`
启用有 bug 的 `internal-regs-ra`（正确性挂，不能用）	~79 TFLOPS	§19 讨论的 Pass #346 bug 路径，报告值但正确性 `max_err=350`

73.8T 是 muBLAS 78.5T 的 94%。这 6% 差距不是"S4000 硬件不如 NVIDIA 某某卡"——S4000 硬件峰值 100T，muBLAS 用的手写汇编已经证明硬件能跑到 78.5T。差距在我这条 C++ WMMA API 编译路径自身的抽象层限制上——VLIW 打包已经 100%，C++ 双缓冲不可行（§21 的 5 种方案全败），手动 fragment 加载被硬件 bank 约束封死。

不是"代码还能更聪明"，是"C++ / WMMA API 的抽象层能表达的最优代码已经到这里"。要再往上推有两条路，都在工具链侧：

修 §19 的 Pass #346 bug → 73T → 79T
修 §20 的 mma.h:761 bug → 解锁 512 线程 + 16 累加器 + 汇编双缓冲

任意一条被修，team03_gemm_work.mu 不改一行代码就能过 78.5T。更远的突破要等 MTX 1.0 低级 IR、load_matrix_sync 增加 swizzle 参数、MUTLASS 成熟（需要 MCC 4.3.4+），或者 MP31+ 下一代硬件。

23. 微基准教训：31 cyc vs 403 cyc 的故事

前面每一步优化——stride-33、tile 尺寸、L2 遍历、flag 搜索——都建立在"我相信自己测的数字"这个前提上。但项目重建早期我犯过一次很贵的错，浪费了差不多两天时间。

事情是这样：重建版 1（stride-33）之后我想建一个性能模型预测 BK 取不同值时的性能。模型需要 MMA 延迟作为关键参数。我没跑微基准——直接从 NVIDIA A100 的经验外推：A100 的 mma.m16n8k16 大概 4 周期延迟，MP22 的 32×32×16 tile 比 m16n8k16 大 8 倍，我猜 MMA.323216 应该在 30 周期左右。

基于 30 周期这个数字建模型：预测 BK=16 能跑 135 TFLOPS、BK=32 能跑 110 TFLOPS、BK=64 能跑 95 TFLOPS。推论：BK=16 最好，应该把 tile 的 K 维度降到 16。

然后我花了一天时间重写代码用 BK=16，加倍 barrier 数量、改 prefetch 寄存器数量、重调 smem 布局。跑出来：

BK=16 实测: 42.5 TFLOPS
BK=32 实测: 57.0 TFLOPS
BK=64 实测: 22.0 TFLOPS (寄存器溢出)

预测和实测差 2-3 倍。模型预测 BK=16 最好，实测 BK=16 最差（除了 BK=64 的溢出灾难）。

我一开始以为是模型里的其他参数估错——smem 带宽估错、barrier 开销估错、ILP 假设估错。折腾了一天多试各种修正都对不上。最后决定"算了，老老实实把 MMA 延迟测一下"。

写了一个最朴素的依赖链微基准——2000 次 mma_sync，每次都用前一次的累加器输出作为下一次的输入，这样流水线无法重叠，测到的就是单条 MMA 从发射到结果可用的完整延迟：

__global__ void bench_mma_latency(float* out) {
    fragment<accumulator, 32, 32, 16, float> acc;
    fragment<matrix_a, 32, 32, 16, half, col_major> a_frag;
    fragment<matrix_b, 32, 32, 16, half, col_major> b_frag;
    // ... 初始化 ...

    clock_t t0 = clock();
    #pragma unroll
    for (int i = 0; i < 2000; i++) {
        mma_sync(acc, a_frag, b_frag, acc);   // serial dependency chain
    }
    clock_t t1 = clock();
    // ... 存 (t1 - t0) / 2000 ...
}

跑出来：403 cycle。比我猜的 31 cycle 高 13 倍。

这里要马上澄清一个容易被误读的地方：403 cycle 是 MMA.323216 这一条指令的"延迟"（从发射到结果可用），不是"吞吐"。两者对性能的影响完全不同：

延迟决定单条指令依赖链的最短执行时间，必须通过"并行发射多条独立 MMA"来掩盖
吞吐决定 GPU 每秒能执行多少条 MMA，这是性能的真实上限

MMA.323216 的计算量是 32 × 32 × 16 × 2 = 32,768 次乘加。同样计算量的 NVIDIA mma.m16n8k16 是 16 × 8 × 16 × 2 = 4,096 次乘加——MP22 的单条 MMA 包含的计算量是 NVIDIA 的 8 倍。所以 403 cycle 的延迟不是"比 NVIDIA 慢 100 倍"，而是"一条指令干 8 倍活"自然需要更长的 pipeline。有效每单位计算延迟 403/8 ≈ 50 cycle，和 NVIDIA 可比。

更关键的是 MP22 的 MMA 吞吐：硬件规格峰值是 100 TFLOPS（S4000 的 FP16 Tensor 规格）。403 cycle 的延迟只要有足够多的独立 MMA 填充 pipeline 就能被掩盖——我的 1024 线程 / 8 物理 warp / 每 warp 8 个独立累加器架构（acc[2][4]）已经把 pipeline 填得很满，实测 73.8T 是从下往上逼近这个 100T 硬件上限的位置。muBLAS 的 78.5T 靠手写汇编又往前推了一步。1024 线程 / 8 物理 warp 配置能跑到 73T 的原因——不是靠"减少 MMA 延迟"，是靠"并发足够多 MMA 把延迟藏在 warp 调度后面"。

回到性能模型的问题：我之前猜 31 cycle 错了一个量级，直接把 BK=16 的预测从 42.5T 拉到 135T——因为低估了"要掩盖这个延迟需要多少并行 MMA"。带回真实的 403 cycle 之后，BK=16 的预测从 135T 降到 42.5T——和实测完全吻合。BK=32 预测 72.9T，实测 73.1T，误差 0.3%。BK=64 的"22T"是寄存器溢出造成的，纯粹的算法预测是 60T 左右，溢出惩罚把它拖到 22T。

之前一切对不上，不是模型结构有问题，是输入参数错了一个量级。

然后我立刻补跑了整套微基准，花了半天时间把 MP22 的关键常数测齐：

参数	值	测量方法
MMA.323216 延迟	403 cyc	2000 次串行 MMA 依赖链
MMA 吞吐（每 warp 1 MMA in-flight，全芯片下限估算）	~57 TFLOPS	延迟 403 cyc × 每 warp 仅 1 个 in-flight MMA × 8 warp/SM × 56 SM 的理论下限
MMA 吞吐（硬件规格峰值）	100 TFLOPS	FP16 Tensor 规格上限，需要每 warp 多个独立累加器填 pipeline 才能逼近
`__syncthreads`（8 warp）	152 cyc	有/无 barrier 对比
smem 顺序读带宽	27 B/cyc/SM	stride-1 sweep
smem stride-8 penalty	1.67×	PAD_A=8 造成
Global DMA 带宽	748 GB/s	float4 顺序读，对应 S4000 官方规格 768 GB/s 的 97% 有效值
DMA 延迟	467 cyc	pointer-chase
INT 吞吐	24 ops/cyc/SM	独立 INT 链
FP 吞吐	79 ops/cyc/SM	独立 FP 链
INT + FP 混合	39 ops/cyc	两种同时，非 24+79=103

测量条件（所有微基准的统一环境，为了复现）：

条件	值
硬件	单卡 MTT S4000，独占运行，无其他 kernel 抢占
SDK	MUSA SDK 3.1.0，mcc clang 14.0.0
编译 flag	`--offload-arch=mp_22 -O3 -march=native` + `/usr/local/bin/mcc` wrapper 的 7 个 flag
GPU 状态	核心频率官方规格 1.5 GHz（`MTT S4000产品规格书.pdf`），`mthreads-gmi -q` 报告 Max graphics clock 1600 MHz、`musaGetDeviceProperties.clockRate` 返回 1669.99 MHz；运行时在 DVFS 控制下浮动。每次测量前跑 `gpu_health_check.sh` 确认健康度 > 70T
Warmup	每个微基准前跑 10 次 warmup、丢弃前 3 次结果
测量次数	每个数字 50 次重复，取中位数；标准差 < 2% 才接受
随机数据	A/B 使用 `rand() % 100 - 50` 缩放到 `±0.5` 的 half，避免零数据 L2 假象
DVFS 状态	Docker 容器内无 `CAP_SYS_ADMIN` 权限，DVFS 不能禁用。小尺寸微基准（单 kernel 耗时 < 3ms）不触发降频，保持 boost 频率；大尺寸 8k GEMM 在 `musaDeviceSynchronize()` 空闲间隙会触发 DVFS 降频，boost/throttle 交替（参见下一小节）

最后一行特别关键——它说 INT 和 FP 在 VLIW 层不完全并行，地址计算会和 MMA 争发射 slot。这正是为什么 §22.2 里我们和 muBLAS 的差距最终定位到"5 个地址计算冗余的 VLIW word"。

带回性能模型：

cyc_per_iter = max(instr × IPC_inv, MMA × 403 / warps) + 403 + barriers × 148
IPC_inv(8w)  = -9.76 + 2.15 / (MMA_count / instr_count)

用 team03_gemm_work.mu 主循环的实际指令分布（16 MMA + 24 LDMA.RD + 5 DMA + 57 INT + 69 FOP + 45 BIT + 5 CTRL + 2 AP = 223 指令）代入：BK=32 预测 72.9 TFLOPS、实测 73.1 TFLOPS；BK=16 预测 42.5 TFLOPS、实测 42.5 TFLOPS。模型误差 < 0.5%。

⚠️ 公式的局限性,必须说明:IPC_inv = -9.76 + 2.15 / ratio 是我用 BK=16 和 BK=32 这两个数据点做的经验线性拟合(把 IPC_inv 看成 1/ratio 的线性函数)。-9.76 这个负常数项在物理上没有意义——它只是拟合出来的截距,**不代表"指令吞吐有 9.76 cyc 的下限"**这种说法。

这个公式只在我测的范围(MMA/instr ratio ≈ 0.07-0.10,对应 BK=16~32 的工作负载形态)有效。外推有几种情况会出错:

MMA 密度极高(比如手写汇编纯 MMA 循环,ratio → 0.5):公式给出 IPC_inv = -9.76 + 4.3 = -5.46,负值,无意义

MMA 密度极低(比如纯 epilogue 写回,ratio → 0.01):公式给出 IPC_inv = -9.76 + 215 = 205 cyc/instr,虚高,远超硬件能力

不同 fragment 尺寸(比如 16×16×16 而不是 32×32×16):MMA 延迟可能不是 403 cyc,基础假设破坏

正确的态度:把这个公式当成"对当前 256×256×32 内核家族的局部插值",不要当成 MP22 通用的性能模型。要建一个真正能外推的模型,需要更多数据点(覆盖不同 ratio 范围),并用 roofline 风格的分段函数(memory-bound / MMA-latency-bound / instruction-throughput-bound 三段)而不是线性拟合。我没有做到这一步——这个公式只是用来验证"403 cyc 的 MMA 延迟数字是对的"的一个最小可行模型,不是文章的核心贡献。

教训：在陌生硬件上，第一天的前 2 小时就应该全部花在微基准上，不要先写 kernel。一个错误的延迟数字可以让后面所有推理和优化方向全部偏掉。这个教训排我自己项目反思的"做错的"第一条。

23.1 关于 DVFS 和 73.8T 的测量条件

73.8T 这个数字需要一个重要的测量条件说明，否则容易被不同环境的测量结果打架。

DVFS 行为：S4000 在 Docker 容器里（没有 CAP_SYS_ADMIN 权限）没法禁用 DVFS。8k GEMM 单次调用耗时约 25-30 ms，在 musaDeviceSynchronize() 带来的空闲间隙里 GPU 会触发 DVFS 降频——我观察到的模式是高频状态（接近官方规格 1.5 GHz 或 gmi 报告的 1600 MHz max）和低频状态（我估计在 1.2-1.4 GHz 区间，没有精确测量每次 kernel 对应的瞬时频率）交替出现，大约每 2-3 次调用循环一次。这对单次测量的影响很大：高频状态下单次 8k 可以跑到 ~80 T，低频状态下掉到 ~55 T。

73.8T 是怎么得到的:我用 scripts/stable_bench.sh 做 20 次重复测量,每次独立的 mcc 编译和 kernel launch,取中位数作为"代表性数字"。我在开发期间(2026-04-03 ~ 04-08)做过几十轮这样的测量,看到的典型分布是:中位数 72-74 TFLOPS,p25 ≈ 73 T,最佳单次 ≈ 75 T,最差单次 ≈ 69 T。73.8 TFLOPS 是那段时间 p25 附近的稳定值——不是中位数,是"DVFS 处于 boost 状态 + 容器刚热起来"那一部分运行的典型值。

给复现读者的一个说明:本仓库 ship 的 reports/stable_bench_20runs.log 是 2026-04-12 fresh clone 从零跑出的 20 次测量,那次中位数是 72.1 T(最佳 74.9 T、最差 69.2 T、CV 2.4%、未检测到 DVFS 强震荡)。73.8 T 和 72.1 T 的差距完全在 DVFS 带来的 ±3% 日常浮动范围内——两者都是正确的复现结果,只是代表不同的统计量(历史 p25 vs 今日中位数)。reports/README.md 里有完整数据和解读。

所以正确的读法是:"73.8 TFLOPS" 是历史最佳稳定运行的代表值,"72-74 TFLOPS" 是中位数区间,两者等价描述"我这条优化路径能达到的性能"。读者在自己的 S4000 上跑 stable_bench 20 后应该落在这个区间里,如果明显偏离(比如 < 65 或 > 78),参考 reports/README.md 的排查指引。

和 muBLAS 78.5T 的对比是在相同 DVFS 状态下的。我单独跑了 bench_mublas 基准 wrapper，用同样的 20 次中位数方法测 muBLAS 的 hgemm 在相同 8k 尺寸下的性能，得到 78.5 T 中位数（同样带 DVFS 噪声）。所以 "73.8T / 78.5T = 94%" 这个比值是在相同测量环境下的苹果对苹果比较，不是和官方 86,012.6 GFLOPS 基线数字的直接比较。

官方 86T 基线怎么来的：官方基线 86 TFLOPS 是在 DVFS 禁用（核心时钟锁定到规格最大频率）的环境下测的——这在 Docker 容器里做不到，需要主机级权限或者硬件开发板。我在容器里能做到的最佳对比是 muBLAS 78.5 T，所以文章里的性能对比全部基于这个数字。如果你在 DVFS 禁用的环境下测，73.8 T 的绝对数字会变高（我估计 ~80-85 T，因为不会被降频），muBLAS 也会相应变高到接近 86 T——相对比值 94% 应该保持。

这个 DVFS 噪声是 MUSA 3.1.0 + S4000 + Docker 环境的通用 confound，不是我内核代码的问题。摩尔线程的 ISA 手册和官方评测文档建议可以提一下这点，帮其他做对比测量的开发者避免被 DVFS 影响了结论。

24. 复现指南

环境：MUSA SDK 3.1.0、MTT S4000、Docker 容器（DVFS 不能禁用）。下面每条命令都在本项目仓库 scripts/ 目录下，本次重建过程中都实测验证过。

第 0 步（强制）：安装 mcc wrapper

要达到 73.8 TFLOPS 必须先装 wrapper——默认 mcc -O3 -march=native 只能跑到 57 TFLOPS。wrapper 预置了 7 个 -mllvm 调度/fence/load-store 优化 flag，对用户完全透明：

sudo cp scripts/mcc_wrapper.sh /usr/local/bin/mcc
sudo chmod +x /usr/local/bin/mcc

这 7 个 flag 的搜索过程在 §16-§18。wrapper 只 exec 原 mcc 并预置 flag，任何通过 PATH 调 mcc 的命令（grader / Makefile / 本文其他脚本）都会自动受益。删除 /usr/local/bin/mcc 即恢复原 mcc（真实 mcc 在 /usr/local/musa/bin/mcc，不受影响）。

实测对比：

配置	8k 绝对性能
默认 mcc（`/usr/local/musa/bin/mcc`）	57.2 TFLOPS
带 wrapper（`/usr/local/bin/mcc` 注入 7 个 flag）	72-74 TFLOPS（随 DVFS 波动）

第 1 步：session 开始的标准三件套

# 1a) GPU 健康检查 (~10 秒)
bash scripts/gpu_health_check.sh
# 健康时报 ~73 TFLOPS；<40T 说明 GPU 进入降级状态，用
#   mthreads-gmi -r 1 -i 0
# 重置，有时需要等 1-2 分钟冷却

# 1b) 完整评测循环：5 个正确性用例 + 8k 性能 + 官方评分
bash scripts/run_official_loop.sh
# 内部调用 run_standard_grader.sh + run_absolute_perf.sh + parse_summary.py，
# 结果汇总到 logs/latest_summary.json；
# 读者可以对照文章数据验证自己的复现结果

# 1c) 编译器资源报告 (~3 秒)
bash scripts/compiler_report.sh src/team03_gemm_work.mu
# 预期输出:
#   Temp regs 126/thread, SMEM 37396 bytes, private_memory 8 bytes,
#   instruction count 956, 1 block/SM (16% occupancy)

开发迭代工具：

# 测试循环: GPU 健康检查 + 标准 grader + 绝对性能 (~90 秒)
bash scripts/test_cycle.sh src/team03_gemm_work.mu

# 稳定性能测量，带百分位数和 DVFS 噪声分离 (~30 秒)
bash scripts/stable_bench.sh src/team03_gemm_work.mu

# 安全编辑: 备份 → 改 → 测 → 失败手动回滚
bash scripts/safe_edit.sh src/team03_gemm_work.mu backup  # 备份到 kernel_backups/
# ... 手动修改 src/team03_gemm_work.mu ...
bash scripts/safe_edit.sh src/team03_gemm_work.mu test    # 跑 test_cycle
bash scripts/safe_edit.sh src/team03_gemm_work.mu revert  # 回滚到 last_good

最小复现 8k bug：

mcc --offload-arch=mp_22 -O3 \
    reproduction/test_subwarp_diverge.mu -lmusart -o /tmp/test_sd
/tmp/test_sd
# 期望（K=32，随机数据）:
#   UNIFORM:    0/65536 errors, max~=0.001
#   DIVERGENT:  ~64477/65536 errors, max~=2.5

mcc wrapper 原始脚本内容（对应仓库 scripts/mcc_wrapper.sh，复现时直接 sudo cp 到 /usr/local/bin/mcc 即可）：

#!/bin/bash
REAL_MCC=/usr/local/musa/bin/mcc
exec "$REAL_MCC" \
    -mllvm --mtgpu-enable-postra-sched \
    -mllvm --mtgpu-enable-max-ilp-scheduling-strategy \
    -mllvm --FenceSetting=mixed \
    -mllvm -mtgpu-load-store-opt \
    -mllvm -mtgpu-load-cluster-mutation \
    -mllvm -mtgpu-memory-sched-mutation \
    -mllvm -mtgpu-tiny-offset-hint \
    "$@"

整个 6 天过程中我维护了两份本地文件：一份 append-only 的实验笔记（每次修改的时间戳 + 正确性 + 性能 + GPU 健康度），一份"尝试了什么 → 结果如何 → 为什么失败"的决策记录。项目末期累积了 60+ 条"死胡同"记录，每条一句话就能告诉后来者这条路走不通。记录每一次错比记录每一次对更有复用价值。

25. 最后

从 4 月 3 日下午发现 8k fail 那一刻到 4 月 8 日把 src/team03_gemm_work.mu 推到 73.8 TFLOPS / 100 分，一共 6 天。

回头看做对的事情：立了"正确性优先于性能"的硬规则；在动大改动前先跑微基准、建性能模型；找到物理 warp 粒度错配这条没人写过的陷阱；读懂了 mcc 的 -mllvm flag 机制并做了系统组合搜索；每次尝试都往本地的实验笔记追加一行，让"错过的路"能被下次直接跳过。

做错的事情更贵：最初没有亲自质疑 "一个 warp 是 32 线程" 这个 NVIDIA 惯性假设——这个假设在 70+ 版迭代里从来没被挑出来，直到 8k 正确性测试暴露它；早期建性能模型时没先跑微基准就凭经验外推 MMA 延迟，用错的数字建模型浪费了一天多；被 muBLAS 的双缓冲诱惑浪费了 13 小时在 5 种 C++ 双缓冲方案上，最后发现是 BIT.LUT.XOR 指令只能手写汇编做到的事情，C++ 抽象层根本够不着。

73.8 TFLOPS 的定位要说清楚——避免被误读。

它不是"S4000 硬件极限"——硬件理论峰值 100 TFLOPS，muBLAS 用手写汇编已经证明了硬件能跑到 78.5T。

它也不是"超越 muBLAS"——相同测量环境下我是 muBLAS 78.5T 的 94%，差了 6%。

它是我这条 C++/WMMA API 编译路径的经验上限——用 mcc 3.1.0 + load_matrix_sync API + 编译器自动生成代码能达到的最好数字。§22 的三条证据支撑了这个定位：VLIW 165/165 满打包、C++ 双缓冲 5 种方案全败、Fragment 寄存器约束封死手动加载路径。这 6% 的差距精确定位在 5 个地址计算冗余的 VLIW word 上——muBLAS 用单条 BIT.LUT.XOR 手写汇编指令消掉了这 5 个 word，而 C++ 编译器面对动态缓冲区索引时做不到这种程度的优化。

这 6% 差距不在代码里，在工具链里。有两条路可以跨过去，任意一条都能让这份 team03_gemm_work.mu 不改一行代码就过 78.5T：

修 §19 的 machine-scheduler + internal-regs-ra 的 WMMA 正确性冲突——在调度器的寄存器替换决策里把 MMA.* / WMMA.* 操作数从 Internal bypass 替换候选集里排除（具体实现路径请摩尔线程编译器团队按真实微架构语义确定）。修好之后 --mtgpu-enable-internal-regs-ra 可以安全启用，73 → 79 TFLOPS，一步跨过 muBLAS
核实 §20 的 /usr/local/musa/include/crt/mma.h:761 的 __frag_base<float, 32> 是否可以改成 __frag_base<float, 8>——如果 32 这个数字背后没有设计意图，改成 8 就可以让 512 线程 + 16 累加器架构的寄存器预算恢复、打开汇编级双缓冲的可能性

以及三条 API 改进建议：

load_matrix_sync 增加 swizzle 参数——load_matrix_sync(frag, ptr, stride, swizzle_bits)，让 XOR swizzle 在函数内部完成，PAD=0 双缓冲因此可行，预期再 +7%
inline asm 寄存器约束——支持 "r"（通用寄存器）、"s"（共享寄存器）约束，解锁手写汇编路径
发布 MTGPU ISA 参考手册——指令延迟/吞吐量表、VLIW 打包规则、Fence 编码语义、寄存器 bank 结构。当前这些数字全部靠微基准反推，发布官方手册后任何人都可以在 1 天内建立精确性能模型

更远一点的突破要等 MTX 1.0 低级 IR（让手写汇编路径可行、可以直接用单条 BIT.LUT.XOR 实现双缓冲基地址切换）、MUTLASS 成熟（需要 MCC 4.3.4+，当前 3.1.0 不支持），或者 MP31+ 下一代硬件（更大 smem 让双缓冲 PAD=0 成为默认配置）。

6 天之前我还以为自己写的 GEMM 比 muBLAS 快。6 天之后我知道：那个"快"是 1/4 工作量 × 4× 公式放大造出来的幻觉，真实性能是触到当前工具链抽象层的极限。对于正在用国产 GPU 写高性能代码的开发者，这 6 天里我最想告诉你的只有一条——在陌生硬件上，先质疑"warp 是 32 线程"这样的常识再写代码。NVIDIA 的铁律不是所有 GPU 的铁律；硬件的执行粒度比你想象的更具体、更不可绕开。

致谢：感谢摩尔线程 GEMM 矩阵挑战赛提供 MTT S4000 硬件平台与评测环境。期待 MTX 1.0 发布与 MUSA 生态继续迭代。

MTT S4000 FP16 GEMM 优化实战:逼近 mcc 编译器的物理上限

1. 起点：摩尔线程的题目和我手里的两份示例​

2. 第 1 版：64×64×32 + 4 warp（~5 TFLOPS）​

3. 第 2 版：128×128×32 + 512 线程（~25 TFLOPS）​

4. 第 3 版：手写双缓冲流水线（~50 TFLOPS）​

5. 第 4 版：256×256×32 Mega Tile（~70 TFLOPS）​

6. 第 5 版：MP22 原生 32×32×16 WMMA，提交 ~90 TFLOPS​

7. v69-v78 的微调实验（性能在 88-90T 之间波动）​

8. 加测 8k 正确性那天：崩塌​

9. 重建之前先立规矩​

10. 排查：五条怀疑全部不命中​

11. 物理 warp = 128 的假设和最小复现​

12. 重建版 0：物理 warp = 128 修复（0 → 22 TFLOPS）​

13. 重建版 1：stride-33 bank conflict 消除（22 → 48 TFLOPS，+118%）​

14. 重建版 2：M-first 蛇形 CTA 排序（48 → 54 TFLOPS）​

15. 重建版 3：256×256×32 大 tile + 1024 线程（54 → 57 TFLOPS）​

16. 重建版 4：postRA + ILP 调度 flag（57 → 67 TFLOPS，+18%）​

17. 重建版 5：FenceSetting=mixed（67 → 73 TFLOPS，+9%）​

18. 重建版 6：load-store-opt 等 4 个新 flag（73 → 73.8 TFLOPS，+1%）​

19. 没走通的路 1：internal-regs-ra 与 machine-scheduler 的 WMMA 正确性冲突​

20. 没走通的路 2：mma.h:761 头文件 bug​

21. 其他踩过的坑汇总（踩坑清单）​

22. 73.8T 是我这个工具链路径的经验上限​

22.1 VLIW 165/165 满打包​

22.2 muBLAS 的完整逆向数据​

22.3 Fragment 寄存器约束（§21 讲过）​

23. 微基准教训：31 cyc vs 403 cyc 的故事​

23.1 关于 DVFS 和 73.8T 的测量条件​

24. 复现指南​

第 0 步（强制）：安装 mcc wrapper​

第 1 步：session 开始的标准三件套​

25. 最后​

1. 起点：摩尔线程的题目和我手里的两份示例

2. 第 1 版：64×64×32 + 4 warp（~5 TFLOPS）

3. 第 2 版：128×128×32 + 512 线程（~25 TFLOPS）

4. 第 3 版：手写双缓冲流水线（~50 TFLOPS）

5. 第 4 版：256×256×32 Mega Tile（~70 TFLOPS）

6. 第 5 版：MP22 原生 32×32×16 WMMA，提交 ~90 TFLOPS

7. v69-v78 的微调实验（性能在 88-90T 之间波动）

8. 加测 8k 正确性那天：崩塌

9. 重建之前先立规矩

10. 排查：五条怀疑全部不命中

11. 物理 warp = 128 的假设和最小复现

12. 重建版 0：物理 warp = 128 修复（0 → 22 TFLOPS）

13. 重建版 1：stride-33 bank conflict 消除（22 → 48 TFLOPS，+118%）

14. 重建版 2：M-first 蛇形 CTA 排序（48 → 54 TFLOPS）

15. 重建版 3：256×256×32 大 tile + 1024 线程（54 → 57 TFLOPS）

16. 重建版 4：postRA + ILP 调度 flag（57 → 67 TFLOPS，+18%）

17. 重建版 5：FenceSetting=mixed（67 → 73 TFLOPS，+9%）

18. 重建版 6：load-store-opt 等 4 个新 flag（73 → 73.8 TFLOPS，+1%）

19. 没走通的路 1：`internal-regs-ra` 与 `machine-scheduler` 的 WMMA 正确性冲突

20. 没走通的路 2：mma.h:761 头文件 bug

21. 其他踩过的坑汇总（踩坑清单）

22. 73.8T 是我这个工具链路径的经验上限

22.1 VLIW 165/165 满打包

22.2 muBLAS 的完整逆向数据

22.3 Fragment 寄存器约束（§21 讲过）

23. 微基准教训：31 cyc vs 403 cyc 的故事

23.1 关于 DVFS 和 73.8T 的测量条件

24. 复现指南

第 0 步（强制）：安装 mcc wrapper

第 1 步：session 开始的标准三件套

25. 最后