From 0e87786f89e2f03978fe2893edf9b5c6ed964889 Mon Sep 17 00:00:00 2001 From: "Xinyi Wu (i26343) - Application Ecology" Date: Mon, 6 Jul 2026 18:49:41 +0800 Subject: [PATCH 1/2] =?UTF-8?q?=E4=BF=AE=E6=94=B9fused=20moe=E6=95=99?= =?UTF-8?q?=E7=A8=8B?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- ...Benchmark 验证到 XPU-OJ 接口提交.md | 883 +++++++++++ .../MCTLASS_Fused MoE 算子优化.md | 1400 ----------------- 2 files changed, 883 insertions(+), 1400 deletions(-) create mode 100644 基于AI Agent开发范式的国产GPU大模型推理算子库优化/Fused MoE 算子入门：从 Benchmark 验证到 XPU-OJ 接口提交.md delete mode 100644 基于AI Agent开发范式的国产GPU大模型推理算子库优化/MCTLASS_Fused MoE 算子优化.md diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/Fused MoE 算子入门：从 Benchmark 验证到 XPU-OJ 接口提交.md b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/Fused MoE 算子入门：从 Benchmark 验证到 XPU-OJ 接口提交.md new file mode 100644 index 0000000..3b35deb --- /dev/null +++ b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/Fused MoE 算子入门：从 Benchmark 验证到 XPU-OJ 接口提交.md @@ -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*"`查找绝对路径
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 | 三种接口供选择：
* CUDA C（`extern "C"`）

* TileLang（Python `@jit`）

* 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 平台后使用分配到的账号进行登录 + + [![image7](https://origin.picgo.net/2026/07/06/image7818fb2b42fb44170.png)](https://www.picgo.net/image/image7.4dk6yw) + +2. 进入比赛页面：点击顶部导航栏【比赛】，选择【进行中】，找到对应比赛进入。 + + [![image8](https://origin.picgo.net/2026/07/06/image8708159862585f494.png)](https://www.picgo.net/image/image8.4dkxr6) + + 3、进入题目页面：本算子对应比赛题目6：`Fused MoE i8 tn`，点击进入题目页面： + + [![image9](https://origin.picgo.net/2026/07/06/image99038c852a7318e8d.png)](https://www.picgo.net/image/image9.4doS64) + + 完成上述步骤可进入如下题目页面： + + * 左侧：题目描述，下滑可查看 CUDA Maca、Triton 和 TileLang 三种语言的接口约定、输入输出格式、示例、数据范围、正确性要求以及提示； + + * 右侧：提交区域，输入编写的`run_kernel(...)`后在下方选择对应的语言即可提交。还可以通过上方导航栏【我的提交】查看历史提交。 + + +[![image10](https://origin.picgo.net/2026/07/06/image1019aed4cf4a3df57d.png)](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. 等待结果：评测时间与题目测试点数量、队列状态和平台负载有关，通常需要等待数十秒到数分钟。以平台实际返回为准。 + + +[![image11](https://origin.picgo.net/2026/07/06/image119978f6695cc5a838.png)](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），便于横向对比。 + + + [![image12](https://origin.picgo.net/2026/07/06/image121919db8452535abb.png)](https://www.picgo.net/image/image12.4doNRi) + + 点击【排行榜】进入如下页面： + + [![image13](https://origin.picgo.net/2026/07/06/image139a158e24a8c895b2.png)](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. 提交优化任务：点击右下角 \[优化\] 按钮，系统将提交任务并进入 \[生成中\] 状态 + + +![image.png](https://alidocs.oss-cn-zhangjiakou.aliyuncs.com/res/r4mlQ5b7084Ndlxo/img/bcc1067e-76be-4e63-b9af-e06390351785.png) + +完成上述步骤将看到如下界面： + +![image.png](https://alidocs.oss-cn-zhangjiakou.aliyuncs.com/res/r4mlQ5b7084Ndlxo/img/51758246-cf96-439a-9a9d-dc7e3bd24041.png) + +### 步骤2：任务查看与结果管理 + +**目标：**在新建优化任务后可追踪任务进度，获取优化结果 + +**操作：** + +1. 查看任务列表：点击左侧【任务查看】，可看到所有提交的优化任务 + + * 任务状态：排队中、环境初始化、算子预编译、精度验证、性能调优、已完成、失败 + + * 任务信息：任务名称、进度、创建时间、适配硬件 + + * 操作按钮：查看详情、删除任务 + + + ![image.png](https://alidocs.oss-cn-zhangjiakou.aliyuncs.com/res/r4mlQ5b7084Ndlxo/img/0102c936-62ba-4226-b6f0-94be689dc08f.png) + +2. 追踪任务进度：当前任务状态为【运行中】时，点击任务列表中的【查看详情】按钮，追踪任务进度： + + * 左侧：原始算子代码（输入的 `input_code.py`） + + * 右侧：任务进度条，包含以下阶段： + + 1. 环境初始化：准备目标硬件编译环境 + + 2. 算子预编译：验证算子代码是可正常编译 + + 3. 精度验证：验证优化后算子输出与原始算子误差的可接受范围 + + 4. 性能调优：按设定的演化轮次迭代优化算子性能 + + + * 顶部：任务名称、创建/更新时间、适配硬件、当前轮次进度 + + +![image.png](https://alidocs.oss-cn-zhangjiakou.aliyuncs.com/res/r4mlQ5b7084Ndlxo/img/f7fb1794-0795-441e-88e8-5e3723e97568.png) + +3. 获取优化结果：当前任务状态为【已完成】时，可在详情页查看优化结果： + + * 优化后算子代码支持一键复制 + + * 算子加速比（基准耗时 / 优化后耗时）、性能数据（如延迟、吞吐量） + + * 可点击【Diff 对比】查看优化前后代码差异，理解性能提升逻辑 + + + ![image.png](https://alidocs.oss-cn-zhangjiakou.aliyuncs.com/res/r4mlQ5b7084Ndlxo/img/8a305199-417d-46cb-9607-f7595d64c059.png) + +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. 硬件适配问题 + + * 提交任务前确认目标硬件支持的算子类型； + + * 优化失败时，可尝试更换适配硬件，或调整算子实现逻辑。 \ No newline at end of file diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/MCTLASS_Fused MoE 算子优化.md b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/MCTLASS_Fused MoE 算子优化.md deleted file mode 100644 index c29c08d..0000000 --- a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/MCTLASS_Fused MoE 算子优化.md +++ /dev/null @@ -1,1400 +0,0 @@ -# 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 推理基本概念。 - - -## 四、前置准备 - -**开始实战前，请确认你已经完成以下准备：** - -**环境准备：** - -* 已进入赛事专属镜像环境。 - - -**工具准备：** - -* 已准备 Agent 工具； - -* 已配置 Token / API Key； - -* 已确认 Agent 可以正常调用模型。 - - -**代码准备：** - -* 已获取 Fused MoE 源码。 - - -## 五、项目实践：从 Benchmark 验证到 XPU-OJ 提交 - -**项目目标：**先跑通 Fused MoE 算子的 benchmark 脚本，建立性能基线；再基于同一题目语义完成 XPU-OJ 冒烟提交，确认评测环境能够正确调用 `run_kernel(...)`，为后续正确性修复和性能优化打基础。 - -### 步骤 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` | - -### 步骤 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 -``` - - -### 步骤 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*"`查找绝对路径
2、将该路径赋值给 `LIBPYTHON_PATH` | -| `recompile with -fPIC` | 编译未开启位置无关代码 | 确保 `mxcc`/ `g++`编译参数中有 `-fPIC` | -| `permission denied` | 无脚本执行权限 | `chmod +x scripts/*.sh` | -| `undefined reference to Py_...` | Python 版本不匹配 | 确认编译脚本中`PYTHON_BIN`路径与当前运行的 Python 环境完全一致 | - -### 步骤 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 硬件限制 | - -### 步骤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：理解 XPU-OJ 提交要求 - -**目标：**在完成前文的本地验证后，理解 XPU-OJ 的提交接口、评测方式和 Candidate 管理方式。 - -在完成前文的本地验证后，本节将带你把实现提交到 XPU-OJ，并确认评测环境能够正确调用 `run_kernel(...)`。 - -完成本节后，你应能够完成一次最小 OJ 提交，查看评测结果，并据此进入正确性修复或后续性能优化。 - -本节的冒烟提交只用于验证函数接口、索引逻辑和提交流程；性能优化请在正确性通过后再进行。 - - -##### 账号准备 - -XPU-OJ 账号由组委会统一发放。登录入口： - - -https://xpuoj.com/ - - -如果登录后看不到比赛或题目，请联系助教或赛事运营确认账号是否已经加入对应比赛或用户组。 - -#### OJ 基础概念 - -##### 什么是 OJ - -OJ 可以理解为“自动评测机”。 - -你提交代码后，OJ 会自动完成： - -1. 加载你的代码； - -2. 构造测试输入； - -3. 调用你的 `run_kernel(...)`； - -4. 生成参考答案； - -5. 对比你的输出和参考输出； - -6. 返回评测状态、耗时、内存和分数。 - - -所以，OJ 不是让你提交 benchmark 日志，也不是让你提交本地运行截图，而是让你提交一份符合接口约定的代码。 - -##### 什么是 Candidate - -Candidate 就是一次可复现的候选方案。 - -建议每一轮都记录： - -| 记录项 | 示例 | -| --- | --- | -| 候选编号 | candidate-001 | -| 代码文件 | `oj/problem_1_fused_moe/solution001.cu` | -| 本地检查结果 | local check passed | -| OJ 结果 | WA / RE / AC | -| 备注 | 初始冒烟版，只验证接口 | - -这样后续多次打榜时，不会忘记哪一版代码对应哪一次提交结果。 - -### 步骤 7：以 CUDA MACA 语言为例实现冒烟代码 - -本节以 CUDA MACA 语言为例介绍冒烟代码的编写方式，选手提交阶段可自行选择 Triton、CUDA MACA 或 TileLang 作为实现语言。 - - -#### 接口约定 - -你必须在提交的 CUDA 源码中提供如下 C 符号，函数名、参数类型、顺序必须完全一致，并使用 `extern "C"` 防止 name mangling： - -```cpp -#include -#include - -extern "C" void run_kernel( - const int8_t* a, - const int8_t* b_col_major, - const float* scale_a, - const float* scale_b, - const float* moe_weights, - const int32_t* token_ids, - const int32_t* expert_ids, - int64_t topk, - __nv_bfloat16* out -); -``` - -#### 题目语义与索引规则 - -本题计算的是 `fused_moe_i8_tn` 形式的 W8A8 MoE GEMM。核心公式是： - -```text -out[r, n] = - sum_k(a[token(r), k] * b_col_major[expert(r), n, k]) - * scale_a[token(r)] - * scale_b[expert(r), n] - * moe_weights[r] -``` - -两个索引最容易写错： - -```text -token(r) = token_ids[r] / topk -expert(r) = expert_ids[r / 128] -``` - -注意： - -* `token_ids` 不是直接拿来当 `a` 的行号，要先除以 `topk`； - -* `expert_ids` 不是每一行一个 expert，而是每 128 行一个 expert； - -* `b_col_major` 的布局是 `[expert, n, k]`，不是 `[expert, k, n]`； - -* 最终结果必须写回传入的 `out`。 - - -#### 准备 OJ 提交文件 - -在终端中创建目录： - -```bash -cd /data/fusedmoe_v2.1 -mkdir -p oj/problem_1_fused_moe -``` - -新建 CUDA 源码文件： - -```bash -touch oj/problem_1_fused_moe/solution001.cu -``` - -注意：冒烟版的目标只是确认 CUDA MACA 提交接口、索引和 OJ 提交流程，不追求性能最优。 - -#### CUDA MACA 冒烟代码 - -将下面代码保存到： - -```text -oj/problem_1_fused_moe/solution001.cu -``` - -```cpp -#include -#include -#include - -// xcore1000's CUDA-compatible compiler does not expose NVIDIA's __dp4a. -// This is a correctness-first replacement: each int32 stores four signed -// int8 values in little-endian byte order. -__device__ inline int32_t signed_byte(uint32_t x) { - x &= 0xffu; - return (int32_t)(x ^ 0x80u) - 128; -} - -__device__ inline int32_t dp4a_compat(int32_t a, int32_t b, int32_t acc) { - uint32_t ua = (uint32_t)a; - uint32_t ub = (uint32_t)b; - acc += signed_byte(ua) * signed_byte(ub); - acc += signed_byte(ua >> 8) * signed_byte(ub >> 8); - acc += signed_byte(ua >> 16) * signed_byte(ub >> 16); - acc += signed_byte(ua >> 24) * signed_byte(ub >> 24); - return acc; -} - -__global__ void w8a8_moe_gemm_kernel( - const int8_t* __restrict__ a, - const int8_t* __restrict__ b_col_major, - const float* __restrict__ scale_a, - const float* __restrict__ scale_b, - const float* __restrict__ moe_weights, - const int32_t* __restrict__ token_ids, - const int32_t* __restrict__ expert_ids, - int K, int N, int topk, - __nv_bfloat16* __restrict__ out) -{ - int n_base = blockIdx.x * 128; - int m_base = blockIdx.y * 128; - int expert = expert_ids[blockIdx.y]; - - int tid = threadIdx.x; - int warp_id = tid / 32; - int lane_id = tid & 31; - - int warp_y = warp_id / 2; - int warp_x = warp_id & 1; - int my = lane_id / 8; - int mx = lane_id & 7; - - int m_idx[8]; - int n_idx[8]; -#pragma unroll - for (int i = 0; i < 8; ++i) { - m_idx[i] = warp_y * 32 + my + i * 4; - } -#pragma unroll - for (int j = 0; j < 8; ++j) { - n_idx[j] = warp_x * 64 + mx + j * 8; - } - - __shared__ int32_t smem_A[2][128 * 17]; - __shared__ int32_t smem_B[2][128 * 17]; - - int32_t accum[8][8] = {0}; - -#pragma unroll - for (int step = 0; step < 2; ++step) { - int load_idx = step * 256 + tid; - int row = load_idx / 4; - int col_int4 = load_idx & 3; - - int r = m_base + row; - int token = token_ids[r] / topk; - int64_t a_idx = (int64_t)token * K; - int4 va = ((const int4*)(a + a_idx))[col_int4]; - - int sa = row * 17 + col_int4 * 4; - smem_A[0][sa + 0] = va.x; - smem_A[0][sa + 1] = va.y; - smem_A[0][sa + 2] = va.z; - smem_A[0][sa + 3] = va.w; - - int64_t b_idx = (int64_t)expert * N * K + (int64_t)(n_base + row) * K; - int4 vb = ((const int4*)(b_col_major + b_idx))[col_int4]; - - int sb = row * 17 + col_int4 * 4; - smem_B[0][sb + 0] = vb.x; - smem_B[0][sb + 1] = vb.y; - smem_B[0][sb + 2] = vb.z; - smem_B[0][sb + 3] = vb.w; - } - __syncthreads(); - - for (int k_outer = 0; k_outer < K; k_outer += 64) { - int comp_buf = (k_outer / 64) & 1; - int load_buf = 1 - comp_buf; - int next_k = k_outer + 64; - - if (next_k < K) { -#pragma unroll - for (int step = 0; step < 2; ++step) { - int load_idx = step * 256 + tid; - int row = load_idx / 4; - int col_int4 = load_idx & 3; - - int r = m_base + row; - int token = token_ids[r] / topk; - int64_t a_idx = (int64_t)token * K + next_k; - int4 va = ((const int4*)(a + a_idx))[col_int4]; - - int sa = row * 17 + col_int4 * 4; - smem_A[load_buf][sa + 0] = va.x; - smem_A[load_buf][sa + 1] = va.y; - smem_A[load_buf][sa + 2] = va.z; - smem_A[load_buf][sa + 3] = va.w; - - int64_t b_idx = (int64_t)expert * N * K + - (int64_t)(n_base + row) * K + next_k; - int4 vb = ((const int4*)(b_col_major + b_idx))[col_int4]; - - int sb = row * 17 + col_int4 * 4; - smem_B[load_buf][sb + 0] = vb.x; - smem_B[load_buf][sb + 1] = vb.y; - smem_B[load_buf][sb + 2] = vb.z; - smem_B[load_buf][sb + 3] = vb.w; - } - } - -#pragma unroll - for (int k_step = 0; k_step < 16; ++k_step) { - int32_t reg_A[8]; - int32_t reg_B[8]; - -#pragma unroll - for (int i = 0; i < 8; ++i) { - reg_A[i] = smem_A[comp_buf][m_idx[i] * 17 + k_step]; - } -#pragma unroll - for (int j = 0; j < 8; ++j) { - reg_B[j] = smem_B[comp_buf][n_idx[j] * 17 + k_step]; - } - -#pragma unroll - for (int i = 0; i < 8; ++i) { -#pragma unroll - for (int j = 0; j < 8; ++j) { - accum[i][j] = dp4a_compat(reg_A[i], reg_B[j], accum[i][j]); - } - } - } - __syncthreads(); - } - - float scale_row[8]; -#pragma unroll - for (int i = 0; i < 8; ++i) { - int r = m_base + m_idx[i]; - int token = token_ids[r] / topk; - scale_row[i] = scale_a[token] * moe_weights[r]; - } - - float scale_col[8]; -#pragma unroll - for (int j = 0; j < 8; ++j) { - int n = n_base + n_idx[j]; - scale_col[j] = scale_b[(int64_t)expert * N + n]; - } - -#pragma unroll - for (int i = 0; i < 8; ++i) { - int r = m_base + m_idx[i]; -#pragma unroll - for (int j = 0; j < 8; ++j) { - int n = n_base + n_idx[j]; - float v = (float)accum[i][j] * scale_row[i] * scale_col[j]; - out[(int64_t)r * N + n] = __float2bfloat16(v); - } - } -} - -static size_t device_allocation_size(const void* p) { - mcDrvDeviceptr_t base = 0; - size_t size = 0; - (void)wcuMemGetAddressRange(&base, &size, (mcDrvDeviceptr_t)(uintptr_t)p); - return size; -} - -extern "C" void run_kernel( - const int8_t* a, - const int8_t* b_col_major, - const float* scale_a, - const float* scale_b, - const float* moe_weights, - const int32_t* token_ids, - const int32_t* expert_ids, - int64_t topk, - __nv_bfloat16* out) -{ - size_t b_size = device_allocation_size(b_col_major); - size_t out_size = device_allocation_size(out); - - int N = 7168; - int K = 2048; - if (b_size > 5000000000ULL) { - N = 4096; - K = 7168; - } - - int EM = 4096; - if (out_size > 128ULL * 1024ULL * 1024ULL) { - EM = 32768; - } else if (out_size == 0) { - // Last-resort fallback if allocation-size probing is unavailable. - int32_t host_tokens[4096]; - cudaMemcpy(host_tokens, token_ids, sizeof(host_tokens), cudaMemcpyDeviceToHost); - int max_token_id = 0; - for (int i = 0; i < 4096; ++i) { - if (host_tokens[i] > max_token_id) { - max_token_id = host_tokens[i]; - } - } - if (max_token_id >= 4096) { - EM = 32768; - } - } - - dim3 block(256); - dim3 grid(N / 128, EM / 128); - w8a8_moe_gemm_kernel<<>>( - a, b_col_major, scale_a, scale_b, moe_weights, - token_ids, expert_ids, K, N, (int)topk, out); -} -``` - -保存前建议人工检查 5 个点： - -| 检查项 | 应该满足 | -| --- | --- | -| C 符号 | 必须是 `extern "C" void run_kernel(...)` | -| 参数顺序 | 必须是 `a, b_col_major, scale_a, scale_b, moe_weights, token_ids, expert_ids, topk, out` | -| token 索引 | 必须使用 `token_ids[r] / topk` | -| expert 索引 | 必须使用 `expert_ids[r / 128]` 或等价的每 128 行一个 expert 逻辑 | -| 输出方式 | 必须原地写入 `out`，不要返回新指针 | - -### 步骤 8：提交并解读 XPU-OJ 结果 - -#### 提交到 XPU-OJ - -1. 打开 XPU-OJ： https://xpuoj.com/ - -2. 使用组委会发放的账号登录； - -3. 进入比赛页面； [![image6](https://origin.picgo.net/2026/06/23/image6047f2ac4bd2a0f08.png)](https://www.picgo.net/image/image6.4ScJM4) - -4. 找到题目： ```text1. Fused MoE i8 tn``` - -5. 点击题目进入详情页； - -6. 在提交区域选择本题支持的语言，例如： ```textCUDA / CUDA MACA``` - - -```plaintext -具体名称以页面下拉框为准。 -``` - -1. 将 `solution001.cu` 中的源码复制到提交框； - -2. 点击提交； - -3. 等待评测结果返回。 - - -#### 查看 OJ 结果 - -提交后，进入： - -```text -我的提交 -``` - -常见状态含义如下： - -| 状态 | 含义 | 下一步 | -| --- | --- | --- | -| Accepted / AC | 正确性通过 | 可以继续优化性能 | -| Wrong Answer / WA | 输出数值不对 | 检查索引、shape、dtype、缩放系数 | -| Runtime Error / RE | 运行时报错 | 点开详情看报错栈 | -| Compile Error / CE | 编译或加载失败 | 检查头文件、语法、`extern "C"` 符号和函数签名 | -| Time Limit Exceeded / TLE | 超时 | 说明代码太慢，需要做 kernel 优化 | - -**测试点结果怎么看：** - -点开提交详情后，优先查看每个测试点是否通过正确性校验，再查看耗时和性能指标。一般按以下顺序判断： - -1. 先看是否通过： - - * 如果出现 `passed`、`pass=true` 或类似通过标记，说明该测试点数值正确； - - * 如果出现 `FAILED`、`Wrong Answer`、`allclose failed`、`Runtime Error`，说明当前实现还没有通过该测试点，不能只看性能时间。 - -2. 再看耗时： - - * `time_ms`：当前提交在该测试点上的运行时间，数值越小越好； - - * `speedup`：相对基线的加速比，数值越大越好； - - * `score_ratio`：当前测试点得分比例，越接近 1 说明越接近该测试点满分。 - -3. 最后看不同测试点之间的差异： - - * 如果某个测试点正确性失败，应先定位该测试点的 shape、topk、dtype、索引规则或容差要求； - - * 如果某个测试点正确但明显更慢，说明该场景下可能存在访存、重复计算或并行度不足问题； - - * 优化时应优先关注失败测试点和耗时占比较高的测试点。 - - -**分数怎么看：** - -XPU-OJ 结果通常需要同时关注正确性和性能分数： - -* `pass=true`：表示该提交通过正确性校验； - -* `pass=false`：表示该提交未通过，通常不会获得有效性能分； - -* `time_ms`：当前提交的运行时间，越小越好； - -* `speedup`：相对基线的加速比，越大越好； - -* `score_ratio`：当前测试点得分比例，越接近 1 说明越接近该测试点满分； - -* `allclose failed`：说明输出与参考结果误差超过阈值，应优先修正确性； - -* `Runtime Error`：说明编译、运行、越界访问或环境限制出错，应先解决运行问题。 - - -评分理解顺序： - -1. 先看 `pass`：不通过时先修正确性； - -2. 再看 `time_ms`：通过后比较耗时； - -3. 再看 `speedup`：判断相对基线是否有提升； - -4. 最后看 `score_ratio`：判断当前优化距离满分还有多远。 - -需要注意：OJ 分数通常不是只由单个测试点决定，而是多个测试点综合计算。因此优化时不要只盯一个最快 case，应优先解决失败测试点和耗时占比较高的测试点。 - -**是否需要优化怎么看：** - -读完 OJ 结果后，可以按下面的顺序判断下一步： - -| 现象 | 说明 | 下一步 | -| --- | --- | --- | -| 提交结果不通过 | 算子语义、索引、shape、dtype、scale、反量化或输出写回存在问题 | 先修正确性 | -| 正确性通过但 `time_ms` 很高 | 初版实现可用，但并行度、访存或计算复用不足 | 进入性能优化 | -| 某个测试点特别慢 | 该场景可能存在重复计算、负载不均或 block 配置不合适 | 针对该测试点单独分析 | -| 多次提交耗时抖动很大 | 可能受 GPU 占用、预热不足或评测波动影响 | 多提交几次或回到本地 benchmark 复查 | - -如果当前版本还没有通过正确性，不建议马上做性能优化。先让冒烟版本跑通，再进入下一轮 Candidate。 - -**怎么优化：** - -在正确性通过后，可以按以下方向逐步优化： - -1. 减少重复计算： - - * 检查同一个 token 或 expert 是否被重复计算； - - * 对 topk2 / topk3 场景，可考虑复用相同 token 的中间结果； - - * 避免每个输出元素都重复加载相同的输入向量。 - -2. 优化访存模式： - - * 尽量让连续线程访问连续内存； - - * 减少非合并访存； - - * 对频繁使用的权重、scale、token id、expert id 做局部缓存； - - * 避免不必要的全局内存读写。 - -3. 提高并行度： - - * 将输出矩阵按行、列或 expert 维度切分； - - * 对小 batch、小 token 场景，适当增加 block 数量； - - * 避免一个 kernel 只有很少 block，导致 GPU 利用率不足。 - -4. 优化计算粒度： - - * 合理设置 `BLOCK_M`、`BLOCK_N`、`BLOCK_K`； - - * 让 tile 大小匹配硬件并行能力和 shared memory 限制； - - * 对不同 topk 或不同 shape 可以采用不同 kernel 配置。 - -5. 对比基线逐步验证： - - * 每次只改一个优化点； - - * 修改后先跑正确性测试； - - * 正确后再跑 benchmark； - - * 记录每次 `time_ms`、`speedup`、`score_ratio` 变化，避免无效优化。 - -具体的 Candidate 保存和下一轮 Agent 优化方式，见步骤 9。 - -如果看到 `0 pts`，通常表示本次提交没有拿到分数。原因可能是： - -* 样例没过； - -* 测试点没过； - -* 代码运行时报错； - -* 代码超时； - -* 输出与参考答案超过容差。 - - -如果看到用时和内存都是 `0`，很多时候说明代码在正式计时前就失败了，例如 C 符号不匹配、编译失败或 kernel launch 失败。 - -#### 理解 OJ 评测流程 - -一次 OJ 提交通常会经历下面这些步骤： - -1. 选手提交代码； - -2. 平台按所选语言加载代码； - -3. 评测程序构造输入 tensor； - -4. 调用选手代码里的 `run_kernel(...)`； - -5. 选手代码把结果写入 `out`； - -6. 评测程序生成参考结果； - -7. 比较 `out` 和参考结果； - -8. 正确性通过后统计运行耗时； - -9. 根据题目规则计算分数； - -10. 在排行榜或提交记录中更新结果。 - - -本题的正确性校验口径是： - -```python -torch.allclose(out_target.float(), out_ref.float(), rtol=0.0, atol=1e-2) -``` - -也就是说，OJ 允许很小的数值误差，但不是随便差一点都能过。 - -### 步骤 9：保存 Candidate 并进入优化 - -#### 保存 Candidate - -建议每一次能跑的版本都用 Git 保存。 - -```bash -cd /data/fusedmoe_v2.1 -git status --short -git add oj/problem_1_fused_moe/solution001.cu -git commit -m "candidate 001 fused moe i8 tn oj smoke" -git tag candidate-001-oj-smoke -``` - -查看最近候选版本： - -```bash -git log --oneline --decorate -5 -``` - -如果下一轮要继续优化，可以复制一份新文件： - -```bash -cp oj/problem_1_fused_moe/solution001.cu oj/problem_1_fused_moe/solution002.cu -``` - -然后让 Agent 基于 `solution002.cu` 继续改。 - -#### 使用 Agent 定位问题与优化 - -本模块中，Agent 主要用来做三件事： - -1. 读题目接口； - -2. 生成最小可提交代码； - -3. 根据 OJ 报错定位问题。 - - -建议不要一开始就让 Agent “直接写最快版本”。更稳的流程是： - -```text -第一步：先写一个能过正确性的最小版本。 -第二步：提交 OJ，看是否 AC。 -第三步：AC 后再优化性能。 -``` - -可以使用下面的 Prompt： - -```text -我正在做 XPU-OJ 的 Fused MoE GEMM 题。 - -请只做一件事：根据题目接口写一个最小正确的 CUDA MACA run_kernel 冒烟版本。 - -要求： -1. C 符号和函数签名必须完全一致： - extern "C" void run_kernel( - const int8_t* a, - const int8_t* b_col_major, - const float* scale_a, - const float* scale_b, - const float* moe_weights, - const int32_t* token_ids, - const int32_t* expert_ids, - int64_t topk, - __nv_bfloat16* out - ) -2. token(r) = token_ids[r] / topk -3. expert(r) = expert_ids[r / 128] -4. b_col_major 的布局是 [expert, n, k] -5. 结果必须原地写入 out -6. 不要做性能优化 -7. 不要依赖外部文件 -8. 请输出完整可复制提交的 CUDA 源码 -``` - -如果 OJ 返回 `Wrong Answer`，可以继续问： - -```text -OJ 返回 Wrong Answer。 - -请不要重写整份代码，先根据下面四点检查可能原因： -1. token_ids 是否正确除以 topk； -2. expert_ids 是否按每 128 行一个 expert 使用； -3. b_col_major 是否按 [expert, n, k] 读取； -4. 是否把结果写入 out，且 dtype 与 out 保持一致。 - -请给出最小修改建议。 -``` - -如果 OJ 返回 `Runtime Error`，可以问： - -```text -OJ 返回 Runtime Error。 - -这是错误日志：[粘贴错误日志] - -请先判断是 extern "C" 符号、头文件、编译选项、dtype、shape、越界访问还是 GPU kernel 调用问题。 -只给出最小修复方案。 -``` - -### 常见问题 - -#### Q1：为什么本地能跑，OJ 上却 Runtime Error？ - -本地环境和 OJ 编译环境不完全一样。OJ 可能对头文件、CUDA/MACA 编译参数、外部文件访问或动态链接行为有限制。 - -常见例子： - -```cpp -extern "C" void run_kernel(...) { - ... -} -``` - -常见问题包括：没有使用 `extern "C"` 导致符号名被 C++ name mangling 改写、缺少必要头文件、调用了当前 OJ 编译环境不支持的 CUDA intrinsic，或 kernel launch 参数越界。 - -处理方式：先保证接口符号和参数类型完全一致，再根据错误日志逐项缩小范围。 - -#### Q2：为什么 OJ 是 Wrong Answer？ - -优先检查四个点： - -1. `token_ids[r]` 是否先除以 `topk`； - -2. `expert_ids` 是否按 `r / 128` 取； - -3. `b_col_major` 是否按 `[expert, n, k]` 理解； - -4. 结果是否写回 `out`，而不是写到临时 buffer 后没有拷回。 - - -#### Q3：为什么冒烟代码很慢？ - -冒烟代码的目标是确认接口正确，不是追求性能。 - -如果它能过正确性，但耗时很高，这是正常的。下一步才是在 CUDA MACA 源码中继续优化访存、并行度、tile 配置和计算复用。 - -#### Q4：50 分、10 分是什么意思？ - -不同比赛和题目的评分规则可能不同。一般可以先这样理解： - -* 正确性没过时，通常拿不到有效分数； - -* 正确性通过后，平台会继续根据耗时或加速比计算分数； - -* 具体分数含义以 XPU-OJ 当前题目的评分说明为准。 - - -#### Q5：榜单怎么看？ - -先看自己的提交是否通过正确性，再看耗时和分数。 - -建议记录： - -| Candidate | OJ 状态 | 用时 | 分数 | 备注 | -| --- | --- | --- | --- | --- | -| candidate-001 | AC / WA / RE | 以页面为准 | 以页面为准 | 冒烟版 | -| candidate-002 | AC / WA / RE | 以页面为准 | 以页面为准 | 第一轮优化 | - -不要只看单次结果。每轮都记录，后面才知道 Agent 的修改到底有没有带来收益。 - -### 从 Benchmark 验证到参赛作品的路径回顾 - -建议按下面顺序推进： - -1. 跑通 benchmark 脚本，理解算子输入输出； - -2. 阅读 XPU-OJ 题目页面，确认 `run_kernel(...)` 接口； - -3. 提交冒烟代码，确认 OJ 链路正常； - -4. 如果冒烟代码 WA / RE，先修正确性； - -5. 正确性通过后，再让 Agent 在 CUDA MACA 版本上继续优化 kernel； - -6. 每一轮提交都保存 candidate、prompt、代码 diff 和 OJ 结果； - -7. 用 OJ 分数和耗时判断优化是否有效。 - - -```text -Benchmark 验证代码用来学习，OJ 用来评分，Candidate 用来管理每一轮结果。 -``` - - -## 六、后续实践：Kernel Swift 智能算子迁移系统自动调优 - -系统链接：[https://deeplink.org.cn/kernelswift/task](https://deeplink.org.cn/kernelswift/task) - -**项目目标：**基于 KernelSwift 智能算子迁移系统，对 Fused MoE 算子进行在线自动调优。通过输入算子的 PyTorch 代码，一键生成适配沐曦硬件的高性能实现，高效完成算子优化与全流程追踪。 - -### 6.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. 提交优化任务：点击右下角 \[优化\] 按钮，系统将提交任务并进入 \[生成中\] 状态 - - - [![image1](https://origin.picgo.net/2026/06/23/image1d46e08e5a17fd767.png)](https://www.picgo.net/image/image1.4SHrb4) - -完成上述步骤将看到如下界面： - - [![image2](https://origin.picgo.net/2026/06/23/image268924dc11f138788.png)](https://www.picgo.net/image/image2.4SHscu) - -### 6.2 任务查看与结果管理 - -**目标：**在新建优化任务后可追踪任务进度，获取优化结果 - -**操作：** - -1. 查看任务列表：点击左侧【任务查看】，可看到所有提交的优化任务 - - * 任务状态：排队中、环境初始化、算子预编译、精度验证、性能调优、已完成、失败 - - * 任务信息：任务名称、进度、创建时间、适配硬件 - - * 操作按钮：查看详情、删除任务 - - - [![image3](https://origin.picgo.net/2026/06/23/image33091601c9a68bd18.png)](https://www.picgo.net/image/image3.4SHDeY) - -2. 追踪任务进度：当前任务状态为【运行中】时，点击任务列表中的【查看详情】按钮，追踪任务进度： - - * 左侧：原始算子代码（输入的 `input_code.py`） - - * 右侧：任务进度条，包含以下阶段： - - 1. 环境初始化：准备目标硬件编译环境 - - 2. 算子预编译：验证算子代码是可正常编译 - - 3. 精度验证：验证优化后算子输出与原始算子误差的可接受范围 - - 4. 性能调优：按设定的演化轮次迭代优化算子性能 - - - * 顶部：任务名称、创建/更新时间、适配硬件、当前轮次进度 - - - [![image4](https://origin.picgo.net/2026/06/23/image403daf417d165a79f.png)](https://www.picgo.net/image/image4.4SHVpp) - -3. 获取优化结果：当前任务状态为【已完成】时，可在详情页查看优化结果： - - * 优化后算子代码支持一键复制 - - * 算子加速比（基准耗时 / 优化后耗时）、性能数据（如延迟、吞吐量） - - * 可点击【Diff 对比】查看优化前后代码差异，理解性能提升逻辑 - - - [![image5](https://origin.picgo.net/2026/06/23/image58f2b2ea36dad2ef0.png)](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. **KernelSwift 系统搜索算子** - - - ```plaintext - 请帮我在算子广场检索 fused_moe 算子 - ``` - -## 八、常见问题与注意事项 - -### 8.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​的条件下进行，降低误差。 - - -### 8.2 Kernel Swift 智能算子迁移系统自动调优项目 - -1. 代码规范问题 - - 输入代码需符合以下标准格式： - - * `class Model`，表示待优化的算子实现； - - * `def get_init_inputs`，表示 module init 的输入测试样例； - - * `def get_inputs`，表示 module forward 的输入测试样例。 - -2. 性能优化建议 - - * 对于复杂算子，可适当提高最大演化轮次（如 100-200），获得更高加速比； - - * 优先选择算子广场中已有优化案例的算子类型，降低适配失败概率。 - -3. 硬件适配问题 - - * 提交任务前确认目标硬件支持的算子类型； - - * 优化失败时，可尝试更换适配硬件，或调整算子实现逻辑。 - From 4b4a56eda8c151ccedfee51cd175c50ae3febdbf Mon Sep 17 00:00:00 2001 From: "Xinyi Wu (i26343) - Application Ecology" Date: Mon, 6 Jul 2026 18:55:12 +0800 Subject: [PATCH 2/2] =?UTF-8?q?=E4=BF=AE=E6=94=B9fused=5Fmoe=E6=95=99?= =?UTF-8?q?=E7=A8=8B?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit --- ...：从 Benchmark 验证到 XPU-OJ 接口提交.md | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/Fused MoE 算子入门：从 Benchmark 验证到 XPU-OJ 接口提交.md b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/Fused MoE 算子入门：从 Benchmark 验证到 XPU-OJ 接口提交.md index 3b35deb..35a3666 100644 --- a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/Fused MoE 算子入门：从 Benchmark 验证到 XPU-OJ 接口提交.md +++ b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/Fused MoE 算子入门：从 Benchmark 验证到 XPU-OJ 接口提交.md @@ -660,11 +660,11 @@ $S(T\_k) = \frac{100}{1 + \left(\frac{1}{0.5} - 1\right) \cdot \frac{T\_ 5. 提交优化任务：点击右下角 \[优化\] 按钮，系统将提交任务并进入 \[生成中\] 状态 -![image.png](https://alidocs.oss-cn-zhangjiakou.aliyuncs.com/res/r4mlQ5b7084Ndlxo/img/bcc1067e-76be-4e63-b9af-e06390351785.png) +[![image1](https://origin.picgo.net/2026/06/23/image1d46e08e5a17fd767.png)](https://www.picgo.net/image/image1.4SHrb4) 完成上述步骤将看到如下界面： -![image.png](https://alidocs.oss-cn-zhangjiakou.aliyuncs.com/res/r4mlQ5b7084Ndlxo/img/51758246-cf96-439a-9a9d-dc7e3bd24041.png) +[![image2](https://origin.picgo.net/2026/06/23/image268924dc11f138788.png)](https://www.picgo.net/image/image2.4SHscu) ### 步骤2：任务查看与结果管理 @@ -681,7 +681,7 @@ $S(T\_k) = \frac{100}{1 + \left(\frac{1}{0.5} - 1\right) \cdot \frac{T\_ * 操作按钮：查看详情、删除任务 - ![image.png](https://alidocs.oss-cn-zhangjiakou.aliyuncs.com/res/r4mlQ5b7084Ndlxo/img/0102c936-62ba-4226-b6f0-94be689dc08f.png) + [![image3](https://origin.picgo.net/2026/06/23/image33091601c9a68bd18.png)](https://www.picgo.net/image/image3.4SHDeY) 2. 追踪任务进度：当前任务状态为【运行中】时，点击任务列表中的【查看详情】按钮，追踪任务进度： @@ -701,7 +701,7 @@ $S(T\_k) = \frac{100}{1 + \left(\frac{1}{0.5} - 1\right) \cdot \frac{T\_ * 顶部：任务名称、创建/更新时间、适配硬件、当前轮次进度 -![image.png](https://alidocs.oss-cn-zhangjiakou.aliyuncs.com/res/r4mlQ5b7084Ndlxo/img/f7fb1794-0795-441e-88e8-5e3723e97568.png) +[![image4](https://origin.picgo.net/2026/06/23/image403daf417d165a79f.png)](https://www.picgo.net/image/image4.4SHVpp) 3. 获取优化结果：当前任务状态为【已完成】时，可在详情页查看优化结果： @@ -712,7 +712,7 @@ $S(T\_k) = \frac{100}{1 + \left(\frac{1}{0.5} - 1\right) \cdot \frac{T\_ * 可点击【Diff 对比】查看优化前后代码差异，理解性能提升逻辑 - ![image.png](https://alidocs.oss-cn-zhangjiakou.aliyuncs.com/res/r4mlQ5b7084Ndlxo/img/8a305199-417d-46cb-9607-f7595d64c059.png) + [![image5](https://origin.picgo.net/2026/06/23/image58f2b2ea36dad2ef0.png)](https://www.picgo.net/image/image5.4ScbBr) 4. 任务异常处理