forked from metax-maca/op_optimization
修改赛题2fused_moe教程
This commit is contained in:
commit
2e3aa374cd
|
|
@ -0,0 +1,883 @@
|
|||
# Fused MoE 算子入门：从 Benchmark 验证到 XPU-OJ 接口提交
|
||||
|
||||
## 一、教程定位
|
||||
|
||||
本教程是赛题二 **Fused MoE** 任务的 “benchmark 性能基线与 XPU-OJ 提交衔接”模块，主要帮助学员跑通目标算子的 benchmark 脚本，理解原库 API、输入输出结构、性能指标和性能基线结果，并进一步读懂 XPU-OJ 题目包中的接口约定、测试数据、参考输出和精度要求。
|
||||
|
||||
需要特别说明：
|
||||
|
||||
* 本教程不提供可直接提交的标准答案代码。
|
||||
|
||||
* 本教程仅提供冒烟级 starter 示例代码，用于验证环境、语言、提交链路和 `run_kernel(...)` 接口。
|
||||
|
||||
* benchmark 脚本用于建立性能基线，不是最终提交物。
|
||||
|
||||
* XPU-OJ 题包中的 `baseline()` 属于 OJ 后台参考实现，用于生成 `output_ref`，不是选手提交代码。
|
||||
|
||||
* 选手最终需要自行实现 `run_kernel(...)`，并在正确性通过后继续优化性能。
|
||||
|
||||
|
||||
完成本教程后，学员应能够跑通 benchmark 脚本，记录性能基线结果，读懂 XPU-OJ 题包，理解 OJ 的测试输入与参考实现，并完成一次冒烟级 OJ 提交。
|
||||
|
||||
## 二、学习目标
|
||||
|
||||
完成本模块后，你将能够：
|
||||
|
||||
1. 理解 Fused MoE 推理算子的基本作用、输入输出和典型应用场景；
|
||||
|
||||
2. 跑通对应 benchmark 脚本，并记录性能基线结果；
|
||||
|
||||
3. 学习如何基于 Trition 与 MXMACA C++ 编写 Fused MOE 算子；
|
||||
|
||||
4. 完成数值正确性测试，即验证 reference 计算、pybind 计算、Triton 计算这三种方式计算结果是否数值完全一致。
|
||||
|
||||
* reference：基于 PyTorch 架构在 CPU 上运行的**数值基准**实现；
|
||||
|
||||
* pybind：将 MXMACA C++ 算子编译并封装为 Python 可调用的动态库，**实现复杂且迁移成本高**；
|
||||
|
||||
* Triton：基于 Python 编写的高效 GPU Kernel，可利用 Agent 自动调优，**开发效率高、易于迁移**；
|
||||
|
||||
* 要求 pybind 和 Triton 结果均与 reference 一致，鼓励参赛选手持续调优 Triton ，使其性能逼近甚至超越 pybind 性能。
|
||||
|
||||
5. 区分 benchmark 性能基线、OJ 参考实现和选手提交代码；
|
||||
|
||||
6. 读懂对应 XPU-OJ 题包中的题目描述、接口约定、数据范围和精度要求；
|
||||
|
||||
7. 完成一次冒烟级 `run_kernel(...)` 提交，确认 OJ 链路、语言环境和接口调用正常；
|
||||
|
||||
8. 使用 AI Agent 辅助阅读题包、生成初版实现、定位错误并规划性能优化方向。
|
||||
|
||||
|
||||
## 三、适用对象
|
||||
|
||||
**本模块适合以下人员：**
|
||||
|
||||
* 参与基于 AI Agent 开发范式的国产 GPU 大模型推理算子库优化比赛的学生；
|
||||
|
||||
* 对 GPU 推理算子性能优化感兴趣的开发者；
|
||||
|
||||
* 需要了解 Fused MoE 推理性能的研究人员。
|
||||
|
||||
|
||||
**学习本模块前，需掌握以下基础知识：**
|
||||
|
||||
* Python、C++ 编程基础；
|
||||
|
||||
* PyTorch 基础；
|
||||
|
||||
* GPU 推理基本概念。
|
||||
|
||||
|
||||
## 四、前置准备
|
||||
|
||||
**开始实战前，请确认你已经完成以下准备：**
|
||||
|
||||
### 4.1 环境准备
|
||||
|
||||
* 已获取 GPU 资源并进入赛事专属镜像环境。
|
||||
|
||||
|
||||
若未完成可参考：[基于AI Agent开发范式的国产GPU大模型推理算子库优化/模力方舟Agent部署准备教程.md-MetaX-MACA/揭榜挂帅-沐曦赛题](https://www.gitlink.org.cn/metax-maca/op_optimization/tree/master/%E5%9F%BA%E4%BA%8EAI%20Agent%E5%BC%80%E5%8F%91%E8%8C%83%E5%BC%8F%E7%9A%84%E5%9B%BD%E4%BA%A7GPU%E5%A4%A7%E6%A8%A1%E5%9E%8B%E6%8E%A8%E7%90%86%E7%AE%97%E5%AD%90%E5%BA%93%E4%BC%98%E5%8C%96%2F%E6%A8%A1%E5%8A%9B%E6%96%B9%E8%88%9FAgent%E9%83%A8%E7%BD%B2%E5%87%86%E5%A4%87%E6%95%99%E7%A8%8B.md)
|
||||
|
||||
### 4.2 工具准备
|
||||
|
||||
* 已准备 Agent 工具；
|
||||
|
||||
* 已配置 Token / API Key；
|
||||
|
||||
* 已确认 Agent 可以正常调用模型。
|
||||
|
||||
|
||||
若未完成可参考：[基于AI Agent开发范式的国产GPU大模型推理算子库优化/模力方舟Agent部署准备教程.md-MetaX-MACA/揭榜挂帅-沐曦赛题](https://www.gitlink.org.cn/metax-maca/op_optimization/tree/master/%E5%9F%BA%E4%BA%8EAI%20Agent%E5%BC%80%E5%8F%91%E8%8C%83%E5%BC%8F%E7%9A%84%E5%9B%BD%E4%BA%A7GPU%E5%A4%A7%E6%A8%A1%E5%9E%8B%E6%8E%A8%E7%90%86%E7%AE%97%E5%AD%90%E5%BA%93%E4%BC%98%E5%8C%96%2F%E6%A8%A1%E5%8A%9B%E6%96%B9%E8%88%9FAgent%E9%83%A8%E7%BD%B2%E5%87%86%E5%A4%87%E6%95%99%E7%A8%8B.md)
|
||||
|
||||
### 4.3 代码准备
|
||||
|
||||
* 已获取 Fused MoE 源码，包括 benchmark 测试脚本和供参考的冒烟代码。
|
||||
|
||||
|
||||
源码目录：[MetaX-MACA/揭榜挂帅-沐曦赛题 | GitLink](https://www.gitlink.org.cn/metax-maca/op_optimization/tree/master/%E5%9F%BA%E4%BA%8EAI%20Agent%E5%BC%80%E5%8F%91%E8%8C%83%E5%BC%8F%E7%9A%84%E5%9B%BD%E4%BA%A7GPU%E5%A4%A7%E6%A8%A1%E5%9E%8B%E6%8E%A8%E7%90%86%E7%AE%97%E5%AD%90%E5%BA%93%E4%BC%98%E5%8C%96/operator_task_package/fused_moe_task_package)
|
||||
|
||||
### 4.4 账号准备
|
||||
|
||||
* 已获取 XPU-OJ 账号。
|
||||
|
||||
|
||||
XPU-OJ 账号由组委会统一发放，参赛者无需自行注册。
|
||||
|
||||
如果登录后看不到题目，请联系助教或赛事运营确认账号是否已加入对应比赛 / 用户组。
|
||||
|
||||
## 五、知识预备
|
||||
|
||||
### 5.1 混合专家模型 MoE
|
||||
|
||||
参考链接：[https://huggingface.co/blog/zh/moe](https://huggingface.co/blog/zh/moe)
|
||||
|
||||
混合专家模型（Mixed Expert Models，简称 MoE）是一种稀疏激活的模型结构。与稠密模型不同，MoE 在前向计算时只激活部分参数，从而在相近的计算预算下，获得更大的模型容量和更快的收敛速度。这也是当前大模型扩展参数规模的主流路径之一。
|
||||
|
||||
MoE 的主要优势体现在预训练和训练阶段的算力效率上，但在推理阶段带来了新的**挑战**：
|
||||
|
||||
* **显存占用高**：尽管每个 Token 只激活部分专家，但所有专家的参数仍需常驻显存。例如 Mixtral 8×7B 的实际参数量接近 47B，而非 8×7B 的简单叠加，因为除 FFN 外的大多数参数在各专家间是共享的。
|
||||
|
||||
* **计算不均衡**：不同专家接收的 Token 数量可能不同，导致计算负载不均，容易形成局部瓶颈。
|
||||
|
||||
* **访存压力大**：专家权重、Token 路由、Permute / Unpermute 都会带来额外的显存读写，而不仅仅是计算量的减少。
|
||||
|
||||
|
||||
因此，在推理阶段对 MoE 算子进行系统级优化具有非常重要的意义，在不改变模型行为和数值精度的前提下，通过更合理的调度、融合与访存优化，缓解稀疏性带来的碎片化和内存压力。
|
||||
|
||||
### 5.2 INT8 模型量化
|
||||
|
||||
参考链接：[https://www.cnblogs.com/chentiao/p/18315901](https://www.cnblogs.com/chentiao/p/18315901)
|
||||
|
||||
INT8 量化是一种用于减少模型大小和计算复杂度的方法，特别是在深度学习模型中。它通过将浮点数（通常是 FP32）转换为 8 位整数 (INT 8)，从而减少内存使用和提高计算效率。
|
||||
|
||||
## 六、项目实践1-从 Benchmark 验证到 XPU-OJ
|
||||
|
||||
**项目目标：**先跑通 Fused MoE 算子的 benchmark 脚本，建立性能基线；再基于同一题目语义完成 XPU-OJ 冒烟提交，确认评测环境能够正确调用 `run_kernel(...)`，为后续正确性修复和性能优化打基础。
|
||||
|
||||
### 6.1 在赛事镜像中验证 Fused MoE Benchmark
|
||||
|
||||
#### Step 1：检查运行环境
|
||||
|
||||
**目标：**确认当前环境满足本模块运行要求，包括编译器、MXMACA 工具链及 Python 依赖库。
|
||||
|
||||
**操作：**检查 Python、编译工具、MXMACA 编译器及关键 Python 包（numpy、torch、triton）是否存在。
|
||||
|
||||
**命令示例：**
|
||||
|
||||
```apl
|
||||
python --version # 检查Python版本，Python ≥ 3.8
|
||||
g++ --version # 确认 C++ 编译器存在
|
||||
which mxcc # 确认 MACA 编译器存在
|
||||
|
||||
|
||||
# 检查 Python 依赖
|
||||
python - << 'EOF'
|
||||
import sys
|
||||
deps = ["numpy", "torch", "triton"]
|
||||
missing = []
|
||||
for d in deps:
|
||||
try:
|
||||
__import__(d)
|
||||
except ImportError:
|
||||
missing.append(d)
|
||||
if missing:
|
||||
print(f"[ERROR] Missing packages: {missing}")
|
||||
sys.exit(1)
|
||||
else:
|
||||
print("[OK] numpy, torch, triton are installed.")
|
||||
EOF
|
||||
```
|
||||
|
||||
**预期结果：**
|
||||
|
||||
* Python 3.12.11
|
||||
|
||||
* g++ (Ubuntu 13.3.0-6ubuntu2~24.04.1) 13.3.0
|
||||
|
||||
* /opt/maca/mxgpu\_llvm/bin/mxcc
|
||||
|
||||
* \[OK\] numpy, torch, triton are installed.
|
||||
|
||||
|
||||
**常见问题：**
|
||||
|
||||
| **报错** | **原因** | **解决办法** |
|
||||
| --- | --- | --- |
|
||||
| `g++：command not found` | 未安装 C++ 编译工具 | `apt update && apt install -y build-essential` |
|
||||
| `Python 3.6.x/ Python 3.7.x` | Python 版本过低 | `conda install python=3.12` （推荐3.10+) |
|
||||
| `ModuleNotFoundError: numpy` | 当前 Python 缺少依赖 | `pip install numpy torch triton` |
|
||||
|
||||
#### Step 2：进入项目目录
|
||||
|
||||
**目标：**进入本模块所需的源码目录。
|
||||
|
||||
**操作：**切换到指定项目路径。
|
||||
|
||||
**命令示例：**
|
||||
|
||||
```apl
|
||||
# 克隆代码仓库
|
||||
git clone https://gitlink.org.cn/metax-maca/op_optimization.git
|
||||
# 切换到fused moe目录下benchmark项目
|
||||
cd '.\op_optimization\基于AI Agent开发范式的国产GPU大模型推理算子库优化\operator_task_package\fused_moe_task_package\benchmark'
|
||||
```
|
||||
|
||||
#### Step 3：pybind 编译
|
||||
|
||||
**目标：**将用 C++ 编写的 fused\_moe 算子编译为 Python 可调用的 pybind 模块。
|
||||
|
||||
**操作：**运行 `fused_moe/scripts/build_fused_moe_i8_tn_pybind.sh` 脚本
|
||||
|
||||
**命令示例：**
|
||||
|
||||
```apl
|
||||
bash scripts/build_fused_moe_i8_tn_pybind.sh
|
||||
```
|
||||
|
||||
切换 Python 环境命令示例：
|
||||
|
||||
```apl
|
||||
PYTHON_BIN=/path/to/python bash scripts/build_fused_moe_i8_tn_pybind.sh
|
||||
```
|
||||
|
||||
**预期结果：**
|
||||
|
||||
编译成功无报错，终端显示：
|
||||
|
||||
* \[SUCCESS\] /root/Project/fused\_moe/standalone/fused\_moe\_i8\_tn/build/fused\_moe\_i8\_tn\_ pybind.so
|
||||
|
||||
且成功生成 `fused_moe/standalone/fused_moe_i8_tn/build/fused_moe_i8_tn_pybind.cpython-310-x86_64-linux-gnu.so` 文件
|
||||
|
||||
**常见问题：**
|
||||
|
||||
| **报错** | **原因** | **解决办法** |
|
||||
| --- | --- | --- |
|
||||
| `Python.h: No such file or directory` | Python 头文件路径未找到 | 确认 `PYTHON_BIN` 路径正确，脚本自动探测 `sysconfig.get_path('include')` |
|
||||
| `libpython3.x.so: cannot find` | 链接时找不到 Python 库 | 1、执行 `find $CONDA_PREFIX -name "libpython3*.so*"`查找绝对路径<br>2、将该路径赋值给 `LIBPYTHON_PATH` |
|
||||
| `recompile with -fPIC` | 编译未开启位置无关代码 | 确保 `mxcc`/ `g++`编译参数中有 `-fPIC` |
|
||||
| `permission denied` | 无脚本执行权限 | `chmod +x scripts/*.sh` |
|
||||
| `undefined reference to Py_...` | Python 版本不匹配 | 确认编译脚本中`PYTHON_BIN`路径与当前运行的 Python 环境完全一致 |
|
||||
|
||||
#### Step 4：正确性验证
|
||||
|
||||
**目标：**验证 reference 计算、pybind 计算、Triton 计算这三种方式计算结果的数值是否一致。
|
||||
|
||||
**操作：**运行 `fused_moe/scripts/run_fused_moe_i8_tn_pybind_test.sh` 脚本
|
||||
|
||||
**命令示例：**
|
||||
|
||||
```apl
|
||||
bash scripts/run_fused_moe_i8_tn_pybind_test.sh --backend all # 运行全部计算方式
|
||||
|
||||
# --backend：选择计算方式
|
||||
# 只测 pybind：
|
||||
bash scripts/run_fused_moe_i8_tn_pybind_test.sh --backend pybind
|
||||
# 只测 triton：
|
||||
bash scripts/run_fused_moe_i8_tn_pybind_test.sh --backend triton
|
||||
# 只测 reference:
|
||||
bash scripts/run_fused_moe_i8_tn_pybind_test.sh --backend reference
|
||||
```
|
||||
|
||||
**预期结果：**
|
||||
|
||||
编译成功无报错，输出示例如下：
|
||||
|
||||
> pybind:fused\_moe\_i8\_tn\_topk1 passed: rows=256, cols=128, sample C\[0\]=0.69531, C\[last\]=-0.44531
|
||||
|
||||
> pybind:fused\_moe\_i8\_tn\_topk2 passed: rows=512, cols=128, sample C\[0\]=-0.57813, C\[last\]=-0.49805
|
||||
|
||||
> pybind:fused\_moe\_i8\_tn\_topk3 passed: rows=384, cols=128, sample C\[0\]=-1.08594, C\[last\]=-0.33594
|
||||
|
||||
> reference:fused\_moe\_i8\_tn\_topk1 passed: rows=256, cols=128, sample C\[0\]=0.6934, C\[last\]=-0.4451
|
||||
|
||||
> reference:fused\_moe\_i8\_tn\_topk2 passed: rows=512, cols=128, sample C\[0\]=-0.5768, C\[last\]=-0.4975
|
||||
|
||||
> reference:fused\_moe\_i8\_tn\_topk3 passed: rows=384, cols=128, sample C\[0\]=-1.0875, C\[last\]=-0.3362
|
||||
|
||||
> triton:fused\_moe\_i8\_tn\_topk1 passed: rows=256, cols=128, sample C\[0\]=0.69337, C\[last\]=-0.44513
|
||||
|
||||
> triton:fused\_moe\_i8\_tn\_topk2 passed: rows=512, cols=128, sample C\[0\]=-0.57678, C\[last\]=-0.49749
|
||||
|
||||
> triton:fused\_moe\_i8\_tn\_topk3 passed: rows=384, cols=128, sample C\[0\]=-1.08748, C\[last\]=-0.33618
|
||||
|
||||
**结果解释：**
|
||||
|
||||
* “pybind/reference/Triton”：三种计算方式；
|
||||
|
||||
* “fused\_moe\_i8\_tn\_topk1/2/3 passed”：测试算子通过数值校验，数值误差在允许范围内且无明显异常，否则会报错 FAILED；
|
||||
|
||||
* ”rows=... , cols=...“：输出 Tensor 的形状；
|
||||
|
||||
* ”sample C\[0\]=... , C\[last\]=...“：首尾采样值，用于辅助定位数值偏差，不作为精度判定依据。
|
||||
|
||||
|
||||
**常见问题：**
|
||||
|
||||
| **报错** | **原因** | **解决办法** |
|
||||
| --- | --- | --- |
|
||||
| `ModuleNotFoundError: fused_moe_i8_tn_pybind` | pybind 模块未编译或未加入 `PYTHONPATH` | 回到步骤 3，确认 `.so`已生成；执行 `export PYTHONPATH=/root/Project/fused_moe:$PYTHONPATH` |
|
||||
| `FAILED: max abs diff too large` | 数值误差超过阈值 | 检查 scale 是否应用位置错误；确认 TopK 索引与权重是否一致 |
|
||||
| `FAILED: shape mismatch` | 输出张量形状不一致 | 检查 Token Permute / Unpermute 逻辑；确认 expert 维度对齐 |
|
||||
| `FAILED: NaN or Inf detected` | 溢出或未初始化内存 | 检查 INT8 乘加是否溢出；确认 GEMM 输出是否反量化 |
|
||||
| 终端长时间无输出 | Kernel 死锁或 Launch 失败 | 减小测试 shape；检查是否触发 MACA 硬件限制 |
|
||||
|
||||
#### Step 5：性能测试
|
||||
|
||||
**目标：**输出 benchmark 结果对比表
|
||||
|
||||
**操作：**运行 `fused_moe/scripts/run_fused_moe_i8_tn_benchmark.sh` 脚本
|
||||
|
||||
**命令示例：**
|
||||
|
||||
```apl
|
||||
bash scripts/run_fused_moe_i8_tn_benchmark.sh --backend all --warmup 5 --iters 20
|
||||
# --backend：选择计算方式
|
||||
# --warmup：设置预热次数
|
||||
# --iters：设置迭代次数
|
||||
```
|
||||
|
||||
**预期结果：**
|
||||
|
||||
编译成功无报错，输出示例如下：
|
||||
|
||||
> pybind:fused\_moe\_i8\_tn\_topk1 benchmark: avg\_ms=0.308978, TOPS=0.027149, warmup=5, iters=20
|
||||
|
||||
> pybind:fused\_moe\_i8\_tn\_topk2 benchmark: avg\_ms=0.304500, TOPS=0.055098, warmup=5, iters=20
|
||||
|
||||
> pybind:fused\_moe\_i8\_tn\_topk3 benchmark: avg\_ms=0.297775, TOPS=0.042256, warmup=5, iters=20
|
||||
|
||||
> reference:fused\_moe\_i8\_tn\_topk1 benchmark: avg\_ms=1685.43, TOPS=0.000005, warmup=5, iters=20
|
||||
|
||||
> reference:fused\_moe\_i8\_tn\_topk2 benchmark: avg\_ms=3384.52, TOPS=0.000005, warmup=5, iters=20
|
||||
|
||||
> reference:fused\_moe\_i8\_tn\_topk3 benchmark: avg\_ms=2532.14, TOPS=0.000005, warmup=5, iters=20
|
||||
|
||||
> triton:fused\_moe\_i8\_tn\_topk1 benchmark: avg\_ms=19.013421, TOPS=0.000441, warmup=5, iters=20
|
||||
|
||||
> triton:fused\_moe\_i8\_tn\_topk2 benchmark: avg\_ms=16.745914, TOPS=0.001002, warmup=5, iters=20
|
||||
|
||||
> triton:fused\_moe\_i8\_tn\_topk3 benchmark: avg\_ms=19.630328, TOPS=0.000641, warmup=5, iters=20
|
||||
|
||||
**结果解释：**
|
||||
|
||||
* “pybind/reference/Triton”：三种计算方式；
|
||||
|
||||
* “fused\_moe\_i8\_tn\_topk1/2/3”：分别对应选择前 1 / 2 / 3 个专家场景下的 MoE 算子；
|
||||
|
||||
* “avg\_ms”：平均算子执行耗时（毫秒），这里不计算预热时间，只计算正式迭代的时间；
|
||||
|
||||
* “TOPS”：Tera Operations Per Second，本次 MoE 算子的总运算量 / 实际耗时；
|
||||
|
||||
* “warmup=5, iters=20”：预热轮数和正式迭代数。
|
||||
|
||||
|
||||
**常见错误：**
|
||||
|
||||
| **报错** | **原因** | **解决办法** |
|
||||
| --- | --- | --- |
|
||||
| `ModuleNotFoundError: fused_moe_i8_tn_pybind` | pybind 模块未编译或未加入 `PYTHONPATH` | 回到步骤 3，确认 `.so`已生成；执行 `export PYTHONPATH=/root/Project/fused_moe:$PYTHONPATH` |
|
||||
| 终端长时间无输出 | Kernel 死锁或 MACA 驱动异常 | 减小测试 shape；重启容器或设备 |
|
||||
| avg\_ms 异常抖动（±50%） | 其他进程占用 GPU | 关闭其他占用显存的进程，单机单任务运行 |
|
||||
|
||||
### 6.2 在 XPU-OJ 平台进行提交
|
||||
|
||||
平台链接：[https://xpuoj.com/](https://xpuoj.com/)
|
||||
|
||||
#### Step 6：从 Benchmark 到 XPU-OJ 提交
|
||||
|
||||
**目标：**理解 Benchmark 和 XPU-OJ 在线评测任务的不同，完成从 Benchmark 到 XPU-OJ 提交的转换。
|
||||
|
||||
**操作：**
|
||||
|
||||
1、厘清 Benchmark 和 XPU-OJ 的区别：
|
||||
|
||||
赛事镜像中的 Benchmark 脚本用于理解目标算子的调用方式、输入输出 shape 和性能基线；XPU-OJ 题包用于定义最终评测接口、数据范围、参考输出和精度要求。
|
||||
|
||||
| **维度** | **Benchmark 脚本** | **XPU-OJ 提交** |
|
||||
| --- | --- | --- |
|
||||
| **目的** | 理解算子接口、建立性能基线 | 统一环境下的正确性+性能评测 |
|
||||
| **接口形式** | Python API | 三种接口供选择：<br>* CUDA C（`extern "C"`）<br> <br>* TileLang（Python `@jit`）<br> <br>* Triton（Python `@triton.jit`） |
|
||||
| **函数签名** | `backend_fn(...)` | `run_kernel(...)` |
|
||||
| **数据范围** | 多种 head\_dim / batch\_size / seq\_len 组合 | 固定参数范围（以题包为准） |
|
||||
| **验证** | 无自动正确性校验，人工对比输出数值 | 强制通过 `torch.allclose(rtol=2e-2, atol=5e-3)` |
|
||||
| **输出** | 终端直接输出 | 排行榜得分 |
|
||||
|
||||
2、理解完成 benchmark 验证并成功建立性能基线后，需要完成以下转换：
|
||||
|
||||
1. 从 benchmark 脚本中理解目标 API；
|
||||
|
||||
2. 在对应 OJ 平台【题目描述】中查看 `run_kernel(...)` 接口；
|
||||
|
||||
3. 对照 OJ 平台【题目描述】中的输入 shape、数据范围和精度要求；
|
||||
|
||||
4. 编写自己的 `run_kernel(...)`；
|
||||
|
||||
5. 在 OJ 平台提交`run_kernel(...)`，先通过正确性；
|
||||
|
||||
6. 正确性通过后，再对比 benchmark 耗时 / OJ 耗时继续优化。
|
||||
|
||||
|
||||
#### Step 7：登录 XPU-OJ 平台并进入题目页面
|
||||
|
||||
**目标：**访问 XPU-OJ 评测平台，找到 `fused_moe_i8_tn`算子题目，熟悉题目页面布局。
|
||||
|
||||
**操作：**
|
||||
|
||||
1. 进入 XPU-OJ 平台后使用分配到的账号进行登录
|
||||
|
||||
[](https://www.picgo.net/image/image7.4dk6yw)
|
||||
|
||||
2. 进入比赛页面：点击顶部导航栏【比赛】，选择【进行中】，找到对应比赛进入。
|
||||
|
||||
[](https://www.picgo.net/image/image8.4dkxr6)
|
||||
|
||||
3、进入题目页面：本算子对应比赛题目6：`Fused MoE i8 tn`，点击进入题目页面：
|
||||
|
||||
[](https://www.picgo.net/image/image9.4doS64)
|
||||
|
||||
完成上述步骤可进入如下题目页面：
|
||||
|
||||
* 左侧：题目描述，下滑可查看 CUDA Maca、Triton 和 TileLang 三种语言的接口约定、输入输出格式、示例、数据范围、正确性要求以及提示；
|
||||
|
||||
* 右侧：提交区域，输入编写的`run_kernel(...)`后在下方选择对应的语言即可提交。还可以通过上方导航栏【我的提交】查看历史提交。
|
||||
|
||||
|
||||
[](https://www.picgo.net/image/image10.4doWIj)
|
||||
|
||||
#### Step 8：理解 XPU-OJ 评测接口和精度要求
|
||||
|
||||
**目标：**明确提交代码的接口规范、函数签名及评测判分标准，避免因接口不匹配或理解偏差导致反复提交失败。
|
||||
|
||||
**操作：**
|
||||
|
||||
1. 在题目页面中找到"接口约定"部分，确认你选择的提交语言（CUDA C / TileLang / Triton），仔细阅读对应的函数签名，**确认参数类型和顺序完全一致**。
|
||||
|
||||
2. 对照 Benchmark 的接口（8 个参数、无 `out`），注意 XPU-OJ 的接口**多了一个** `**out**`**参数**（共 9 个），你的代码必须**原地写回结果到** `**out**`，不能只 `return`。
|
||||
|
||||
3. 阅读数据范围与提示部分，记住以下关键约束：
|
||||
|
||||
* `topk`恒为 8，`num_experts`恒为 256
|
||||
|
||||
* `EM = num_tokens × 8`，且 `EM`必须是 128 的倍数
|
||||
|
||||
* `N`、`K`由 case 携带：Gate-up 为 (4096, 7168)，Down 为 (7168, 2048)
|
||||
|
||||
* `b_col_major`布局是 `[expert, n, k]`，不是 `[expert, k, n]`
|
||||
|
||||
* `expert_ids`每 128 行一个 tile：`expert(r) = expert_ids[r // 128]`
|
||||
|
||||
* `a`和 `scale_a`已按 routed row 展开，直接用 `a[r, :]`和 `scale_a[r]`即可
|
||||
|
||||
4. 找到正确性要求中的评测口径，确认精度要求：
|
||||
|
||||
* 容差：`rtol=2e-2, atol=5e-3`
|
||||
|
||||
* 通过率：`matched_ratio >= 0.99`（至少 99% 的元素在容差范围内）
|
||||
|
||||
|
||||
#### Step 9：提交 OJ 冒烟代码
|
||||
|
||||
**目标：**提交冒烟代码，确认 OJ 提交链路、语言环境和 `run_kernel(...)` 接口可用。
|
||||
|
||||
**操作：**
|
||||
|
||||
1. 粘贴代码：在题目页面右侧的提交区域输入编写的`run_kernel(...)` ；
|
||||
|
||||
2. 选择语言：在提交界面语言下拉框中选择对应的开发语言（本任务支持 CUDA Mac、Triton 和 TileLang，教程示例对应 CUDA Maca）；
|
||||
|
||||
3. 执行提交：点击【提交】按钮，系统将自动进入评测队列，出现如下界面；
|
||||
|
||||
4. 等待结果：评测时间与题目测试点数量、队列状态和平台负载有关，通常需要等待数十秒到数分钟。以平台实际返回为准。
|
||||
|
||||
|
||||
[](https://www.picgo.net/image/image11.4do0yf)
|
||||
|
||||
**预期结果：**
|
||||
|
||||
* 提交成功后，系统会自动运行评测程序。
|
||||
|
||||
* 首先进行正确性校验，如果输出结果与参考实现差异超过容差，则标记为 `Wrong Answer`。
|
||||
|
||||
* 正确性通过后，进行性能评测，计算加速比和得分。
|
||||
|
||||
|
||||
#### Step 10：分析评测结果与评分机制
|
||||
|
||||
**目标：**深入理解 OJ 评测报告的各项指标含义，结合官方评分规则（参考图片），分析当前代码的性能瓶颈与得分潜力。
|
||||
|
||||
**操作：**
|
||||
|
||||
1. 查看结果详情：在提交记录中查看状态、总得分、耗时、内存及 SPJ Report（单测试点检查器信息）。
|
||||
|
||||
2. 解读关键指标：
|
||||
|
||||
* Config：测试场景参数（如 batch, seqlen, heads 等）；
|
||||
|
||||
* Baseline：官方基准实现耗时（对应 50分）；
|
||||
|
||||
* User kernel：你的代码实际耗时；
|
||||
|
||||
* Speedup vs base：加速比（Baseline / User kernel）；
|
||||
|
||||
* Score ratio：得分比例（0~1），映射为 0~100 分；
|
||||
|
||||
* Pass：功能正确性（OK 表示通过，FAIL 表示错误）。
|
||||
|
||||
3. 评分规则：
|
||||
|
||||
* 正确性优先：未通过正确性测试或稳定性测试的作品，客观评测得分记为 0 分；
|
||||
|
||||
* OJ 平台对单测试点的评分遵循以下公式：
|
||||
|
||||
|
||||
$S(T\_k) = \frac{100}{1 + \left(\frac{1}{0.5} - 1\right) \cdot \frac{T\_k - T\_h}{T\_b - T\_h}}$
|
||||
|
||||
其中，$T\_k$代表你的 kernel 平均时间；$T\_b$代表 PyTorch baseline 平均时间（对应 50 分）；$T\_h$代表硬件理论下限 。
|
||||
|
||||
#### Step 13：榜单查看与优化方向
|
||||
|
||||
**目标：**掌握榜单查看方法，了解自身排名位置，并制定针对性的代码优化策略。
|
||||
|
||||
**操作：**
|
||||
|
||||
1. 查看榜单：进入题目榜单页面，观察：
|
||||
|
||||
* 总得分：各题目得分总和，决定最终排名。
|
||||
|
||||
* 个人排名：页面顶部显示“我的排名”与“我的总分”。
|
||||
|
||||
* 各题得分：表格中每列对应一个任务（FlashInfer / FlashAttention / MCTLASS Fused MoE），便于横向对比。
|
||||
|
||||
|
||||
[](https://www.picgo.net/image/image12.4doNRi)
|
||||
|
||||
点击【排行榜】进入如下页面：
|
||||
|
||||
[](https://www.picgo.net/image/image13.4doU7W)
|
||||
|
||||
2. 制定优化方向：
|
||||
|
||||
|
||||
| **优化方向** | **具体说明** |
|
||||
| --- | --- |
|
||||
| 算子融合 | 将矩阵乘、scale、softmax 等步骤合并为单个内核，减少 HBM 往返。 |
|
||||
| 并行策略调整 | 在 Decode 阶段采用分块或 Split-K 思路，提升长 KV 序列并行度。 |
|
||||
| 在线 Softmax | 引入局部最大值与局部求和动态缩放，保证数值稳定并减少访存。 |
|
||||
| 显存访问合并 | 保证相邻线程访问相邻地址，按 HeadDim 等连续维度向量化加载。 |
|
||||
| 利用内存层级 | 将频繁更新的标量放入寄存器，块内复用数据放入共享内存。 |
|
||||
| 软件流水线 | 在当前分块计算时预取下一分块数据，隐藏内存加载延迟。 |
|
||||
| 减少冗余计算 | 提取循环不变量，处理变长序列时减少复杂分支。 |
|
||||
|
||||
## 七、项目实践2-Kernel Swift 智能算子迁移系统自动调优
|
||||
|
||||
系统链接：[https://deeplink.org.cn/kernelswift/task](https://deeplink.org.cn/kernelswift/task)
|
||||
|
||||
**项目目标：**基于 KernelSwift 智能算子迁移系统，对 Fused MoE 算子进行在线自动调优。通过输入算子的 PyTorch 代码，一键生成适配沐曦硬件的高性能实现，高效完成算子优化与全流程追踪。
|
||||
|
||||
### 步骤1：复用算子广场的Fused MoE 算子进行二次优化
|
||||
|
||||
**目标：**通过提交算子广场的 fused\_moe 算子代码发起自动优化流程，实现二次优化
|
||||
|
||||
**操作：**
|
||||
|
||||
1. 进入算子广场：点击左侧导航栏 【算子广场】，进入算子列表页
|
||||
|
||||
搜索 fused\_moe 算子，复制 `input_code.py` 代码，也可直接复制以下代码：
|
||||
|
||||
```python
|
||||
import torch
|
||||
import torch.nn as nn
|
||||
import torch.nn.functional as F
|
||||
|
||||
|
||||
class Model(nn.Module):
|
||||
"""
|
||||
Reference PyTorch MoE forward (no fused kernels).
|
||||
Expects inputs:
|
||||
hidden_states: (M, in_size)
|
||||
w1: (E, hidden_size, in_size) where hidden_size = 2 * up_dim
|
||||
w2: (E, out_size, up_dim)
|
||||
topk_weights: (M, top_k)
|
||||
topk_idx: (M, top_k)
|
||||
top_k: int
|
||||
renormalize: bool
|
||||
"""
|
||||
|
||||
def __init__(self):
|
||||
super().__init__()
|
||||
|
||||
def forward(
|
||||
self,
|
||||
hidden_states: torch.Tensor,
|
||||
w1: torch.Tensor,
|
||||
w2: torch.Tensor,
|
||||
topk_weights: torch.Tensor,
|
||||
topk_idx: torch.Tensor,
|
||||
top_k: int,
|
||||
renormalize: bool = True,
|
||||
) -> torch.Tensor:
|
||||
if renormalize:
|
||||
topk_weights = topk_weights / topk_weights.sum(dim=-1, keepdim=True)
|
||||
|
||||
seq_len = hidden_states.size(0)
|
||||
out_size = w2.size(1)
|
||||
output = hidden_states.new_zeros(seq_len, out_size)
|
||||
num_experts = w1.size(0)
|
||||
|
||||
# Accumulate expert contributions
|
||||
for eid in range(num_experts):
|
||||
token_idx, k_idx = torch.where(topk_idx == eid)
|
||||
if token_idx.numel() == 0:
|
||||
continue
|
||||
gate_proj, up_proj = w1[eid].chunk(2, dim=0)
|
||||
down_proj = w2[eid]
|
||||
tmp = F.linear(hidden_states[token_idx], gate_proj)
|
||||
tmp = F.silu(tmp) * F.linear(hidden_states[token_idx], up_proj)
|
||||
tmp = F.linear(tmp, down_proj)
|
||||
tmp = tmp * topk_weights[token_idx, k_idx, None]
|
||||
output.index_add_(0, token_idx, tmp.to(output.dtype))
|
||||
return output
|
||||
|
||||
|
||||
# Hyperparameters
|
||||
seq_len = 128
|
||||
in_size = 128
|
||||
hidden_size = 256 # 2 * up_dim
|
||||
out_size = 128
|
||||
num_experts = 32
|
||||
top_k = 4
|
||||
|
||||
dtype = torch.float16
|
||||
|
||||
def get_inputs():
|
||||
hidden_states = (torch.rand(seq_len, in_size, dtype=dtype) - 0.5) / 2
|
||||
w1 = (torch.rand(num_experts, hidden_size, in_size, dtype=dtype) - 0.5) / 2
|
||||
w2 = (torch.rand(num_experts, out_size, hidden_size//2, dtype=dtype) - 0.5) / 2
|
||||
routing_logits = (torch.rand(seq_len, num_experts, dtype=dtype) - 0.5) / 2
|
||||
routing_weights = torch.softmax(routing_logits, dim=-1, dtype=torch.float32)
|
||||
topk_weights, topk_idx = torch.topk(routing_weights, top_k, dim=-1)
|
||||
return [hidden_states, w1, w2, topk_weights, topk_idx, top_k, True]
|
||||
|
||||
def get_init_inputs():
|
||||
return []
|
||||
```
|
||||
|
||||
2. 进入新建任务页：点击左侧导航栏【新建任务】 ，进入算子提交页面。
|
||||
|
||||
3. 编写算子代码：在 `model.py` 编辑器中输入刚刚复制的 fused\_moe 算子代码。
|
||||
|
||||
如果想自行编写算子代码，需严格遵循标准格式规范：输入代码必须包含 `class Model` 定义算子实现，`get_init_inputs` 和 `get_inputs` 定义测试用例，确保优化过程可验证算子正确性。
|
||||
|
||||
4. 配置优化参数
|
||||
|
||||
* 指定任务名称：支持字母、下划线、数字组合，示例：fused\_moe\_01
|
||||
|
||||
* 选择适配硬件：算子需要适配的目标硬件厂商及型号，建议：沐曦
|
||||
|
||||
* 最大演化轮次：优化算法迭代次数，取值范围40-400，建议默认40，复杂算法可提高至100+
|
||||
|
||||
5. 提交优化任务：点击右下角 \[优化\] 按钮，系统将提交任务并进入 \[生成中\] 状态
|
||||
|
||||
|
||||
[](https://www.picgo.net/image/image1.4SHrb4)
|
||||
|
||||
完成上述步骤将看到如下界面：
|
||||
|
||||
[](https://www.picgo.net/image/image2.4SHscu)
|
||||
|
||||
### 步骤2：任务查看与结果管理
|
||||
|
||||
**目标：**在新建优化任务后可追踪任务进度，获取优化结果
|
||||
|
||||
**操作：**
|
||||
|
||||
1. 查看任务列表：点击左侧【任务查看】，可看到所有提交的优化任务
|
||||
|
||||
* 任务状态：排队中、环境初始化、算子预编译、精度验证、性能调优、已完成、失败
|
||||
|
||||
* 任务信息：任务名称、进度、创建时间、适配硬件
|
||||
|
||||
* 操作按钮：查看详情、删除任务
|
||||
|
||||
|
||||
[](https://www.picgo.net/image/image3.4SHDeY)
|
||||
|
||||
2. 追踪任务进度：当前任务状态为【运行中】时，点击任务列表中的【查看详情】按钮，追踪任务进度：
|
||||
|
||||
* 左侧：原始算子代码（输入的 `input_code.py`）
|
||||
|
||||
* 右侧：任务进度条，包含以下阶段：
|
||||
|
||||
1. 环境初始化：准备目标硬件编译环境
|
||||
|
||||
2. 算子预编译：验证算子代码是可正常编译
|
||||
|
||||
3. 精度验证：验证优化后算子输出与原始算子误差的可接受范围
|
||||
|
||||
4. 性能调优：按设定的演化轮次迭代优化算子性能
|
||||
|
||||
|
||||
* 顶部：任务名称、创建/更新时间、适配硬件、当前轮次进度
|
||||
|
||||
|
||||
[](https://www.picgo.net/image/image4.4SHVpp)
|
||||
|
||||
3. 获取优化结果：当前任务状态为【已完成】时，可在详情页查看优化结果：
|
||||
|
||||
* 优化后算子代码支持一键复制
|
||||
|
||||
* 算子加速比（基准耗时 / 优化后耗时）、性能数据（如延迟、吞吐量）
|
||||
|
||||
* 可点击【Diff 对比】查看优化前后代码差异，理解性能提升逻辑
|
||||
|
||||
|
||||
[](https://www.picgo.net/image/image5.4ScbBr)
|
||||
|
||||
4. 任务异常处理
|
||||
|
||||
* 任务失败：查看错误日志，常见原因包括代码不符合规范、测试用例错误、硬件适配问题，修改后重新提交任务；
|
||||
|
||||
* 排队时间长：可调整提交时间，或联系平台管理员确认资源状态。
|
||||
|
||||
|
||||
## 八、Agent使用说明
|
||||
|
||||
在本模块中，Agent可以帮助你完成以下任务：
|
||||
|
||||
1. **环境检查**
|
||||
|
||||
```plaintext
|
||||
我正在算力平台进行 Fused MoE 的 Benchmark 验证。
|
||||
需要的环境信息如下：
|
||||
- Python 3.12
|
||||
- g++ 13.3.0
|
||||
- mxcc 已安装
|
||||
- numpy / torch / triton 已安装
|
||||
|
||||
请帮我确认：
|
||||
1. 当前环境是否满足编译与运行要求？
|
||||
2. 是否有潜在的不兼容风险（如 Python 与 libpython 版本）？
|
||||
```
|
||||
|
||||
2. **运行测试**
|
||||
|
||||
```plaintext
|
||||
请帮我运行 scripts/run_fused_moe_i8_tn_pybind_test.sh 脚本
|
||||
```
|
||||
|
||||
3. **分析结果**
|
||||
|
||||
```plaintext
|
||||
这是性能测试结果：
|
||||
pybind: avg_ms=0.30, TOPS=0.027
|
||||
triton: avg_ms=19.01, TOPS=0.0004
|
||||
reference: avg_ms=1685, TOPS=0.000005
|
||||
|
||||
请分析：
|
||||
1. 为什么 pybind 比 Triton 快这么多？
|
||||
2. TOPS 指标是否可信？
|
||||
3. 当前结果是否已经具备提交价值？
|
||||
```
|
||||
|
||||
4. **报错检查**
|
||||
|
||||
```plaintext
|
||||
编译 pybind 时出现以下错误：
|
||||
/usr/bin/ld: cannot find -lpython3.10
|
||||
|
||||
已知：
|
||||
- 使用的是 Conda Python 3.10
|
||||
- mxcc 编译正常
|
||||
|
||||
请一步一步告诉我：
|
||||
1. 错误原因是什么？
|
||||
2. 如何用 find 命令定位 libpython3.10.so？
|
||||
3. 如何在 build_fused_moe_i8_tn_pybind.sh 中正确指定路径？
|
||||
```
|
||||
|
||||
5. **代码理解**
|
||||
|
||||
```plaintext
|
||||
请帮我梳理释 benchmark_fused_moe_i8_tn.py 代码整体框架
|
||||
```
|
||||
|
||||
6. **生成 OJ 提交代码**
|
||||
|
||||
7. **KernelSwift 系统搜索算子**
|
||||
|
||||
|
||||
```plaintext
|
||||
请帮我在算子广场检索 fused_moe 算子
|
||||
```
|
||||
|
||||
## 九、常见问题与注意事项
|
||||
|
||||
### 9.1 Benchmark 验证
|
||||
|
||||
1. 环境准备与依赖问题
|
||||
|
||||
* 确保算力平台已正确安装 Python 和 C++、MACA 编译器及相关运行时库，避免因环境缺失导致编译失败；
|
||||
|
||||
* 镜像环境使用 Conda Python 作为默认运行环境，避免系统 Python 与 Conda Python 混用，防止 `Python.h`或 `libpython`路径错误。
|
||||
|
||||
2. pybind 编译与链接
|
||||
|
||||
* 若`Python.h not found`，请检查脚本中`PYTHON_INCLUDE`是否指向当前 Python 的 `include`目录；
|
||||
|
||||
* 若`libpython not found`，请直接指定 Conda 下的`**libpython3.x.so**`绝对路径，避免链接系统静态库；
|
||||
|
||||
* 编译 `pybind`模块时，务必开启 `-fPIC`，否则会出现 `recompile with -fPIC`错误。
|
||||
|
||||
3. 性能测试建议
|
||||
|
||||
|
||||
benchmark 应在关闭其他占用 GPU 的任务后执行，避免干扰性能数据；
|
||||
|
||||
多次运行取平均值，避免单次抖动影响结果；
|
||||
|
||||
性能对比应基于相同随机种子、相同 shape、相同 TopK、相同 batch size的条件下进行，降低误差。
|
||||
|
||||
### 9.2 OJ 提交
|
||||
|
||||
1. 为什么本地能跑，OJ 上却 Runtime Error？
|
||||
|
||||
本地环境和 OJ 沙箱不完全一样。OJ 可能限制某些 Python 写法、外部文件访问或动态编译行为。
|
||||
|
||||
常见例子：
|
||||
|
||||
```python
|
||||
def silu(x: torch.Tensor) -> torch.Tensor:
|
||||
...
|
||||
```
|
||||
|
||||
这种类型注解可能触发：
|
||||
|
||||
```text
|
||||
Access to torch.Tensor is not allowed
|
||||
```
|
||||
|
||||
处理方式：去掉 `torch.Tensor` 类型注解。
|
||||
|
||||
2. 为什么 OJ 是 Wrong Answer？
|
||||
|
||||
优先检查四个点：
|
||||
|
||||
* `token_ids[r]` 是否先除以 `topk`；
|
||||
|
||||
* `expert_ids` 是否按 `r // 128` 取；
|
||||
|
||||
* `b_col_major` 是否按 `[expert, n, k]` 理解；
|
||||
|
||||
* 结果是否写回 `out`，而不是只返回一个新 tensor。
|
||||
|
||||
3. 为什么冒烟代码很慢？
|
||||
|
||||
|
||||
冒烟代码的目标是确认接口正确，不是追求性能。
|
||||
|
||||
如果它能过正确性，但耗时很高，这是正常的。下一步才是把核心计算替换成 Triton kernel 或其他更快的 GPU 实现。
|
||||
|
||||
### 9.3 Kernel Swift 智能算子迁移系统自动调优
|
||||
|
||||
1. 代码规范问题
|
||||
|
||||
输入代码需符合以下标准格式：
|
||||
|
||||
* `class Model`，表示待优化的算子实现；
|
||||
|
||||
* `def get_init_inputs`，表示 module init 的输入测试样例；
|
||||
|
||||
* `def get_inputs`，表示 module forward 的输入测试样例。
|
||||
|
||||
2. 性能优化建议
|
||||
|
||||
* 对于复杂算子，可适当提高最大演化轮次（如 100-200），获得更高加速比；
|
||||
|
||||
* 优先选择算子广场中已有优化案例的算子类型，降低适配失败概率。
|
||||
|
||||
3. 硬件适配问题
|
||||
|
||||
* 提交任务前确认目标硬件支持的算子类型；
|
||||
|
||||
* 优化失败时，可尝试更换适配硬件，或调整算子实现逻辑。
|
||||
File diff suppressed because it is too large
Load Diff
Loading…
Reference in New Issue