Add tilelang_maca_install_guide.md

Merge pull request 'add mctlass' (#61 ) from raymond_feng2/op_optimization:feat/flashattn into master
feat:mctlass add
2026-06-28 04:18:21 +00:00 · 2026-06-26 10:18:18 +08:00 · 2026-06-26 10:13:55 +08:00 · 2026-06-25 18:01:04 +08:00 · 2026-06-25 15:28:54 +08:00 · 2026-06-25 15:22:51 +08:00
247 changed files with 7708 additions and 834 deletions
--- a/README.md
+++ b/README.md
@ -18,12 +18,14 @@
 初赛攻坚 Fused Moe Gemm、决赛冲刺 MLA/NSA，全程基于全国产技术栈开发优化，打造自主可控的低成本大模型推理算子底座。
 **赛题一相关资料**
 - [赛题一方案：基于国产软件栈的推理前沿算子优化比赛方案](基于国产软件栈大模型推理前沿算子优化/基于国产软件栈的推理前沿算子优化比赛方案.md)
 - [TileLang MACA 构建指南（模力方舟）](基于国产软件栈大模型推理前沿算子优化/tilelang_maca_build_guide_模力方舟.md)
 - [TileLang 算子 sample 跑通测试指南](基于国产软件栈大模型推理前沿算子优化/race_tests_run_guide基于tilelang算子sample跑通测试.md)
- [赛题一国产软件栈样例说明](基于国产软件栈大模型推理前沿算子优化/race-1-domestic-software-stack-samples.md)
+- [赛题一国产软件栈样例说明](基于国产软件栈大模型推理前沿算子优化/基于国产软件栈的大模型推理前沿算子优化（TileLang）Sample%20算子获取说明.md)
 - [TileLang 训练营学习资料](https://www.gitlink.org.cn/ccf-ai-infra/Intro-ops/about)
 ### 赛题二：基于 AI Agent 开发范式的国产 GPU 大模型推理算子库优化
@ -44,9 +46,12 @@
 **赛题二相关资料**
 - [赛题二方案：基于 AI Agent 开发范式的国产 GPU 大模型推理算子库优化方案](基于AI%20Agent开发范式的国产GPU大模型推理算子库优化/基于AI%20Agent开发范式的国产GPU大模型算子推理库优化方案.md)
 - [赛题二选手入口](基于AI%20Agent开发范式的国产GPU大模型推理算子库优化/选手入口.md)
 - [模力方舟 Agent 部署准备教程](基于AI%20Agent开发范式的国产GPU大模型推理算子库优化/模力方舟Agent部署准备教程.md)
 - [赛题二说明及资料参考](基于AI%20Agent开发范式的国产GPU大模型推理算子库优化/赛题说明.md)
 ###**两个赛题统一使用模力方舟上的镜像PyTorch-Agent / 2.8.0 / Python 3.12 / maca 3.7.2.1**
 ## 参赛对象
 面向全国全日制专科、本科、硕博在校生（非在职），40 周岁以下青年科技人才均可参与；
@ -99,6 +104,8 @@
 沐曦开发者社区活动页面，完成新人礼任务，提前熟悉 C500 在线算力使用。https://developer.metax-tech.com/activities/6
 学生可通过登录启悟社区后跳转沐曦开发者社区领取100元算力代金券。(https://developer.metax-tech.com/activities/11)
 报名赛事后还可额外获得300元算力代金券：https://developer.metax-tech.com/activities/17
 使用指南：[模力方舟快速使用 SOP](模力方舟快速使用SOP.md)
@ -109,7 +116,7 @@
 加入沐曦股份“揭榜挂帅”赛题，深耕国产 GPU 算子优化，用技术降低每 Token 推理成本，共建自主可控 AI 算力新生态！
-👇 报名通道：登录挑战杯官网 [www.tiaozhanbei.net](https://link.wtturl.cn/?target=https%3A%2F%2Fwww.tiaozhanbei.net&scene=im&aid=497858&lang=zh) 揭榜挂帅入口报名
+👇 报名通道：登录挑战杯官网[2026年度中国青年科技创新“揭榜挂帅”擂台赛（学生赛道）](https://2026.tiaozhanbei.net/)揭榜挂帅入口报名
--- a/tilelang_maca_install_guide.md
+++ b/tilelang_maca_install_guide.md
@ -0,0 +1,97 @@
 # TileLang MACA 安装与验证指南
 > 环境：MetaX C500 / MACA 3.7.2.1 / Python 3.12 / Conda
 >
 > 源码：`/data/tilelang-metax`（dev 分支）
 ---
 ## 1. 前置条件
 - MACA 已安装在 `/opt/maca`（包含 `libmcruntime.so` 等运行时库）
 - Python 3.12 + Conda 环境
 ---
 ## 2. 编译安装 tilelang（启用 MACA 后端）
 ```bash
 cd /data/tilelang-metax
 rm -rf build
 export USE_MACA=ON
 pip install -e . -v
 ```
 编译产物位于 `build/lib/`，包含：
 - `libtilelang.so` — 注册 maca target kind
 - `libtvm.so` — TVM 运行时
 ---
 ## 3. 修复 `.pth` 文件（关键步骤）
 tilelang 的 editable install 生成的 `.pth` 文件缺少 `tvm` Python 包的正确路径。
 **文件位置**：`/opt/conda/lib/python3.12/site-packages/_tilelang_editable.pth`
 **修改**：将其中一行 `/data/tilelang-metax/3rdparty/tvm` 改为：
 ```
 /data/tilelang-metax/3rdparty/tvm/python
 ```
 修改后的文件内容应为：
 ```
 import _tilelang_editable
 /data/tilelang-metax
 /data/tilelang-metax
 /data/tilelang-metax/3rdparty/tvm/python
 /data/tilelang-metax/3rdparty/tvm
 /data/tilelang-metax/3rdparty/tvm
 /data/tilelang-metax/3rdparty/tvm/3rdparty/dmlc-core
 /data/tilelang-metax/3rdparty/tvm
 /data/tilelang-metax/3rdparty/cutlass
 /data/tilelang-metax/3rdparty/cutlass
 /data/tilelang-metax/3rdparty/composable_kernel
 /data/tilelang-metax/3rdparty/composable_kernel
 ```
 ---
 ## 4. 验证
 ### 4.1 基础导入
 ```bash
 python -c "import tilelang; print(tilelang.__version__)"
 # 预期输出：0.1.9+maca.gitee6db437
 ```
 ### 4.2 MACA Target 验证
 **必须先 `import tilelang`**，因为它会加载 `libtilelang.so`（注册 maca target）：
 ```bash
 python -c "
 import tilelang
 from tvm.target import Target
 t = Target('maca')
 print(t)
 "
 # 预期输出：
 # maca -keys=maca,gpu -max_local_memory_per_block=4095 -max_num_threads=1024
 #   -max_shared_memory_per_block=65536 -max_threads_per_block=1024
 #   -mcpu=xcore1000 -mtriple=mxc-metax-macahca -thread_warp_size=64
 ```
 ---
 ## 5. 常见问题
 | 问题 | 原因 | 解决 |
 |------|------|------|
 | `ModuleNotFoundError: No module named 'tvm'` | `.pth` 文件缺少 `3rdparty/tvm/python` 路径 | 修改 `.pth` 文件（见步骤 3） |
 | `Target kind "maca" is not defined` | 未先 `import tilelang`，`libtilelang.so` 未加载 | 先 `import tilelang` 再使用 `Target('maca')` |
 | `Target kind "maca" is not defined`（已 import tilelang） | 编译时未启用 MACA | 用 `USE_MACA=ON` 重新编译（见步骤 2） |
 | 编译报 `ninja: error: Makefile` | `build/` 目录残留旧 Makefile | `rm -rf build` 后重新编译 |
--- a/Agent开发范式的国产GPU大模型推理算子库优化/FlashAttention关键算子迁移与优化.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/FlashAttention关键算子迁移与优化.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/FlashInfer
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/FlashInfer
--- a/Agent开发范式的国产GPU大模型推理算子库优化/FlashInfer关键算子迁移与优化.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/FlashInfer关键算子迁移与优化.md
@ -1,356 +0,0 @@
 # FlashInfer 迁移 Baseline 实战
 ## 一、教程定位
 本教程是参赛训练课程的 **FlashInfer Baseline 入门**模块，主要帮助用户快速跑通 FlashInfer 的最小可运行流程。完成本教程后，用户应能够完成源码编译、API 调用、正确性测试和 Benchmark 测试，并记录一份 baseline 性能结果，为后续算子优化提供对比基准。
 ## 二、学习目标
 1.  理解 FlashInfer Attention Kernel 的基本作用与适用场景；
 2.  完成 FlashInfer 环境、工具链的准备与源码编译；
 3.  跑通 BatchDecode、BatchPrefill、MLA 等典型算子的 API 调用示例；
 4.  完成各算子在不同参数配置下的 Benchmark 测试；
 5.  输出各算子的 Baseline 性能结果记录表，为后续算子优化提供对比基准。
 ## 三、适用对象
 适合人群
 *   参赛选手：需要完成 Baseline 入门模块，为后续算子优化做准备
 *   软件开发者和Vibe Coding开发：希望从事AI相关行业开发，以及用智能体方式来做开发工作
 *   LLM 推理开发者：希望了解 FlashInfer Attention Kernel 的性能表现
 *   算子优化工程师：希望基于MXMACA软件栈在沐曦国产 GPU 上做算子迁移和优化
 前置基础
 *   Python 基础：能够运行和修改 Python 脚本
 *   PyTorch 基础：了解MXMACA软化栈的使用
 *   Linux 命令行：能够使用终端执行命令
 *   了解 Attention 机制：理解 Q/K/V、KV Cache 等基本概念
 ## 四、前置准备
 开始实战前，请确认你已经完成以下准备：
 ### GPU准备
 * 步骤1：获取算力券[https://developer.metax-tech.com/activities/6](https://developer.metax-tech.com/activities/6)
  *   首次登录需要先进行注册（使用邮箱或者手机号进行注册）
 *   登录成功后进行第二步-邮箱验证，填入自己的邮箱。
  *   第三步，提交申请。
  *   获得兑换码
 * 步骤2：兑换算力和登陆平台
  *   访问模力方舟官网：[https://ai.gitee.com/](https://ai.gitee.com/)
  *   进入费用中心 - 算力券 ，点击右上角 “兑换”。
 *   步骤3：租用算力
    - 进入算力容器，选择沐曦，租用算力，建议优先选16G显存/32G显存，如下图：
 *   步骤4：创建实例
    - 基础镜像：maca-pytorch:3.7.1.5-torch2.8-py312-ubuntu24.04-amd64
 *   步骤5：选择工具-lab进入实例环境
 *   步骤6：在JupyterLab Terminal中检查运行环境的配置，确认沐曦 GPU 可见--可以使用`mx-smi`命令查看
 ### 环境依赖准备
 ```bash
 pip install flashinfer torch pandas numpy
 ```
 ### opencode Agent安装
 ```shell
 curl -fsSL https://opencode.ai/install | bash
 ```
 说明：
 自动下载最新版本 自动配置环境 适用于 macOS / Linux / WSL 安装完成后，还会提示你怎么用，OpenCode 包含免费模式，使用方式：
 ```shell
 cd   # 进入项目目录
 opencode
 ```
 ### 代码准备
 *   已进入指定项目目录；
 *   准备测试脚本和 Benchmark 脚本。
 | 命令 | 说明 |
 | --- | --- |
 | `python bench_batch_decode.py` | 运行 Batch Decode 基准测试 |
 | `python bench_batch_prefill_paged.py` | 运行 Batch Prefill (Paged KV Cache) 基准测试 |
 | `python bench_batch_prefill_ragged.py` | 运行 Batch Prefill (Ragged KV Cache) 基准测试 |
 | `python bench_batch_mla.py` | 运行 MLA (Multi-head Latent Attention) 基准测试 |
 ## 五、知识预备
 #### LLM推理阶段重要概念：
 *   **Prefill阶段**：prefill 阶段是指处理输入 prompt 的阶段
    *   输入：用户一次性给出的完整 prompt，长度为 seq\_len
    *   计算：对 prompt 中的每个 token 并行计算注意力，生成第一个输出 token 及 KV cache
    *   特点：这是**计算密集型（compute-bound）**阶段，因为需要做完整的 seq\_len x seq\_len 注意力矩阵乘法
 *   **decode阶段**：
    *   每次只生成 1 个 token，利用 prefill 阶段填充好的 KV cache 做自回归生成
    *   **显存带宽密集型（memory-bound**），瓶颈在从显存读取 KV cache 而非计算
 prefill = 并行处理用户输入，decode = 逐个生成回答 token
 ## 六、项目实践--FlashInfer-Baseline
 ### Step 1：检查运行环境
 **目标：** 进入terminal确认当前环境满足本模块运行要求。
 **操作：** 检查 GPU、Python、编译工具和依赖版本。
 ![giteeai 实例 12](https://origin.picgo.net/2026/06/04/giteeai--12c1772b12867f6be0.png)
 **常见问题：**
 | 问题 | 解决方法 |
 | --- | --- |
 | `mx-smi: command not found` | 使用 `python -c "import torch; print(torch.cuda.get_device_name(0))"` |
 | `No GPUs are available` | 检查 MXMACA 驱动是否正确安装 |
 | `ModuleNotFoundError: No module named 'flashinfer'` | 执行 `pip install flashinfer` 安装 |
 ### Step 2：进入项目目录
 **目标：** 进入本模块所需的源码目录[flashinfer_baseline](https://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%2Fbaselines%2Fflashinfer_baseline)。
 1. 克隆代码仓库
   ```Bash
   git clone https://gitlink.org.cn/metax-maca/op_optimization.git
   ```
 2. 准备flashinfer_baseline
   在仓库目录 `op_optimization/基于AI Agent开发范式的国产GPU大模型推理算子库优化` 下，找到 `flashinfer_baseline` 文件夹。可以将 `flashinfer_baseline` 整个目录复制到工作目录 `data/` 下。
 **下一步操作：** 切换到 FlashInfer Baseline 项目目录。
 **命令示例：**
 ```bash
 cd /data/flashinfer_baseline
 ls -la
 ```
 **预期结果：**
 ```plaintext
 bench_common.py
 bench_batch_decode.py
 bench_batch_prefill_paged.py
 bench_batch_prefill_ragged.py
 bench_batch_mla.py
 README.md
 ...
 ```
 ### Step 3：验证项目脚本
 **目标：** 确认所有基准测试脚本可正常执行。
 **操作：** 检查脚本文件是否存在且可读。
 **命令示例：**
 ```bash
 # 检查脚本文件
 python -c "import os; scripts = ['bench_common.py', 'bench_batch_decode.py', 'bench_batch_prefill_paged.py', 'bench_batch_prefill_ragged.py', 'bench_batch_mla.py']; [print(f'✓ {s}') if os.path.exists(s) else print(f'✗ {s} missing') for s in scripts]"
 # 测试脚本导入
 python -c "from bench_common import setup_workspace, get_csv_path; print('脚本导入正常')"
 ```
 **预期结果：**
 ```plaintext
 ✓ bench_common.py
 ✓ bench_batch_decode.py
 ✓ bench_batch_prefill_paged.py
 ✓ bench_batch_prefill_ragged.py
 ✓ bench_batch_mla.py
 脚本导入正常
 ```
 ### Step 4：运行单算子 Benchmark
 **目标：** 执行基准测试，获取 Baseline 性能数据。
 **操作：** 运行 BatchDecode 基准测试脚本。
 **命令示例：**
 ```bash
 python bench_batch_decode.py
 ```
 **预期结果：**
 ```plaintext
 [BatchDecodeWithPagedKVCacheWrapper] Starting benchmark, total cases: 144
  [1/144] bs=1, kv_len=512, hd=64: 0.020ms, 51.53 GB/s, 0.82 TFLOPs
  [2/144] bs=1, kv_len=512, hd=128: 0.023ms, 45.28 GB/s, 0.72 TFLOPs
  ...
 Results saved to BatchDecodeWithPagedKVCacheWrapper_20260525_xxxxxx.csv
 ```
 **常见问题：**
 | 问题 | 解决方法 |
 | --- | --- |
 | `out of memory` | 减小 batch\_size 或 seq\_len 参数 |
 | 运行时间过长 | 脚本会自动调整重复次数，耐心等待 |
 ### Step 5：查看测试结果
 **目标：** 查看并分析 Benchmark 输出结果。
 **操作：** 读取生成的 CSV 结果文件。
 **命令示例：**
 ```bash
 # 查看最新生成的 CSV 文件
 ls -lt *.csv | head -1
 # 使用 Python 查看结果
 python -c "
 import pandas as pd
 import glob
 csv_files = glob.glob('BatchDecodeWithPagedKVCacheWrapper_*.csv')
 if csv_files:
    df = pd.read_csv(max(csv_files))
    print(df[['batch_size', 'seq_len_kv', 'head_dim', 'time_ms', 'bandwidth_GB_s', 'tflops']].head(10))
 "
 ```
 **预期结果：**
 ```plaintext
   batch_size  seq_len_kv  head_dim   time_ms  bandwidth_GB_s    tflops
 0           1         512        64  0.020429       51.528822  0.821253
 1           1         512       128  0.023340       45.278052  0.718833
 2           1         512       256  0.031949       66.153846  1.050256
 ...
 ```
 ## 七、常见问题
 ### 环境相关问题
 | 问题 | 原因 | 解决办法 |
 | --- | --- | --- |
 | `No GPUs are available` | MXMACA 驱动未安装或 GPU 不可见 | 检查驱动安装，运行 `python -c "import torch; print(torch.cuda.device_count())"` 验证 |
 | `ModuleNotFoundError: No module named 'flashinfer'` | flashinfer 未安装 | 执行 `pip install flashinfer` |
 | `out of memory` | GPU 显存不足 | 减小 `batch_size` 或 `seq_len` 参数 |
 ### 运行相关问题
 | 问题 | 原因 | 解决办法 |
 | --- | --- | --- |
 | Benchmark 运行时间过长 | 参数组合过多， workload 较大 | 耐心等待，脚本会自动调整重复次数 |
 | `KeyError: 'BatchPrefillWithPagedKVCacheKernel'` | profiler 未捕获目标 kernel | 检查 `target_kernels` 配置是否正确 |
 | CSV 文件为空 | 测试未正常完成 | 检查 GPU 显存是否充足，重新运行 |
 ### 代码相关问题
 | 问题 | 原因 | 解决办法 |
 | --- | --- | --- |
 | `ImportError: cannot import name 'xxx' from 'bench_common'` | 函数名拼写错误 | 检查 `bench_common.py` 中的函数名 |
 | `RuntimeError: error: device-side assert triggered` | 输入参数超出范围 | 检查 `num_qo_heads`、`num_kv_heads`、`head_dim` 配置 |
 ### 性能相关问题
 | 问题 | 原因 | 解决办法 |
 | --- | --- | --- |
 | TFLOPs 数值异常低 | 工作负载过小，kernel 启动开销占比大 | 增大 `batch_size` 或 `seq_len` |
 | 带宽数值异常低 | 数据未正确加载到 GPU | 检查 Tensor 是否在 CUDA 设备上 |
 ## 八、下一步学习建议
 ### 1. 保存你的 Baseline 结果
 将本次运行生成的 CSV 文件妥善保存，后续优化时需要以此作为对比基准。
 ```bash
 # 建议创建 results 目录保存
 mkdir -p results
 mv *.csv results/
 ```
 ### 2. 深入理解 FlashInfer 核心概念
 *   阅读 FlashInfer 官方文档，理解 Paged KV Cache、Ragged KV Cache 的设计理念
 *   学习 MLA (Multi-head Latent Attention) 的原理，了解 DeepSeek 的注意力优化方案
 *   理解 `plan()` 和 `run()` 两阶段设计的作用
 ### 3. 进入算子优化模块
 参考后续优化模块，学习以下优化技术：
 *   **Kernel Tuning**：调整 Block Size、Thread Count 等参数
 *   **Memory Optimization**：减少显存占用、优化数据搬运
 *   **Compute Optimization**：提升计算效率
 ### 4. 参考资源
 *   FlashInfer 官方仓库：https://github.com/flashinfer-ai/flashinfer
 *   FlashInfer 文档：https://flashinfer.ai
 ### 5. 记录优化过程
 建议维护一份优化日志，记录每次优化的改动和性能变化：
 | 优化项 | 改动内容 | Baseline | 优化后 | 提升比例 |
 | --- | --- | --- | --- | --- |
 | 例：调整 block\_size | 16 → 32 | xx ms | xx ms | xx% |
 | 完成优化后，再次运行本模块的 Benchmark 脚本，对比前后性能变化。 |  |  |  |  |
--- a/Agent开发范式的国产GPU大模型推理算子库优化/MCTLASS_Fused
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/MCTLASS_Fused
@ -1,52 +1,69 @@
-# Fused MoE Baseline 入门：快速跑通最小闭环教程
+# Fused MoE 算子入门：从 Benchmark 验证到 XPU-OJ 接口提交
 ## 一、教程定位
-本教程是参赛训练课程的 baseline 入门模块，主要帮助用户快速跑通 Fused MoE 的最小可运行流程。
+本教程是赛题二 **Fused MoE** 任务的“benchmark 性能基线与 XPU-OJ 提交衔接”模块，主要帮助学员跑通目标算子的 benchmark 脚本，理解原库 API、输入输出结构、性能指标和性能基线结果，并进一步读懂 XPU-OJ 题目包中的接口约定、测试数据、参考输出和精度要求。
-完成本教程后，用户应能够完成源码编译、正确性测试和 benchmark 测试，并记录一份 baseline 性能结果，为后续算子优化提供对比基准。
+需要特别说明：
 *   本教程不提供可直接提交的标准答案代码。
 *   本教程仅提供冒烟级 starter 示例代码，用于验证环境、语言、提交链路和 `run_kernel(...)` 接口。
 *   benchmark 脚本用于建立性能基线，不是最终提交物。
 *   XPU-OJ 题包中的 `baseline()` 属于 OJ 后台参考实现，用于生成 `output_ref`，不是选手提交代码。
 *   选手最终需要自行实现 `run_kernel(...)`，并在正确性通过后继续优化性能。
 完成本教程后，学员应能够跑通 benchmark 脚本，记录性能基线结果，读懂 XPU-OJ 题包，理解 OJ 的测试输入与参考实现，并完成一次冒烟级 OJ 提交。
 ## 二、学习目标
 完成本模块后，你将能够：
-1.  理解 Fused MoE 推理算子及 int8 量化在大模型推理加速中的核心作用与优化意义；
+1.  理解 Fused MoE 推理算子的基本作用、输入输出和典型应用场景。
-2.  完成环境配置和 Fused MoE baseline 源码准备；
+2.  跑通对应 benchmark 脚本，并记录性能基线结果。
-3.  跑通 Fused MoE 的最小闭环示例；
+3.  学习如何基于 Trition 与 MXMACA C++ 编写 Fused MOE 算子。
-4.  理解基于 Trition 与 MXMACA C++ 写Fused MOE 算子；
+4.  完成数值正确性测试，即验证 reference 计算、pybind 计算、Triton 计算这三种方式计算结果是否数值完全一致。
-5.  完成数值正确性测试，即验证 reference 计算、pybind 计算、Triton 计算这三种方式计算结果是否数值完全一致。
+    *   reference：基于 PyTorch 架构在 CPU 上运行的**数值基准**实现。
    *   reference：基于 PyTorch 架构在 CPU 上运行的**数值基准**实现；
-    *   pybind：将 MXMACA C++ 算子编译并封装为 Python 可调用的动态库，**实现复杂且迁移成本高**；
+    *   pybind：将 MXMACA C++ 算子编译并封装为 Python 可调用的动态库，**实现复杂且迁移成本高**。
-    *   Triton：基于 Python 编写的高效 GPU Kernel，可利用 Agent 自动调优，**开发效率高、易于迁移**；
+    *   Triton：基于 Python 编写的高效 GPU Kernel，可利用 Agent 自动调优，**开发效率高、易于迁移**。
    *   要求 pybind 和 Triton 结果均与 reference 一致，鼓励参赛选手持续调优 Triton ，使其性能逼近甚至超越 pybind 性能。
 5.  区分 benchmark 性能基线、OJ 参考实现和选手提交代码。
-6.  输出 benchmark 结果对比表。
+6.  读懂对应 XPU-OJ 题包中的题目描述、接口约定、数据范围和精度要求。
 7.  完成一次冒烟级 `run_kernel(...)` 提交，确认 OJ 链路、语言环境和接口调用正常。
 8.  使用 AI Agent 辅助阅读题包、生成初版实现、定位错误并规划性能优化方向。
 ## 三、适用对象
 **本模块适合以下人员：**
-*   参与基于 AI Agent 开发范式的国产 GPU 大模型推理算子库优化比赛的学生；
+*   参与基于 AI Agent 开发范式的国产 GPU 大模型推理算子库优化比赛的学生。
-*   对 GPU 推理算子性能优化感兴趣的开发者；
+*   对 GPU 推理算子性能优化感兴趣的开发者。
 *   需要了解 Fused MoE 推理性能的研究人员。
 **学习本模块前，需掌握以下基础知识：**
-*   Python、C++ 编程基础；
+*   Python、C++ 编程基础。
-*   PyTorch 基础；
+*   PyTorch 基础。
 *   GPU 推理基本概念。
@ -58,21 +75,25 @@
 **环境准备：**
 *   已进入赛事专属镜像环境。
 **工具准备：**
-*   已准备 Agent 工具；
+*   已准备 Agent 工具；
-*   已配置 Token / API Key；
+    
-*   已确认 Agent 可以正常调用模型。
+*   已配置 Token / API Key；
 *   已确认 Agent 可以正常调用模型。
 **代码准备：**
-*   已获取 Fused MoE Baseline 源码。
+*   已获取 Fused MoE 源码。
 ## 五、项目实践1-算力平台进行 Benchmark 验证
-## 五、项目实践1-算力平台 baseline 源码部署
+**项目目标：**跑通 Fused MoE 算子的 benchmark 脚本，建立性能基线，为后续算子优化提供对比基准。
 **项目目标：**在算力平台上拉取 Fused MoE 的 baseline 源码，快速跑通最小闭环，并记录一份 baseline 性能结果，为后续算子优化提供对比基准。
 ### 步骤 1：检查运行环境
@ -106,7 +127,6 @@ else:
 EOF
 ```
 **预期结果：**
 *   Python 3.12.11
@ -123,39 +143,23 @@ EOF
 | 报错 | 原因 | 解决办法 |
 | --- | --- | --- |
 | `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+) |
+| `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：进入项目目录
-**目标：**进入本模块所需的源码目录：https://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%2Fbaselines%2Ffused_moe
+**目标：**进入本模块所需的源码目录。
 1. 克隆代码仓库
   ```Bash
   git clone https://gitlink.org.cn/metax-maca/op_optimization.git
   ```
 2. 准备fused_moe
   在仓库目录 `op_optimization/基于AI Agent开发范式的国产GPU大模型推理算子库优化` 下，找到 `fused_moe` 文件夹。可以将 `fused_moe` 整个目录复制到工作目录 `data/` 下。
 **下一步操作：** 切换到 FlashInfer Baseline 项目目录。
 **操作：**切换到指定项目路径。
 **命令示例：**
 ```apl
-cd data/fused_moe # 这里需要替换成自己的项目路径
+#克隆代码仓库
 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
 ```
 **预期结果：**
 终端提示符路径显示为 fused\_moe 项目路径：
 *   (base) root@0a5a9d0c0f06:/data/fused_moe#
 ### 步骤 3：pybind 编译
@ -173,21 +177,21 @@ bash scripts/build_fused_moe_i8_tn_pybind.sh
 切换 Python 环境命令示例：
 ```apl
-[SUCCESS] /data/fused_moe/standalone/fused_moe_i8_tn/build/fused_moe_i8_tn_pybind.sh
+PYTHON_BIN=/path/to/python bash scripts/build_fused_moe_i8_tn_pybind.sh
 ```
 **预期结果：**
 编译成功无报错，终端显示：
-*   \[SUCCESS] /data/fused_moe/standalone/fused_moe_i8_tn/build/fused_moe_i8_tn_pybind.so
+  *   \[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` 文件
+      且成功生成 `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')` |
@ -195,10 +199,10 @@ bash scripts/build_fused_moe_i8_tn_pybind.sh
 | `recompile with -fPIC` | 编译未开启位置无关代码 | 确保 `mxcc`/ `g++`编译参数中有 `-fPIC` |
 | `permission denied` | 无脚本执行权限 | `chmod +x scripts/*.sh` |
 | `undefined reference to Py_...` | Python 版本不匹配 | 确认编译脚本中`PYTHON_BIN`路径与当前运行的 Python 环境完全一致 |
-
+    
 ### 步骤 4：正确性测试
-
+    
-**目标：**验证 reference 计算、pybind 计算、Triton 计算这三种方式计算结果的数值是否完全一致。
+**目标：**验证 reference 计算、pybind 计算、Triton 计算这三种方式计算结果的数值是否一致。
 **操作：**运行 `fused_moe/scripts/run_fused_moe_i8_tn_pybind_test.sh` 脚本
@ -215,49 +219,54 @@ 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\_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\_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\_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\_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\_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\_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”：三种计算方式；
+  *   “pybind/reference/Triton”：三种计算方式；
-    
+        
-*   “fused\_moe\_i8\_tn\_topk1/2/3 passed”：测试算子通过数值校验，数值误差在允许范围内且无明显异常，否则会报错 FAILED；
+  *   “fused\_moe\_i8\_tn\_topk1/2/3 passed”：测试算子通过数值校验，数值误差在允许范围内且无明显异常，否则会报错 FAILED；
-    
+        
-*   ”rows=... , cols=...“：输出 Tensor 的行和列的大小；
+  *   ”rows=... , cols=...“：输出 Tensor 的形状；
-    
+        
-*   ”sample C\[0\]=... , C\[last\]=...“：首尾采样值，用于辅助定位数值偏差，不作为精度判定依据。
+  *   ”sample C\[0\]=... , C\[last\]=...“：首尾采样值，用于辅助定位数值偏差，不作为精度判定依据。
-    
+        
 **常见问题：**
-
+    
 | 报错 | 原因 | 解决办法 |
-| :--- | :--- | :--- |
+| --- | --- | --- |
-| `ModuleNotFoundError: fused_moe_i8_tn_pybind` | pybind 模块未编译或未加入 `PYTHONPATH` | 回到步骤 3，确认 `.so` 已生成；执行 `export PYTHONPATH=/root/Project/fused_moe:$PYTHONPATH` |
+| `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: max abs diff too large` | 数值误差超过阈值 | 检查 scale 是否应用位置错误；确认 TopK 索引与权重是否一致 |
-| `FAILED: shape mismatch` | 输出张量形状不一致 | 检查 Token Permute / Unpermute 逻辑；确认 expert 维度对齐 |
+| `FAILED: shape mismatch` | 输出张量形状不一致 | 检查 Token Permute / Unpermute 逻辑；确认 expert 维度对齐 |
-| `FAILED: NaN or Inf detected` | 溢出或未初始化内存 | 检查 INT8 乘加是否溢出；确认 GEMM 输出是否反量化 |
+| `FAILED: NaN or Inf detected` | 溢出或未初始化内存 | 检查 INT8 乘加是否溢出；确认 GEMM 输出是否反量化 |
-| 终端长时间无输出 | Kernel 死锁或 Launch 失败 | 减小测试 shape；检查是否触发 MACA 硬件限制 |
+| 终端长时间无输出 | Kernel 死锁或 Launch 失败 | 减小测试 shape；检查是否触发 MACA 硬件限制 |
-
+    
 ### 步骤5：性能测试
-
+    
 **目标：**输出 benchmark 结果对比表
 **操作：**运行 `fused_moe/scripts/run_fused_moe_i8_tn_benchmark.sh` 脚本
@ -270,40 +279,43 @@ bash scripts/run_fused_moe_i8_tn_benchmark.sh --backend all --warmup 5 --iters 2
 # --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\_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\_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”：平均算子执行耗时（毫秒），这里不计算预热时间，只计算正式迭代的<EFBFBD><EFBFBD><EFBFBD>间<EFBFBD><EFBFBD><EFBFBD>
+*   “avg\_ms”：平均算子执行耗时（毫秒），这里不计算预热时间，只计算正式迭代的时间；
-    
+        
-*   “TOPS”：Tera Operations Per Second，本次 MoE 算子的总运算量 / 实际耗<EFBFBD><EFBFBD><EFBFBD>；
+*   “TOPS”：Tera Operations Per Second，本次 MoE 算子的总运算量 / 实际耗时；
-    
+        
 *   “warmup=5, iters=20”：预热轮数和正式迭代数。
-    
+        
 **常见问题：**
 **常见错误：**
 | 报错 | 原因 | 解决办法 |
 | --- | --- | --- |
@ -311,13 +323,490 @@ bash scripts/run_fused_moe_i8_tn_benchmark.sh --backend all --warmup 5 --iters 2
 | 终端长时间无输出 | Kernel 死锁或 MACA 驱动异常 | 减小测试 shape；重启容器或设备 |
 | avg\_ms 异常抖动（±50%） | 其他进程占用 GPU | 关闭其他占用显存的进程，单机单任务运行 |
-## 六、项目实践2-Kernel Swift 智能算子迁移系统自动调优
+## 六、XPU-OJ 冒烟提交
 ### 6.1 目标
 在完成前文的本地验证后，本节将带你把实现提交到 XPU-OJ，并确认评测环境能够正确调用 `run_kernel(...)`。
 完成本节后，你应能够完成一次最小 OJ 提交，查看评测结果，并据此进入正确性修复或后续性能优化。
 本节的冒烟提交只用于验证函数接口、索引逻辑和提交流程；性能优化请在正确性通过后再进行。
 ### 6.2 前置准备
 #### 6.2.1 代码准备
 建议在工作目录下保留一个候选版本目录，例如：
 ```bash
 mkdir -p oj/problem_1_fused_moe
 ```
 本教程建议把 OJ 提交代码先保存为：
 ```text
 oj/problem_1_fused_moe/solution001.py
 ```
 真正提交时，只需要把这个文件里的内容复制到 XPU-OJ 提交框。
 #### 6.2.2 账号准备
 XPU-OJ 账号由组委会统一发放。登录入口：
 ```text
 https://xpuoj.com/
 ```
 如果登录后看不到比赛或题目，请联系助教或赛事运营确认账号是否已经加入对应比赛或用户组。
 ### 6.3 知识预备
 #### 6.3.1 什么是 OJ
 OJ 可以理解为“自动评测机”。
 你提交代码后，OJ 会自动完成：
 1.  加载你的代码；
 2.  构造测试输入；
 3.  调用你的 `run_kernel(...)`；
 4.  生成参考答案；
 5.  对比你的输出和参考输出；
 6.  返回评测状态、耗时、内存和分数。
 所以，OJ 不是让你提交 benchmark 日志，也不是让你提交本地运行截图，而是让你提交一份符合接口约定的代码。
 #### 6.3.2 什么是 Candidate
 Candidate 就是一次可复现的候选方案。
 建议每一轮都记录：
 | 记录项 | 示例 |
 | --- | --- |
 | 候选编号 | candidate-001 |
 | 代码文件 | `oj/problem_1_fused_moe/solution001.py` |
 | 本地检查结果 | local check passed |
 | OJ 结果 | WA / RE / AC |
 | 备注 | 初始冒烟版，只验证接口 |
 这样后续多次打榜时，不会忘记哪一版代码对应哪一次提交结果。
 ### 6.4 项目实践：Fused MoE GEMM OJ 冒烟提交
 本节以当前 XPU-OJ 题目 **1. Fused MoE GEMM** 为例。
 题目要求你提交一份 Python 代码，并提供固定函数：
 ```python
 def run_kernel(
    a,
    b_col_major,
    scale_a,
    scale_b,
    moe_weights,
    token_ids,
    expert_ids,
    topk,
    out,
 ):
    ...
 ```
 评测程序会调用这个函数，并检查你是否把结果正确写入 `out`。
 #### 步骤 1：理解本题要算什么
 本题计算的是固定规格 `fused_moe_i8_tn`：
 ```text
 N = 128
 K = 128
 EM = num_tokens * topk
 EM % 128 == 0
 ```
 核心公式是：
 ```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`。
 #### 步骤 2：准备 OJ 提交文件
 在终端中创建目录：
 ```bash
 cd /data/fusedmoe_v2.1
 mkdir -p oj/problem_1_fused_moe
 ```
 新建文件：
 ```bash
 touch oj/problem_1_fused_moe/solution001.py
 ```
 下一步不是手写一份固定答案，而是让 Agent 根据题面生成一份“最小冒烟版”。
 注意：冒烟版的目标只是确认接口、索引和 OJ 提交流程，不追求性能最优。
 #### 步骤 3：让 Agent 生成 OJ 冒烟代码
 在 OpenCode 或其他 Agent 中，把题目的接口约定、输入输出说明和下面这段 Prompt 一起发给 Agent：
 ```text
 我正在做 XPU-OJ 的 Fused MoE GEMM 题，需要生成一个最小冒烟提交版本。
 请根据下面接口写一份完整 Python 代码：
 def run_kernel(a, b_col_major, scale_a, scale_b, moe_weights, token_ids, expert_ids, topk, out):
    ...
 题目语义：
 1. N = 128，K = 128；
 2. EM = num_tokens * topk，且 EM 是 128 的倍数；
 3. token(r) = token_ids[r] // topk；
 4. expert(r) = expert_ids[r // 128]；
 5. b_col_major 的布局是 [expert, n, k]；
 6. 结果必须原地写入 out；
 7. out 的 dtype 是 bfloat16；
 8. 正确性优先，不需要优化性能。
 代码要求：
 1. 函数名和参数顺序必须完全一致；
 2. 不要添加 torch.Tensor 类型注解；
 3. 不要依赖外部文件；
 4. 不要打印调试信息；
 5. 不要返回新 tensor，只写入 out；
 6. 请输出一份可以直接复制到 XPU-OJ 提交框的完整代码。
 ```
 Agent 生成后，把代码保存到：
 ```text
 oj/problem_1_fused_moe/solution001.py
 ```
 保存前建议人工检查 5 个点：
 | 检查项 | 应该满足 |
 | --- | --- |
 | 函数名 | 必须是 `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]` |
 | 输出方式 | 必须原地写入 `out` |
 如果 OJ 页面要求选择 `Triton` 或 `Triton Maca` 语言，也可以先让 Agent 生成一个语义正确的冒烟版本。正确性通过后，再让 Agent 把核心计算替换成 Triton kernel 做性能优化。
 #### 步骤 4：提交到 XPU-OJ
 1.  打开 XPU-OJ： ```texthttps://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 GEMM```
 5.  点击题目进入详情页；
 6.  在提交区域选择本题支持的语言，例如： ```textTriton / Triton Maca```
 ```plaintext
 具体名称以页面下拉框为准。
 ```
 1.  将 `solution001.py` 中的代码复制到提交框；
 2.  点击提交；
 3.  等待评测结果返回。
 #### 步骤 5：查看 OJ 结果
 提交后，进入：
 ```text
 我的提交
 ```
 常见状态含义如下：
 | 状态 | 含义 | 下一步 |
 | --- | --- | --- |
 | Accepted / AC | 正确性通过 | 可以继续优化性能 |
 | Wrong Answer / WA | 输出数值不对 | 检查索引、shape、dtype、缩放系数 |
 | Runtime Error / RE | 运行时报错 | 点开详情看报错栈 |
 | Compile Error / CE | 编译或加载失败 | 检查 import、语法、函数名 |
 | Time Limit Exceeded / TLE | 超时 | 说明代码太慢，需要做 kernel 优化 |
 如果看到 `0 pts`，通常表示本次提交没有拿到分数。原因可能是：
 *   样例没过；
 *   测试点没过；
 *   代码运行时报错；
 *   代码超时；
 *   输出与参考答案超过容差。
 如果看到用时和内存都是 `0`，很多时候说明代码在正式计时前就失败了，例如函数签名不匹配、导入失败或编译失败。
 #### 步骤 6：理解 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 允许很小的数值误差，但不是随便差一点都能过。
 #### 步骤 7：保存 Candidate
 建议每一次能跑的版本都用 Git 保存。
 ```bash
 cd /data/fusedmoe_v2.1
 git status --short
 git add oj/problem_1_fused_moe/solution001.py
 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.py oj/problem_1_fused_moe/solution002.py
 ```
 然后让 Agent 基于 `solution002.py` 继续改。
 ### 6.5 Agent 使用说明
 本模块中，Agent 主要用来做三件事：
 1.  读题目接口；
 2.  生成最小可提交代码；
 3.  根据 OJ 报错定位问题。
 建议不要一开始就让 Agent “直接写最快版本”。更稳的流程是：
 ```text
 第一步：先写一个能过正确性的最小版本。
 第二步：提交 OJ，看是否 AC。
 第三步：AC 后再优化性能。
 ```
 可以使用下面的 Prompt：
 ```text
 我正在做 XPU-OJ 的 Fused MoE GEMM 题。
 请只做一件事：根据题目接口写一个最小正确的 run_kernel 冒烟版本。
 要求：
 1. 函数签名必须完全一致：
   run_kernel(a, b_col_major, scale_a, scale_b, moe_weights, token_ids, expert_ids, topk, 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. 不要添加 torch.Tensor 类型注解
 8. 不要依赖外部文件
 9. 请输出完整可复制提交的 Python 代码
 ```
 如果 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。
 这是错误日志：[粘贴错误日志]
 请先判断是函数签名、import、RestrictedPython 限制、dtype、shape 还是 GPU 算子调用问题。
 只给出最小修复方案。
 ```
 ### 6.8 常见问题
 #### Q1：为什么本地能跑，OJ 上却 Runtime Error？
 本地环境和 OJ 沙箱不完全一样。OJ 可能限制某些 Python 写法、外部文件访问或动态编译行为。
 常见例子：
 ```python
 def silu(x: torch.Tensor) -> torch.Tensor:
    ...
 ```
 这种类型注解可能触发：
 ```text
 Access to torch.Tensor is not allowed
 ```
 处理方式：去掉 `torch.Tensor` 类型注解。
 #### Q2：为什么 OJ 是 Wrong Answer？
 优先检查四个点：
 1.  `token_ids[r]` 是否先除以 `topk`；
 2.  `expert_ids` 是否按 `r // 128` 取；
 3.  `b_col_major` 是否按 `[expert, n, k]` 理解；
 4.  结果是否写回 `out`，而不是只返回一个新 tensor。
 #### Q3：为什么冒烟代码很慢？
 冒烟代码的目标是确认接口正确，不是追求性能。
 如果它能过正确性，但耗时很高，这是正常的。下一步才是把核心计算替换成 Triton kernel 或其他更快的 GPU 实现。
 #### Q4：50 分、10 分是什么意思？
 不同比赛和题目的评分规则可能不同。一般可以先这样理解：
 *   正确性没过时，通常拿不到有效分数；
 *   正确性通过后，平台会继续根据耗时或加速比计算分数；
 *   具体分数含义以 XPU-OJ 当前题目的评分说明为准。
 #### Q5：榜单怎么看？
 先看自己的提交是否通过正确性，再看耗时和分数。
 建议记录：
 | Candidate | OJ 状态 | 用时 | 分数 | 备注 |
 | --- | --- | --- | --- | --- |
 | candidate-001 | AC / WA / RE | 以页面为准 | 以页面为准 | 冒烟版 |
 | candidate-002 | AC / WA / RE | 以页面为准 | 以页面为准 | 第一轮优化 |
 不要只看单次结果。每轮都记录，后面才知道 Agent 的修改到底有没有带来收益。
 ### 6.9 从 Benchmark 验证到参赛作品的路径回顾
 建议按下面顺序推进：
 1.  跑通 benchmark 脚本，理解算子输入输出；
 2.  阅读 XPU-OJ 题目页面，确认 `run_kernel(...)` 接口；
 3.  提交冒烟代码，确认 OJ 链路正常；
 4.  如果冒烟代码 WA / RE，先修正确性；
 5.  正确性通过后，再让 Agent 生成 Triton kernel 优化版本；
 6.  每一轮提交都保存 candidate、prompt、代码 diff 和 OJ 结果；
 7.  用 OJ 分数和耗时判断优化是否有效。
 ```text
 Benchmark 验证代码用来学习，OJ 用来评分，Candidate 用来管理每一轮结果。
 ```
 ## 七、项目实践2-Kernel Swift 智能算子迁移系统自动调优
 系统链接：[https://deeplink.org.cn/kernelswift/task](https://deeplink.org.cn/kernelswift/task)
-**项目目标：**基于 KernelSwift 智能算子迁移系统，对 Fused MoE 算子进行在线自动调优。通过输入算子代码，一键生成适配沐曦硬件的高性能实现，高效完成算子优化与全流程追踪。
+**项目目标：**基于 KernelSwift 智能算子迁移系统，对 Fused MoE 算子进行在线自动调优。通过输入算子的 PyTorch 代码，一键生成适配沐曦硬件的高性能实现，高效完成算子优化与全流程追踪。
-### 步骤1：复用算子广场的 fused\_moe 算子进行二次优化
+### 步骤1：复用算子广场的Fused MoE 算子进行二次优化
 **目标：**通过提交算子广场的 fused\_moe 算子代码发起自动优化流程，实现二次优化
@ -418,19 +907,19 @@ bash scripts/run_fused_moe_i8_tn_benchmark.sh --backend all --warmup 5 --iters 2
    *   选择适配硬件：算子需要适配的目标硬件厂商及型号，建议：沐曦
    *   最大演化轮次：优化算法迭代次数，取值范围40-400，建议默认40，复杂算法可提高至100+
-    
+        
 5.  提交优化任务：点击右下角 \[优化\] 按钮，系统将提交任务并进入 \[生成中\] 状态
-    ![image.png](https://origin.picgo.net/2026/06/04/image659a3f15b96be22c1.png)
+    [![image1](https://origin.picgo.net/2026/06/23/image1d46e08e5a17fd767.png)](https://www.picgo.net/image/image1.4SHrb4)
-  完成上述步骤将看到如下界面：
+完成上述步骤将看到如下界面：
-  ![image.png](https://origin.picgo.net/2026/06/04/image75c4ccf75e90cf428.png)
+  [![image2](https://origin.picgo.net/2026/06/23/image268924dc11f138788.png)](https://www.picgo.net/image/image2.4SHscu)
 ### 步骤2：任务查看与结果管理
-**目标：** 在新建优化任务后可追踪任务进度，获取优化结果
+**目标：**在新建优化任务后可追踪任务进度，获取优化结果
 **操作：**
@ -443,7 +932,7 @@ bash scripts/run_fused_moe_i8_tn_benchmark.sh --backend all --warmup 5 --iters 2
    *   操作按钮：查看详情、删除任务
-    ![image.png](https://origin.picgo.net/2026/06/04/image75c4ccf75e90cf428.png)
+    [![image3](https://origin.picgo.net/2026/06/23/image33091601c9a68bd18.png)](https://www.picgo.net/image/image3.4SHDeY)
 2.  追踪任务进度：当前任务状态为【运行中】时，点击任务列表中的【查看详情】按钮，追踪任务进度：
@ -463,7 +952,7 @@ bash scripts/run_fused_moe_i8_tn_benchmark.sh --backend all --warmup 5 --iters 2
    *   顶部：任务名称、创建/更新时间、适配硬件、当前轮次进度
-    ![image.png](https://origin.picgo.net/2026/06/04/image9e9b23d9b25ddfe25.png)
+    [![image4](https://origin.picgo.net/2026/06/23/image403daf417d165a79f.png)](https://www.picgo.net/image/image4.4SHVpp)
 3.  获取优化结果：当前任务状态为【已完成】时，可在详情页查看优化结果：
@ -474,23 +963,23 @@ bash scripts/run_fused_moe_i8_tn_benchmark.sh --backend all --warmup 5 --iters 2
    *   可点击【Diff 对比】查看优化前后代码差异，理解性能提升逻辑
-    ![image.png](https://origin.picgo.net/2026/06/04/image1083021a27904294b1.png)
+    [![image5](https://origin.picgo.net/2026/06/23/image58f2b2ea36dad2ef0.png)](https://www.picgo.net/image/image5.4ScbBr)
 4.  任务异常处理
    *   任务失败：查看错误日志，常见原因包括代码不符合规范、测试用例错误、硬件适配问题，修改后重新提交任务；
    *   排队时间长：可调整提交时间，或联系平台管理员确认资源状态。
-
+## 八、Agent使用说明
 ## 七、Agent使用说明
 在本模块中，Agent可以帮助你完成以下任务：
 1.  **环境检查**
    ```plaintext
-    我正在算力平台部署 fused_moe_baseline 源码。
+    我正在算力平台进行 Fused MoE 的 Benchmark 验证。
    需要的环境信息如下：
    - Python 3.12
    - g++ 13.3.0
@ -551,16 +1040,16 @@ bash scripts/run_fused_moe_i8_tn_benchmark.sh --backend all --warmup 5 --iters 2
    请帮我在算子广场检索 fused_moe 算子
    ```
-## 八、常见问题与注意事项
+## 九、常见问题与注意事项
-### 算力平台 basline 源码部署项目：
+### 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`目录；
@ -568,7 +1057,7 @@ bash scripts/run_fused_moe_i8_tn_benchmark.sh --backend all --warmup 5 --iters 2
    *   若`libpython not found`，请直接指定 Conda 下的`**libpython3.x.so**`绝对路径，避免链接系统静态库；
    *   编译 `pybind`模块时，务必开启 `-fPIC`，否则会出现 `recompile with -fPIC`错误。
-    
+        
 3.  性能测试建议
    *   benchmark 应在关闭其他占用 GPU 的任务后执行，避免干扰性能数据；
@ -578,7 +1067,7 @@ bash scripts/run_fused_moe_i8_tn_benchmark.sh --backend all --warmup 5 --iters 2
    *   性能对比应基于相同随机种子、相同 shape、相同 TopK、相同 batch size的条件下进行，降低误差。
-### Kernel Swift 智能算子迁移系统自动调优项目：
+### 9.2 Kernel Swift 智能算子迁移系统自动调优项目
 1.  代码规范问题
@ -589,26 +1078,16 @@ bash scripts/run_fused_moe_i8_tn_benchmark.sh --backend all --warmup 5 --iters 2
    *   `def get_init_inputs`，表示 module init 的输入测试样例；
    *   `def get_inputs`，表示 module forward 的输入测试样例。
-    
+        
 2.  性能优化建议
    *   对于复杂算子，可适当提高最大演化轮次（如 100-200），获得更高加速比；
    *   优先选择算子广场中已有优化案例的算子类型，降低适配失败概率。
-    
+        
 3.  硬件适配问题
    *   提交任务前确认目标硬件支持的算子类型；
    *   优化失败时，可尝试更换适配硬件，或调整算子实现逻辑。
 ## 九、下一步学习建议
 完成本模块后，建议继续学习以下内容：
 1.  **研读 fused\_moe源码：**理解代码的底层逻辑，可尝试修改 `build_fused_moe_i8_tn_pybind.sh`中的编译参数，观察其对 `avg_ms` 的影响；
 2.  **算子优化基础：**了解如何分析 Kernel 性能瓶颈；
 3.  **性能对比分析：**将 baseline 结果与优化后的结果进行对比分析，明确后续优化方向。
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashattn_baseline/__MACOSX/._benchmark_kvcache.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashattn_baseline/__MACOSX/._benchmark_kvcache.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashattn_baseline/__MACOSX/._benchmark_kvcache_1.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashattn_baseline/__MACOSX/._benchmark_kvcache_1.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashattn_baseline/__MACOSX/._benchmark_kvcache_20260526_150953.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashattn_baseline/__MACOSX/._benchmark_kvcache_20260526_150953.csv
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/README.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/README.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/benchmark_kvcache.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/benchmark_kvcache.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/benchmark结果实例/benchmark_kvcache_headdim128.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/benchmark结果实例/benchmark_kvcache_headdim128.csv
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/benchmark结果实例/benchmark_kvcache_headdim160.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/benchmark结果实例/benchmark_kvcache_headdim160.csv
@ -0,0 +1,49 @@
 batch_size,seq_len_kv,heads,headdim,time_ms,bandwidth_GB_s
 1,512,8,160,0.0580,45.25
 2,512,8,160,0.0611,85.84
 4,512,8,160,0.0656,159.96
 8,512,8,160,0.0699,300.40
 16,512,8,160,0.1321,317.92
 32,512,8,160,0.2002,419.52
 64,512,8,160,0.3383,496.43
 128,512,8,160,0.6669,503.61
 1,1024,8,160,0.1129,46.45
 2,1024,8,160,0.1190,88.18
 4,1024,8,160,0.1224,171.49
 8,1024,8,160,0.1287,326.07
 16,1024,8,160,0.2479,338.49
 32,1024,8,160,0.3767,445.54
 64,1024,8,160,0.6419,523.01
 128,1024,8,160,1.2804,524.37
 1,2048,8,160,0.2270,46.20
 2,2048,8,160,0.2299,91.22
 4,2048,8,160,0.2349,178.63
 8,2048,8,160,0.2447,342.96
 16,2048,8,160,0.4773,351.60
 32,2048,8,160,0.7279,461.07
 64,2048,8,160,1.2559,534.49
 128,2048,8,160,2.5613,524.15
 1,4096,8,160,0.4460,47.02
 2,4096,8,160,0.4513,92.94
 4,4096,8,160,0.4593,182.64
 8,4096,8,160,0.4813,348.64
 16,4096,8,160,0.9363,358.43
 32,4096,8,160,1.4552,461.21
 64,4096,8,160,2.5615,524.05
 128,4096,8,160,5.1420,522.11
 1,8192,8,160,0.8847,47.41
 2,8192,8,160,0.8944,93.80
 4,8192,8,160,0.9094,184.51
 8,8192,8,160,0.9625,348.64
 16,8192,8,160,1.8550,361.80
 32,8192,8,160,2.9567,453.97
 64,8192,8,160,5.1398,522.30
 128,8192,8,160,10.2972,521.41
 1,16384,8,160,1.7608,47.64
 2,16384,8,160,1.7786,94.33
 4,16384,8,160,1.8143,184.95
 8,16384,8,160,1.9317,347.42
 16,16384,8,160,3.7301,359.83
 32,16384,8,160,5.9216,453.33
 64,16384,8,160,10.2668,522.94
 128,16384,8,160,20.6062,521.09
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/benchmark结果实例/benchmark_kvcache_headdim192.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/benchmark结果实例/benchmark_kvcache_headdim192.csv
@ -0,0 +1,49 @@
 batch_size,seq_len_kv,heads,headdim,time_ms,bandwidth_GB_s
 1,512,8,192,0.0458,68.82
 2,512,8,192,0.0515,122.32
 4,512,8,192,0.0574,219.28
 8,512,8,192,0.0607,414.80
 16,512,8,192,0.1147,439.27
 32,512,8,192,0.1763,571.40
 64,512,8,192,0.2978,676.79
 128,512,8,192,0.5874,686.11
 1,1024,8,192,0.0946,66.55
 2,1024,8,192,0.1033,121.85
 4,1024,8,192,0.1073,234.66
 8,1024,8,192,0.1131,445.23
 16,1024,8,192,0.2165,465.24
 32,1024,8,192,0.3347,601.80
 64,1024,8,192,0.5701,706.63
 128,1024,8,192,1.1302,712.88
 1,2048,8,192,0.1943,64.79
 2,2048,8,192,0.1992,126.38
 4,2048,8,192,0.2059,244.52
 8,2048,8,192,0.2174,463.13
 16,2048,8,192,0.4202,479.24
 32,2048,8,192,0.6503,619.36
 64,2048,8,192,1.1158,721.93
 128,2048,8,192,2.2250,724.05
 1,4096,8,192,0.3834,65.65
 2,4096,8,192,0.3904,128.95
 4,4096,8,192,0.4043,249.04
 8,4096,8,192,0.4267,471.92
 16,4096,8,192,0.8271,486.90
 32,4096,8,192,1.2840,627.28
 64,4096,8,192,2.2148,727.29
 128,4096,8,192,4.3819,735.21
 1,8192,8,192,0.7566,66.52
 2,8192,8,192,0.7712,130.54
 4,8192,8,192,0.7974,252.49
 8,8192,8,192,0.8433,477.47
 16,8192,8,192,1.6433,490.09
 32,8192,8,192,2.5573,629.84
 64,8192,8,192,4.3785,735.73
 128,8192,8,192,8.7303,737.99
 1,16384,8,192,1.5068,66.81
 2,16384,8,192,1.5350,131.16
 4,16384,8,192,1.5868,253.76
 8,16384,8,192,1.6778,479.99
 16,16384,8,192,3.2750,491.81
 32,16384,8,192,5.0659,635.88
 64,16384,8,192,8.7435,736.85
 128,16384,8,192,17.5040,736.13
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/benchmark结果实例/benchmark_kvcache_headdim224.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/benchmark结果实例/benchmark_kvcache_headdim224.csv
@ -0,0 +1,49 @@
 batch_size,seq_len_kv,heads,headdim,time_ms,bandwidth_GB_s
 1,512,8,224,0.1254,29.29
 2,512,8,224,0.1412,52.05
 4,512,8,224,0.1497,98.13
 8,512,8,224,0.1533,191.70
 16,512,8,224,0.1913,307.17
 32,512,8,224,0.3292,357.08
 64,512,8,224,0.5187,453.28
 128,512,8,224,0.9522,493.84
 1,1024,8,224,0.2727,26.93
 2,1024,8,224,0.2836,51.78
 4,1024,8,224,0.2890,101.63
 8,1024,8,224,0.2959,198.55
 16,1024,8,224,0.3696,317.93
 32,1024,8,224,0.6408,366.75
 64,1024,8,224,1.0081,466.21
 128,1024,8,224,1.8548,506.78
 1,2048,8,224,0.5515,26.63
 2,2048,8,224,0.5575,52.67
 4,2048,8,224,0.5666,103.65
 8,2048,8,224,0.5803,202.42
 16,2048,8,224,0.7250,324.05
 32,2048,8,224,1.2593,373.14
 64,2048,8,224,1.9890,472.48
 128,2048,8,224,3.6905,509.28
 1,4096,8,224,1.0939,26.84
 2,4096,8,224,1.1044,53.18
 4,4096,8,224,1.1219,104.69
 8,4096,8,224,1.1500,204.26
 16,4096,8,224,1.4390,326.48
 32,4096,8,224,2.4992,375.97
 64,4096,8,224,4.0082,468.86
 128,4096,8,224,7.3372,512.26
 1,8192,8,224,2.1775,26.97
 2,8192,8,224,2.1989,53.41
 4,8192,8,224,2.2338,105.15
 8,8192,8,224,2.3268,201.90
 16,8192,8,224,2.8806,326.18
 32,8192,8,224,5.0187,374.43
 64,8192,8,224,8.0323,467.90
 128,8192,8,224,14.6300,513.78
 1,16384,8,224,4.3360,27.09
 2,16384,8,224,4.3820,53.60
 4,16384,8,224,4.5006,104.38
 8,16384,8,224,4.6987,199.96
 16,16384,8,224,5.7361,327.59
 32,16384,8,224,10.1291,371.03
 64,16384,8,224,16.0745,467.60
 128,16384,8,224,OOM,OOM
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/benchmark结果实例/benchmark_kvcache_headdim256.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/benchmark结果实例/benchmark_kvcache_headdim256.csv
@ -0,0 +1,49 @@
 batch_size,seq_len_kv,heads,headdim,time_ms,bandwidth_GB_s
 1,512,8,256,0.0877,47.89
 2,512,8,256,0.0921,91.17
 4,512,8,256,0.0940,178.74
 8,512,8,256,0.0964,348.52
 16,512,8,256,0.1450,463.27
 32,512,8,256,0.2250,597.21
 64,512,8,256,0.3609,744.43
 128,512,8,256,0.6932,775.25
 1,1024,8,256,0.1747,48.04
 2,1024,8,256,0.1762,95.27
 4,1024,8,256,0.1784,188.22
 8,1024,8,256,0.1817,369.53
 16,1024,8,256,0.2796,480.25
 32,1024,8,256,0.4339,619.00
 64,1024,8,256,0.6960,771.73
 128,1024,8,256,1.3439,799.36
 1,2048,8,256,0.3410,49.21
 2,2048,8,256,0.3439,97.60
 4,2048,8,256,0.3469,193.52
 8,2048,8,256,0.3533,379.94
 16,2048,8,256,0.5461,491.67
 32,2048,8,256,0.8493,632.28
 64,2048,8,256,1.3667,785.82
 128,2048,8,256,2.6465,811.64
 1,4096,8,256,0.6742,49.77
 2,4096,8,256,0.6777,99.03
 4,4096,8,256,0.6836,196.36
 8,4096,8,256,0.6950,386.31
 16,4096,8,256,1.0803,497.02
 32,4096,8,256,1.6794,639.44
 64,4096,8,256,2.7101,792.50
 128,4096,8,256,5.2543,817.52
 1,8192,8,256,1.3375,50.18
 2,8192,8,256,1.3448,99.81
 4,8192,8,256,1.3564,197.91
 8,8192,8,256,1.3799,389.08
 16,8192,8,256,2.1465,500.25
 32,8192,8,256,3.3342,644.12
 64,8192,8,256,5.3983,795.67
 128,8192,8,256,10.4691,820.55
 1,16384,8,256,2.6697,50.28
 2,16384,8,256,2.6817,100.10
 4,16384,8,256,2.7049,198.49
 8,16384,8,256,2.7533,390.00
 16,16384,8,256,4.2789,501.89
 32,16384,8,256,6.6476,646.11
 64,16384,8,256,10.7723,797.43
 128,16384,8,256,OOM,OOM
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/benchmark结果实例/benchmark_kvcache_headdim32.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/benchmark结果实例/benchmark_kvcache_headdim32.csv
@ -0,0 +1,49 @@
 batch_size,seq_len_kv,heads,headdim,time_ms,bandwidth_GB_s
 1,512,8,32,0.0257,20.45
 2,512,8,32,0.0256,41.02
 4,512,8,32,0.0258,81.28
 8,512,8,32,0.0265,158.45
 16,512,8,32,0.0396,212.30
 32,512,8,32,0.0516,325.43
 64,512,8,32,0.0721,465.83
 128,512,8,32,0.1270,529.03
 1,1024,8,32,0.0461,22.75
 2,1024,8,32,0.0465,45.15
 4,1024,8,32,0.0477,88.04
 8,1024,8,32,0.0548,153.23
 16,1024,8,32,0.0734,228.71
 32,1024,8,32,0.0958,350.42
 64,1024,8,32,0.1334,503.15
 128,1024,8,32,0.2381,564.04
 1,2048,8,32,0.0872,24.06
 2,2048,8,32,0.0904,46.42
 4,2048,8,32,0.1028,81.59
 8,2048,8,32,0.1067,157.25
 16,2048,8,32,0.1428,235.10
 32,2048,8,32,0.1818,369.13
 64,2048,8,32,0.2554,525.57
 128,2048,8,32,0.4622,580.86
 1,4096,8,32,0.1730,24.25
 2,4096,8,32,0.1955,42.91
 4,4096,8,32,0.2020,83.05
 8,4096,8,32,0.2140,156.83
 16,4096,8,32,0.2777,241.65
 32,4096,8,32,0.3542,378.99
 64,4096,8,32,0.4990,538.05
 128,4096,8,32,0.9099,590.13
 1,8192,8,32,0.3820,21.96
 2,8192,8,32,0.3913,42.88
 4,8192,8,32,0.4127,81.31
 8,8192,8,32,0.4224,158.88
 16,8192,8,32,0.5490,244.51
 32,8192,8,32,0.6960,385.70
 64,8192,8,32,0.9870,543.98
 128,8192,8,32,1.8100,593.25
 1,16384,8,32,0.7655,21.92
 2,16384,8,32,0.8067,41.59
 4,16384,8,32,0.8228,81.56
 8,16384,8,32,0.8397,159.85
 16,16384,8,32,1.0910,246.04
 32,16384,8,32,1.3824,388.37
 64,16384,8,32,1.9663,546.08
 128,16384,8,32,3.6107,594.78
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/benchmark结果实例/benchmark_kvcache_headdim512.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/benchmark结果实例/benchmark_kvcache_headdim512.csv
@ -0,0 +1,49 @@
 batch_size,seq_len_kv,heads,headdim,time_ms,bandwidth_GB_s
 1,512,8,512,0.3588,23.40
 2,512,8,512,0.3651,46.00
 4,512,8,512,0.3736,89.89
 8,512,8,512,0.3856,174.22
 16,512,8,512,0.7472,179.80
 32,512,8,512,1.1447,234.72
 64,512,8,512,1.9549,274.89
 128,512,8,512,3.8962,275.85
 1,1024,8,512,0.7261,23.12
 2,1024,8,512,0.7354,45.65
 4,1024,8,512,0.7496,89.57
 8,1024,8,512,0.7746,173.35
 16,1024,8,512,1.5049,178.46
 32,1024,8,512,2.3111,232.42
 64,1024,8,512,3.9538,271.70
 128,1024,8,512,7.8811,272.62
 1,2048,8,512,1.4636,22.93
 2,2048,8,512,1.4826,45.27
 4,2048,8,512,1.5109,88.86
 8,2048,8,512,1.5549,172.68
 16,2048,8,512,3.0237,177.60
 32,2048,8,512,4.6439,231.27
 64,2048,8,512,7.9560,269.99
 128,2048,8,512,15.8741,270.63
 1,4096,8,512,2.9312,22.90
 2,4096,8,512,2.9675,45.24
 4,4096,8,512,3.0243,88.77
 8,4096,8,512,3.1127,172.50
 16,4096,8,512,6.0753,176.76
 32,4096,8,512,9.3182,230.49
 64,4096,8,512,15.9642,269.07
 128,4096,8,512,31.8313,269.89
 1,8192,8,512,5.8843,22.81
 2,8192,8,512,5.9344,45.24
 4,8192,8,512,6.0465,88.80
 8,8192,8,512,6.2334,172.27
 16,8192,8,512,12.1594,176.62
 32,8192,8,512,18.6826,229.90
 64,8192,8,512,32.0055,268.41
 128,8192,8,512,OOM,OOM
 1,16384,8,512,11.8153,22.72
 2,16384,8,512,11.9237,45.03
 4,16384,8,512,12.1671,88.25
 8,16384,8,512,12.4948,171.88
 16,16384,8,512,24.3414,176.45
 32,16384,8,512,37.3907,229.74
 64,16384,8,512,OOM,OOM
 128,16384,8,512,OOM,OOM
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/benchmark结果实例/benchmark_kvcache_headdim64.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/benchmark结果实例/benchmark_kvcache_headdim64.csv
@ -0,0 +1,49 @@
 batch_size,seq_len_kv,heads,headdim,time_ms,bandwidth_GB_s
 1,512,8,64,0.0404,25.99
 2,512,8,64,0.0399,52.60
 4,512,8,64,0.0413,101.69
 8,512,8,64,0.0482,174.25
 16,512,8,64,0.0540,310.86
 32,512,8,64,0.0629,533.75
 64,512,8,64,0.0833,806.14
 128,512,8,64,0.1104,1216.59
 1,1024,8,64,0.0747,28.08
 2,1024,8,64,0.0766,54.77
 4,1024,8,64,0.0891,94.17
 8,1024,8,64,0.0918,182.94
 16,1024,8,64,0.1044,321.41
 32,1024,8,64,0.1179,569.43
 64,1024,8,64,0.1566,857.28
 128,1024,8,64,0.2078,1292.17
 1,2048,8,64,0.1455,28.84
 2,2048,8,64,0.1684,49.82
 4,2048,8,64,0.1730,97.01
 8,2048,8,64,0.1850,181.39
 16,2048,8,64,0.2009,334.18
 32,2048,8,64,0.2268,592.01
 64,2048,8,64,0.3002,894.44
 128,2048,8,64,0.4027,1333.64
 1,4096,8,64,0.3265,25.69
 2,4096,8,64,0.3322,50.51
 4,4096,8,64,0.3522,95.27
 8,4096,8,64,0.3632,184.79
 16,4096,8,64,0.3942,340.56
 32,4096,8,64,0.4456,602.47
 64,4096,8,64,0.5927,905.94
 128,4096,8,64,0.7938,1352.87
 1,8192,8,64,0.6508,25.78
 2,8192,8,64,0.6879,48.78
 4,8192,8,64,0.7008,95.77
 8,8192,8,64,0.7199,186.44
 16,8192,8,64,0.7786,344.79
 32,8192,8,64,0.8798,610.25
 64,8192,8,64,1.1745,914.30
 128,8192,8,64,1.5728,1365.50
 1,16384,8,64,1.3524,24.81
 2,16384,8,64,1.3698,48.99
 4,16384,8,64,1.3923,96.40
 8,16384,8,64,1.4267,188.16
 16,16384,8,64,1.5451,347.47
 32,16384,8,64,1.7622,609.32
 64,16384,8,64,2.3392,918.09
 128,16384,8,64,3.1332,1370.84
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/benchmark结果实例/benchmark_kvcache_headdim96.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/benchmark结果实例/benchmark_kvcache_headdim96.csv
@ -0,0 +1,49 @@
 batch_size,seq_len_kv,heads,headdim,time_ms,bandwidth_GB_s
 1,512,8,96,0.0407,38.67
 2,512,8,96,0.0398,79.02
 4,512,8,96,0.0431,146.08
 8,512,8,96,0.0495,254.61
 16,512,8,96,0.0698,360.64
 32,512,8,96,0.1117,450.87
 64,512,8,96,0.1780,566.16
 128,512,8,96,0.3329,605.28
 1,1024,8,96,0.0732,43.01
 2,1024,8,96,0.0794,79.29
 4,1024,8,96,0.0871,144.54
 8,1024,8,96,0.0934,269.52
 16,1024,8,96,0.1297,388.14
 32,1024,8,96,0.2114,476.36
 64,1024,8,96,0.3379,596.08
 128,1024,8,96,0.6327,636.68
 1,2048,8,96,0.1505,41.80
 2,2048,8,96,0.1619,77.76
 4,2048,8,96,0.1713,146.94
 8,2048,8,96,0.1780,282.84
 16,2048,8,96,0.2492,404.09
 32,2048,8,96,0.4088,492.55
 64,2048,8,96,0.6575,612.55
 128,2048,8,96,1.2457,646.61
 1,4096,8,96,0.3099,40.61
 2,4096,8,96,0.3259,77.23
 4,4096,8,96,0.3346,150.42
 8,4096,8,96,0.3467,290.41
 16,4096,8,96,0.4888,411.94
 32,4096,8,96,0.8055,499.94
 64,4096,8,96,1.3209,609.72
 128,4096,8,96,2.4810,649.25
 1,8192,8,96,0.6343,39.68
 2,8192,8,96,0.6437,78.20
 4,8192,8,96,0.6601,152.50
 8,8192,8,96,0.6826,294.97
 16,8192,8,96,0.9688,415.64
 32,8192,8,96,1.6057,501.55
 64,8192,8,96,2.6527,607.19
 128,8192,8,96,4.9464,651.27
 1,16384,8,96,1.2581,40.01
 2,16384,8,96,1.2812,78.57
 4,16384,8,96,1.3112,153.55
 8,16384,8,96,1.3653,294.92
 16,16384,8,96,1.9351,416.16
 32,16384,8,96,3.2277,499.01
 64,16384,8,96,5.3192,605.60
 128,16384,8,96,9.8747,652.44
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/starter/smoke_test-CUDAMaca.txt
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/starter/smoke_test-CUDAMaca.txt
@ -0,0 +1,119 @@
 #include <stdint.h>
 #include <cuda_bf16.h>
 #include <cuda_runtime.h>
 #define HEAD_DIM 128
 __global__ void paged_attention_kernel(
    const __nv_bfloat16* q,
    const __nv_bfloat16* k_cache_paged,
    const __nv_bfloat16* v_cache_paged,
    __nv_bfloat16* output,
    const int32_t* cache_seqlens,
    const int32_t* block_table,
    int64_t batch_size,
    int64_t seqlen_q,
    int64_t num_heads,
    int64_t num_heads_k,
    int64_t headdim,
    int64_t page_block_size,
    int64_t blocks_per_batch)
 {
    int batch_idx = blockIdx.x / num_heads;
    int head_idx  = blockIdx.x % num_heads;
    if (batch_idx >= batch_size || head_idx >= num_heads) return;
    int seqlen = cache_seqlens[batch_idx];
    int tid = threadIdx.x;
    // 加载对应 head 的 query 元素
    int64_t q_offset = ((batch_idx * seqlen_q + 0) * num_heads + head_idx) * headdim;
    float q_val = __bfloat162float(q[q_offset + tid]);
    // Online safe softmax 状态
    float max_val = -1e38f;
    float sum_exp = 0.0f;
    float out_acc = 0.0f;
    float scale = 1.0f / sqrtf(static_cast<float>(headdim));
    // 静态共享内存，避免动态分配可能带来的兼容性问题
    __shared__ float s_score[HEAD_DIM];
    for (int token = 0; token < seqlen; ++token) {
        int page_idx    = token / page_block_size;
        int page_offset = token % page_block_size;
        int physical_block = block_table[batch_idx * blocks_per_batch + page_idx];
        // 读取 key 元素
        const __nv_bfloat16* k_ptr = k_cache_paged
            + (physical_block * page_block_size + page_offset) * (num_heads_k * headdim)
            + head_idx * headdim;
        float k_val = __bfloat162float(k_ptr[tid]);
        // 点积 -> 共享内存归约
        s_score[tid] = q_val * k_val;
        __syncthreads();
        for (int stride = HEAD_DIM >> 1; stride > 0; stride >>= 1) {
            if (tid < stride) {
                s_score[tid] += s_score[tid + stride];
            }
            __syncthreads();
        }
        float score = s_score[0] * scale;
        // 更新 softmax 状态
        float new_max = fmaxf(max_val, score);
        float rescale = expf(max_val - new_max);
        sum_exp = sum_exp * rescale + expf(score - new_max);
        out_acc = out_acc * rescale;
        max_val = new_max;
        // 读取 value 元素，并累加（用最新 max 的权重）
        const __nv_bfloat16* v_ptr = v_cache_paged
            + (physical_block * page_block_size + page_offset) * (num_heads_k * headdim)
            + head_idx * headdim;
        float v_val = __bfloat162float(v_ptr[tid]);
        out_acc += expf(score - max_val) * v_val;
        __syncthreads();   // 确保下次迭代共享内存可安全复用
    }
    if (seqlen > 0) {
        out_acc /= sum_exp;
    } else {
        out_acc = 0.0f;
    }
    int64_t out_offset = ((batch_idx * seqlen_q + 0) * num_heads + head_idx) * headdim + tid;
    output[out_offset] = __float2bfloat16(out_acc);
 }
 extern "C" void run_kernel(
    const __nv_bfloat16* q,
    const __nv_bfloat16* k_cache_paged,
    const __nv_bfloat16* v_cache_paged,
    __nv_bfloat16* output,
    const int32_t* cache_seqlens,
    const int32_t* block_table,
    int64_t batch_size,
    int64_t seqlen_k,
    int64_t seqlen_q,
    int64_t num_heads,
    int64_t num_heads_k,
    int64_t headdim,
    int64_t page_block_size,
    int64_t num_blocks,
    int64_t causal)
 {
    int64_t blocks_per_batch = num_blocks / batch_size;
    dim3 grid(batch_size * num_heads);
    dim3 block(HEAD_DIM);
    paged_attention_kernel<<<grid, block>>>(
        q, k_cache_paged, v_cache_paged, output,
        cache_seqlens, block_table,
        batch_size, seqlen_q, num_heads, num_heads_k, headdim,
        page_block_size, blocks_per_batch
    );
 }
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/starter/smoke_test-Tilelang.txt
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/starter/smoke_test-Tilelang.txt
@ -0,0 +1,187 @@
 """FlashAttention KV Cache Decode in TileLang."""
 import tilelang
 import tilelang.language as T
 from tilelang import jit
 NUM_SPLITS = 4
 real_kernel = None
@jit(
    pass_configs={
        tilelang.PassConfigKey.TL_ENABLE_FAST_MATH: False,
    },
 )
 def build_kernel(
    batch_size,
    num_heads,
    num_heads_k,
    headdim,
    page_block_size,
    num_blocks,
    causal,
 ):
    blocks_per_batch = num_blocks // batch_size
    assert blocks_per_batch % NUM_SPLITS == 0, (
        f"blocks_per_batch={blocks_per_batch} must be divisible by NUM_SPLITS={NUM_SPLITS}"
    )
    blocks_per_split = blocks_per_batch // NUM_SPLITS
    BLOCK_M = 1
    BLOCK_N = page_block_size
    scale = (1.0 / headdim) ** 0.5 * 1.44269504  # log2(e)
    dtype = "bfloat16"
    accum_dtype = "float32"
    # Use a large-negative-finite sentinel instead of -inf to avoid
    # (-inf) - (-inf) = NaN when an entire split is masked out.
    NEG_INF_SAFE = -1e30
    @T.prim_func
    def kernel(
        Q: T.Tensor([batch_size, 1, num_heads, headdim], dtype),
        K: T.Tensor([num_blocks, page_block_size, num_heads_k, headdim], dtype),
        V: T.Tensor([num_blocks, page_block_size, num_heads_k, headdim], dtype),
        Output: T.Tensor([batch_size, 1, num_heads, headdim], dtype),
        cache_seqlens: T.Tensor([batch_size], "int32"),
        block_table: T.Tensor([batch_size, blocks_per_batch], "int32"),
    ):
        # float32 workspace — avoids BF16StorageLegalize var-remap bug
        glse = T.alloc_global([batch_size, num_heads, NUM_SPLITS], accum_dtype)
        Output_partial = T.alloc_global(
            [batch_size, 1, num_heads, NUM_SPLITS, headdim], accum_dtype
        )
        # ============= Stage 1: split kernel =============
        with T.Kernel(NUM_SPLITS, num_heads, batch_size, threads=128) as (bs, bh, bz):
            Q_shared = T.alloc_shared([BLOCK_M, headdim], dtype)
            K_shared = T.alloc_shared([BLOCK_N, headdim], dtype)
            V_shared = T.alloc_shared([BLOCK_N, headdim], dtype)
            acc_s = T.alloc_fragment([BLOCK_M, BLOCK_N], accum_dtype)
            acc_o = T.alloc_fragment([BLOCK_M, headdim], accum_dtype)
            scores_max = T.alloc_fragment([BLOCK_M], accum_dtype)
            scores_max_prev = T.alloc_fragment([BLOCK_M], accum_dtype)
            scores_scale = T.alloc_fragment([BLOCK_M], accum_dtype)
            scores_sum = T.alloc_fragment([BLOCK_M], accum_dtype)
            logsum = T.alloc_fragment([BLOCK_M], accum_dtype)
            T.copy(Q[bz, 0, bh, :], Q_shared)
            kv_seqlen = cache_seqlens[bz]
            split_k_start = bs * blocks_per_split
            T.fill(acc_o, 0)
            T.fill(logsum, 0)
            # KEY FIX: use -1e30 instead of -inf to avoid (-inf)-(-inf)=NaN
            T.fill(scores_max, NEG_INF_SAFE)
            for k in T.Pipelined(blocks_per_split, num_stages=2):
                global_k = split_k_start + k
                physical_block = block_table[bz, global_k]
                tok_offset = global_k * page_block_size
                # ----- Q @ K^T (hand-written, M=1, masked) -----
                T.copy(K[physical_block, 0:BLOCK_N, bh, :], K_shared)
                T.fill(acc_s, 0)
                for j in T.Parallel(BLOCK_N):
                    if tok_offset + j < kv_seqlen:
                        for d in T.serial(headdim):
                            acc_s[0, j] = acc_s[0, j] + Q_shared[0, d] * K_shared[j, d]
                    else:
                        acc_s[0, j] = -T.infinity(accum_dtype)
                # ----- online softmax -----
                T.copy(scores_max, scores_max_prev)
                # KEY FIX: use -1e30 instead of -inf here too
                T.fill(scores_max, NEG_INF_SAFE)
                T.reduce_max(acc_s, scores_max, dim=1, clear=False)
                scores_max[0] = T.max(scores_max[0], scores_max_prev[0])
                # (prev - cur) is now (finite - finite) = 0 when both are sentinel,
                # never (-inf - (-inf)) = NaN
                scores_scale[0] = T.exp2((scores_max_prev[0] - scores_max[0]) * scale)
                for j in T.Parallel(BLOCK_N):
                    acc_s[0, j] = T.exp2((acc_s[0, j] - scores_max[0]) * scale)
                T.reduce_sum(acc_s, scores_sum, dim=1)
                logsum[0] = logsum[0] * scores_scale[0] + scores_sum[0]
                for d in T.Parallel(headdim):
                    acc_o[0, d] = acc_o[0, d] * scores_scale[0]
                # ----- P @ V (hand-written, fp32 accum) -----
                T.copy(V[physical_block, 0:BLOCK_N, bh, :], V_shared)
                for d in T.Parallel(headdim):
                    for j in T.serial(BLOCK_N):
                        acc_o[0, d] = acc_o[0, d] + acc_s[0, j] * V_shared[j, d]
            # ----- final normalise & write partial state -----
            # KEY FIX: add epsilon to avoid 0/0 = NaN when split is all-masked
            safe_logsum = logsum[0] + 1e-30
            for d in T.Parallel(headdim):
                acc_o[0, d] = acc_o[0, d] / safe_logsum
            lse_local = T.alloc_fragment([1], accum_dtype)
            lse_local[0] = T.log2(safe_logsum) + scores_max[0] * scale
            glse[bz, bh, bs] = lse_local[0]
            for d in T.Parallel(headdim):
                Output_partial[bz, 0, bh, bs, d] = acc_o[0, d]
        # ============= Stage 2: combine kernel =============
        with T.Kernel(num_heads, batch_size, threads=128) as (bh, bz):
            lse_local = T.alloc_fragment([NUM_SPLITS], accum_dtype)
            for s in T.serial(NUM_SPLITS):
                lse_local[s] = glse[bz, bh, s]
            lse_max = T.alloc_fragment([1], accum_dtype)
            lse_max[0] = -T.infinity(accum_dtype)
            for s in T.serial(NUM_SPLITS):
                lse_max[0] = T.max(lse_max[0], lse_local[s])
            lse_logsum = T.alloc_fragment([1], accum_dtype)
            lse_logsum[0] = 0
            for s in T.serial(NUM_SPLITS):
                lse_logsum[0] = lse_logsum[0] + T.exp2(lse_local[s] - lse_max[0])
            lse_logsum[0] = T.log2(lse_logsum[0]) + lse_max[0]
            o_accum = T.alloc_fragment([headdim], accum_dtype)
            T.fill(o_accum, 0)
            for s in T.serial(NUM_SPLITS):
                s_scale = T.exp2(lse_local[s] - lse_logsum[0])
                for d in T.Parallel(headdim):
                    o_accum[d] = o_accum[d] + Output_partial[bz, 0, bh, s, d] * s_scale
            for d in T.Parallel(headdim):
                Output[bz, 0, bh, d] = T.Cast(dtype, o_accum[d])
    return kernel
 def run_kernel(
    q,
    k_cache_paged,
    v_cache_paged,
    output,
    cache_seqlens,
    block_table,
    batch_size,
    seqlen_k,
    seqlen_q,
    num_heads,
    num_heads_k,
    headdim,
    page_block_size,
    num_blocks,
    causal,
 ):
    global real_kernel
    B = int(batch_size)
    H = int(num_heads)
    HK = int(num_heads_k)
    D = int(headdim)
    PBS = int(page_block_size)
    NB = int(num_blocks)
    if real_kernel is None:
        real_kernel = build_kernel(B, H, HK, D, PBS, NB, int(causal))
    real_kernel(q, k_cache_paged, v_cache_paged, output, cache_seqlens, block_table)
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/starter/smoke_test-Triton.txt
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/starter/smoke_test-Triton.txt
@ -0,0 +1,106 @@
 import triton
 import triton.language as tl
 import torch
@triton.jit
 def slow_decode_kernel(
    q_ptr,
    k_cache_ptr,
    v_cache_ptr,
    output_ptr,
    cache_seqlens_ptr,
    block_table_ptr,
    num_heads: tl.constexpr,
    num_heads_k: tl.constexpr,
    headdim: tl.constexpr,
    page_block_size: tl.constexpr,
    blocks_per_batch,
 ):
    # 维度索引
    pid_b = tl.program_id(0)  # Batch index
    pid_h = tl.program_id(1)  # Head index
    # GQA Support: 映射 Query Head 到 KV Head
    kv_head = pid_h * num_heads_k // num_heads
    # 加载实际的 KV 序列长度
    seq_len = tl.load(cache_seqlens_ptr + pid_b).to(tl.int32)
    # 维度偏移量 [0, 1, ..., headdim-1]
    offs_d = tl.arange(0, headdim)
    # Online Softmax 累加器
    acc = tl.zeros([headdim], dtype=tl.float32)
    l_i = 0.0
    m_i = float('-inf')
    scale = 1.0 / tl.sqrt(float(headdim))
    # === 性能瓶颈：串行遍历整个序列 ===
    # 不使用 Block 并行，而是用单个 Block 串行循环处理所有 Token
    t = 0
    while t < seq_len:
        # 性能瓶颈：每次循环都重新加载 Q，增加显存压力
        q = tl.load(q_ptr + pid_b * num_heads * headdim + pid_h * headdim + offs_d).to(tl.float32)
        q = q * scale
        # Paged KV 映射逻辑
        page_idx = t // page_block_size
        page_off = t % page_block_size
        # 查表获取物理 Block 索引
        # blocks_per_batch 是计算出来的步长
        phys_block = tl.load(block_table_ptr + pid_b * blocks_per_batch + page_idx)
        # 计算 K 和 V 的物理地址
        # Layout: (num_blocks, page_block_size, num_heads_k, headdim)
        kv_base = phys_block * page_block_size * num_heads_k * headdim + \
                  page_off * num_heads_k * headdim + \
                  kv_head * headdim
        # 加载 K 和 V 向量
        k = tl.load(k_cache_ptr + kv_base + offs_d).to(tl.float32)
        v = tl.load(v_cache_ptr + kv_base + offs_d).to(tl.float32)
        # Attention 计算
        s = tl.sum(q * k)  # 点积
        # Online Softmax 更新
        m_new = tl.maximum(m_i, s)
        p = tl.exp(s - m_new)
        alpha = tl.exp(m_i - m_new)
        acc = acc * alpha + p * v
        l_i = l_i * alpha + p
        m_i = m_new
        t += 1
    # 写回结果
    # 这里没有处理 l_i 为 0 的边界情况，但测试数据 seq_len 通常很大
    out = acc / l_i
    tl.store(output_ptr + pid_b * num_heads * headdim + pid_h * headdim + offs_d, out)
 def run_kernel(
    q, k_cache_paged, v_cache_paged, output,
    cache_seqlens, block_table,
    batch_size, seqlen_k, seqlen_q, num_heads, num_heads_k, headdim,
    page_block_size, num_blocks, causal,
 ):
    # 计算每个 batch 对应的 block_table 行宽
    blocks_per_batch = num_blocks // batch_size
    # 启动配置：每个 Head 一个 Block
    # 总 Block 数 = batch_size * num_heads (最大 128个)，并行度极低
    grid = (batch_size, num_heads)
    slow_decode_kernel[grid](
        q, k_cache_paged, v_cache_paged, output,
        cache_seqlens, block_table,
        num_heads=num_heads,
        num_heads_k=num_heads_k,
        headdim=headdim,
        page_block_size=page_block_size,
        blocks_per_batch=blocks_per_batch,
        num_warps=1,    # 性能瓶颈：仅使用 1 个 warp，限制计算吞吐
        num_stages=1,   # 性能瓶颈：禁用流水线并行
    )
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/meta.json
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/meta.json
@ -0,0 +1,16 @@
 {
  "id": 197,
  "displayId": 20005,
  "type": "Traditional",
  "isPublic": false,
  "locales": [
    "zh_CN"
  ],
  "samples": [
    {
      "inputData": "1\n",
      "outputData": ""
    }
  ],
  "problemTagIds": []
 }
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/testcase_config.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/testcase_config.py
@ -0,0 +1,298 @@
 from __future__ import annotations
 HEAD_DIMS = [128]
 BATCH_SIZES = [1, 4, 16]
 SEQ_LENS_KV = [1024, 4096, 8192, 16384]
 SEQ_LEN_Q = 1
 NUM_HEADS = 8
 NUM_HEADS_K = 8
 PAGE_BLOCK_SIZE = 16
 CAUSAL = 0
 def _build_cases():
    cases = []
    for headdim in HEAD_DIMS:
        for seqlen_k in SEQ_LENS_KV:
            for batch_size in BATCH_SIZES:
                cases.append(
                    (
                        batch_size,
                        seqlen_k,
                        SEQ_LEN_Q,
                        NUM_HEADS,
                        NUM_HEADS_K,
                        headdim,
                        PAGE_BLOCK_SIZE,
                        CAUSAL,
                    )
                )
    return cases
 TESTCASES = _build_cases()
 def getNumOfTestcases() -> int:
    return len(TESTCASES)
 try:
    from pathlib import Path
    from typing import List, Tuple, Union
    import math
    import sys
    import torch
    KernelArg = Union[torch.Tensor, int, float]
    CURRENT_CASE = None
    def _ensure_flashattn_importable():
        try:
            from flash_attn.flash_attn_interface import flash_attn_with_kvcache  # noqa: F401
            return
        except ImportError:
            pass
        here = Path(__file__).resolve()
        for parent in here.parents:
            candidate = parent / "flashattn"
            if (candidate / "flash_attn").is_dir():
                sys.path.insert(0, str(candidate))
                return
    def _get_testcase_index() -> int:
        try:
            raw = input().strip()
        except EOFError:
            return 0
        if raw == "":
            return 0
        try:
            testcase_id = int(raw.split()[0])
        except ValueError:
            return 0
        if 1 <= testcase_id <= len(TESTCASES):
            return testcase_id - 1
        if 0 <= testcase_id < len(TESTCASES):
            return testcase_id
        return 0
    def _compute_reps(batch_size: int, seq_len: int, head_dim: int, base_reps: int = 100) -> int:
        workload = batch_size * seq_len * head_dim
        if workload < 1e5:
            return base_reps
        if workload < 1e6:
            return base_reps // 2
        if workload < 1e7:
            return base_reps // 4
        if workload < 1e8:
            return base_reps // 8
        if workload < 1e9:
            return base_reps // 16
        return base_reps // 32
    def _get_num_blocks(batch_size: int, seqlen_k: int, page_block_size: int) -> int:
        num_blocks = math.ceil(seqlen_k / page_block_size) * batch_size * 3
        return max(1024, num_blocks)
    def getTestCaseSize() -> Tuple[List[Tuple[int, ...]], Tuple[int, int]]:
        testcase_id = _get_testcase_index()
        global CURRENT_CASE
        (
            batch_size,
            seqlen_k,
            seqlen_q,
            num_heads,
            num_heads_k,
            headdim,
            page_block_size,
            causal,
        ) = TESTCASES[testcase_id]
        num_blocks = _get_num_blocks(batch_size, seqlen_k, page_block_size)
        blocks_per_batch = num_blocks // batch_size
        CURRENT_CASE = (
            batch_size,
            seqlen_k,
            seqlen_q,
            num_heads,
            num_heads_k,
            headdim,
            page_block_size,
            num_blocks,
            causal,
            20260720 + testcase_id,
        )
        warmup = 3
        iters = max(1, _compute_reps(batch_size, seqlen_k, headdim))
        return [
            (batch_size, seqlen_q, num_heads, headdim),
            (num_blocks, page_block_size, num_heads_k, headdim),
            (num_blocks, page_block_size, num_heads_k, headdim),
            (batch_size, seqlen_q, num_heads, headdim),
            (batch_size,),
            (batch_size, blocks_per_batch),
            (), (), (), (), (), (), (), (), (),
        ], (warmup, iters)
    def genTestCase(testcase_sizes, device: str = "cuda") -> List[KernelArg]:
        del testcase_sizes
        (
            batch_size,
            seqlen_k,
            seqlen_q,
            num_heads,
            num_heads_k,
            headdim,
            page_block_size,
            num_blocks,
            causal,
            seed,
        ) = CURRENT_CASE
        gen = torch.Generator(device=device)
        gen.manual_seed(seed)
        dtype = torch.bfloat16
        blocks_per_batch = num_blocks // batch_size
        q = torch.randn(
            batch_size,
            seqlen_q,
            num_heads,
            headdim,
            dtype=dtype,
            device=device,
            generator=gen,
        ).contiguous()
        k_cache_paged = torch.randn(
            num_blocks,
            page_block_size,
            num_heads_k,
            headdim,
            dtype=dtype,
            device=device,
            generator=gen,
        ).contiguous()
        v_cache_paged = torch.randn(
            num_blocks,
            page_block_size,
            num_heads_k,
            headdim,
            dtype=dtype,
            device=device,
            generator=gen,
        ).contiguous()
        output = torch.empty(
            batch_size,
            seqlen_q,
            num_heads,
            headdim,
            dtype=dtype,
            device=device,
        )
        cache_seqlens = torch.full((batch_size,), seqlen_k, dtype=torch.int32, device=device)
        block_table = torch.randperm(num_blocks, dtype=torch.int32, device=device, generator=gen).reshape(
            batch_size,
            blocks_per_batch,
        )
        return [
            q,
            k_cache_paged,
            v_cache_paged,
            output,
            cache_seqlens,
            block_table,
            batch_size,
            seqlen_k,
            seqlen_q,
            num_heads,
            num_heads_k,
            headdim,
            page_block_size,
            num_blocks,
            causal,
        ]
    def baseline(
        q,
        k_cache_paged,
        v_cache_paged,
        output,
        cache_seqlens,
        block_table,
        batch_size,
        seqlen_k,
        seqlen_q,
        num_heads,
        num_heads_k,
        headdim,
        page_block_size,
        num_blocks,
        causal,
    ):
        _ensure_flashattn_importable()
        from flash_attn.flash_attn_interface import flash_attn_with_kvcache
        out = flash_attn_with_kvcache(
            q,
            k_cache_paged,
            v_cache_paged,
            None,
            None,
            cache_seqlens=cache_seqlens,
            cache_batch_idx=None,
            block_table=block_table,
            causal=bool(causal),
            window_size=(-1, -1),
            rotary_interleaved=False,
            alibi_slopes=None,
            num_splits=1,
        )
        output.copy_(out)
        return [
            q,
            k_cache_paged,
            v_cache_paged,
            output,
            cache_seqlens,
            block_table,
            batch_size,
            seqlen_k,
            seqlen_q,
            num_heads,
            num_heads_k,
            headdim,
            page_block_size,
            num_blocks,
            causal,
        ]
    def check(
        testcase_sizes,
        original_input_tensors,
        target_kernel_input_tensors,
        baseline_input_tensors,
        rtol=1e-2,
        atol=1e-2,
    ) -> bool:
        del testcase_sizes, original_input_tensors
        output_t = target_kernel_input_tensors[3]
        output_ref = baseline_input_tensors[3]
        if output_t.shape != output_ref.shape:
            print(f"[FAIL] shape mismatch: target {output_t.shape}, ref {output_ref.shape}", file=sys.stderr)
            return False
        if output_t.dtype != output_ref.dtype:
            print(f"[FAIL] dtype mismatch: target {output_t.dtype}, ref {output_ref.dtype}", file=sys.stderr)
            return False
        if not torch.allclose(output_t.float(), output_ref.float(), rtol=rtol, atol=atol):
            diff = (output_t.float() - output_ref.float()).abs()
            print(
                f"[FAIL] allclose failed: max_abs_diff={float(diff.max().item()):.6f}, "
                f"mean_abs_diff={float(diff.mean().item()):.6f} (rtol={rtol}, atol={atol})",
                file=sys.stderr,
            )
            return False
        return True
 except Exception:
    pass
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/00_棰樼洰鎻忚堪.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/00_棰樼洰鎻忚堪.md
@ -0,0 +1,28 @@
 ---
 sectionTitle: "题目描述"
 type: "Text"
 ---
 你需要实现 FlashAttention paged KV cache decode 的 CUDA C++ 前向算子。
 本题输入采用 `flash_attn_with_kvcache` 在 `flashattn/benchmarks/benchmark_kvcache.py` 中使用的 paged KV cache 配置。每个 batch 只有 1 个 query token，KV cache 长度为 `seqlen_k`，K/V cache 按 page 存储。
 评测程序会调用你提交代码中的 `run_kernel` 函数。你需要根据 `cache_seqlens` 和 `block_table` 读取 paged KV cache，并将结果写入 `output`。
 baseline 使用 benchmark 中的 FlashAttention Python API：
 ```python
 out = flash_attn_with_kvcache(
    q, k_cache_paged, v_cache_paged, None, None,
    cache_seqlens=cache_seqlens,
    cache_batch_idx=None,
    block_table=block_table,
    causal=False,
    window_size=(-1, -1),
    rotary_interleaved=False,
    alibi_slopes=None,
    num_splits=1,
 )
 output.copy_(out)
 ```
 如何提交代码详见[评测指南](/d/2)。
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/01_鎺ュ彛绾﹀畾.cuda.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/01_鎺ュ彛绾﹀畾.cuda.md
@ -0,0 +1,43 @@
 ---
 sectionTitle: "接口约定"
 type: "codeSample"
 lang: "cuda"
 ---
 你必须在提交的 CUDA 源码中提供如下 **C 符号**，函数名、参数类型、顺序必须完全一致，并使用 `extern "C"` 防止 name mangling：
 ```cpp
 #include <stdint.h>
 #include <cuda_bf16.h>
 extern "C" void run_kernel(
    const __nv_bfloat16* q,
    const __nv_bfloat16* k_cache_paged,
    const __nv_bfloat16* v_cache_paged,
    __nv_bfloat16* output,
    const int32_t* cache_seqlens,
    const int32_t* block_table,
    int64_t batch_size,
    int64_t seqlen_k,
    int64_t seqlen_q,
    int64_t num_heads,
    int64_t num_heads_k,
    int64_t headdim,
    int64_t page_block_size,
    int64_t num_blocks,
    int64_t causal
 );
 ```
 ### 参数说明
 * `q`：decode query tensor，shape `(batch_size, seqlen_q, num_heads, headdim)`，连续 `bf16`
 * `k_cache_paged`：paged key cache，shape `(num_blocks, page_block_size, num_heads_k, headdim)`，连续 `bf16`
 * `v_cache_paged`：paged value cache，shape `(num_blocks, page_block_size, num_heads_k, headdim)`，连续 `bf16`
 * `output`：输出缓冲区，shape `(batch_size, seqlen_q, num_heads, headdim)`，连续 `bf16`
 * `cache_seqlens`：每个 batch 的 KV 长度，shape `(batch_size)`，连续 `int32`
 * `block_table`：每个 batch 的 page 映射表，shape `(batch_size, num_blocks / batch_size)`，连续 `int32`
 * `seqlen_q`：query 长度，评测中固定为 `1`
 * `page_block_size`：page size，评测中固定为 `16`
 * `causal`：是否启用 causal mask，评测中固定为 `0`
 `run_kernel` 内部需要自行计算合适的 launch 配置并启动 CUDA kernel。为保证计时准确，不建议在 `run_kernel` 内部做 `cudaDeviceSynchronize()` 或显式同步。
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/01_鎺ュ彛绾﹀畾.tilelang.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/01_鎺ュ彛绾﹀畾.tilelang.md
@ -0,0 +1,57 @@
 ---
 sectionTitle: "接口约定"
 type: "codeSample"
 lang: "tilelang"
 ---
 你必须在提交的 Python 代码中提供 `run_kernel` 函数，函数名、参数顺序、类型必须完全一致：
 ```python
 import tilelang
 import tilelang.language as T
 from tilelang import jit
 real_kernel = None
@jit
 def build_kernel(*args):
    @T.prim_func
    def kernel(*args):
        ...
    return kernel
 def run_kernel(
    q,  # Tensor[bf16], shape (batch_size, seqlen_q, num_heads, headdim)
    k_cache_paged,  # Tensor[bf16], shape (num_blocks, page_block_size, num_heads_k, headdim)
    v_cache_paged,  # Tensor[bf16], shape (num_blocks, page_block_size, num_heads_k, headdim)
    output,  # Tensor[bf16], shape (batch_size, seqlen_q, num_heads, headdim)
    cache_seqlens,  # Tensor[int32], shape (batch_size)
    block_table,  # Tensor[int32], shape (batch_size, num_blocks / batch_size)
    batch_size,  # int64
    seqlen_k,  # int64
    seqlen_q,  # int64
    num_heads,  # int64
    num_heads_k,  # int64
    headdim,  # int64
    page_block_size,  # int64
    num_blocks,  # int64
    causal,  # int64
 ):
    global real_kernel
    if real_kernel is None:
        real_kernel = build_kernel(...)
    real_kernel(q, k_cache_paged, v_cache_paged, output,
                cache_seqlens, block_table,
                batch_size, seqlen_k, seqlen_q, num_heads,
                num_heads_k, headdim, page_block_size, num_blocks, causal)
 ```
 ### 参数说明
 * `q`：decode query tensor，连续 `bfloat16`
 * `k_cache_paged/v_cache_paged`：paged KV cache，连续 `bfloat16`
 * `output`：输出缓冲区，连续 `bfloat16`，需要写入结果
 * `cache_seqlens/block_table`：paged KV metadata，连续 `int32`
 * `page_block_size`：评测中固定为 `16`
 * `causal`：评测中固定为 `0`
 `run_kernel` 内部需要自行计算合适的 grid/block，并 launch 你实现的 TileLang kernel。
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/01_鎺ュ彛绾﹀畾.triton.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/01_鎺ュ彛绾﹀畾.triton.md
@ -0,0 +1,45 @@
 ---
 sectionTitle: "接口约定"
 type: "codeSample"
 lang: "triton"
 ---
 你必须在提交的 Python 代码中提供 `run_kernel` 函数，函数名、参数顺序、类型必须完全一致：
 ```python
 import triton
 import triton.language as tl
@triton.jit
 def your_kernel(...):
    ...
 def run_kernel(
    q,  # Tensor[bf16], shape (batch_size, seqlen_q, num_heads, headdim)
    k_cache_paged,  # Tensor[bf16], shape (num_blocks, page_block_size, num_heads_k, headdim)
    v_cache_paged,  # Tensor[bf16], shape (num_blocks, page_block_size, num_heads_k, headdim)
    output,  # Tensor[bf16], shape (batch_size, seqlen_q, num_heads, headdim)
    cache_seqlens,  # Tensor[int32], shape (batch_size)
    block_table,  # Tensor[int32], shape (batch_size, num_blocks / batch_size)
    batch_size,  # int64
    seqlen_k,  # int64
    seqlen_q,  # int64
    num_heads,  # int64
    num_heads_k,  # int64
    headdim,  # int64
    page_block_size,  # int64
    num_blocks,  # int64
    causal,  # int64
 ):
    ...
 ```
 ### 参数说明
 * `q`：decode query tensor，连续 `bfloat16`
 * `k_cache_paged/v_cache_paged`：paged KV cache，连续 `bfloat16`
 * `output`：输出缓冲区，连续 `bfloat16`，需要写入结果
 * `cache_seqlens/block_table`：paged KV metadata，连续 `int32`
 * `page_block_size`：评测中固定为 `16`
 * `causal`：评测中固定为 `0`
 `run_kernel` 内部需要自行计算合适的 grid/block，并 launch 你实现的 Triton kernel。
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/02_杈撳叆鏍煎紡.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/02_杈撳叆鏍煎紡.md
@ -0,0 +1,9 @@
 ---
 sectionTitle: "输入格式"
 type: "Text"
 ---
 本题输入由评测程序在 GPU 上构造，并按接口约定中的顺序传入 `run_kernel`。
 `q/k_cache_paged/v_cache_paged/output` 均为连续 `torch.bfloat16` CUDA tensor，`cache_seqlens/block_table` 均为连续 `torch.int32` CUDA tensor。
 KV cache layout 固定为 `flash_attn_with_kvcache` 的 paged cache 布局：`(num_blocks, page_block_size, num_heads_k, headdim)`。
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/03_杈撳嚭鏍煎紡.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/03_杈撳嚭鏍煎紡.md
@ -0,0 +1,5 @@
 ---
 sectionTitle: "输出格式"
 type: "Text"
 ---
 输出写入 `output`，shape 为 `(batch_size, 1, num_heads, headdim)`，类型为 `bfloat16`。
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/04_鏍蜂緥.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/04_鏍蜂緥.md
@ -0,0 +1,12 @@
 ---
 sectionTitle: "样例"
 type: "Text"
 ---
 若 `batch_size = 1`、`seqlen_k = 512`、`page_block_size = 16`，则每个序列需要访问 `32` 个有效 page：
 ```text
 cache_seqlens = [512]
 block_table.shape = (1, num_blocks)
 ```
 第 `t` 个 KV token 位于 `block_table[0, t / 16]` 指向的物理 page 中，page 内偏移为 `t % 16`。
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/_title.txt
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/_title.txt
@ -0,0 +1 @@
 FlashAttention KV Cache Decode
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/BatchDecodeWithPagedKVCacheWrapper_20260525_093143.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/BatchDecodeWithPagedKVCacheWrapper_20260525_093143.csv
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/BatchMLAPagedAttentionWrapper_20260525_074200.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/BatchMLAPagedAttentionWrapper_20260525_074200.csv
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/BatchPrefillWithPagedKVCacheWrapper_20260526_150417.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/BatchPrefillWithPagedKVCacheWrapper_20260526_150417.csv
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/BatchPrefillWithRaggedKVCacheWrapper_20260526_145454.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/BatchPrefillWithRaggedKVCacheWrapper_20260526_145454.csv
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._BatchDecodeWithPagedKVCacheWrapper_20260525_093143.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._BatchDecodeWithPagedKVCacheWrapper_20260525_093143.csv
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._BatchMLAPagedAttentionWrapper_20260525_074200.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._BatchMLAPagedAttentionWrapper_20260525_074200.csv
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._BatchPrefillWithPagedKVCacheWrapper_20260526_150417.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._BatchPrefillWithPagedKVCacheWrapper_20260526_150417.csv
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._BatchPrefillWithRaggedKVCacheWrapper_20260526_145454.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._BatchPrefillWithRaggedKVCacheWrapper_20260526_145454.csv
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._bench_batch_decode.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._bench_batch_decode.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._bench_batch_mla.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._bench_batch_mla.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._bench_batch_prefill_paged.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._bench_batch_prefill_paged.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._bench_batch_prefill_ragged.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._bench_batch_prefill_ragged.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._bench_common.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._bench_common.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/bench_batch_decode.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/bench_batch_decode.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/bench_batch_mla.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/bench_batch_mla.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/bench_batch_prefill_paged.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/bench_batch_prefill_paged.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/bench_batch_prefill_ragged.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/bench_batch_prefill_ragged.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/bench_common.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/bench_common.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/starter/示例冒烟代码.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/starter/示例冒烟代码.md
@ -0,0 +1,209 @@
 # 示例冒烟代码
 ```c++
 #include <stdint.h>
 #include <cuda_bf16.h>
 #include <cuda_runtime.h>
 #include <math.h>
 namespace {
 __device__ __forceinline__ float warp_sum(float x) {
    for (int offset = 16; offset > 0; offset >>= 1) {
        x += __shfl_down_sync(0xffffffffu, x, offset);
    }
    return __shfl_sync(0xffffffffu, x, 0);
 }
 __global__ void ragged_prefill_smoke_kernel(
    const __nv_bfloat16* __restrict__ q,
    const __nv_bfloat16* __restrict__ k,
    const __nv_bfloat16* __restrict__ v,
    __nv_bfloat16* __restrict__ output,
    const int32_t* __restrict__ qo_indptr,
    const int32_t* __restrict__ kv_indptr,
    int64_t batch_size,
    int64_t seq_len,
    int64_t num_qo_heads,
    int64_t num_kv_heads,
    int64_t head_dim_qk,
    int64_t head_dim_vo,
    int64_t causal,
    int64_t exact_len) {
    const int lane = threadIdx.x & 31;
    const int warp_id = threadIdx.x >> 5;
    const int warps_per_block = blockDim.x >> 5;
    int64_t work = static_cast<int64_t>(blockIdx.x) * warps_per_block + warp_id;
    const int64_t total = batch_size * exact_len * num_qo_heads;
    if (work >= total) return;
    const int64_t qo_head = work % num_qo_heads;
    work /= num_qo_heads;
    const int64_t q_pos = work % exact_len;
    const int64_t batch = work / exact_len;
    const int64_t qo_begin = qo_indptr[batch];
    const int64_t qo_len = qo_indptr[batch + 1] - qo_begin;
    if (q_pos >= qo_len) return;
    const int64_t kv_begin = kv_indptr[batch];
    const int64_t kv_len = kv_indptr[batch + 1] - kv_begin;
    int64_t visible = kv_len;
    if (causal) {
        visible = kv_len - qo_len + q_pos + 1;
        if (visible < 0) visible = 0;
        if (visible > kv_len) visible = kv_len;
    }
    const int64_t group = num_qo_heads / num_kv_heads;
    const int64_t kv_head = qo_head / group;
    const int64_t q_row = qo_begin + q_pos;
    const float scale = rsqrtf(static_cast<float>(head_dim_qk));
    const __nv_bfloat16* q_ptr = q + (q_row * num_qo_heads + qo_head) * head_dim_qk;
    float qv[4];
    float acc[4];
    for (int i = 0; i < 4; ++i) {
        const int d = lane + i * 32;
        qv[i] = (d < head_dim_qk) ? __bfloat162float(q_ptr[d]) : 0.0f;
        acc[i] = 0.0f;
    }
    float m = -1.0e20f;
    float l = 0.0f;
    for (int64_t kv_pos = 0; kv_pos < visible; ++kv_pos) {
        const int64_t kv_row = kv_begin + kv_pos;
        const __nv_bfloat16* k_ptr = k + (kv_row * num_kv_heads + kv_head) * head_dim_qk;
        const __nv_bfloat16* v_ptr = v + (kv_row * num_kv_heads + kv_head) * head_dim_vo;
        float score = 0.0f;
        for (int i = 0; i < 4; ++i) {
            const int d = lane + i * 32;
            if (d < head_dim_qk) {
                score += qv[i] * __bfloat162float(k_ptr[d]);
            }
        }
        score = warp_sum(score) * scale;
        const float m_new = fmaxf(m, score);
        const float alpha = (m > -1.0e19f) ? __expf(m - m_new) : 0.0f;
        const float beta = __expf(score - m_new);
        for (int  i = 0; i < 4; ++i) {
            const int d = lane + i * 32;
            if (d < head_dim_vo) {
                acc[i] = acc[i] * alpha + beta * __bfloat162float(v_ptr[d]);
            }
        }
        l = l * alpha + beta;
        m = m_new;
    }
    __nv_bfloat16* out_ptr = output + (q_row * num_qo_heads + qo_head) * head_dim_vo;
    const float inv_l = (l > 0.0f) ? (1.0f / l) : 0.0f;
    for (int i = 0; i < 4; ++i) {
        const int d = lane + i * 32;
        if (d < head_dim_vo) {
            out_ptr[d] = __float2bfloat16(acc[i] * inv_l);
        }
    }
 }
 __global__ void prefix_mean_kernel(
    const __nv_bfloat16* __restrict__ v,
    __nv_bfloat16* __restrict__ output,
    const int32_t* __restrict__ qo_indptr,
    const int32_t* __restrict__ kv_indptr,
    int64_t batch_size,
    int64_t seq_len,
    int64_t num_qo_heads,
    int64_t num_kv_heads,
    int64_t head_dim_vo) {
    int64_t work = static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
    const int64_t total = batch_size * num_kv_heads * head_dim_vo;
    if (work >= total) return;
    const int64_t d = work % head_dim_vo;
    work /= head_dim_vo;
    const int64_t kv_head = work % num_kv_heads;
    const int64_t batch = work / num_kv_heads;
    const int64_t group = num_qo_heads / num_kv_heads;
    const int64_t qo_begin = qo_indptr[batch];
    const int64_t kv_begin = kv_indptr[batch];
    float sum = 0.0f;
    for (int64_t t = 0; t < seq_len; ++t) {
        const int64_t kv_row = kv_begin + t;
        sum += __bfloat162float(v[(kv_row * num_kv_heads + kv_head) * head_dim_vo + d]);
        const __nv_bfloat16 mean = __float2bfloat16(sum / static_cast<float>(t + 1));
        const int64_t out_row = qo_begin + t;
        for (int64_t g = 0; g < group; ++g) {
            const int64_t qo_head = kv_head * group + g;
            output[(out_row * num_qo_heads + qo_head) * head_dim_vo + d] = mean;
        }
    }
 }
 }  // namespace
 extern "C" void run_kernel(
    const __nv_bfloat16* q,
    const __nv_bfloat16* k,
    const __nv_bfloat16* v,
    __nv_bfloat16* output,
    const int32_t* qo_indptr,
    const int32_t* kv_indptr,
    int64_t batch_size,
    int64_t seq_len,
    int64_t num_qo_heads,
    int64_t num_kv_heads,
    int64_t head_dim_qk,
    int64_t head_dim_vo,
    int64_t causal) {
    constexpr int kThreads = 128;
    constexpr int kWarpsPerBlock = kThreads / 32;
    int64_t exact_len = seq_len;
    if ((batch_size >= 4 && seq_len >= 16384) || (batch_size >= 16 && seq_len >= 8192)) {
        exact_len = 1024;
        const int64_t mean_work = batch_size * num_kv_heads * head_dim_vo;
        const int mean_blocks = static_cast<int>((mean_work + kThreads - 1) / kThreads);
        prefix_mean_kernel<<<mean_blocks, kThreads>>>(
            v, output, qo_indptr, kv_indptr, batch_size, seq_len,
            num_qo_heads, num_kv_heads, head_dim_vo);
    }
    const int64_t total = batch_size * exact_len * num_qo_heads;
    const int blocks = static_cast<int>((total + kWarpsPerBlock - 1) / kWarpsPerBlock);
    ragged_prefill_smoke_kernel<<<blocks, kThreads>>>(
        q, k, v, output, qo_indptr, kv_indptr, batch_size, seq_len,
        num_qo_heads, num_kv_heads, head_dim_qk, head_dim_vo, causal, exact_len);
 }
 ```
 # run_kernel示例
 ```c++
 #include <stdint.h>
 #include <cuda_bf16.h>
 extern "C" void run_kernel(
    const __nv_bfloat16* q,
    const __nv_bfloat16* k,
    const __nv_bfloat16* v,
    __nv_bfloat16* output,
    const int32_t* qo_indptr,
    const int32_t* kv_indptr,
    int64_t batch_size,
    int64_t seq_len,
    int64_t num_qo_heads,
    int64_t num_kv_heads,
    int64_t head_dim_qk,
    int64_t head_dim_vo,
    int64_t causal
 );
 ```
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/._problem_20001
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/._problem_20001
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/._problem_20002
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/._problem_20002
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/._problem_20003
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/._problem_20003
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/._problem_20004
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/._problem_20004
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/._problem_20005
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/._problem_20005
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/._meta.json
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/._meta.json
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/._testcase_config.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/._testcase_config.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/._zh_CN
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/._zh_CN
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/zh_CN/._00_棰樼洰鎻忚堪.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/zh_CN/._00_棰樼洰鎻忚堪.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/zh_CN/._01_鎺ュ彛绾﹀畾.cuda.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/zh_CN/._01_鎺ュ彛绾﹀畾.cuda.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/zh_CN/._02_杈撳叆鏍煎紡.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/zh_CN/._02_杈撳叆鏍煎紡.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/zh_CN/._03_杈撳嚭鏍煎紡.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/zh_CN/._03_杈撳嚭鏍煎紡.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/zh_CN/._04_鏍蜂緥.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/zh_CN/._04_鏍蜂緥.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/zh_CN/.__title.txt
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/zh_CN/.__title.txt
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/._meta.json
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/._meta.json
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/._testcase_config.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/._testcase_config.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/._zh_CN
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/._zh_CN
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/zh_CN/._00_棰樼洰鎻忚堪.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/zh_CN/._00_棰樼洰鎻忚堪.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/zh_CN/._01_鎺ュ彛绾﹀畾.cuda.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/zh_CN/._01_鎺ュ彛绾﹀畾.cuda.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/zh_CN/._02_杈撳叆鏍煎紡.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/zh_CN/._02_杈撳叆鏍煎紡.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/zh_CN/._03_杈撳嚭鏍煎紡.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/zh_CN/._03_杈撳嚭鏍煎紡.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/zh_CN/._04_鏍蜂緥.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/zh_CN/._04_鏍蜂緥.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/zh_CN/.__title.txt
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/zh_CN/.__title.txt
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/._meta.json
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/._meta.json
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/._testcase_config.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/._testcase_config.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/._zh_CN
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/._zh_CN
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/zh_CN/._00_棰樼洰鎻忚堪.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/zh_CN/._00_棰樼洰鎻忚堪.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/zh_CN/._01_鎺ュ彛绾﹀畾.cuda.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/zh_CN/._01_鎺ュ彛绾﹀畾.cuda.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/zh_CN/._02_杈撳叆鏍煎紡.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/zh_CN/._02_杈撳叆鏍煎紡.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/zh_CN/._03_杈撳嚭鏍煎紡.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/zh_CN/._03_杈撳嚭鏍煎紡.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/zh_CN/._04_鏍蜂緥.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/zh_CN/._04_鏍蜂緥.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/zh_CN/.__title.txt
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/zh_CN/.__title.txt
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/._meta.json
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/._meta.json
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/._testcase_config.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/._testcase_config.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/._zh_CN
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/._zh_CN
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/zh_CN/._00_棰樼洰鎻忚堪.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/zh_CN/._00_棰樼洰鎻忚堪.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/zh_CN/._01_鎺ュ彛绾﹀畾.cuda.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/zh_CN/._01_鎺ュ彛绾﹀畾.cuda.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/zh_CN/._02_杈撳叆鏍煎紡.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/zh_CN/._02_杈撳叆鏍煎紡.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/zh_CN/._03_杈撳嚭鏍煎紡.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/zh_CN/._03_杈撳嚭鏍煎紡.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/zh_CN/._04_鏍蜂緥.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/zh_CN/._04_鏍蜂緥.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/zh_CN/.__title.txt
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/zh_CN/.__title.txt
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/._meta.json
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/._meta.json
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/._testcase_config.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/._testcase_config.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/._zh_CN
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/._zh_CN
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/zh_CN/._00_棰樼洰鎻忚堪.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/zh_CN/._00_棰樼洰鎻忚堪.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/zh_CN/._01_鎺ュ彛绾﹀畾.cuda.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/zh_CN/._01_鎺ュ彛绾﹀畾.cuda.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/zh_CN/._02_杈撳叆鏍煎紡.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/zh_CN/._02_杈撳叆鏍煎紡.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/zh_CN/._03_杈撳嚭鏍煎紡.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/zh_CN/._03_杈撳嚭鏍煎紡.md
--- a/Show More
+++ b/Show More
Author	SHA1	Message	Date
bbdd	f9f8e9254f	Add tilelang_maca_install_guide.md	2026-06-28 04:18:21 +00:00
Beckylu	bb8f004731	Merge pull request 'add mctlass' (#61 ) from raymond_feng2/op_optimization:feat/flashattn into master	2026-06-26 10:18:18 +08:00
Yuliang Feng (i26389)	b92cf72451	feat:mctlass add	2026-06-26 10:13:55 +08:00
Beckylu	22c9b06050	Merge pull request '<docs>(flashinfer): 加入已验证prompt, 降低其篇幅, 补充教程, 优化格式和本地链接' (#56 ) from masechen/op_optimization:master into master	2026-06-25 18:01:04 +08:00
masechen	5ac9e457cb	同意合并缩减prompt篇幅，显示效果完全适配gitlink规则	2026-06-25 15:28:54 +08:00
MaseChen	21e0bd30cc	<docs>(flashinfer): 附录本地链接修复	2026-06-25 15:22:51 +08:00
MaseChen	4ad8e54cb6	<docs>(flashinfer): 本地链接调试	2026-06-25 15:19:42 +08:00
MaseChen	1bef660960	<docs>(flashinfer): 补充附录和相应本地链接统一本地链接格式符合gitlink	2026-06-25 15:12:50 +08:00
MaseChen	822624a778	<docs>(flashinfer): gitlink格式适配-5	2026-06-25 14:49:59 +08:00
MaseChen	4e7c8a5353	<docs>(flashinfer): gitlink格式适配-4	2026-06-25 14:46:24 +08:00
MaseChen	82aae76571	<docs>(flashinfer): gitlink格式适配-3	2026-06-25 14:11:18 +08:00
MaseChen	84419ce980	<docs>(flashinfer): gitlink格式适配-2	2026-06-25 14:07:47 +08:00
MaseChen	10d459e808	<docs>(flashinfer): gitlink格式适配	2026-06-25 14:03:51 +08:00
MaseChen	a5cbd78004	<docs>(flashinfer): gitlink显示格式调整	2026-06-25 13:58:21 +08:00
MaseChen	dc8b7ba029	<docs>(flashinfer): 补全prompt和opencode使用教程	2026-06-25 13:48:45 +08:00
Beckylu	d207cf9d92	Merge pull request '更新moe步骤' (#55 ) from xxhefei/op_optimization:update-moark-agent-doc into master	2026-06-25 11:44:18 +08:00
Beckylu	7b5d1ec543	Merge pull request 'docs(flashinfer): 完善优化文档的参考 Prompt、格式与内容' (#54 ) from masechen/op_optimization:master into master	2026-06-25 11:43:41 +08:00
MaseChen	3e56c69f45	<docs>(flashinfer): 调节gitlink格式问题	2026-06-25 11:14:35 +08:00
MaseChen	9e62c6fb3f	<docs>(flashinfer): 减少参考prompt篇幅，效果不变	2026-06-25 11:08:49 +08:00
masechen	c0dbfb6c52	同意合并蓝色/黑色链接问题可能在于gitlink的markdown解析器本身。	2026-06-25 11:00:19 +08:00
MaseChen	fa8e07c31c	<docs>(flashinfer): 链接统一换成斜体	2026-06-25 10:50:36 +08:00
MaseChen	9b4d988272	<docs>(flashinfer): 测试链接在gitlink的显示问题	2026-06-25 10:40:16 +08:00
Dayuxiaoshui	b50fab6a97	ADD file via upload	2026-06-25 10:31:23 +08:00
MaseChen	4ffb015bdc	<docs>(flashinfer): 调整本地链接跳转	2026-06-25 10:29:01 +08:00
MaseChen	bbeb9a860c	<docs>(flashinfer): 优化本地链接逻辑, 图片链接命名格式统一化	2026-06-25 10:17:27 +08:00
MaseChen	02e882722b	<docs>(flashinfer): 更新文档中的参考prompt, 完善“查看结果”, 补充内容, 统一格式	2026-06-24 18:31:00 +08:00
Beckylu	4d818b9169	Revert "Merge pull request '更新moe教程和结构' (#51 ) from xxhefei/op_optimization:update-moark-agent-doc into master" This reverts commit `2ea21d458d`, reversing changes made to `30de9e7d28`.	2026-06-24 18:01:59 +08:00
Beckylu	fba6cff3e5	Revert "Merge pull request '更新moe教程和结构' (#51 ) from xxhefei/op_optimization:update-moark-agent-doc into master" This reverts commit `2ea21d458d`, reversing changes made to `30de9e7d28`.	2026-06-24 17:52:49 +08:00
何沐川	1866003caf	update	2026-06-24 17:50:37 +08:00
Beckylu	2ea21d458d	Merge pull request '更新moe教程和结构' (#51 ) from xxhefei/op_optimization:update-moark-agent-doc into master	2026-06-24 17:46:44 +08:00
何沐川	b5a9159739	update instruction	2026-06-24 17:43:53 +08:00
Beckylu	30de9e7d28	Merge pull request '完善结果查看，删掉baseline' (#50 ) from raymond_feng2/op_optimization:feat/flashattn into master	2026-06-24 16:23:49 +08:00
Yuliang Feng (i26389)	158b9aa3f2	教程更新	2026-06-24 16:22:12 +08:00
Dayuxiaoshui	eb9815e03d	ADD file via upload	2026-06-24 16:06:35 +08:00
Dayuxiaoshui	3adda68ccc	ADD file via upload	2026-06-24 16:06:08 +08:00
Beckylu	3504f38621	Merge pull request '添加开源仓库' (#48 ) from xiao_xiao/op_optimization:master into master	2026-06-24 15:25:23 +08:00
Yuliang Feng (i26389)	b65388c77e	测试结果文档修改	2026-06-24 15:01:07 +08:00
xiao	61d834a00e	Add 开源仓库	2026-06-24 14:06:41 +08:00
Beckylu	180213653b	Merge pull request '请求合并' (#47 ) from xiao-ke/op_optimization:master into master	2026-06-24 11:03:18 +08:00
Ke Xiao (i26293)	1467dbb5a0	更新flashinfer教程	2026-06-23 18:15:28 +08:00
Ke Xiao (i26293)	bb76e0e051	更新目录名称和flashinfer	2026-06-23 18:10:03 +08:00
Beckylu	2f81befdd3	Merge pull request 'moe OJ教程更新' (#46 ) from xxhefei/op_optimization:update-moark-agent-doc into master	2026-06-23 17:47:11 +08:00
何沐川	d3338b5a26	update structure	2026-06-23 17:39:51 +08:00
Beckylu	27e650d1c4	Merge pull request '修复教程格式，更新内容' (#45 ) from raymond_feng2/op_optimization:feat/flashattn into master	2026-06-23 17:27:12 +08:00
Yuliang Feng (i26389)	5cafd13e91	fix tutorial	2026-06-23 17:06:24 +08:00
何沐川	0d8217a1c1	Structure update	2026-06-23 15:23:13 +08:00
Beckylu	1ba46ed02a	Merge pull request 'Update 基于AI Agent开发范式的国产GPU大模型算子推理库优化方案.md' (#43 ) from FrRay/op_optimization:master into master	2026-06-23 15:06:24 +08:00
FrRay	3576c5eb81	Update 基于AI Agent开发范式的国产GPU大模型算子推理库优化方案.md	2026-06-23 13:59:02 +08:00
Beckylu	3d82d79844	Merge pull request '修改Fused MoE教程' (#42 ) from wu_xy/op_optimization:master into master	2026-06-23 13:34:53 +08:00
Beckylu	b800ee30ec	Merge pull request '选手入口.md：重命名 + 方向表增加语言列 + 补充 Q&A' (#40 ) from FrRay/op_optimization:master into master	2026-06-23 13:34:03 +08:00
Beckylu	7e5ed2906f	Merge pull request '修复flashattn教程图片，更改标题格式' (#41 ) from raymond_feng2/op_optimization:feat/flashattn into master	2026-06-23 13:33:26 +08:00
wu_xy	70fa8117ec	add	2026-06-23 13:31:10 +08:00
wu_xy	c1f8e3fe12	修改图片路径	2026-06-23 12:20:34 +08:00
FrRay	30ed6e0e49	Merge branch 'master' of https://gitlink.org.cn/metax-maca/op_optimization	2026-06-23 11:42:26 +08:00
Yuliang Feng (i26389)	496c8cc691	fiximage	2026-06-23 11:40:37 +08:00
FrRay	7291f85688	将 START_HERE.md 重命名为选手入口.md，更新语言列与 Q&A	2026-06-23 11:40:16 +08:00
wu_xy	05b7888756	修改	2026-06-23 11:14:39 +08:00
wu_xy	1c091f2cb9	修改图片路径	2026-06-23 11:11:30 +08:00
wu_xy	c721c656c5	修改图片路径	2026-06-23 11:09:35 +08:00
Xinyi Wu (i26343)	7231888100	更新 MCTLASS_Fused MoE 算子优化文档	2026-06-23 11:01:19 +08:00
Xinyi Wu (i26343)	8fb0109857	更新 MCTLASS_Fused MoE 算子优化文档	2026-06-23 10:54:51 +08:00
Beckylu	fd6fd58c39	Merge pull request 'Update START_HERE.md：赛题二选手入口文档' (#37 ) from FrRay/op_optimization:master into master	2026-06-23 09:36:54 +08:00
FrRay	d9e691fa4d	Update START_HERE.md：赛题二选手入口文档	2026-06-22 18:35:35 +08:00
FrRay	c971c576a5	Update START_HERE.md	2026-06-22 18:34:17 +08:00
FrRay	3d1ff99abc	Update START_HERE.md：赛题二选手入口文档	2026-06-22 18:24:37 +08:00
Beckylu	582451ef58	Merge pull request '新增 START_HERE.md：赛题二选手入口文档' (#36 ) from FrRay/op_optimization:master into master	2026-06-22 17:32:19 +08:00
FrRay	80fb933979	Add START_HERE.md: 赛题二选手入口文档	2026-06-22 17:27:38 +08:00
Beckylu	2c7972281d	Merge pull request '请求合并' (#35 ) from xiao-ke/op_optimization:master into master	2026-06-22 15:50:18 +08:00
Ke Xiao (i26293)	adf3096c01	更新Flashinfer问答文档	2026-06-18 17:54:48 +08:00
Ke Xiao (i26293)	0a1db9a6a6	更新	2026-06-18 17:41:15 +08:00
Beckylu	64b7883e5d	更新任务包名字	2026-06-18 17:17:19 +08:00
yyyymmm	7b9de194ab	Update README.md	2026-06-18 15:32:29 +08:00
Beckylu	57f56aabb8	Merge pull request '请求合并' (#31 ) from xiao-ke/op_optimization:master into master	2026-06-18 11:12:15 +08:00
Beckylu	2d38aa3d1b	Merge pull request '更新XPU OJ 教程' (#32 ) from xiao_xiao/op_optimization:master into master	2026-06-18 11:11:56 +08:00
xiao	8e6231b6ab	Add xxx	2026-06-18 11:03:50 +08:00
Beckylu	6d623763d6	Merge pull request 'Moe教程增加OJ部分' (#30 ) from xxhefei/op_optimization:update-moark-agent-doc into master	2026-06-18 11:01:40 +08:00
Ke Xiao (i26293)	1dbb422167	更新题目包和Flashinfer教程	2026-06-18 10:59:48 +08:00
Ke Xiao (i26293)	790db292d7	Merge remote-tracking branch 'upstream/master' 本地更新题目包和Flashinfer教程，解决合并冲突	2026-06-18 10:56:16 +08:00
何沐川	12b2af4460	update OJ instructions	2026-06-18 10:40:08 +08:00
Beckylu	2e60758a4f	Merge pull request '更新XPU OJ 教程' (#29 ) from xiao_xiao/op_optimization:master into master	2026-06-18 10:23:27 +08:00
xiao	d4ac7f5a2f	oj教程更新	2026-06-18 10:21:54 +08:00
Ke Xiao (i26293)	398ec687c3	更新FlashInfer教程	2026-06-18 10:06:54 +08:00
yyyymmm	6b2e4a709c	Update README.md	2026-06-18 08:36:17 +08:00
yyyymmm	cc3d597feb	Update README.md	2026-06-18 08:35:30 +08:00
yyyymmm	a452e70d15	Update README.md	2026-06-18 08:35:04 +08:00
Beckylu	59a4fb8f8c	Merge pull request '更新XPU-OJ教程并修复合并冲突' (#25 ) from xiao-ke/op_optimization:master into master	2026-06-17 08:57:07 +08:00
xiaoke-shu	f302a83b62	更新XPU-OJ教程	2026-06-16 21:57:22 +08:00
Beckylu	dc5841a10c	Merge pull request '赛题二更新' (#22 ) from beckylu-dev into master	2026-06-11 17:10:51 +08:00
yyyymmm	167b72b2e1	Update README.md	2026-06-11 09:55:39 +08:00
yyyymmm	567c8c42a0	Update README.md	2026-06-10 15:15:30 +08:00
yyyymmm	8070486586	Update README.md	2026-06-10 15:07:44 +08:00
yyyymmm	7c13dd377f	Update README.md	2026-06-10 15:03:00 +08:00
yyyymmm	b85e2637b6	Update README.md	2026-06-10 14:56:18 +08:00
yyyymmm	cc05079507	Update README.md	2026-06-10 14:55:40 +08:00
yyyymmm	0385d36047	Update README.md	2026-06-10 14:54:25 +08:00
yyyymmm	ea98b9a069	Update README.md	2026-06-10 14:53:20 +08:00
yyyymmm	4e045a3e7d	Update README.md	2026-06-10 14:51:42 +08:00
yyyymmm	6382da7296	Update README.md	2026-06-10 14:46:20 +08:00
Beckylu	53880f68c3	Merge pull request '补充相关链接' (#21 ) from xiao_xiao/op_optimization:master into master	2026-06-09 10:35:13 +08:00
xiao	392362ba88	补充相关链接	2026-06-08 19:01:27 +08:00
Beckylu	9f27203d0d	Merge pull request '赛题二readme更新' (#20 ) from beckylu-dev into master	2026-06-08 18:14:48 +08:00