bmm-c500-evidence

TileOPs-Metax BmmFwdOp（torch.bmm 3D-3D 严格批 GEMM）在 MetaX C500 上的适配与优化实验证据。

目标仓库：ccf-ai-infra/TileOPs-Metax
冻结基线：392b5d4c3929d9c054dcaba5d5fd0d6268337412 （summer-camp-2026，认领时）
实现基线：PR #1676（fp16/bf16，commit e37033c，作者 Fto）、 PR #1709（fp8）
实验分支：chen-buhao/TileOPs-Metax optimize/bmm，被测提交 3df5670（基于 62570e9）
认领 Issue：https://gitlink.org.cn/ccf-ai-infra/TileOPs-Metax/issues/139
数据快照：snapshot/

结果速览

官方 harness（ManifestBenchmark，CUPTI kernel 计时、10 warmup、 50 rep × 3 trials、L2 flush、clone 输入），当前 C500 切片：

workload	dtype	基线 ms	优化后 ms	加速比
small-b8-128	fp16	0.013000	0.012800	1.02x
small-b8-128	bf16	0.012406	0.012300	1.01x
square-b8-1k	fp16	0.311439	0.245700	1.27x
square-b8-1k	bf16	0.303923	0.246500	1.23x
square-b16-512	fp16	0.098703	0.088100	1.12x
square-b16-512	bf16	0.089605	0.088500	1.01x
square-b32-256	fp16	0.042102	0.038500	1.09x
square-b32-256	bf16	0.043638	0.039100	1.12x
square-b4-4k	bf16	7.269555	6.090100	1.19x
square-b8-2k	fp16	1.892654	1.612100	1.17x
square-b8-2k	bf16	1.865585	1.613500	1.16x
mha-decode-b64-qk	fp16	0.124370	0.129600	0.96x
mha-decode-b64-qk	bf16	0.130683	0.128300	1.02x
mha-decode-b64-pv	fp16	0.135619	0.104700	1.30x
mha-decode-b64-pv	bf16	0.135501	0.104700	1.29x
moe-prefill-b128	bf16	1.902582	1.540500	1.24x

几何平均 1.132x，算术平均 1.137x，最大 1.29x，最小 0.96x。
两次独立全量运行趋势一致；mha-decode-b64-qk 在 parity 附近波动（本次 fp16 0.96x / bf16 1.02x），是 write-dominated 边界 shape。
详细表、TFLOPS、mcProfiler、环境与方法见 results/benchmark_summary.md。

关键结论

C500 上 BmmFwdOp 不受 arch gate 限制：capability (8,0) → sm=80，在 BmmKernel.supported_archs=[80,89,90] 内；被挡住的是 BmmFp8Kernel。
现有 BmmKernel 的默认 schedule（64x64x64 / ns2 / 128t / c_smem）面向 Hopper；C500 的 64 KB shared memory 上限使 ns≥3 的大 tile 不可行。换成 MACA 调度（ns0、直写 epilogue、128x128x64/256t 或 64x64x64/128t）后取得上述收益。
mcProfiler：square 的全局读 1.310 → 1.119 GB、MTE duty 30.7% → 16.6%； decode 流量不变、MTE duty 30.7% → 19.4%，均无尾波。
TileLang MACA 的 dense GEMM 后端约为 cuBLAS 的 0.3–0.5x（独立 GEMM 对照，测于上一块切片），本优化已接近该上限；不承诺 cuBLAS parity。

README.md                       本文件
LICENSE                         MIT
docs/EXPERIMENT.md              环境、设计、方法、结果、限制
docs/REPRODUCE.md               自包含复现步骤（clone + mktemp）
docs/CLAIM_ISSUE.md             认领 Issue 正文（快照）
evidence/baseline/              基线 16-workload 原始输出
evidence/optimized/             优化后 16-workload 原始日志（含 torch）
evidence/mcprofiler/            优化前后 mcProfiler report.txt + json
evidence/archive_oldslice/      2026-10-24 旧切片原始日志（历史对照）
evidence/tests/                 pytest / validate_manifest / K-tail 日志
results/benchmark_summary.md    最终对比表、TFLOPS、profiler、测试
results/profile_report.md       早期 profile / Go-No-Go 报告（旧切片）
results/review_todos.md         review 记录（三项已修复）
results/competition_audit/      2026-10-24 竞争审计原始 JSON
scripts/                        实验与复现脚本
snapshot/                       代码快照（被测提交 3df5670）

环境

MetaX C500，sGPU 切片：Compute 25%，Vram Quota 16000 MiB （被测时 0000:10:00.0，minor 006）
驱动 3.8.30，MACA 3.7.1.5，BIOS 1.33.5.0
Python 3.12.11，PyTorch 2.8.0+metax3.7.1.3
TileLang /opt/tilelang-metax-v0.1.10，版本 0.1.10+cuda.gitf549117c， determine_target('auto') == maca
设备能力 (8, 0)；get_sm_version() == 80

所有加速比都是同一 sGPU 切片、同一 session 内的 TileOps vs 基线比值；绝对 TFLOPS 是切片值，不能直接外推到整卡。不同物理切片的小 shape 可差 10–20%，因此只比较同一次运行内的相对值。

固定 commit 链接

认领 Issue 中使用固定 commit 链接；仓库内部文档使用相对路径，避免每次更新文档都要替换 SHA。当前建议钉住的证据 commit 见 results/benchmark_summary.md 顶部与被测提交 3df5670。

bmm-c500-evidence

结果速览

关键结论

目录

环境

固定 commit 链接