metaxRace — MetaX C500 GPU 高性能 Kernel 优化项目

MetaX C500 GPU 上的 CUDA / TileLang 高性能计算 Kernel 仓库。聚焦 DeepSeek 风格 MoE (Mixture of Experts)、FlashMLA 和 Native Sparse Attention (NSA) 等核心算子的 TileLang kernel 实现与性能优化。配套 mcProfiler 性能分析工具链与 mxmaca-performance-tuning-guide 性能优化指南。

1. 项目概览

1.1 目标

本仓库的核心目标是在 MetaX（沐曦）C500 GPU 上，使用 TileLang 与 CUDA 编写并优化大语言模型推理中的关键计算 Kernel，包括：

**MoE (Mixture of Experts)**：DeepSeek 风格的分组 Fused MoE 前向 Kernel
FlashMLA：Multi-head Latent Attention 的 FlashAttention 风格实现
**NSA (Native Sparse Attention)**：原生稀疏注意力机制

1.2 技术栈

组件	版本 / 说明
GPU	MetaX C500（曦云）
MACA SDK	3.7.1.5（沐曦 CUDA 兼容层）
Driver	3.8.30
TileLang	0.1.10（基于 Apache TVM 的 Tile DSL）
PyTorch	2.8.0+metax
mcProfiler	3.8.1.4（性能分析工具）
mxcc	MACA C/C++ 编译器

1.3 学习路线

推荐阅读顺序（对 C500 高性能 Kernel 开发者）：

mcoplib — 沐曦工程师的实际 CUDA kernel 实现（最直接的写码参考）
vLLM-metax/csrc — vLLM 适配 C500 的引擎层代码
FlashMLA — MLA 算子的 TileLang + CUDA 实现
TileKernels-Metax — TileLang kernel 集合
race_tests/ — 竞赛测试用例（MoE / MLA / NSA TileLang 实现）
mxmaca-performance-tuning-guide/ — 性能优化指南（含 microbenchmark）

2. 硬件平台：MetaX C500 架构

2.1 MetaX ↔ NVIDIA ↔ AMD 术语对照

MetaX	NVIDIA	AMD	含义
DPC (Data Process Cluster)	GPC	Shader Engine	顶层计算集群，8 个，每 DPC 含 13 个 AP
AP (Accelerator Processor)	SM	CU	执行 thread block 的主要计算单元，共 104 个
PEU	SM sub-partition	SIMD Lane	AP 内的执行单元，4 个/AP，共 416 个
Wave / Wavefront	Warp	Wavefront	调度执行的基本单位，64 线程（CUDA 为 32）
WSM	Shared Memory	LDS	64 KiB/AP，由 4 个 PEU 共享
MMA	Tensor Core	Matrix Core	矩阵乘加单元
MTE	Load/Store 单元	Load/Store Unit	数据搬运单元
STE	Scalar 单元	Scalar Unit	标量执行单元

2.2 硬件层次结构

GPU
 └─ 8 × DPC
     └─ 每个 DPC 包含 13 × AP
         ├─ Wave 调度器（4 组，对应 4 个 PEU）
         ├─ MMA 矩阵计算单元
         ├─ MTE 数据搬运单元
         ├─ STE 标量执行单元
         └─ 存储层次：
             ├─ 向量寄存器堆：512 KiB/AP（2048 × 32-bit/lane）
             ├─ 标量寄存器堆：800 × 32-bit/AP
             ├─ WSM (Shared Memory)：64 KiB/AP
             ├─ VL1 (Vector L1 Cache)：32 KiB/AP（默认关闭）
             ├─ L2 Cache：8 MB（全芯片共享）
             └─ 全局内存 (DRAM)：64 GB，峰值带宽 1843 GB/s

2.3 C500 关键参数速记

参数	值
AP 数量	104 (8 DPC × 13 AP)
PEU 数量	416 (4 PEU/AP)
Wavefront 大小	64 线程
WSM 容量	64 KiB/AP
每 AP 最大 wave 数	32 (2048 线程)
每 AP 最大 block 数	32
每 block 最大线程数	1024
每线程最大向量寄存器	256 × 32-bit
VL1 容量	32 KiB/AP（默认关闭）
L2 容量	8 MB
DRAM 峰值带宽	1843 GB/s
Shared Memory 带宽	128 B/cycle/AP
Shared Memory 延迟	~64 cycles
L2 延迟	~170 cycles
DRAM 延迟	~420 cycles
FP16×FP16→FP32 MMA 粒度	Wave-level M16×N16×K16

2.4 存储层次与数据局部性

C500 的存储子系统采用三层架构，类似认知科学的三层记忆模型：

工作记忆区（寄存器堆）
    ↕ 全双工
短期记忆区（WSM / VL1 / L2 Cache）
    ↕ 半双工
长期记忆区（DRAM 全局内存）

关键特性：

全双工通路：L2↔VL1、VL1↔寄存器、WSM↔寄存器之间可同时双向传输
半双工通路：DRAM↔L2 一个时刻只能单向传输
VL1 默认关闭：C500 某些默认配置关闭 VL1，低 hit rate 不能单独证明 locality 差
Coalescing 模型：64 线程 wave → 4 个 16 线程合并组 → 地址聚合为 128B 对齐的 transaction

2.5 指令发射单元

每周期最多发射 5 条指令（来自不同 warp）：

1 条标量 ALU 或标量访存（STE/LDU）
1 条 MMA 矩阵指令或向量 ALU 计算指令（MTE）
1 条向量全局访存指令（Global Load/Store/Atomics）
1 条共享内存访存指令（BSM）
1 条 MISC 指令（BRA 等）

3. 项目结构

/data/metaxRace/
├── README.md                          ← 本文档
├── init_tilelang.sh                   ← TileLang 安装/注册脚本
├── skills/                            ← 技能目录
│   ├── install_skills.sh
│   └── mcprof-report-skill/           ← mcProfiler 报告分析技能
│       ├── SKILL.md
│       ├── metax-kernel-guide.md      ← C500 Kernel 开发指南
│       └── tests/conformance/         ← 一致性测试
├── doc/                               ← 文档目录
│   ├── record.md                      ← 术语对照与架构速查
│   ├── metax_c500_arch.md             ← C500 架构说明
│   ├── image.png
│   ├── GPU扫盲/
│   │   └── gemm.md                    ← GEMM Tile 切分基础
│   ├── problem_description/           ← OJ 题目描述
│   │   ├── oj_requirement.md          ← 评测系统规范（CUDA/Triton/TileLang）
│   │   └── moe.md                     ← MoE 题目描述
│   ├── mxmaca-performance-tuning-guide/ ← 性能优化指南
│   │   ├── README.md
│   │   ├── guide/                     ← 10 章优化文档
│   │   │   ├── ch1.初探异构编程.vectoradd.md
│   │   │   ├── ch2.曦云C500芯片架构.md
│   │   │   ├── ch3.Kernel编程入门.reduction.md
│   │   │   ├── ch4.Kernel性能建模.sgemm.md
│   │   │   ├── ch5.Kernel性能优化技巧.md
│   │   │   ├── ch6.Kernel性能分析工具.md
│   │   │   ├── ch7.Host代码性能优化.md
│   │   │   ├── ch8.张量编程.hgemm.md
│   │   │   ├── ch9.C500_HW_limitation.md
│   │   │   └── ch10.常用Compiler参数和Driver环境变量.md
│   │   ├── case/                      ← 文档对应示例代码
│   │   │   ├── 1.vector_add/          ← 向量加法
│   │   │   ├── 2.reduction/           ← 归约
│   │   │   ├── 3.sgemv/               ← 矩阵-向量乘
│   │   │   └── 4.sgemm_tt/            ← 矩阵乘法 (18 个版本)
│   │   ├── microbenchmark/            ← 微架构参数测试
│   │   └── requirements.txt
│   ├── kernel-optimize/               ← Kernel 优化笔记
│   │   └── moe/
│   │       ├── moe.md                 ← MoE Kernel 全面优化记录
│   │       ├── analyze.md             ← 流水线与卡顿分析
│   │       ├── plan.md                ← 优化计划
│   │       ├── case.md                ← 案例研究
│   │       ├── fused.md               ← 融合分析
│   │       ├── kernel.md              ← Kernel 细节
│   │       ├── analyze-data-reuse.md  ← 数据复用分析
│   │       ├── optimize.md            ← 优化分析整理版
│   │       └── optimize/
│   │           ├── HANDOFF.md
│   │           ├── bsm_instruction_num.md
│   │           ├── plan.md
│   │           ├── state.md
│   │           └── tl_thread.md
│   └── skill/
│       └── review.md                  ← 代码评审指南
├── example/                           ← 参考实现
│   ├── README.md                      ← 学习路线指南
│   ├── FlashMLA/                      ← FlashMLA 参考实现
│   │   ├── README.md
│   │   ├── setup.py
│   │   ├── flash_mla/
│   │   ├── csrc/                      ← CUDA 源码
│   │   └── tests/
│   ├── mcoplib/                       ← 沐曦算子库源码
│   │   ├── README.md
│   │   ├── kernel/                    ← CUDA kernel
│   │   ├── op/                        ← 算子实现 (CUDA)
│   │   └── benchmark/                 ← 性能测试
│   │       ├── config/                ← 130 个算子 benchmark 配置
│   │       └── runners/               ← 131 个 benchmark runner
│   ├── kernels/                       ← 各种 kernel 参考
│   │   ├── FlashMLA/                  ← FlashMLA 参考
│   │   ├── TileKernels-Metax/         ← TileLang kernel 集
│   │   ├── TileOPs-Metax/             ← TileLang ops
│   │   ├── mcoplib/                   ← mcoplib 符号链接
│   │   └── vLLM-metax/                ← vLLM 适配代码
│   └── mcoplib/                       ← 沐曦官方算子库（完整版）
│       ├── CMakeLists.txt
│       ├── setup.py / pyproject.toml
│       ├── include/                   ← 头文件
│       ├── kernel/                    ← CUDA kernel (14 个文件)
│       ├── op/                        ← 算子入口 (31 个 CUDA 文件)
│       ├── benchmark/                 ← 性能测试框架
│       └── unit_test/
├── race_tests/                        ← 竞赛测试用例
│   ├── moe/                           ← MoE Fused Kernel
│   │   ├── README.md                  ← 使用说明
│   │   ├── doc.md                     ← 算法拆解
│   │   ├── custom_fusedmoe.py         ← TileLang kernel（主入口）
│   │   ├── ref_fusedmoe.py            ← PyTorch 参考实现
│   │   ├── fusedmoe_benchmark.py      ← host 侧编排
│   │   ├── moe_test_configs.json      ← 测试配置
│   │   ├── moe_prof_driver.py         ← Profiling 被测进程
│   │   ├── profile_mcprof.sh          ← 一键 profiling 脚本
│   │   ├── mcprof_summary.py          ← Profiler 报告摘要解析
│   │   ├── run.sh / run_min.sh        ← 运行脚本
│   │   └── submission*/               ← 各版本提交
│   ├── mla/                           ← Multi-head Latent Attention
│   │   ├── test_tilelang_mla.py       ← TileLang MLA Kernel
│   │   └── test_cases_mla_batch_ctx.json
│   └── nsa/                           ← Native Sparse Attention
│       ├── test_tilelang_nsa_fwd.py   ← TileLang NSA Kernel
│       ├── test_cases_nsa_fwd.json
│       └── reference.py               ← PyTorch 参考实现

4. OJ 提交评测系统

4.1 评测流程

OJ 平台支持三种语言提交：CUDA（C/C++ 源码）、Triton、TileLang。

单测试点执行步骤：

生成测试数据：按题目规格生成多组随机张量
预热：连续执行若干次（不计时），稳定 JIT/缓存/调度
测速：连续执行若干次，GPU 端 cupti 测时取平均
校验：输出与 PyTorch baseline allclose 对比
基线计时：跑一次 PyTorch baseline

4.2 评分公式

$S(T_k)=\dfrac{100}{1+\left(\dfrac{1}{s}-1\right)\dfrac{T_k-T_h}{T_b-T_h}}$

$s = 0.5$ ：PyTorch baseline 对应 50 分
$T_{h}$ ：基于 FLOPs 与带宽估算的理论下限
$T_{b}$ ：PyTorch baseline 平均时间
$S > 100$ 时按对数压缩展示，上限 150 分

4.3 TileLang 提交规范

入口函数必须为 run_kernel（无装饰器，按 run_kernel(*args) 调用）
允许的根包：torch、tilelang（含 tilelang.language、tilelang.intrinsics）、math
内核装饰器：@tilelang.jit 包装 + @T.prim_func 内层
参数注解：T.Tensor((M, K), T.float16) 或 T.Tensor[(M, K), T.float16]
禁用：os、sys、subprocess、socket、pickle、exec/eval/compile
PyTorch 接口白名单制，不在白名单的调 TorchProxyError

5. MoE Kernel：核心攻关项目

5.1 题目描述

实现 DeepSeek 风格 MoE 中的 pre-routed fused expert GEMM。

计算链条：

for each expert e:
    gate_logits = stacked_expert_tokens @ gate_w[e]^T    [d_hidden → d_expert]
    up_logits   = stacked_expert_tokens @ up_w[e]^T      [d_hidden → d_expert]
    hidden      = SiLU(gate_logits) * up_logits           [d_expert]
    output      = hidden @ down_w[e]^T                    [d_expert → d_hidden]
    output     *= routed_expert_weights                   [逐行缩放]

5.2 测试用例

ID	n_experts	d_hidden	d_expert	group_sum	padded/M	warmup/iters
cfg0	16	2048	8192	2272	3072/24	5/30
cfg1	32	7168	2048	4544	6144/48	5/20
cfg2	64	7168	2048	9088	12288/96	5/20

重要：routed_expert_weights 为 torch.float32（非 float16），需在 TileLang kernel 中单独声明 dtype。

5.3 两阶段 Kernel 结构

设计上采用 2 次 launch（相对朴素 5 次 launch 的融合）：

Stage 1：gate/up 双 GEMM + SiLU 融合

grid: (M, ceil(d_expert/128))
K 维 = d_hidden, T.Pipelined 循环
每轮: 读 input[128,128] + gate_w[128,128] + up_w[128,128]
      两次 T.gemm, fp32 累加
Epilogue: SiLU(gate) * up → 写回 up_logits

Stage 2：down GEMM + 权重缩放融合

grid: (M, ceil(d_hidden/128))
K 维 = d_expert, T.Pipelined 循环
每轮: 读 up_logits[128,128] + down_w[128,128]
Epilogue: output *= routed_expert_weights[m]

5.4 当前最优参数

参数	值	说明
block_token	128	token 维度 tile
block_dhidden	128	d_hidden 维度 tile
block_dexpert	128	d_expert 维度 tile
threads	512	每 block 线程数
num_stages	1	Pipelined 级数
无 split	—	split-K 实测负收益

5.5 性能现状（cfg0）

配置	ms/call	分数	相对 baseline
小 tile 64/64/ns2/t256	16.06	—	—
大 tile 128/128/ns1/t256	11.30	—	-30%
大 tile t512	9.60	—	-15%
大 tile t1024	~12.0	—	-25%
大 tile t512 + split_k=2 (OJ)	9.91	53.16	-4.3%
大 tile t512 (OJ, 最优)	9.502	54.21	—

5.6 职责划分：Host vs Kernel

步骤	位置
Router（softmax + top-k）	PyTorch（`MoEGate`）
分组排序、padding	PyTorch（host 侧）
gate GEMM、up GEMM	TileLang Kernel（Stage 1）
SiLU、gate·up 乘法	TileLang Kernel（Stage 1 epilogue）
down GEMM、乘专家权重	TileLang Kernel（Stage 2 epilogue）
跨专家 Σ 求和（scatter_reduce）	PyTorch

5.7 融合的四个层面

#	融合点	收益
1	gate + up 双 GEMM 合在 Stage 1，共享同一份 `x` 分块	`x` 只读一次；少一次 launch
2	SiLU + gate·up 乘法融进 Stage 1 epilogue	不物化 `gate_logits` 中间张量
3	乘专家权重融进 Stage 2 epilogue	省一轮读写
4	跨专家分组 GEMM	n 次 launch 变 1 次

5.8 瓶颈诊断（mcProfiler 证据）

cfg0 瓶颈：卡在 Stage 2（K=8192 的 64 轮长依赖链）

指标	Stage 1	Stage 2
TotalCycles 占比	52.5%	47.5%
per-wave vls stall	30K（已兑现）	118K（必要等待）
vls_pipeline_stall / AP active	66.5%	74.7%
MMA duty（真实值）	17.8%	10.1%
blocks / AP	1536/104 ≈ 14.8	384/104 ≈ 3.7

cfg1 瓶颈：卡在 Stage 1（K=7168 的 56 轮 + smem 锁死 1 block/AP）

指标	Stage 1	Stage 2
TotalCycles 占比	52.7%	47.3%
per-wave vls stall	104K（必要等待）	30K（未验证）
MMA duty	18.4%（全场最高）	9.9%
blocks / AP	768/104 ≈ 7.4	2688/104 ≈ 25.8

关键发现：K 链长的一侧（16 轮 → 30K per-wave；56-64 轮 → 104-118K per-wave）构成”结构性必要等待”——所有并行度杠杆（split-K / 小 tile / t1024 / 双缓冲）均实测无效或负收益。

5.9 关于”融合”的澄清

不是一次 launch 的终极融合，而是两次连续 launch 的”两段式融合”。

“gate/up 产出 → down 消费”之间必须经全局内存——因为一个 128-token tile 的 up_logits 是 128×2048×2B=512KB，远超 C500 的 64 KiB WSM。这是所有主流实现都采用的结构。

6. MLA / NSA 算子

6.1 FlashMLA（Multi-head Latent Attention）

文件：race_tests/mla/test_tilelang_mla.py

TileLang 实现的 Flash Attention 风格 MLA kernel，支持 split 与 no-split 两种模式。

核心特性：

Q/KV 分离的 RoPE 位置编码处理（Q_pe / K_pe）
支持 split-KV（num_split > 1）与 combine kernel
GQA (Grouped Query Attention) 支持
T.use_swizzle(10) 改善 L2 命中

参考实现：example/FlashMLA/（含 CUDA 源码与 PyTorch 接口）

6.2 NSA（Native Sparse Attention）

文件：race_tests/nsa/test_tilelang_nsa_fwd.py、race_tests/nsa/reference.py

基于 block-sparse 的原生稀疏注意力机制。

核心特性：

每个 query token 选择 S 个 key block 计算注意力
支持 causal masking
Grouped Query Attention（groups = HQ // H）
与 PyTorch naive_nsa 参考实现对比验证

当前限制：

Shared memory 需求较大，部分配置会超 64 KiB 上限
大配置跳过参考对比（naive_nsa 是 O(N²) Python 实现）

7. 性能优化指南

7.1 优化策略速查表

优化策略	对计算的好处	对访存的好处
最大化占用率	隐藏流水线延迟	隐藏 DRAM 延迟
线程粗化	平摊创建开销	提高共享内存数据复用
合并全局内存访问	—	减少全局内存请求
向量化访存	—	提高实际带宽
减少共享内存体冲突	—	提高 WSM 带宽
减少控制发散	减少无效线程	增加合并访存
使用编译期常量	减少指令数	增加内存级并行
使用 `__restrict__`	减少指令数	增加内存级并行
Warp 特化	避免加载阻塞计算	—
私有化	—	减少原子更新竞争

7.2 优化方法论（数据路径图方法）

分析任意 Kernel 时按以下三步：

判断吞吐瓶颈：检查各执行单元 duty（MMA/MTE/STE/L2C/HBM）
定位延迟站点：识别 stall 构成，找出等待主导的路径段
评估延迟掩盖上限：检查并行度（block 数、wave 数）与串行链长的匹配

典型思考框架：

墙钟时间 ≈ Σ(每 AP 串行 block 数 × 每 block K 循环轮数 × 每轮串行下界)
            - 可被并行掩盖的延迟

7.3 C500 关键优化约束

约束	影响
WSM 64 KiB/AP	决定每 AP 并发 block 数上限（大 tile → 1 block/AP）
VL1 默认关闭	L2 直接响应所有 transaction，VL1 hit rate 不参与诊断
Wave 64 线程	Coalescing 组为 4×16 threads → 4×128B transactions
寄存器 256/lane	大 tile + fp32 累加容易 spill
半双工 DRAM↔L2	读写不可同时进行

7.4 优化案例：MoE Kernel 的经验教训

尝试	结果	教训
小 tile + 双缓冲	-36%	放大 MTE 搬运指令
大 tile + t256	-30%	吞吐减半但出现 spill
大 tile + t512	-15%	消除 spill + per-wave stall 减半
大 tile + t1024	-25%	同步开销爆炸
split-K	-4.3%	原子加 + finalize 纯开销
Stage1 A 先搬 shared	-21%	窄访问换通路（VLS→BSM），stall 塌 8.6 倍
Stage1 block_dexpert=128	-5.3%	削 spill，0 栈帧

8. mcProfiler 性能分析工具

8.1 概述

mcProfiler 3.8.1.4 是 MetaX C500 配套的性能分析工具，基于 Client-Server 架构。

组件	路径	作用
客户端	`/usr/local/bin/mcProfiler`	CLI（Python Flask 客户端）
服务端	`/opt/mcProfiler-ubuntu18.04/profiler_server`	监听 `127.0.0.1:50123`
GUI	`/opt/mcProfiler-ubuntu18.04/gui-profiler-0.1.0.AppImage`	图形界面
配置文件	`/opt/mcProfiler-ubuntu18.04/config/*.pcd`	指标定义

8.2 一键 Profiling 流程

cd /data/metaxRace/race_tests/moe
./profile_mcprof.sh           # 使用默认配置
MOE_PROF_CFG=1 ./profile_mcprof.sh   # 使用 cfg1
PER_KERNEL=1 ./profile_mcprof.sh 0   # 逐 kernel detail

8.3 关键口径修正（重要）

口径	说明
AP busy Duty	分母为 Total Cycles（含 host），被稀释
单元 duty（MMA/MTE/STE/L2C）	分母为 AP active cycles（kernel 窗口），真实值可直接读
计数类（指令数/字节/hit rate）	不受影响，单次直接读
stall 对比必须按 wave 归一	wave 数变化的 A/B 中 stall 占比会因争用假象失真

8.4 常用指标

HEAD: Summary          — Total Instructions/Cycles, AP busy Duty, RoofLine
CE Statistics          — WORKGROUPS, WAVES, Average Wave life cycles
ISU Statistics         — stall cycles layout, DPC comparison
Memory Statistics      — Global/Private/Shared 指令数, L2C/VL1 Hit Rate
Occupancy              — Achieved/Dispatched waves
GPU Throughput         — AP/MMA/MTE/STE/VLS/L2C Duty ratio
Compute workload       — IPC, instruction throughput/efficiency

8.5 已知问题

config/ 按 CWD 解析：必须在 /opt/mcProfiler-ubuntu18.04 下运行 mcProfiler
报告异步生成：perf done 后约 5-6 分钟才生成 report.txt.json
--headnames 900s+ timeout：不做无人值守采集
full-event = 9 batches：任务重跑 9 遍，每遍换一组 perf event
**--single-pass**：不稳定（出现过 108050% duty 异常）

9. 工程文件清单

9.1 竞赛测试

文件	职责
`race_tests/moe/custom_fusedmoe.py`	TileLang Fused MoE Kernel（主入口）
`race_tests/moe/ref_fusedmoe.py`	PyTorch 参考实现
`race_tests/moe/fusedmoe_benchmark.py`	host 侧编排：router、分组、测试
`race_tests/moe/moe_test_configs.json`	功能/性能测试配置
`race_tests/moe/moe_prof_driver.py`	Profiling 被测进程
`race_tests/moe/profile_mcprof.sh`	一键 profiling 编排
`race_tests/moe/mcprof_summary.py`	Profiler 报告摘要解析
`race_tests/mla/test_tilelang_mla.py`	FlashMLA TileLang kernel
`race_tests/nsa/test_tilelang_nsa_fwd.py`	NSA TileLang kernel
`race_tests/nsa/reference.py`	NSA PyTorch 参考实现

9.2 参考实现与算子库

路径	内容
`example/FlashMLA/csrc/`	FlashMLA CUDA 源码（含 kernel/dispatch/run）
`example/mcoplib/kernel/`	14 个 CUDA kernel（MoE gate/RoPE/RMSNorm 等）
`example/mcoplib/op/`	31 个算子入口
`example/mcoplib/benchmark/`	130+ 算子 benchmark 配置
`example/kernels/`	各类 kernel 参考（TileLang CUDA）

9.3 性能分析工具与脚本

文件	职责
`skills/mcprof-report-skill/SKILL.md`	mcProfiler 报告分析技能
`skills/mcprof-report-skill/metax-kernel-guide.md`	C500 Kernel 开发指南
`skills/mcprof-report-skill/tests/conformance/`	mcProfiler 语义一致性测试
`doc/mxmaca-performance-tuning-guide/guide/`	10 章性能优化文档
`doc/mxmaca-performance-tuning-guide/case/`	文档对应示例代码
`doc/mxmaca-performance-tuning-guide/microbenchmark/`	微架构参数测试

10. 开发计划与状态

10.1 McProfiler 报告分析技能状态

轨道	状态	说明
Track A：确定性工具工程	✅ 完成	manifest/normalization/report parser/通用 collector
Track B：mcProfiler 采样语义	✅ 完成	3.8.1.4 conformance matrix 固化
Track C：C500 诊断知识	⚠️ 部分完成	基础事实、来源分级完成；counter 公式/全卡基线待验证

10.2 仍需验证的项目

项目	优先级	说明
Counter 公式（Duty 分母等）	高	需官方定义或更强对照实验
全卡绝对健康区间	高	当前设备为 25% compute sGPU
JIT 编译/初始化额外 launch	中	不能从 warmup 外推
`--headnames` 900s timeout	中	是否稳定复现
诊断阈值噪声门控	中	需 3 次独立 full-event baseline

10.3 后续顺序

在 full C500 环境重跑 conformance 与微基准
验证 duty/shared/Roofline 公式
建立按 workload family 的统计基线
将高置信度模式写入 diagnosis engine
用真实目标 kernel 做闭环复测

附录 A：MoE 优化 A/B 报告索引

run	配置	说明
`run_cfg0_20260810_025910`	小 tile 64/64/ns2/t256	Baseline
`run_cfg0_20260810_032333`	大 tile 128/128/ns1/t256	-30%
`run_cfg0_20260810_033517`	per-kernel detail (t256)	首次区分 stage1/stage2
`run_cfg0_20260810_041838`	大 tile t512	-15% (当前最优)
`run_cfg0_20260810_050659`	per-kernel detail (t512, 无 split)	确认 stage 占比
`run_cfg1_20260810_055801`	cfg1 per-kernel detail	cfg1 瓶颈分析
OJ split_k=2	大 tile t512 + split	-4.3% (否定)

附录 B：mcProfiler 一键命令参考

# 启动服务端
/opt/mcProfiler-ubuntu18.04/profiler_server &

# 手动 perf_exec
cd /opt/mcProfiler-ubuntu18.04
./mcProfiler perf_exec \
  --cmdline '/opt/conda/bin/python /data/metaxRace/race_tests/moe/moe_prof_driver.py --iters 30 --warmup 3' \
  --casename 'moe_fused_gemm' \
  [--per-kernel] [--counts 50] [--single-pass]

# 查看指标
./mcProfiler show_metrics
./mcProfiler version
./mcProfiler help

附录 C：TileLang 安装

bash /data/metaxRace/init_tilelang.sh

安装脚本支持三级回退：

[A] /opt/tilelang-metax-v0.1.10 已编译 → .pth 注册（秒级）
[B] 源码存在 → pip install -e .（触发编译）
[C] 源码不存在 → git clone + install