新增沐曦通用GPU MXMACA编译器内建函数编程指南入口

Merge pull request 'docs: 补充 XPU-OJ MoE 评分异常 FAQ' (#82 ) from yuting2003/op_optimization:codex/xpuoj-faq-notice into master
docs: add XPU-OJ MoE scoring FAQ
2026-08-07 16:35:06 +08:00 · 2026-08-03 14:34:53 +08:00 · 2026-08-03 14:30:14 +08:00 · 2026-08-03 14:27:34 +08:00 · 2026-08-03 14:23:07 +08:00 · 2026-08-03 14:21:38 +08:00
155 changed files with 7546 additions and 3509 deletions
--- a/.gitignore
+++ b/.gitignore
@ -0,0 +1,79 @@
+# =========================
+# macOS
+# =========================
+.DS_Store
+.AppleDouble
+.LSOverride
+Icon?
+._*
+.Spotlight-V100
+.Trashes
+.fseventsd
+
+# =========================
+# IDE / Editor
+# =========================
+.vscode/
+.idea/
+*.swp
+*.swo
+*~
+
+# =========================
+# Python
+# =========================
+__pycache__/
+*.py[cod]
+*.pyo
+*.pyd
+.pytest_cache/
+.mypy_cache/
+.ruff_cache/
+.coverage
+htmlcov/
+.env
+.venv/
+venv/
+env/
+
+# =========================
+# Jupyter
+# =========================
+.ipynb_checkpoints/
+
+# =========================
+# Build / Packaging
+# =========================
+build/
+dist/
+*.egg-info/
+.eggs/
+pip-wheel-metadata/
+
+# =========================
+# C / C++ / CMake
+# =========================
+CMakeFiles/
+CMakeCache.txt
+cmake-build-*/
+Makefile
+*.o
+*.so
+*.dylib
+*.dll
+*.a
+*.lib
+
+# =========================
+# Logs / Temp
+# =========================
+*.log
+*.tmp
+*.temp
+logs/
+tmp/
+
+# =========================
+# OS / Tool caches
+# =========================
+.cache/
--- a/FAQ.md
+++ b/FAQ.md
@ -0,0 +1,329 @@
+# 常见问题 FAQ
+
+> 最后整理：2026-07-13
+
+本文档汇总沐曦“揭榜挂帅”两项赛题的常见问题。使用 `Ctrl+F`（Windows/Linux）或 `⌘F`（macOS）搜索关键词。
+
+环境版本、评测配置、时间安排可能调整。请以仓库 README、XPU-OJ 公告、比赛群通知和容器中的实际版本为准。发现内容过期或未覆盖的问题，请[提交 Issue](https://gitlink.org.cn/metax-maca/op_optimization/issues)。
+
+## 快速导航
+
+- [重要入口与咨询方式](#entry)：XPU-OJ 开放情况、问题反馈渠道和赛事联系人。
+- [XPU-OJ、提交与排行榜](#xpuoj)：账号申请、提交环境、测试样例、评测硬件和排名指标。
+- [赛题一：TileLang 与 Fused MoE](#track-one)：提交限制、算子实现、Baseline、性能优化和决赛加分规则。
+- [赛题二：AI Agent 与推理算子库](#track-two)：任务范围、Agent 参与证明、Baseline、测试参数和上游代码引用。
+- [评测规则与通用技术问题](#evaluation)：技术资料发布、正确性与稳定性要求，以及 FAQ 内容纠错。
+- [环境、镜像与算力资源](#environment)：比赛镜像、MACA 与 PyTorch 版本、开发工具、算力券和资源申请。
+- [报名、组队与资格审核](#registration)：参赛资格、跨校组队、指导教师、材料盖章和审核流程。
+
+<a id="entry"></a>
+
+## 重要入口与咨询方式
+
+<a id="q-xpuoj-open"></a>
+
+### ❓ 问题 1：XPU-OJ 平台是否已经上线？
+
+**回答：** XPU-OJ 已开放。账号申领流程和使用指南见[赛事 XPU-OJ 账号申领说明](赛事XPUOJ账号申领说明.md)。
+
+<a id="q-contact"></a>
+
+### ❓ 问题 2：两个赛题的联系人不同，遇到问题应该联系谁？
+
+**回答：** 请先通过 [GitLink Issue](https://gitlink.org.cn/metax-maca/op_optimization/issues) 提问，维护者会把可复用的答案更新到本文档。比赛群用于接收赛事通知和临时信息。需要单独沟通时，请按对应比赛方案联系章老师或杨老师。
+
+## XPU-OJ、提交与排行榜
+
+<a id="q-moe-score-decrease"></a>
+
+### ❓ 问题 1：XPUOJ测评 MoE 耗时减少了但是分数反而降低了
+
+**回答：**
+
+针对近期部分同学反馈的“XPU.OJ”第三方评测系统中基线（baseline）不稳定的问题，我们高度重视，并已第一时间组织排查与测试。在此，我们对因此给大家带来的困扰深表歉意，也衷心感谢各位同学提出的宝贵意见。
+目前，相关问题已修复完毕。为确保评测的公平性与准确性，我们将对现有榜单进行清空处理。历史提交记录仍可查看，但后续排名将统一以基线修复后重新提交的算子成绩为准。
+
+比赛期间，我们将持续关注系统运行状态，也欢迎大家继续向我们反馈建议。
+祝大家比赛顺利，取得理想成绩！
+
+<a id="q-xpuoj-environment"></a>
+
+### ❓ 问题 2：XPU-OJ 与模力方舟的运行环境一致吗？
+
+**回答：** 排行榜评测环境与模力方舟开发环境保持一致。版本调整时以 XPU-OJ 公告为准。
+
+<a id="q-xpuoj-account"></a>
+
+### ❓ 问题 3：如何申请 XPU-OJ 账号？
+
+**回答：** 请按照[赛事 XPU-OJ 账号申领说明](赛事XPUOJ账号申领说明.md)提交申请。账号发放进度以赛事通知和回复邮件为准。
+
+<a id="q-official-ranking"></a>
+
+### ❓ 问题 4：赛题一 MoE 初赛排名以哪个入口为准？
+
+**回答：** 正式排名和初筛结果以 XPU-OJ 的评测结果为准。Sample benchmark 用于本地功能验证、调试和性能对比，不作为正式榜单依据。
+
+<a id="q-test-cases"></a>
+
+### ❓ 问题 5：MACA C++、Triton 和 TileLang 是否分别设榜？
+
+**回答：** 不按语言分别设榜。每个任务支持 MACA C++、Triton 和 TileLang，团队可以使用一种或多种语言提交。排行榜采用通过正确性和稳定性测试后的最高成绩。
+
+<a id="q-ranking-metric"></a>
+
+### ❓ 问题 6：排行榜使用 latency、speedup 还是综合 score？
+
+**回答：** 当前 XPU-OJ 以 speedup 作为核心排名指标。赛事方调整计算方式时，以 XPU-OJ 公告为准。
+
+<a id="track-one"></a>
+
+## 赛题一：TileLang 与 Fused MoE
+
+<a id="q-track-one-submission"></a>
+
+### ❓ 问题 1：赛题一的提交要求有哪些调整？
+
+**回答：** 正式提交禁止使用 `MACA Maca running` 方式，也不能使用 PyTorch 实现算子。参赛者需使用 TileLang 实现并提交。
+
+<a id="q-ops-reference"></a>
+
+### ❓ 问题 2：OPS 目录中的 TileLang、CUDA、CUTLASS 和 MACA 代码有什么用途？
+
+**回答：** 这些代码用于解释算子的实现原理和设计思路，可作为 TileLang 实现的参考。
+
+<a id="q-baseline-modification"></a>
+
+### ❓ 问题 3：官方 Baseline 可以修改到什么范围？
+
+**回答：** 参赛者可以重新设计和优化算子实现，但需保持与统一 Workload 测试框架的接口兼容。
+
+<a id="q-gemm-optimization"></a>
+
+### ❓ 问题 4：GEMM 计算中可以引入其他优化策略吗？
+
+**回答：** 可以，前提是实现符合赛题规则和评测要求。
+
+<a id="q-benchmark-modification"></a>
+
+### ❓ 问题 5：可以优化 `fusedmoe_benchmark` 吗？
+
+**回答：** 本地修改 benchmark 不会提高正式成绩，XPU-OJ 使用赛事方的评测框架。参赛者应把优化工作放在规定的算子实现和允许修改的接口上。
+
+<a id="q-forward-modification"></a>
+
+### ❓ 问题 6：可以修改 `fusedmoe_benchmark.py` 中的 MoE forward 吗？
+
+**回答：** 正式成绩以 XPU-OJ 的独立评测为准。本地修改 forward 不能替代对提交算子的优化，也不会改变赛事方的评测代码。
+
+<a id="q-async-copy"></a>
+
+### ❓ 问题 7：赛题一允许使用异步拷贝吗？
+
+**回答：** 不允许。当前比赛规则禁用异步拷贝。
+
+<a id="q-baseline-performance"></a>
+
+### ❓ 问题 8：Fused MoE 初赛成绩如何影响决赛？
+
+**回答：** 当前规则只给初赛前 10 名决赛加分，第 11 名及之后不获得额外初赛加分。
+
+<a id="track-two"></a>
+
+## 赛题二：AI Agent 与推理算子库
+
+<a id="q-track-two-update"></a>
+
+### ❓ 问题 1：赛题二的内容有什么调整？
+
+**回答：** “Agent 推理算子库优化 - FlashAttention KV Cache Decode”新增 `mctlass/cute` 要求。参赛者需基于 `mctlass/cute` 实现或优化对应任务。
+
+<a id="q-track-two-selection"></a>
+
+### ❓ 问题 2：赛题二可以选择几个任务？
+
+**回答：** 参赛团队可以从 FlashInfer、FlashAttention、MCTLASS/Fused MoE 等方向选择一项或多项。提交多个任务时，每个任务按赛事规则取有效最高成绩。支持语言包括 Triton、MXMACA C++ 和 TileLang。
+
+<a id="q-agent-proof"></a>
+
+### ❓ 问题 3：如何证明 AI Agent 参与了优化过程？
+
+**回答：** 参赛团队应保留 Agent 配置、Skill 文件、关键提示词、操作日志、代码变更记录、测试结果和复现实验步骤，等待赛事方发布核验细则。
+
+<a id="q-agent-baseline"></a>
+
+### ❓ 问题 4：Agent 赛题的性能 baseline 使用哪个版本？
+
+**回答：** 当前评测使用赛事方提供的 baseline，标准环境为 `PyTorch-Agent / 2.8.0 / Python 3.12 / MACA 3.7.1.5`。赛事方变更 baseline 或评测方式时会发布通知。
+
+<a id="q-mla-dimensions"></a>
+
+### ❓ 问题 5：MLA 的 `QK dim = 576, VO dim = 512` 与 `race_tests` 参数冲突吗？
+
+**回答：** 不冲突。`race_tests` 中的 `dim=512, pe_dim=64` 对应 `QK dim = 576, V dim = 512`。
+
+<a id="q-nsa-ranking"></a>
+
+### ❓ 问题 6：NSA 的 109 个测试 case 如何计算榜单成绩？
+
+**回答：** XPU-OJ 通过统一接口统计测试总运行时间，再根据 baseline 计算整体 speedup。
+
+<a id="q-upstream-code"></a>
+
+### ❓ 问题 7：可以引用或修改 FlashInfer、FlashAttention 等上游代码吗？
+
+**回答：** 可以。参赛团队需遵守上游项目许可证，保留版权和许可证声明，并在提交材料中说明引用范围、迁移工作和自主优化内容。
+
+<a id="evaluation"></a>
+
+## 环境、镜像与算力资源
+
+<a id="q-environment-image"></a>
+
+### ❓ 问题 1：比赛使用哪个在线算力环境？
+
+**回答：** 比赛使用模力方舟沐曦算力专区。仓库 README 当前标注的统一镜像为：
+
+```text
+PyTorch-Agent / 2.8.0 / Python 3.12 / MACA 3.7.1.5
+```
+
+创建实例和连接环境的步骤见[模力方舟快速使用 SOP](模力方舟快速使用SOP.md)。
+
+<a id="q-download-maca"></a>
+
+### ❓ 问题 2：如何获取 MACA 镜像或安装包？
+
+**回答：** 参赛者可以在[模力方舟沐曦算力专区](https://ai.gitee.com/compute/metax)选择 `PyTorch-Agent` 镜像。需要单独获取软件包时，请前往[沐曦开发者社区软件中心](https://developer.metax-tech.com/softnova/docker?chip_name=%E6%9B%A6%E4%BA%91C500%E7%B3%BB%E5%88%97&package_kind=AI&dimension=docker)，并选择与比赛标准环境一致的 MACA `3.7.1.5` 版本。
+
+<a id="q-download-pytorch"></a>
+
+### ❓ 问题 3：比赛使用的 PyTorch 镜像可以下载吗？
+
+**回答：** 可以。请在[沐曦开发者社区](https://developer.metax-tech.com/)或其 [PyTorch 镜像列表](https://developer.metax-tech.com/softnova/docker?chip_name=%E6%9B%A6%E4%BA%91C500%E7%B3%BB%E5%88%97&package_kind=AI&dimension=docker&deliver_type=%E5%88%86%E5%B1%82%E5%8C%85&ai_frame=pytorch)中查询。下载前请核对比赛镜像的 Python、PyTorch 和 MACA 版本。
+
+<a id="q-version-mismatch"></a>
+
+### ❓ 问题 4：页面标注的 MACA 版本与容器内版本不一致怎么办？
+
+**回答：** 比赛标准版本为 MACA `3.7.1.5`。
+
+<a id="q-pytorch-source"></a>
+
+### ❓ 问题 5：Linux 版本的 mcprofiler 是否可用？
+
+**回答：** mcprofiler 的 Linux 版本已经打包进模力方舟上的 pytorch-agent 比赛镜像。
+
+<a id="q-compute-coupons"></a>
+
+### ❓ 问题 6：算力券按团队还是按个人领取？
+
+**回答：** 新人礼和启悟社区算力券按学生个人发放，符合条件的团队成员均可领取。团队主申请人的额度用完后，其他成员可以继续申请资源。
+
+- [沐曦开发者社区新人礼](https://developer.metax-tech.com/activities/6)
+- [启悟社区学生算力券](https://developer.metax-tech.com/activities/11)
+- [赛事算力券活动](https://developer.metax-tech.com/activities/17)
+
+<a id="q-more-compute"></a>
+
+### ❓ 问题 7：算力额度不足时可以追加申请吗？
+
+**回答：** 可以先领取上述活动中的算力券。仍需额外资源时，请发送需求邮件至 `opensource@metax-tech.com`。
+
+<a id="q-commercial-agent-cost"></a>
+
+<a id="registration"></a>
+
+## 报名、组队与资格审核
+
+<a id="q-register-both"></a>
+
+### ❓ 问题 1：同一团队或个人可以同时参加两个赛题吗？
+
+**回答：** 可以。同一团队或个人可以同时报名两个赛题。
+
+<a id="q-register-multiple-tracks"></a>
+
+### ❓ 问题 2：同一名学生可以报名不同赛道的不同赛题吗？
+
+**回答：** 可以。赛事不统一限制学生报名不同赛道或不同赛题，但同一作品不得用相同核心技术内容重复申报不同赛题。
+
+<a id="q-new-graduate"></a>
+
+### ❓ 问题 3：本科应届毕业、尚未正式入学的研一新生可以报名吗？
+
+**回答：** 可以。参赛者可联系原本科学校完成认证手续，并以本科生身份报名。
+
+<a id="q-advisor-required"></a>
+
+### ❓ 问题 4：参赛必须配备指导教师吗？
+
+**回答：** 不强制。填写指导教师时，每支队伍可以设置 1 至 3 名指导教师。
+
+<a id="q-advisor-team-limit"></a>
+
+### ❓ 问题 5：一名指导教师最多可以指导几支队伍？
+
+**回答：** 赛事暂未设置统一的硬性上限。指导教师应根据可投入的时间控制队伍数量。
+
+<a id="q-cross-school-stamp"></a>
+
+### ❓ 问题 6：跨校组队时，报名表应该由哪所学校盖章？
+
+**回答：** 资格审批阶段，每名参赛者需到本人学校的校团委或院团委完成盖章确认。后续材料由团队牵头学生统一整理和提交。
+
+<a id="q-upload-stamped-form"></a>
+
+### ❓ 问题 7：提交报名后还可以补充已盖章的报名表扫描件吗？
+
+**回答：** 审核人员发现材料缺少盖章时，会退回申请。团队补齐材料后可以重新提交。尚未完成盖章的团队应先与学院、校团委或学校相关部门确认办理方式。
+
+<a id="q-stamp-department"></a>
+
+### ❓ 问题 8：资格审查材料应该加盖哪个部门的公章？
+
+**回答：** 各高校的管理口径不同。教务处、学生处等学籍或学生管理部门通常可以办理，参赛团队应以本校校团委或相关管理部门的要求为准。
+
+<a id="q-no-youth-league"></a>
+
+### ❓ 问题 9：学校未设校团委，可以用院系公章替代吗？
+
+**回答：** 赛事原则上要求校级部门公章。学校未设校团委时，可以联系校级学工、双创或教务部门盖章，并提交情况说明。院系公章不能直接替代校级部门公章。
+
+<a id="q-student-status-proof"></a>
+
+### ❓ 问题 10：学校无法配合盖章，可以用学籍证明替代吗？
+
+**回答：** 不可以。参赛团队应使用报名系统导出的报名表，并按要求完成学校盖章。
+
+<a id="q-public-notice"></a>
+
+### ❓ 问题 11：公示材料需要包含哪些内容？
+
+**回答：** 请参考赛事工作群发布的参考文本，并按学校要求调整。跨校团队涉及的学校应分别在学校官网公示，公示渠道原则上使用学校官网。
+
+<a id="q-review-deadline"></a>
+
+### ❓ 问题 12：报名审核需要在报名截止日前完成吗？
+
+**回答：** 原则上需要。往届出现过系统延后关闭的情况，但本届参赛团队不应据此推迟材料提交或审核。
+
+<a id="q-review-flow"></a>
+
+### ❓ 问题 13：报名材料的审核顺序是什么？
+
+**回答：** 学生提交材料后，学校校团委先审核；学校审核通过后，企业再审核。企业审核通过即视为报名成功。
+
+<a id="q-school-review-account"></a>
+
+### ❓ 问题 14：后台显示“校团委审核”，但学校不了解审核事项，怎么办？
+
+**回答：** 校团委需在报名系统内完成审核。省级团委通常会向各高校团委发放账号和密码。学校未收到或不了解安排时，请学校联系省级团委确认。
+
+## FAQ 维护约定
+
+1. 参赛者通过 Issue 提交问题。
+2. 维护者确认答案后更新本文档。
+3. 每个答案保留稳定锚点；需要时附上来源 Issue 或公告。
+4. 维护者在原 Issue 中回复 FAQ 锚点链接，并关闭已经解决的问题。
+5. 涉及版本、日期、评测参数的答案应标注确认日期。
--- a/README.md
+++ b/README.md
@ -1,5 +1,22 @@
 # 降低Token 成本，攻坚国产推理生态｜沐曦两大赛题登陆 2026 揭榜挂帅擂台赛，邀青年共破局！

+## 常用入口
+
+- [常见问题 FAQ](FAQ.md)
+- [沐曦通用GPU MXMACA编译器内建函数编程指南](https://developer.metax-tech.com/api/client/document/preview/1395/index.html)
+
+## XPU-OJ 基线修复及榜单调整通知
+
+针对近期部分同学反馈的“XPU.OJ”第三方评测系统中基线（baseline）不稳定的问题，我们高度重视，并已第一时间组织排查与测试。在此，我们对因此给大家带来的困扰深表歉意，也衷心感谢各位同学提出的宝贵意见。
+目前，相关问题已修复完毕。为确保评测的公平性与准确性，我们将对现有榜单进行清空处理。历史提交记录仍可查看，但后续排名将统一以基线修复后重新提交的算子成绩为准。
+
+比赛期间，我们将持续关注系统运行状态，也欢迎大家继续向我们反馈建议。
+祝大家比赛顺利，取得理想成绩！
+
+- XPU-OJ 地址：[https://xpuoj.com/](https://xpuoj.com/)
+- 账号申领说明：[赛事 XPU-OJ 账号申领说明](赛事XPUOJ账号申领说明.md)
+- 账号申领邮箱：`opensource@metax-tech.com`
+
 2026 年度中国青年科技创新「揭榜挂帅」擂台赛正式启幕。沐曦股份重磅发布两大 AI 算力硬核榜题，聚焦国产 GPU 大模型推理算子优化，以硬核赛事搭建科研攻关平台，邀全国青年学子、科研人才揭榜攻坚，用技术重构推理效率，用创新拉低每 Token 算力成本！

 ## 两大重磅赛题 直击推理成本核心痛点
@ -18,16 +35,18 @@

 初赛攻坚 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 大模型推理算子库优化

-大模型推理具有高并发、长序列、高调用频次等特点，FlashInfer、FlashAttention、Fused MoE 等核心算子直接决定模型服务的吞吐、延迟与显存开销，影响单 Token 综合推理成本。
+大模型推理具有高并发、长序列、高<EFBFBD><EFBFBD><EFBFBD>用频次等特点，FlashInfer、FlashAttention、Fused MoE 等核心算子直接决定模型服务的吞吐、延迟与显存开销，影响单 Token 综合推理成本。

 本赛题面向沐曦国产 GPU 及 MXMACA 软件栈，鼓励参赛团队构建或使用 AI Agent / Skill 工作流，围绕推理算子库开展代码理解、算子迁移、性能分析、Kernel 优化、自动调优、Benchmark 验证和多轮迭代，探索“Agent 驱动算子优化”的新型开发范式。

@ -44,9 +63,12 @@
 **赛题二相关资料**

 - [赛题二方案：基于 AI Agent 开发范式的国产 GPU 大模型推理算子库优化方案](基于AI%20Agent开发范式的国产GPU大模型推理算子库优化/基于AI%20Agent开发范式的国产GPU大模型算子推理库优化方案.md)
- [模力方舟 Agent 环境准备教程](基于AI%20Agent开发范式的国产GPU大模型推理算子库优化/模力方舟Agent环境准备教程.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.1.5**
+
 ## 参赛对象

 面向全国全日制专科、本科、硕博在校生（非在职），40 周岁以下青年科技人才均可参与；
@ -99,6 +121,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 +133,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/Agent开发范式的国产GPU大模型推理算子库优化/Agent说明.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/Agent说明.md
@ -1,42 +1,43 @@
 # Agent 使用说明
+本说明涵盖三个算子的 Agent 辅助能力（FlashInfer、FlashAttention、Fused MoE）以及使用 Agent 时的推荐方式和注意事项。

-## 一、FlashInfer：
+## 一、模块功能速查
+### 1.FlashInferFlashInfer：
 在本模块中，Agent 可以帮助你完成以下任务：

-### 环境检查
+#### 环境检查

 ```plaintext
 请帮我检查当前环境是否满足 FlashInfer 运行要求，包括 GPU、Python、PyTorch 和 flashinfer 依赖。
 ```

-### 运行测试
+#### 运行测试

 ```plaintext
 请帮我运行 bench_batch_decode.py 脚本，执行 BatchDecode 的基准测试。
 ```

-### 分析结果
+#### 分析结果

 ```plaintext
 请帮我读取最新的 CSV 结果文件，分析各参数配置下的性能表现，找出带宽最高和 TFLOPs 最高的配置。
 ```

-### 问题排查
+#### 问题排查

 ```plaintext
 运行时报错 out of memory，请帮我分析原因并给出解决方案。
 ```

-### 代码理解
+#### 代码理解

 ```plaintext
 请帮我解释 bench_common.py 中 run_with_profiler 函数的工作原理。
 ```

-## 二、FlashAttention：
+### 2. FlashAttention：
 在本模块中，Agent 可用于以下场景：
-
-### Prompt 模板
+#### Prompt 模板

 **环境验证：**

@ -62,10 +63,10 @@

 ```
 ---
-## 三、Fused MoE：
+### 3. Fused MoE：
 在本模块中，Agent可以帮助你完成以下任务：

-### 环境检查
+#### 环境检查
    
    ```plaintext
    我正在算力平台部署 fused_moe_baseline 源码。
@ -80,13 +81,13 @@
    2. 是否有潜在的不兼容风险（如 Python 与 libpython 版本）？
    ```
    
-### 运行测试
+#### 运行测试
    
    ```plaintext
    请帮我运行 scripts/run_fused_moe_i8_tn_pybind_test.sh 脚本 
    ```
    
-### 分析结果
+#### 分析结果
    
    ```plaintext
    这是性能测试结果：
@ -100,7 +101,7 @@
    3. 当前结果是否已经具备提交价值？
    ```
    
-### 报错检查
+#### 报错检查
    
    ```plaintext
    编译 pybind 时出现以下错误：
@ -116,15 +117,34 @@
    3. 如何在 build_fused_moe_i8_tn_pybind.sh 中正确指定路径？
    ```
    
-### 代码理解
+#### 代码理解
    
    ```plaintext
    请帮我梳理释 benchmark_fused_moe_i8_tn.py 代码整体框架
    ```
    
-### KernelSwift 系统搜索算子
+#### KernelSwift 系统搜索算子
    

 ```plaintext
 请帮我在算子广场检索 fused_moe 算子
-```
+```
+## 二、推荐工作方式
+使用 Agent 时，不要一开始就让它“直接优化到最快”。推荐节奏是：
+
+先让它检查目录。
+
+再让它跑通原始 baseline。
+
+然后让它只做一处小修改。
+
+每次修改后必须 build、test、benchmark。
+
+每轮都记录结果。
+
+## 三、注意事项
+使用 Agent 不要相信没有命令输出支持的环境判断。
+
+不要让 Agent 同时修改很多文件。
+
+不要让 Agent 修改测试标准。
--- a/Agent开发范式的国产GPU大模型推理算子库优化/FlashAttention_Baseline入门.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/FlashAttention_Baseline入门.md
@ -1,829 +0,0 @@
-# FlashAttention Baseline 入门：从环境验证到 KV-Cache Benchmark 结果记录（2）
-
-## 一、教程定位
-
-本教程是参赛训练课程的 FlashAttention Baseline 入门模块，主要帮助用户快速跑通 FlashAttention 的 KV-Cache 推理性能基准测试最小可运行流程。
-
-完成本教程后，用户应能够：
-
-*   完成环境验证和依赖检查
-    
-*   理解并配置基准测试参数
-    
-*   明确Baseline的解读，理解为什么KV-Cache是性能瓶颈
-    
-*   运行 KV-Cache Benchmark 测试
-    
-*   输出一份 baseline 性能结果记录表，为后续算子优化提供对比基准
-    
-
-> Baseline 解读：为什么本教程基于 KV Cache 做性能基线？
-
-#### 什么是 Baseline（性能基线）
-
-> Baseline是在引入任何优化代码前，系统处于初始可用状态时的参考数据。它是衡量后续所有优化收益的“锚点”。记录内容通常包括：执行时间、吞吐量、带宽、显存占用。
-
-#### 为什么需要 Baseline
-
-> Baseline 是未做任何优化前的参考性能指标，它可以回答：
-
-*   当前性能处于什么水平？
-    
-*   是否存在明显性能瓶颈？
-    
-*   后续优化是否有效？
-    
-*   性能提升有多大？
-    
-
-> 没有 Baseline，则无法量化优化效果。
-
-> **注意**：Baseline ≠ Benchmark。
-
-*   **Benchmark** 是测量性能的手段（脚本、工具）。
-    
-*   **Baseline** 是测量得到的具体结果数值。
-    
-
-> 本教程运行的 `benchmark_kvcache.py` 是一个 Benchmark 脚本，它输出的 CSV 文件就是Baseline。
-
-#### 为什么FlashAttention里面基线是针对KV Cache做benchmark性能验证？
-
-> 因为KV Cache是最核心的性能瓶颈，尤其是在大模型推理的解码阶段。
-
-##### 瓶颈从计算转移到显存访问
-
-> Transformer 推理分为两个截然不同的阶段：
-
-| > 阶段 | > 特点 | > 计算模式 | > 主要限制 |
-| --- | --- | --- | --- |
-| > **Prefill（预填充）** | > 一次性处理全部输入 token | > 矩阵运算密集，Tensor Core 利用率高 | > **计算受限（Compute Bound）** |
-| > **Decode（解码）** | > 逐个生成新 token，每步只算一个 token | > 每次都要**读取全部历史 KV Cache** | > **显存访问受限（Memory Bound）** |
-
-*   Prefill 阶段计算量大但形态规整，通常能较好利用 GPU 算力，不是主要瓶颈。
-    
-*   Decode 阶段占推理过程的大部分时间（尤其是长上下文交互），每个 token 的生成都需要搬运整个 KV Cache。随着序列增长，显存访问开销占比越来越高。
-    
-
-##### KV‑Cache 的访存密集型特征
-
-> 在 Decode 阶段，flash\_attn\_with\_kvcache 内核的工作是：
-
-*   读取当前 token 的 Q 向量（很小）
-    
-*   **反复读取整个历史的 KV Cache**（很大，线性增长）
-    
-*   执行 FlashAttention 计算，然后将新的 K/V 追加写入缓存。
-    
-
-> 95% 以上的时间花在读取 KV Cache 上。因此，Decode 阶段的性能完全由显存带宽决定，而不是 GPU 算力。
-
-##### KV Cache 显存容量直接限制并发能力
-
-> 大模型推理服务需要同时处理多个请求（batch\_size）。每个请求都有自己的 KV‑Cache，显存总占用量与 `batch_size × 序列长度` 成正比。因此，KV Cache 的显存开销直接决定了系统可以同时服务多少用户。
-
-> 通过 Benchmark 对不同 `batch_size × seq_len_kv` 组合进行压力测试，可以：
-
-*   找出 **OOM 边界**：哪些参数组合会导致显存溢出，无法运行
-    
-*   量化每个请求的平均显存开销
-    
-*   为后续 **分页 KV‑Cache（PagedAttention）** 等优化提供基线对比。
-    
-
-##### 为什么选择 flash\_attn\_with\_kvcache作为测试对象
-
-> 它是 FlashAttention 专门为推理阶段设计的核心算子。它融合了高效的 KV Cache 读取与更新逻辑、FlashAttention 的分块、重计算技术以节省显存
-
-> 推理场景下精确的带宽优化。选择它作为 Benchmark 对象，可以直接回答：
-
-*   GPU带宽利用率是多少？是否接近理论峰值？
-    
-*   哪些参数组合（batch\_size, seq\_len\_kv, headdim）达到峰值性能？
-    
-*   OOM边界在哪里？
-    
-*   后续优化（如分页缓存、算子融合）是否有效？
-    
-
-#### 为什么要记录 Baseline 结果
-
-> Baseline（基线）是优化前的参考性能数据，记录 baseline 的意义在于：
-
-*   **量化优化收益**：优化后对比 baseline，计算加速比（speedup = baseline\_time / optimized\_time）
-    
-*   **防止性能回退**：代码变更后重跑 benchmark，确认没有引入性能退化
-    
-*   **建立测试矩阵**：记录不同参数组合下的 baseline，全面了解性能特征
-    
-
-> 在本教程中，baseline 结果以 CSV 文件保存，包含每种 `batch_size × seq_len_kv` 配置的执行时间和带宽，为后续算子优化提供对比基准。
-
---
-
-## 二、学习目标
-
-完成本模块后，你将能够：
-
-1.  理解 FlashAttention `flash_attn_with_kvcache` 核函数的基本作用与应用场景；
-    
-2.  利用预装专属镜像，完成沐曦 GPU 硬件环境的快速验证；
-    
-3.  深入理解 Baseline 的概念，掌握性能测试与正确性测试的联系与区别；
-    
-4.  跑通 KV-Cache Benchmark 基准测试脚本；
-    
-5.  完成多种 `batch_size × seq_len_kv` 组合的性能测试；
-    
-6.  输出带宽性能结果 CSV 文件并进行结果分析。
-    
-
---
-
-## 三、适用对象
-
-本模块适合以下人员：
-
-*   参与 AI 基础设施竞赛的参赛者
-    
-*   对 GPU 算子性能优化感兴趣的开发者
-    
-*   需要了解 FlashAttention KV-Cache 推理性能的研究人员
-    
-
-**基础知识要求：**
-
-*   了解 Python 编程基础
-    
-*   了解 PyTorch 基本用法
-    
-*   了解 GPU 推理的基本概念
-    
-
---
-
-## 四、前置准备
-
-开始实战前，请确认你已经完成以下准备：
-
-### 环境准备
-
-#### 开发环境设置
-
-1.  **在沐曦开发者社区领取算力券**
-    
-    *   领取链接：[https://developer.metax-tech.com/activities/6](https://developer.metax-tech.com/activities/6)
-        
-        *   登录平台
-            
-        
-        ![platform login](https://origin.picgo.net/2026/06/04/platform-login626620122b08424d.png)
-        
-        *   首次登录需要先进行注册（使用邮箱或者手机号进行注册）
-            
-        
-        ![platform registration](https://origin.picgo.net/2026/06/04/platform-registrationdb267074af39bf4c.png)
-        
-        *   登录成功后进行第二步-邮箱验证，填入自己的邮箱。
-            
-        
-        ![email verification](https://origin.picgo.net/2026/06/04/email-verificationc3f391bb747318e0.png)
-        
-        *   第三步，提交申请。
-            
-        
-        ![submit application](https://origin.picgo.net/2026/06/04/submit-application3bf7ac4724e13ae8.png)
-        
-        *   获得兑换码
-            
-    
-    ![get redeem code](https://origin.picgo.net/2026/06/04/get-redeem-code3f6a20e5f9cbbd38.png)
-    
-2.  **在模力方舟平台兑换算力券**
-    
-    *   平台链接：[https://ai.gitee.com/](https://ai.gitee.com/)
-        
-    *   1.登录模力方舟平台
-        
-    
-    ![ai.gitee login](https://origin.picgo.net/2026/06/04/ai.gitee-login7d9fe2b5e35e3a92.png)
-    
-    *   2.进入费用中心 - 算力券 ， 点击右上角“兑换”
-        
-    
-    ![redeem compute voucher](https://origin.picgo.net/2026/06/04/redeem-compute-voucher0eb15e2f3f9b7bbd.png)
-    
-3.  **租用算力**
-    
-    *   模力方舟算力市场链接：https://ai.gitee.com/compute
-        
-    *   选择沐曦芯片厂商，并根据项目要求选择相应的配置。
-        
-    
-    ![rent compute](https://origin.picgo.net/2026/06/04/rent-compute1197cc6d884ce429.png)
-    
-4.  **创建实例**
-    
-
-专属镜像文件：
-
-![create instance1](https://origin.picgo.net/2026/06/04/create-instance10ab33dd1b7e14727.png)
-
-![create instance2](https://origin.picgo.net/2026/06/04/create-instance23666efb720fefa60.png)
-
-![create instance3](https://origin.picgo.net/2026/06/04/create-instance3f18b4323644d3447.png)
-
-  进入算力容器，刚创建的实例默认开机状态，点击工具-lab开始项目创作。
-
-**重要说明：**由于本次使用的是预装的专属镜像，环境中已经默认安装并配置好了 PyTorch、FlashAttention、einops 等所有依赖包。因此在启动实例后，无需再进行繁琐的依赖库版本验证即可直接进入测试环节。
-
-### 代码准备
-
-*   已获取基准测试脚本 `benchmark_kvcache.py`
-    
-*   已进入项目目录 `/data/flashattn_baseline`
-    
-
---
-
-## 五、知识预备
-
-### 名词解释
-
-| 术语 | 说明 |
-| --- | --- |
-| **KV-Cache** | Key-Value Cache，Transformer 推理时缓存历史 token 的 Key 和 Value 向量，避免重复计算 |
-| **Paged KV-Cache** | 将 KV-Cache 分页管理，提高显存利用率，类似操作系统的虚拟内存分页机制 |
-| **flash\_attn\_with\_kvcache** | FlashAttention 提供的带 KV-Cache 支持的注意力计算核函数 |
-| **batch\_size** | 批大小，一次处理的样本数量 |
-| **seq\_len\_kv** | KV 序列长度，KV-Cache 中缓存的历史 token 数量 |
-| **headdim** | Head Dimension，注意力头的维度 |
-| **带宽 (Bandwidth)** | 显存带宽，单位 GB/s，衡量 GPU 读写显存的速度 |
-
-### 核心概念详解
-
-#### 什么是正确性测试与性能测试
-
-*   **正确性测试 (Correctness Testing)：**解决“算得对不对”的问题。它的目标是验证当前算子的输出结果，在数学精度上是否与标准参考实现完全一致。这是所有测试的绝对前提底线。
-    
-*   **性能测试 (Performance Testing)：**解决“跑得快不快”的问题。它的目标是在验证正确性的基础上，测量算子在特定硬件上的执行耗时、吞吐量和有效带宽利用率。本教程执行的 Benchmark 脚本，正是一个纯粹的性能测试。
-    
-*   **二者区别：**
-    
-
-| 维度 | 性能测试 | 正确性测试 |
-| --- | --- | --- |
-| 测试目标 | 测量速度、带宽 | 验证输出结果 |
-| 关注输出 | 否 | 是 |
-| 关注效率 | 是 | 否 |
-| 是否需要 Baseline | 是 | 不一定 |
-| 是否受实现不同而影响 | 大 | 是（精度不同可能影响） |
-
-**二者联系：**
-
-在实际开发中：
-
-正确性测试（先）
-
-      ↓
-
-建立 Baseline
-
-      ↓
-
-性能分析
-
-      ↓
-
-优化实现
-
-      ↓
-
-性能测试对比 Baseline
-
-      ↓
-
-回归正确性验证（保证没变坏）
-
-在算子优化迭代中，每一次修改底层代码，都必须**先通过正确性测试**确立功能基准，**再运行性能测试**对比性能基准 Baseline，确保速度的提升绝不是以牺牲结果正确性为代价。
-
-#### 什么是 Benchmark（基准测试）
-
-Benchmark 是一种标准化的性能测量方法，通过在固定条件下反复运行同一任务，获取可重复、可对比的性能指标。在 GPU 算子优化场景中，benchmark 的作用是：
-
-*   **建立性能基线**：在优化前记录原始性能数据，作为后续对比的参照
-    
-*   **量化优化效果**：优化后运行同样的 benchmark，直接对比时间/带宽变化
-    
-*   **发现性能瓶颈**：通过不同参数组合的测试结果，定位性能拐点
-    
-
-> 参考：[MLPerf Benchmark 介绍](https://mlcommons.org/benchmarks/)
-
-#### 什么是 batch\_size、seq\_len、headdim
-
-这三个参数共同决定了注意力计算的**工作量**和**显存占用**：
-
-*   **batch\_size（批大小）**：一次推理同时处理的样本数量。batch\_size 越大，GPU 并行度越高，但显存占用也线性增长。在 KV-Cache 场景中，batch\_size 对应同时服务的请求数。
-    
-*   **seq\_len / seq\_len\_kv（序列长度）**：序列中 token 的数量。seq\_len\_kv 特指 KV-Cache 中已缓存的历史 token 数量。序列越长，注意力计算的计算量呈 O(n²) 增长（但 FlashAttention 将其优化为 O(n) 显存），KV-Cache 的显存占用则呈 O(n) 线性增长。
-    
-*   **headdim（注意力头维度）**：每个注意力头的向量维度。常见的有 64、128、256。headdim 越大，单个 token 的 Key/Value 向量越宽，KV-Cache 的显存占用与 headdim 成正比。
-    
-
-三者与显存占用的关系：
-
-```Plain
-KV-Cache 显存 ≈ batch_size × seq_len_kv × num_heads_k × headdim × 2(K+V) × bytes_per_elem
-```
-
-> 参考：[Attention Is All You Need (Vaswani et al., 2017)](https://arxiv.org/abs/1706.03762)
-
-#### 什么是 Kernel 执行时间
-
-Kernel（核函数）是运行在 GPU 上的并行计算函数。Kernel 执行时间指从 GPU 开始执行该核函数到执行完毕所花费的时间，通常以**毫秒 (ms)** 为单位。
-
-测量方式有两种：
-
-*   **CPU 端计时**：使用 `torch.cuda.synchronize()` + `time.time()`，包含 GPU 调度开销，时间偏大
-    
-*   **GPU 端计时**：使用 CUDA Event 或 profiler，精度更高，直接测量 GPU 上的实际执行时间
-    
-
-本教程使用 GPU 端同步计时。
-
-> 参考：[PyTorch CUDA Semantics](https://pytorch.org/docs/stable/notes/cuda.html)
-
-#### 什么是有效带宽
-
-有效带宽（Effective Bandwidth）是衡量 kernel 实际利用显存带宽效率的指标，计算公式为：
-
-```Plain
-有效带宽 (GB/s) = 数据传输量 (GB) / kernel 执行时间 (s)
-
-```
-
-GPU 显存带宽是有限的（例如沐曦 C500 的理论峰值带宽），有效带宽越接近理论峰值，说明 kernel 对显存带宽的利用率越高。对于**访存密集型**算子（如 KV-Cache 注意力），有效带宽是衡量优化效果的核心指标。
-
-*   有效带宽 **接近理论峰值** → kernel 已接近最优，优化空间有限
-    
-*   有效带宽 **远低于理论峰值** → 存在优化空间（如内存访问不合并、bank conflict 等）
-    
-
-> 参考：[CUDA C++ Programming Guide - Performance Guidelines](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#performance-guidelines)
-
-#### 为什么要 Warmup / Repeat
-
-GPU 程序的首次运行往往比后续运行慢，原因包括：
-
-*   **JIT 编译**：部分框架会延迟编译 kernel 代码
-    
-*   **缓存冷启动**：GPU L2 Cache、TLB 等初始状态为空
-    
-*   **频率爬升**：GPU 需要时间从低功耗状态切换到高频率状态
-    
-
-因此，benchmark 流程通常分为两步：
-
-1.  **Warmup（预热）**：先运行若干次（如 10 次），不记录时间，让 GPU 进入稳定状态
-    
-2.  **Repeat（重复测量）**：正式运行多次（如 100 次），记录每次时间，取统计值（均值/中位数）
-    
-
-重复测量可以消除随机波动，获得更可靠的性能数据。次数越多，结果越稳定，但耗时也越长。
-
-> 参考：[PyTorch Benchmark Utils](https://pytorch.org/tutorials/recipes/recipes/benchmark.html)
-
-#### CUDA Stream 与同步
-
-CUDA 采用异步执行模型，CPU 提交 kernel 到 GPU 后不等待完成就继续执行。`torch.cuda.synchronize()` 会阻塞 CPU 直到 GPU 上所有已提交的任务完成，这是精确计时的前提。
-
-> 参考：[CUDA Streams](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#asynchronous-concurrent-execution)
-
-#### 数据类型（dtype）对性能的影响
-
-不同数据类型占用的字节数不同，直接影响显存带宽需求和计算吞吐：
-
-| 数据类型 | 字节数 | 说明 |
-| --- | --- | --- |
-| float32 | 4 | 单精度浮点，精度最高 |
-| float16 | 2 | 半精度浮点，精度足够且带宽减半 |
-| bfloat16 | 2 | Brain Float 16，动态范围与 float32 相同，训练/推理常用 |
-
-本教程使用 `bfloat16`，在精度和性能之间取得平衡。
-
-> 参考：[Mixed Precision Training (Micikevicius et al., 2018)](https://arxiv.org/abs/1710.03740)
-
-#### Paged KV-Cache 与 Block Table
-
-传统 KV-Cache 为每个请求预分配连续显存，容易造成碎片和浪费。Paged KV-Cache（灵感来自操作系统虚拟内存）将显存分成固定大小的 page/block，通过 **block\_table** 映射逻辑位置到物理位置：
-
-*   **page\_block\_size**：每个 block 包含的 token 数量（本教程默认 16）
-    
-*   **block\_table**：索引张量，记录每个 batch 的 KV-Cache 页面映射关系
-    
-*   **优势**：减少显存碎片，支持动态分配，提高多请求并发效率
-    
-
-> 参考：[Efficient Memory Management for Large Language Model Serving with PagedAttention (Kwon et al., 2023)](https://arxiv.org/abs/2309.06180)
-
-#### OOM（Out of Memory）
-
-OOM 表示 GPU 显存不足，无法完成当前计算。常见原因：
-
-*   batch\_size 或 seq\_len\_kv 过大，超出显存容量
-    
-*   同时存在多个占用显存的进程
-    
-*   未释放的中间变量占用显存
-    
-
-应对策略：减小 batch\_size/seq\_len\_kv、使用更小的 dtype（如 bfloat16 替代 float32）、使用梯度检查点等。
-
-> 参考：[PyTorch CUDA Memory Management](https://pytorch.org/docs/stable/notes/cuda.html#memory-management)
-
-#### Tensor Core 与矩阵乘法加速
-
-现代 GPU（包括沐曦 C500）配备 Tensor Core 单元，专门加速矩阵乘法运算。Attention 计算中的 Q×K^T 和 Attn×V 都是矩阵乘法，能够受益于 Tensor Core 加速。Tensor Core 对数据类型和矩阵维度有对齐要求（通常要求维度为 8 或 16 的倍数），这也是 headdim 通常取 64/128/256 的原因之一。
-
-> 参考：[NVIDIA Tensor Core Technology](https://developer.nvidia.com/tensor-cores)
-
-### 相关链接
-
-*   [FlashAttention 官方仓库](https://github.com/Dao-AILab/flash-attention)
-    
-*   [FlashAttention API 文档](https://github.com/Dao-AILab/flash-attention/blob/main/flash_attn/flash_attn_interface.py)
-    
-*   [FlashAttention 论文 (Dao et al., 2022)](https://arxiv.org/abs/2205.14135)
-    
-*   [FlashAttention-2 论文 (Dao, 2023)](https://arxiv.org/abs/2307.08691)
-    
-*   [PyTorch CUDA 编程最佳实践](https://pytorch.org/docs/stable/notes/cuda.html)
-    
-
---
-
-## 六、项目实践：FlashAttention KV-Cache Benchmark
-
-### 项目目标
-
-对 FlashAttention 的 paged KV-cache 推理核函数（`flash_attn_with_kvcache`）进行自动化性能基准测试，覆盖多种 `batch_size × seq_len_kv` 组合，输出执行时间和有效显存带宽。
-
-基准测试脚本 `benchmark_kvcache.py` 的最小闭环包括：
-
-1.  准备 paged KV-cache 张量和 block table。
-    
-2.  通过 `torch.profiler` 对 kernel 进行计时。
-    
-3.  计算有效显存带宽（GB/s）。
-    
-4.  将结果写入带时间戳的 CSV 文件。
-    
-
---
-
-### 步骤 0：进入创建的实例环境
-
-模力方舟链接：https://ai.gitee.com/fwlhecko/dashboard/compute/instances
-
-选择工具-lab进入实例环境
-
-![lab enter instance environment](https://origin.picgo.net/2026/06/04/lab-enter-instance-environment67995d391603a686.png)
-
-### 步骤 1：检查运行环境
-
-**目标：** 确认当前环境满足本模块运行要求。
-
-在JupyterLab Terminal中检查运行环境的配置。
-
-![jupyterlab terminal check](https://origin.picgo.net/2026/06/04/jupyterlab-terminal-check143d7d42426453cf.png)
-
-**操作：** 检查 GPU 状态、Python 版本和依赖版本。
-
-**命令示例：**
-
-```Bash
-# 检查沐曦 GPU 状态
-mx-smi
-
-# 检查 Python 版本
-python --version
-
-# 检查 PyTorch 是否能识别 GPU
-python -c "import torch; print(f'GPU available: {torch.cuda.is_available()}'); print(f'GPU count: {torch.cuda.device_count()}')"
-
-# 检查依赖版本
-python -c "import torch; print(f'PyTorch {torch.__version__}')"
-python -c "import flash_attn; print(f'flash-attn {flash_attn.__version__}')"
-python -c "import einops; print('einops OK')"
-
-```
-
-**预期结果：**
-
-*   `mx-smi` 显示沐曦 GPU 信息
-    
-
-![result mx smi](https://origin.picgo.net/2026/06/04/result-mx-smif86a3bed6681382e.png)
-
-*   Python 版本 = 3.8
-    
-
-![result python](https://origin.picgo.net/2026/06/04/result-pythond15f856ddb84c649.png)
-
-*   `torch.cuda.is_available()` 返回 `True`
-    
-
-![result gpu available](https://origin.picgo.net/2026/06/04/result-gpu-available44e049addf5638fb.png)
-
-*   所有依赖版本符合要求
-    
-
-![result dependency version](https://origin.picgo.net/2026/06/04/result-dependency-versionc17845898ee89434.png)
-
-**常见问题：**
-
-| 问题 | 解决方法 |
-| --- | --- |
-| `torch.cuda.is_available()` 返回 `False` | 检查 MXMACA 环境变量是否正确配置 |
-| `ModuleNotFoundError: No module named &#39;flash_attn &#39;` | 确认 flash-attn 已安装且版本= 2.6 |
-| `mx-smi` 命令不存在 | 确认已配置沐曦 GPU 驱动环境 |
-
---
-
-### 步骤 2：进入项目目录
-
-**目标：** 进入本模块所需的项目目录。
-
-1.  将flashattn\_baseline压缩文件进行解压，并将解压文件加入JupyterLab。
-    
-
-![flashattn baseline extract](https://origin.picgo.net/2026/06/04/flashattn-baseline-extract66d40632f4253973.png)
-
-**操作：** 切换到基准测试脚本所在目录。
-
-**命令示例：**
-
-```Bash
-cd flashattn_baseline
-ls -la
-
-```
-
-**预期结果：**
-
-```Plain
-total xx
-drwxr-xr-x 2 root root 4096 Jun  1 09:00 __MACOSX
-
- rw-r--r-- 1 root root 5232 Jun  1 09:00 benchmark_kvcache.py
-
-
- rw-r--r-- 1 root root 1440 Jun  1 09:00 benchmark_kvcache_20260526_150953.csv
-
-...
-
-```
---
-
-### 步骤 3：配置基准测试参数
-
-**目标：** 根据测试需求配置基准测试参数。
-
-**操作：** 在 `benchmark_kvcache.py` 的 `main()` 函数中配置基准测试参数。
-
-**参数说明：**
-
-| 参数 | 默认值 | 说明 |
-| --- | --- | --- |
-| `headdims` | `[256]` | head dimension，可改为 `[128]` 等 |
-| `page_block_size` | `16` | paged KV-cache 的 block 大小 |
-| `batch_sizes` | `[1, 2, 4, 8, 16, 32, 64, 128]` | 批大小扫描范围 |
-| `seq_lens_kv` | `[512, 1024, 2048, 4096, 8192, 16384]` | KV 序列长度扫描范围 |
-| `num_heads` | `8` | query head 数量 |
-| `num_heads_k` | `8` | KV head 数量 |
-| `seqlen_q` | `1` | query 序列长度（单 token 推理） |
-| `dtype` | `torch.bfloat16` | 数据类型 |
-| `causal` | `False` | 是否启用 causal mask |
-| `warmup` | `10` | 预热迭代次数 |
-| `repeat` | `100` | 正式 profiling 迭代次数 |
-
-**配置示例：**
-
-如需测试 `headdim=128`，修改对应列表：
-
-```Python
-headdims = [128]
-
-```
-
-如需启用 causal mask：
-
-```Python
-causal = True
-
-```
---
-
-### 步骤 4：运行基准测试
-
-**目标：** 运行基准测试脚本，收集性能数据。
-
-**操作：** 执行基准测试脚本并观察输出。
-
-**命令示例：**
-
-```Bash
-cd flashattn_baseline
-python benchmark_kvcache.py
-
-```
-
-**预期结果：**
-
-脚本运行时会在终端实时打印结果表格：
-
-```Plain
-batch_size   seq_len_kv  heads  headdim    time_ms  bandwidth_GB_s
---------------------------------------------------------------------------
-         1          512      8      256     0.0948           44.27
-         2          512      8      256     0.0988           84.95
-       ...
-
-```
-
-同时生成带时间戳的 CSV 文件，命名格式为 `benchmark_kvcache_YYYYMMDD_HHMMSS.csv`。
-
-**常见问题：**
-
-| 问题 | 解决方法 |
-| --- | --- |
-| 出现 `OOM` 标记 | 该配置超出 GPU 显存容量，可减小 batch\_size 或 seq\_len\_kv |
-| 脚本运行缓慢 | 减少 `repeat` 次数或缩小扫描范围 |
-
---
-
-### 步骤 5：查看与分析结果
-
-**目标：** 理解输出结果格式，分析性能数据。
-
-**CSV 输出格式：**
-
-| 列名 | 说明 |
-| --- | --- |
-| `batch_size` | 批大小 |
-| `seq_len_kv` | KV 序列长度 |
-| `heads` | head 数量 |
-| `headdim` | head 维度 |
-| `time_ms` | kernel 执行时间（毫秒） |
-| `bandwidth_GB_s` | 有效显存带宽（GB/s） |
-
-如果某个配置因显存不足而失败，对应的 `time_ms` 和 `bandwidth_GB_s` 列会标记为 `OOM`。
-
-**带宽计算公式：**
-
-```Plain
-total_bytes = q_bytes + kv_bytes
-q_bytes   = batch_size × seqlen_q × num_heads × headdim × bytes_per_elem
-kv_bytes  = batch_size × seqlen_k × num_heads_k × headdim × bytes_per_elem × 2
-bandwidth = (total_bytes / 1e9) / (time_ms / 1e3)   [GB/s]
-
-```
-
-其中 `bytes_per_elem` 在 `bfloat16` 下为 2，`float32` 下为 4。
-
-**性能观察：**
-
-*   **小 batch 时带宽较低**：batch\_size=1 时 kernel 无法充分利用 GPU 并行度，带宽通常 &lt; 100 GB/s
-    
-*   **大 batch + 长序列时带宽较高**：batch\_size=128 时可接近 GPU 显存带宽上限
-    
-*   **headdim 增大时 OOM 风险增加**：headdim=256 的显存占用是 headdim=128 的两倍，大 batch + 长序列更容易 OOM
-    
-
-**测试结果示例：**
-
-#### headdim=128（2026-05-26）
-
-所有 48 个配置均成功运行，峰值带宽约 1251 GB/s。
-
-| batch\_size | seq\_len\_kv | time\_ms | bandwidth\_GB\_s |
-| --- | --- | --- | --- |
-| 1 | 512 | 0.0322 | 65.27 |
-| 128 | 512 | 0.2453 | 1095.50 |
-| 1 | 16384 | 0.8356 | 80.32 |
-| 128 | 16384 | 6.8668 | 1250.98 |
-
-#### headdim=256（2026-05-27）
-
-48 个配置中有 3 个因显存不足（OOM）而失败，峰值带宽约 807 GB/s。
-
-| batch\_size | seq\_len\_kv | time\_ms | bandwidth\_GB\_s |
-| --- | --- | --- | --- |
-| 1 | 512 | 0.0948 | 44.27 |
-| 128 | 4096 | 5.3221 | 807.10 |
-| 128 | 8192 | OOM | OOM |
-| 64 | 16384 | OOM | OOM |
-| 128 | 16384 | OOM | OOM |
-
---
-
-### 步骤 6：自定义扩展（可选）
-
-如需进一步测试，可参考以下扩展方法：
-
-**修改 headdim：**
-
-```Python
-headdims = [128, 256]
-
-```
-
-**启用 causal mask：**
-
-```Python
-causal = True
-
-```
-
-**调整 profiling 精度：**
-
-```Python
-warmup = 20
-repeat = 200
-
-```
-
-**启用详细 profiler 输出：**
-
-```Python
-ms = run_with_profiler(run_fn, warmup=warmup, reps=repeat, print_result=True, target_kernels=["flash"])
-
-```
---
-
-## 七、Agent 使用说明
-
-在本模块中，Agent 可用于以下场景：
-
-### Prompt 模板
-
-**环境验证：**
-
-```Plain
-请帮我验证当前环境是否满足 FlashAttention KV-Cache Benchmark 的运行要求，包括：
-1. 沐曦 GPU 是否可见
-2. PyTorch 版本和 CUDA 支持
-3. flash-attn 和 einops 是否已安装
-
-```
-
-**参数配置建议：**
-
-```Plain
-我需要测试 headdim=128 和 headdim=256 的性能差异，请帮我推荐合适的 batch_sizes 和 seq_lens_kv 扫描范围。
-
-```
-
-**结果分析：**
-
-```Plain
-请帮我分析这份 benchmark 结果 CSV 文件，找出峰值带宽配置和 OOM 边界。
-
-```
---
-
-## 八、常见问题
-
-### Q1: 运行时提示 `ModuleNotFoundError: No module named &#39;flash_attn &#39;`
-
-**原因：** flash-attn 未安装或版本不兼容。 **解决：** 确认已安装 flash-attn = 2.6，可使用 `pip show flash-attn` 检查。
-
-### Q2: 所有配置都显示 OOM
-
-**原因：** GPU 显存不足或 batch\\_size 设置过大。 **解决：** 减小 `batch_sizes` 和 `seq_lens_kv` 的范围重新测试。
-
-### Q3: `mx-smi` 命令无输出或报错
-
-**原因：** 沐曦 GPU 驱动未正确安装或环境变量未配置。 **解决：** 确认已正确配置 MXMACA 环境，检查 `/usr/local/maca` 目录是否存在。
-
-### Q4: 带宽数值异常低
-
-**原因：** 可能是 warmup 不足或 GPU 未达到稳态。 **解决：** 增加 `warmup` 次数，如 `warmup = 20`。
-
---
-
-## 九、下一步学习建议
-
-完成本模块后，建议继续学习以下内容：
-
-1.  **算子优化基础** — 了解如何分析 kernel 性能瓶颈
-    
-2.  **FlashAttention 源码解析** — 深入理解 `flash_attn_with_kvcache` 的实现原理
-    
-3.  **自定义 kernel 开发** — 学习如何编写和优化沐曦 GPU 上的算子
-    
-4.  **性能对比分析** — 将 baseline 结果与优化后结果进行对比
-    
-
---
--- a/Agent开发范式的国产GPU大模型推理算子库优化/FlashAttention关键算子迁移与优化.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/FlashAttention关键算子迁移与优化.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/FlashInfer
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/FlashInfer
--- a/Agent开发范式的国产GPU大模型推理算子库优化/FlashInfer
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/FlashInfer
@ -1,349 +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)。
-
-**操作：** 切换到 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大模型推理算子库优化/Fused
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/Fused
@ -1,604 +0,0 @@
-# Fused MoE Baseline 入门：快速跑通最小闭环教程
-
-## 一、教程定位
-
-本教程是参赛训练课程的 baseline 入门模块，主要帮助用户快速跑通 Fused MoE 的最小可运行流程。
-
-完成本教程后，用户应能够完成源码编译、正确性测试和 benchmark 测试，并记录一份 baseline 性能结果，为后续算子优化提供对比基准。
-
-## 二、学习目标
-
-完成本模块后，你将能够：
-
-1.  理解 Fused MoE 推理算子及 int8 量化在大模型推理加速中的核心作用与优化意义；
-    
-2.  完成环境配置和 Fused MoE baseline 源码准备；
-    
-3.  跑通 Fused MoE 的最小闭环示例；
-    
-4.  理解基于 Trition 与 MXMACA C++ 写Fused MOE 算子；
-    
-5.  完成数值正确性测试，即验证 reference 计算、pybind 计算、Triton 计算这三种方式计算结果是否数值完全一致。
-    
-    *   reference：基于 PyTorch 架构在 CPU 上运行的**数值基准**实现；
-        
-    *   pybind：将 MXMACA C++ 算子编译并封装为 Python 可调用的动态库，**实现复杂且迁移成本高**；
-        
-    *   Triton：基于 Python 编写的高效 GPU Kernel，可利用 Agent 自动调优，**开发效率高、易于迁移**；
-        
-    *   要求 pybind 和 Triton 结果均与 reference 一致，鼓励参赛选手持续调优 Triton ，使其性能逼近甚至超越 pybind 性能。
-        
-6.  输出 benchmark 结果对比表。
-    
-
-## 三、适用对象
-
-**本模块适合以下人员：**
-
-*   参与基于 AI Agent 开发范式的国产 GPU 大模型推理算子库优化比赛的学生；
-    
-*   对 GPU 推理算子性能优化感兴趣的开发者；
-    
-*   需要了解 Fused MoE 推理性能的研究人员。
-    
-
-**学习本模块前，需掌握以下基础知识：**
-
-*   Python、C++ 编程基础；
-    
-*   PyTorch 基础；
-    
-*   GPU 推理基本概念。
-    
-
-## 四、前置准备
-
-**开始实战前，请确认你已经完成以下准备：**
-
-**环境准备：**
-
-*   已进入赛事专属镜像环境。
-
-**工具准备：**
-
-*   已准备 Agent 工具；
-*   已配置 Token / API Key；
-*   已确认 Agent 可以正常调用模型。
-
-**代码准备：**
-
-*   已获取 Fused MoE Baseline 源码。
-
-
-## 五、项目实践1-算力平台 baseline 源码部署
-
-**项目目标：**在算力平台上拉取 Fused MoE 的 baseline 源码，快速跑通最小闭环，并记录一份 baseline 性能结果，为后续算子优化提供对比基准。
-
-### 步骤 1：检查运行环境
-
-**目标：**确认当前环境满足本模块运行要求，包括编译器、MXMACA 工具链及 Python 依赖库。
-
-**操作：**检查 Python、编译工具、MXMACA 编译器及关键 Python 包（numpy、torch、triton）是否存在。
-
-**命令示例：**
-
-```apl
-python --version # 检查Python版本，Python ≥ 3.8
-g++ --version # 确认 C++ 编译器存在
-which mxcc # 确认 MACA 编译器存在
-
-
-# 检查 Python 依赖
-python - << 'EOF'
-import sys
-deps = ["numpy", "torch", "triton"]
-missing = []
-for d in deps:
-    try:
-        __import__(d)
-    except ImportError:
-        missing.append(d)
-if missing:
-    print(f"[ERROR] Missing packages: {missing}")
-    sys.exit(1)
-else:
-    print("[OK] numpy, torch, triton are installed.")
-EOF
-```
-
-
-**预期结果：**
-
-*   Python 3.12.11
-    
-*   g++ (Ubuntu 13.3.0-6ubuntu2~24.04.1) 13.3.0
-    
-*   /opt/maca/mxgpu\_llvm/bin/mxcc
-    
-*   \[OK\] numpy, torch, triton are installed.
-    
-
-**常见问题：**
-
-| 报错 | 原因 | 解决办法 |
-| --- | --- | --- |
-| `g++：command not found` | 未安装 C++ 编译工具 | `apt update && apt install -y build-essential` |
-| `Python 3.6.x/ Python 3.7.x` | Python 版本过低 | `conda install python=3.12` （推荐3.10+) |
-| `ModuleNotFoundError: numpy` | 当前 Python 缺少依赖 | `pip install numpy torch triton` |
-
-
-### 步骤 2：进入项目目录
-
-**目标：**进入本模块所需的源码目录。
-
-**操作：**切换到指定项目路径。
-
-**命令示例：**
-
-```apl
-cd /root/Project/fused_moe # 这里需要替换成自己的项目路径
-```
-
-
-**预期结果：**
-
-终端提示符路径显示为 fused\_moe 项目路径：
-
-*   (base) root@0a5a9d0c0f06:~/Project/fused\_moe#
-    
-
-### 步骤 3：pybind 编译
-
-**目标：**将用 C++ 编写的 fused\_moe 算子编译为 Python 可调用的 pybind 模块。
-
-**操作：**运行 `fused_moe/scripts/build_fused_moe_i8_tn_pybind.sh` 脚本
-
-**命令示例：**
-
-```apl
-bash scripts/build_fused_moe_i8_tn_pybind.sh
-```
-
-切换 Python 环境命令示例：
-
-```apl
-PYTHON_BIN=/path/to/python bash scripts/build_fused_moe_i8_tn_pybind.sh
-```
-
-
-**预期结果：**
-
-编译成功无报错，终端显示：
-
-*   \[SUCCESS\] /root/Project/fused\_moe/standalone/fused\_moe\_i8\_tn/build/fused\_moe\_i8\_tn\_ pybind.so
-    
-    且成功生成 `fused_moe/standalone/fused_moe_i8_tn/build/fused_moe_i8_tn_pybind.cpython-310-x86_64-linux-gnu.so` 文件
-    
-
-**常见问题：**
-
-
-| 报错 | 原因 | 解决办法 |
-| --- | --- | --- |
-| `Python.h: No such file or directory` | Python 头文件路径未找到 | 确认 `PYTHON_BIN` 路径正确，脚本自动探测 `sysconfig.get_path('include')` |
-| `libpython3.x.so: cannot find` | 链接时找不到 Python 库 | 1、执行 `find $CONDA_PREFIX -name "libpython3*.so*"`查找绝对路径<br>2、将该路径赋值给 `LIBPYTHON_PATH` |
-| `recompile with -fPIC` | 编译未开启位置无关代码 | 确保 `mxcc`/ `g++`编译参数中有 `-fPIC` |
-| `permission denied` | 无脚本执行权限 | `chmod +x scripts/*.sh` |
-| `undefined reference to Py_...` | Python 版本不匹配 | 确认编译脚本中`PYTHON_BIN`路径与当前运行的 Python 环境完全一致 |
-    
-### 步骤 4：正确性测试
-    
-**目标：**验证 reference 计算、pybind 计算、Triton 计算这三种方式计算结果的数值是否完全一致。
-    
-**操作：**运行 `fused_moe/scripts/run_fused_moe_i8_tn_pybind_test.sh` 脚本
-    
-**命令示例：**
-    
-```apl
-bash scripts/run_fused_moe_i8_tn_pybind_test.sh --backend all # 运行全部计算方式
-    
-# --backend：选择计算方式
-# 只测 pybind：
-bash scripts/run_fused_moe_i8_tn_pybind_test.sh --backend pybind
-# 只测 triton：
-bash scripts/run_fused_moe_i8_tn_pybind_test.sh --backend triton
-# 只测 reference:
-bash scripts/run_fused_moe_i8_tn_pybind_test.sh --backend reference
-```
-
-
-**预期结果：**
-    
-编译成功无报错，输出示例如下：
-    
-> pybind:fused\_moe\_i8\_tn\_topk1 passed: rows=256, cols=128, sample C\[0\]=0.69531, C\[last\]=-0.44531  
-> pybind:fused\_moe\_i8\_tn\_topk2 passed: rows=512, cols=128, sample C\[0\]=-0.57813, C\[last\]=-0.49805 
-> pybind:fused\_moe\_i8\_tn\_topk3 passed: rows=384, cols=128, sample C\[0\]=-1.08594, C\[last\]=-0.33594
-    
-> reference:fused\_moe\_i8\_tn\_topk1 passed: rows=256, cols=128, sample C\[0\]=0.6934, C\[last\]=-0.4451 
-> reference:fused\_moe\_i8\_tn\_topk2 passed: rows=512, cols=128, sample C\[0\]=-0.5768, C\[last\]=-0.4975  
-> reference:fused\_moe\_i8\_tn\_topk3 passed: rows=384, cols=128, sample C\[0\]=-1.0875, C\[last\]=-0.3362
-    
-> triton:fused\_moe\_i8\_tn\_topk1 passed: rows=256, cols=128, sample C\[0\]=0.69337, C\[last\]=-0.44513  
-> triton:fused\_moe\_i8\_tn\_topk2 passed: rows=512, cols=128, sample C\[0\]=-0.57678, C\[last\]=-0.49749  
-> triton:fused\_moe\_i8\_tn\_topk3 passed: rows=384, cols=128, sample C\[0\]=-1.08748, C\[last\]=-0.33618
-    
-
-**结果解释：**
-    
-*   “pybind/reference/Triton”：三种计算方式；
-        
-*   “fused\_moe\_i8\_tn\_topk1/2/3 passed”：测试算子通过数值校验，数值误差在允许范围内且无明显异常，否则会报错 FAILED；
-        
-*   ”rows=... , cols=...“：输出 Tensor 的行和列的大小；
-        
-*   ”sample C\[0\]=... , C\[last\]=...“：首尾采样值，用于辅助定位数值偏差，不作为精度判定依据。
-    
-
-**常见问题：**
-
-
-| 报错 | 原因 | 解决办法 |
-| :--- | :--- | :--- |
-| `ModuleNotFoundError: fused_moe_i8_tn_pybind` | pybind 模块未编译或未加入 `PYTHONPATH` | 回到步骤 3，确认 `.so` 已生成；执行 `export PYTHONPATH=/root/Project/fused_moe:$PYTHONPATH` |
-| `FAILED: max abs diff too large` | 数值误差超过阈值 | 检查 scale 是否应用位置错误；确认 TopK 索引与权重是否一致 |
-| `FAILED: shape mismatch` | 输出张量形状不一致 | 检查 Token Permute / Unpermute 逻辑；确认 expert 维度对齐 |
-| `FAILED: NaN or Inf detected` | 溢出或未初始化内存 | 检查 INT8 乘加是否溢出；确认 GEMM 输出是否反量化 |
-| 终端长时间无输出 | Kernel 死锁或 Launch 失败 | 减小测试 shape；检查是否触发 MACA 硬件限制 |
-    
-### 步骤5：性能测试
-    
-**目标：**输出 benchmark 结果对比表
-    
-**操作：**运行 `fused_moe/scripts/run_fused_moe_i8_tn_benchmark.sh` 脚本
-    
-**命令示例：**
-    
-```apl
-bash scripts/run_fused_moe_i8_tn_benchmark.sh --backend all --warmup 5 --iters 20 
-# --backend：选择计算方式
-# --warmup：设置预热次数
-# --iters：设置迭代次数
-```
-    
-    
-**预期结果：**
-    
-编译成功无报错，输出示例如下：
-    
-> pybind:fused\_moe\_i8\_tn\_topk1 benchmark: avg\_ms=0.308978, TOPS=0.027149, warmup=5, iters=20
-> pybind:fused\_moe\_i8\_tn\_topk2 benchmark: avg\_ms=0.304500, TOPS=0.055098, warmup=5, iters=20
-> pybind:fused\_moe\_i8\_tn\_topk3 benchmark: avg\_ms=0.297775, TOPS=0.042256, warmup=5, iters=20
-    
-> reference:fused\_moe\_i8\_tn\_topk1 benchmark: avg\_ms=1685.43, TOPS=0.000005, warmup=5, iters=20
-> reference:fused\_moe\_i8\_tn\_topk2 benchmark: avg\_ms=3384.52, TOPS=0.000005, warmup=5, iters=20 
-> reference:fused\_moe\_i8\_tn\_topk3 benchmark: avg\_ms=2532.14, TOPS=0.000005, warmup=5, iters=20
-    
-> triton:fused\_moe\_i8\_tn\_topk1 benchmark: avg\_ms=19.013421, TOPS=0.000441, warmup=5, iters=20 
-> triton:fused\_moe\_i8\_tn\_topk2 benchmark: avg\_ms=16.745914, TOPS=0.001002, warmup=5, iters=20
-> triton:fused\_moe\_i8\_tn\_topk3 benchmark: avg\_ms=19.630328, TOPS=0.000641, warmup=5, iters=20
-    
-    
-**结果解释：**
-    
-*   “pybind/reference/Triton”：三种计算方式；
-        
-*   “fused\_moe\_i8\_tn\_topk1/2/3”：分别对应选择前 1 / 2 / 3 个专家场景下的 MoE 算子；
-        
-*   “avg\_ms”：平均算子执行耗时（毫秒），这里不计算预热时间，只计算正式迭代的<E4BBA3><E79A84><EFBFBD>间<EFBFBD><E997B4><EFBFBD>
-        
-*   “TOPS”：Tera Operations Per Second，本次 MoE 算子的总运算量 / 实际耗时；
-        
-*   “warmup=5, iters=20”：预热轮数和正式迭代数。
-    
-
-**常见问题：**
-
-
-| 报错 | 原因 | 解决办法 |
-| --- | --- | --- |
-| `ModuleNotFoundError: fused_moe_i8_tn_pybind` | pybind 模块未编译或未加入 `PYTHONPATH` | 回到步骤 3，确认 `.so`已生成；执行 `export PYTHONPATH=/root/Project/fused_moe:$PYTHONPATH` |
-| 终端长时间无输出 | Kernel 死锁或 MACA 驱动异常 | 减小测试 shape；重启容器或设备 |
-| avg\_ms 异常抖动（±50%） | 其他进程占用 GPU | 关闭其他占用显存的进程，单机单任务运行 |
-
-## 六、项目实践2-Kernel Swift 智能算子迁移系统自动调优
-
-系统链接：[https://deeplink.org.cn/kernelswift/task](https://deeplink.org.cn/kernelswift/task)
-
-**项目目标：**基于 KernelSwift 智能算子迁移系统，对 Fused MoE 算子进行在线自动调优。通过输入算子代码，一键生成适配沐曦硬件的高性能实现，高效完成算子优化与全流程追踪。
-
-### 步骤1：复用算子广场的 fused\_moe 算子进行二次优化
-
-**目标：**通过提交算子广场的 fused\_moe 算子代码发起自动优化流程，实现二次优化
-
-**操作：**
-
-1.  进入算子广场：点击左侧导航栏 【算子广场】，进入算子列表页
-    
-    搜索 fused\_moe 算子，复制 `input_code.py` 代码，也可直接复制以下代码：
-    
-    ```python
-    import torch
-    import torch.nn as nn
-    import torch.nn.functional as F
-    
-    
-    class Model(nn.Module):
-        """
-        Reference PyTorch MoE forward (no fused kernels).
-        Expects inputs:
-            hidden_states: (M, in_size)
-            w1: (E, hidden_size, in_size) where hidden_size = 2 * up_dim
-            w2: (E, out_size, up_dim)
-            topk_weights: (M, top_k)
-            topk_idx: (M, top_k)
-            top_k: int
-            renormalize: bool
-        """
-    
-        def __init__(self):
-            super().__init__()
-    
-        def forward(
-            self,
-            hidden_states: torch.Tensor,
-            w1: torch.Tensor,
-            w2: torch.Tensor,
-            topk_weights: torch.Tensor,
-            topk_idx: torch.Tensor,
-            top_k: int,
-            renormalize: bool = True,
-        ) -> torch.Tensor:
-            if renormalize:
-                topk_weights = topk_weights / topk_weights.sum(dim=-1, keepdim=True)
-    
-            seq_len = hidden_states.size(0)
-            out_size = w2.size(1)
-            output = hidden_states.new_zeros(seq_len, out_size)
-            num_experts = w1.size(0)
-    
-            # Accumulate expert contributions
-            for eid in range(num_experts):
-                token_idx, k_idx = torch.where(topk_idx == eid)
-                if token_idx.numel() == 0:
-                    continue
-                gate_proj, up_proj = w1[eid].chunk(2, dim=0)
-                down_proj = w2[eid]
-                tmp = F.linear(hidden_states[token_idx], gate_proj)
-                tmp = F.silu(tmp) * F.linear(hidden_states[token_idx], up_proj)
-                tmp = F.linear(tmp, down_proj)
-                tmp = tmp * topk_weights[token_idx, k_idx, None]
-                output.index_add_(0, token_idx, tmp.to(output.dtype))
-            return output
-    
-    
-    # Hyperparameters
-    seq_len = 128
-    in_size = 128
-    hidden_size = 256  # 2 * up_dim
-    out_size = 128
-    num_experts = 32
-    top_k = 4
-    
-    dtype = torch.float16
-    
-    def get_inputs():
-        hidden_states = (torch.rand(seq_len, in_size, dtype=dtype) - 0.5) / 2
-        w1 = (torch.rand(num_experts, hidden_size, in_size, dtype=dtype) - 0.5) / 2
-        w2 = (torch.rand(num_experts, out_size, hidden_size//2, dtype=dtype) - 0.5) / 2
-        routing_logits = (torch.rand(seq_len, num_experts, dtype=dtype) - 0.5) / 2
-        routing_weights = torch.softmax(routing_logits, dim=-1, dtype=torch.float32)
-        topk_weights, topk_idx = torch.topk(routing_weights, top_k, dim=-1)
-        return [hidden_states, w1, w2, topk_weights, topk_idx, top_k, True]
-    
-    def get_init_inputs():
-        return []
-    ```
-    
-2.  进入新建任务页：点击左侧导航栏【新建任务】 ，进入算子提交页面。
-    
-3.  编写算子代码：在 `model.py` 编辑器中输入刚刚复制的 fused\_moe 算子代码。
-    
-    如果想自行编写算子代码，需严格遵循标准格式规范：输入代码必须包含 `class Model` 定义算子实现，`get_init_inputs` 和 `get_inputs` 定义测试用例，确保优化过程可验证算子正确性。
-    
-4.  配置优化参数
-    
-    *   指定任务名称：支持字母、下划线、数字组合，示例：fused\_moe\_01
-        
-    *   选择适配硬件：算子需要适配的目标硬件厂商及型号，建议：沐曦
-        
-    *   最大演化轮次：优化算法迭代次数，取值范围40-400，建议默认40，复杂算法可提高至100+
-        
-5.  提交优化任务：点击右下角 \[优化\] 按钮，系统将提交任务并进入 \[生成中\] 状态
-    
-
-    ![image.png](https://origin.picgo.net/2026/06/04/image659a3f15b96be22c1.png)
-
-  完成上述步骤将看到如下界面：
-
-  ![image.png](https://origin.picgo.net/2026/06/04/image75c4ccf75e90cf428.png)
-
-### 步骤2：任务查看与结果管理
-
-**目标：** 在新建优化任务后可追踪任务进度，获取优化结果
-
-**操作：**
-
-1.  查看任务列表：点击左侧【任务查看】，可看到所有提交的优化任务
-    
-    *   任务状态：排队中、环境初始化、算子预编译、精度验证、性能调优、已完成、失败
-        
-    *   任务信息：任务名称、进度、创建时间、适配硬件
-        
-    *   操作按钮：查看详情、删除任务
-        
-    
-    ![image.png](https://origin.picgo.net/2026/06/04/image75c4ccf75e90cf428.png)
-    
-2.  追踪任务进度：当前任务状态为【运行中】时，点击任务列表中的【查看详情】按钮，追踪任务进度：
-    
-    *   左侧：原始算子代码（输入的 `input_code.py`）
-        
-    *   右侧：任务进度条，包含以下阶段：
-        
-        1.  环境初始化：准备目标硬件编译环境
-            
-        2.  算子预编译：验证算子代码是可正常编译
-            
-        3.  精度验证：验证优化后算子输出与原始算子误差的可接受范围
-            
-        4.  性能调优：按设定的演化轮次迭代优化算子性能
-            
-    
-    *   顶部：任务名称、创建/更新时间、适配硬件、当前轮次进度
-        
-
-    ![image.png](https://origin.picgo.net/2026/06/04/image9e9b23d9b25ddfe25.png)
-
-3.  获取优化结果：当前任务状态为【已完成】时，可在详情页查看优化结果：
-    
-    *   优化后算子代码支持一键复制
-        
-    *   算子加速比（基准耗时 / 优化后耗时）、性能数据（如延迟、吞吐量）
-        
-    *   可点击【Diff 对比】查看优化前后代码差异，理解性能提升逻辑
-        
-    
-    ![image.png](https://origin.picgo.net/2026/06/04/image1083021a27904294b1.png)
-    
-4.  任务异常处理
-    
-    *   任务失败：查看错误日志，常见原因包括代码不符合规范、测试用例错误、硬件适配问题，修改后重新提交任务；
-        
-    *   排队时间长：可调整提交时间，或联系平台管理员确认资源状态。
-
-
-## 七、Agent使用说明
-
-在本模块中，Agent可以帮助你完成以下任务：
-
-1.  **环境检查**
-    
-    ```plaintext
-    我正在算力平台部署 fused_moe_baseline 源码。
-    需要的环境信息如下：
-    - Python 3.12
-    - g++ 13.3.0
-    - mxcc 已安装
-    - numpy / torch / triton 已安装
-    
-    请帮我确认：
-    1. 当前环境是否满足编译与运行要求？
-    2. 是否有潜在的不兼容风险（如 Python 与 libpython 版本）？
-    ```
-    
-2.  **运行测试**
-    
-    ```plaintext
-    请帮我运行 scripts/run_fused_moe_i8_tn_pybind_test.sh 脚本 
-    ```
-    
-3.  **分析结果**
-    
-    ```plaintext
-    这是性能测试结果：
-    pybind:    avg_ms=0.30,  TOPS=0.027
-    triton:    avg_ms=19.01, TOPS=0.0004
-    reference: avg_ms=1685,  TOPS=0.000005
-    
-    请分析：
-    1. 为什么 pybind 比 Triton 快这么多？
-    2. TOPS 指标是否可信？
-    3. 当前结果是否已经具备提交价值？
-    ```
-    
-4.  **报错检查**
-    
-    ```plaintext
-    编译 pybind 时出现以下错误：
-    /usr/bin/ld: cannot find -lpython3.10
-    
-    已知：
-    - 使用的是 Conda Python 3.10
-    - mxcc 编译正常
-    
-    请一步一步告诉我：
-    1. 错误原因是什么？
-    2. 如何用 find 命令定位 libpython3.10.so？
-    3. 如何在 build_fused_moe_i8_tn_pybind.sh 中正确指定路径？
-    ```
-    
-5.  **代码理解**
-    
-    ```plaintext
-    请帮我梳理释 benchmark_fused_moe_i8_tn.py 代码整体框架
-    ```
-    
-6.  **KernelSwift 系统搜索算子**
-    
-
-    ```plaintext
-    请帮我在算子广场检索 fused_moe 算子
-    ```
-
-## 八、常见问题与注意事项
-
-### 算力平台 basline 源码部署项目：
-
-1.  环境准备与依赖问题
-    
-    *   确保算力平台已正确安装 Python 和 C++、MACA 编译器及相关运行时库，避免因环境缺失导致编译失败；
-        
-    *   镜像环境使用 Conda Python 作为默认运行环境，避免系统 Python 与 Conda Python 混用，防止 `Python.h`或 `libpython`路径错误。
-        
-2.  pybind 编译与链接
-    
-    *   若`Python.h not found`，请检查脚本中`PYTHON_INCLUDE`是否指向当前 Python 的 `include`目录；
-        
-    *   若`libpython not found`，请直接指定 Conda 下的`**libpython3.x.so**`绝对路径，避免链接系统静态库；
-        
-    *   编译 `pybind`模块时，务必开启 `-fPIC`，否则会出现 `recompile with -fPIC`错误。
-        
-3.  性能测试建议
-    
-    *   benchmark 应在关闭其他占用 GPU 的任务后执行，避免干扰性能数据；
-        
-    *   多次运行取平均值，避免单次抖动影响结果；
-        
-    *   性能对比应基于相同随机种子、相同 shape、相同 TopK、相同 batch size的条件下进行，降低误差。
-        
-
-### Kernel Swift 智能算子迁移系统自动调优项目：
-
-1.  代码规范问题
-    
-    输入代码需符合以下标准格式：
-    
-    *   `class Model`，表示待优化的算子实现；
-        
-    *   `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大模型推理算子库优化/Fused
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/Fused
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/README.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/README.md
@ -1,3 +0,0 @@
-# README
-
-此目录将持续更新baselines。
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashinfer_baseline/BatchDecodeWithPagedKVCacheWrapper_20260525_093143.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashinfer_baseline/BatchDecodeWithPagedKVCacheWrapper_20260525_093143.csv
@ -1,145 +0,0 @@
-api,batch_size,seq_len_q,seq_len_kv,num_qo_heads,num_kv_heads,head_dim,time_ms,bandwidth_GB_s,tflops
-BatchDecodeWithPagedKVCacheWrapper,1,1,512,32,8,64,0.02042879999999998,51.528822055137894,0.8212531328320811
-BatchDecodeWithPagedKVCacheWrapper,1,1,512,32,4,128,0.02333952000000001,45.27805199078642,0.718832949435121
-BatchDecodeWithPagedKVCacheWrapper,1,1,512,32,4,256,0.0319488,66.15384615384615,1.0502564102564103
-BatchDecodeWithPagedKVCacheWrapper,1,1,1024,32,8,64,0.023262719999999973,90.32684054143292,1.4424122372620245
-BatchDecodeWithPagedKVCacheWrapper,1,1,1024,32,4,128,0.025041919999999992,84.07278675117566,1.3399304845634845
-BatchDecodeWithPagedKVCacheWrapper,1,1,1024,32,4,256,0.033387520000000004,126.11562643766291,2.0099984664928687
-BatchDecodeWithPagedKVCacheWrapper,1,1,2048,32,8,64,0.028298240000000037,148.36258368011562,2.371485435136599
-BatchDecodeWithPagedKVCacheWrapper,1,1,2048,32,4,128,0.027745280000000008,151.4670603432367,2.418748846650673
-BatchDecodeWithPagedKVCacheWrapper,1,1,2048,32,4,256,0.03723775999999999,225.7115358174069,3.604344837068611
-BatchDecodeWithPagedKVCacheWrapper,1,1,4096,32,8,64,0.03886591999999997,215.93992886312756,3.4533526544592306
-BatchDecodeWithPagedKVCacheWrapper,1,1,4096,32,4,128,0.03426815999999998,245.03212311370103,3.916689078141344
-BatchDecodeWithPagedKVCacheWrapper,1,1,4096,32,4,256,0.066048,254.26356589147287,4.064248062015504
-BatchDecodeWithPagedKVCacheWrapper,1,1,8192,32,8,64,0.052495359999999984,319.6722910367698,5.1135082414902975
-BatchDecodeWithPagedKVCacheWrapper,1,1,8192,32,4,128,0.04628480000000001,362.6548672566371,5.799646017699114
-BatchDecodeWithPagedKVCacheWrapper,1,1,8192,32,4,256,0.08975359999999999,374.0330861380491,5.981608670849972
-BatchDecodeWithPagedKVCacheWrapper,1,1,16384,32,8,64,0.08625152000000001,389.0775258221536,6.224480588863825
-BatchDecodeWithPagedKVCacheWrapper,1,1,16384,32,4,128,0.0638464,525.6776263031276,8.408789093825181
-BatchDecodeWithPagedKVCacheWrapper,1,1,16384,32,4,256,0.13059071999999994,514.0123892417473,8.222190857053247
-BatchDecodeWithPagedKVCacheWrapper,2,1,512,32,8,64,0.02342912,89.86013986013987,1.4321678321678322
-BatchDecodeWithPagedKVCacheWrapper,2,1,512,32,4,128,0.02486784,84.99073502161829,1.3493102738315832
-BatchDecodeWithPagedKVCacheWrapper,2,1,512,32,4,256,0.03340287999999998,126.54812998160644,2.009074187614961
-BatchDecodeWithPagedKVCacheWrapper,2,1,1024,32,8,64,0.02839040000000001,148.02524797114512,2.3637871956717755
-BatchDecodeWithPagedKVCacheWrapper,2,1,1024,32,4,128,0.028165120000000012,149.5000908925649,2.382694055626249
-BatchDecodeWithPagedKVCacheWrapper,2,1,1024,32,4,256,0.03740160000000001,225.16084873374396,3.5885557837097872
-BatchDecodeWithPagedKVCacheWrapper,2,1,2048,32,8,64,0.03881984000000001,216.30176734370872,3.457451859667633
-BatchDecodeWithPagedKVCacheWrapper,2,1,2048,32,4,128,0.03601408000000001,233.38072220642587,3.7268126243957904
-BatchDecodeWithPagedKVCacheWrapper,2,1,2048,32,4,256,0.06728704000000002,249.82498858621207,3.9894080048698815
-BatchDecodeWithPagedKVCacheWrapper,2,1,4096,32,8,64,0.052490240000000014,319.7815060476004,5.114007023019897
-BatchDecodeWithPagedKVCacheWrapper,2,1,4096,32,4,128,0.04626431999999999,362.9924745462595,5.802213368747235
-BatchDecodeWithPagedKVCacheWrapper,2,1,4096,32,4,256,0.08993791999999999,373.44870773084375,5.969349880450872
-BatchDecodeWithPagedKVCacheWrapper,2,1,8192,32,8,64,0.08536063999999999,393.18618042226495,6.289443378119003
-BatchDecodeWithPagedKVCacheWrapper,2,1,8192,32,4,128,0.0630784,532.2077922077922,8.51116883116883
-BatchDecodeWithPagedKVCacheWrapper,2,1,8192,32,4,256,0.12952576,518.3650881492608,8.289793659577834
-BatchDecodeWithPagedKVCacheWrapper,2,1,16384,32,8,64,0.15207424000000003,441.34401723789637,7.0606423809844445
-BatchDecodeWithPagedKVCacheWrapper,2,1,16384,32,4,128,0.10330112,649.8017446471055,10.394290245836638
-BatchDecodeWithPagedKVCacheWrapper,2,1,16384,32,4,256,0.2281984,588.3060354498541,9.410599057662106
-BatchDecodeWithPagedKVCacheWrapper,4,1,512,32,8,64,0.0283904,148.3137962128043,2.3637871956717764
-BatchDecodeWithPagedKVCacheWrapper,4,1,512,32,4,128,0.028078080000000036,150.5470459518598,2.3900802334062696
-BatchDecodeWithPagedKVCacheWrapper,4,1,512,32,4,256,0.03707903999999999,228.00331400165706,3.619773543220106
-BatchDecodeWithPagedKVCacheWrapper,4,1,1024,32,8,64,0.03844096000000004,218.64677677144357,3.4915290356952546
-BatchDecodeWithPagedKVCacheWrapper,4,1,1024,32,4,128,0.03641856000000004,231.23857725291697,3.6854210600309254
-BatchDecodeWithPagedKVCacheWrapper,4,1,1024,32,4,256,0.06640640000000002,253.63145720894363,4.04231303006939
-BatchDecodeWithPagedKVCacheWrapper,4,1,2048,32,8,64,0.059007999999999984,284.5986984815619,4.5491366594360105
-BatchDecodeWithPagedKVCacheWrapper,4,1,2048,32,4,128,0.04641792000000003,362.1442753143611,5.783013456871825
-BatchDecodeWithPagedKVCacheWrapper,4,1,2048,32,4,256,0.08961023999999998,375.1799794309223,5.991178151068451
-BatchDecodeWithPagedKVCacheWrapper,4,1,4096,32,8,64,0.09185279999999997,365.484949832776,5.84490523968785
-BatchDecodeWithPagedKVCacheWrapper,4,1,4096,32,4,128,0.06349823999999998,528.9469440412838,8.454894371875506
-BatchDecodeWithPagedKVCacheWrapper,4,1,4096,32,4,256,0.1303347200000001,515.3991200502825,8.238340666247638
-BatchDecodeWithPagedKVCacheWrapper,4,1,8192,32,8,64,0.16568319999999992,405.1421508034613,6.480692212608161
-BatchDecodeWithPagedKVCacheWrapper,4,1,8192,32,4,128,0.10290176,652.4828341128471,10.43463031147378
-BatchDecodeWithPagedKVCacheWrapper,4,1,8192,32,4,256,0.22947840000000008,585.1673360107093,9.358107987505575
-BatchDecodeWithPagedKVCacheWrapper,4,1,16384,32,8,64,0.30601215999999987,438.65613706331163,7.017641547316293
-BatchDecodeWithPagedKVCacheWrapper,4,1,16384,32,4,128,0.18384895999999992,730.2216776205863,11.680695109724857
-BatchDecodeWithPagedKVCacheWrapper,4,1,16384,32,4,256,0.4362026666666668,615.5418398787107,9.846265564630505
-BatchDecodeWithPagedKVCacheWrapper,8,1,512,32,8,64,0.038655999999999975,217.85430463576174,3.472105960264903
-BatchDecodeWithPagedKVCacheWrapper,8,1,512,32,4,128,0.03645951999999999,231.87754528858312,3.6812807190001418
-BatchDecodeWithPagedKVCacheWrapper,8,1,512,32,4,256,0.06676480000000001,253.25153374233125,4.020613496932515
-BatchDecodeWithPagedKVCacheWrapper,8,1,1024,32,8,64,0.05858303999999996,286.9428421604617,4.582135990211505
-BatchDecodeWithPagedKVCacheWrapper,8,1,1024,32,4,128,0.04676608000000001,360.14889424129615,5.73996058681848
-BatchDecodeWithPagedKVCacheWrapper,8,1,1024,32,4,256,0.08992768000000002,374.58437713504884,5.970029606012297
-BatchDecodeWithPagedKVCacheWrapper,8,1,2048,32,8,64,0.092416,363.43490304709144,5.8092853185595565
-BatchDecodeWithPagedKVCacheWrapper,8,1,2048,32,4,128,0.07130112000000002,471.5208961654457,7.5296280338934345
-BatchDecodeWithPagedKVCacheWrapper,8,1,2048,32,4,256,0.14862335999999993,452.41835469202175,7.224583161085851
-BatchDecodeWithPagedKVCacheWrapper,8,1,4096,32,8,64,0.16396288000000003,409.4928803397451,6.548688483637271
-BatchDecodeWithPagedKVCacheWrapper,8,1,4096,32,4,128,0.11201536000000002,599.6891854831337,9.585665965810401
-BatchDecodeWithPagedKVCacheWrapper,8,1,4096,32,4,256,0.24935424000000006,538.7869081351894,8.61218019793848
-BatchDecodeWithPagedKVCacheWrapper,8,1,8192,32,8,64,0.3056947200000001,439.16524302415155,7.024928817874248
-BatchDecodeWithPagedKVCacheWrapper,8,1,8192,32,4,128,0.20128768000000002,667.1211273337741,10.668728697156228
-BatchDecodeWithPagedKVCacheWrapper,8,1,8192,32,4,256,0.46690133333333317,575.2104541716168,9.198875628255509
-BatchDecodeWithPagedKVCacheWrapper,8,1,16384,32,8,64,0.5866495999999998,457.6296037702917,7.321179961598886
-BatchDecodeWithPagedKVCacheWrapper,8,1,16384,32,4,128,0.37337600000000015,719.1169009256082,11.503062050051419
-BatchDecodeWithPagedKVCacheWrapper,8,1,16384,32,4,256,0.8934826666666666,601.0211546726517,9.613991308915525
-BatchDecodeWithPagedKVCacheWrapper,16,1,512,32,8,64,0.0698112,241.26145947928126,3.845163182984965
-BatchDecodeWithPagedKVCacheWrapper,16,1,512,32,4,128,0.04724735999999999,357.8673602080625,5.681491114000869
-BatchDecodeWithPagedKVCacheWrapper,16,1,512,32,4,256,0.08954879999999998,377.6329331046313,5.995288736420813
-BatchDecodeWithPagedKVCacheWrapper,16,1,1024,32,8,64,0.12070911999999998,278.52052935188334,4.447641669494402
-BatchDecodeWithPagedKVCacheWrapper,16,1,1024,32,4,128,0.07076864000000004,475.9947909130369,7.586282737664589
-BatchDecodeWithPagedKVCacheWrapper,16,1,1024,32,4,256,0.14710784000000002,457.9702074342197,7.2990115550605585
-BatchDecodeWithPagedKVCacheWrapper,16,1,2048,32,8,64,0.22239232000000014,302.05359609540454,4.8281425545630325
-BatchDecodeWithPagedKVCacheWrapper,16,1,2048,32,4,128,0.11209728000000002,599.8355713894217,9.578660820316067
-BatchDecodeWithPagedKVCacheWrapper,16,1,2048,32,4,256,0.2504192,537.0190145164587,8.575555101206296
-BatchDecodeWithPagedKVCacheWrapper,16,1,4096,32,8,64,0.42098688000000006,318.97256275539985,5.101070247129791
-BatchDecodeWithPagedKVCacheWrapper,16,1,4096,32,4,128,0.2027008,662.7936347562515,10.594352109118466
-BatchDecodeWithPagedKVCacheWrapper,16,1,4096,32,4,256,0.46432,578.6905582356995,9.250015713301172
-BatchDecodeWithPagedKVCacheWrapper,16,1,8192,32,8,64,0.8234496000000004,326.06851955480926,5.215822918609709
-BatchDecodeWithPagedKVCacheWrapper,16,1,8192,32,4,128,0.3726506666666667,720.6924662239522,11.525451797572705
-BatchDecodeWithPagedKVCacheWrapper,16,1,8192,32,4,256,0.8939733333333334,600.8378952392316,9.608714568667223
-BatchDecodeWithPagedKVCacheWrapper,16,1,16384,32,8,64,1.6324906666666663,328.9062896122735,5.261858317107276
-BatchDecodeWithPagedKVCacheWrapper,16,1,16384,32,4,128,0.7114879999999999,754.7590177206082,12.073196725735361
-BatchDecodeWithPagedKVCacheWrapper,16,1,16384,32,4,256,1.742272,616.4387466480549,9.860612570253094
-BatchDecodeWithPagedKVCacheWrapper,32,1,512,32,8,64,0.08406016,400.730905104154,6.386746254111341
-BatchDecodeWithPagedKVCacheWrapper,32,1,512,32,4,128,0.08498175999999999,397.92746113989637,6.317484034220991
-BatchDecodeWithPagedKVCacheWrapper,32,1,512,32,4,256,0.1808896,373.8918765921313,5.935895839230116
-BatchDecodeWithPagedKVCacheWrapper,32,1,1024,32,8,64,0.14712832,457.0155902004454,7.297995545657015
-BatchDecodeWithPagedKVCacheWrapper,32,1,1024,32,4,128,0.14887935999999996,452.5208061077104,7.212160396175805
-BatchDecodeWithPagedKVCacheWrapper,32,1,1024,32,4,256,0.3279462400000001,410.8661712358707,6.548279522887651
-BatchDecodeWithPagedKVCacheWrapper,32,1,2048,32,8,64,0.27223039999999993,493.51137859695325,7.888478465299983
-BatchDecodeWithPagedKVCacheWrapper,32,1,2048,32,4,128,0.27833343999999993,483.1610316029581,7.715507155733786
-BatchDecodeWithPagedKVCacheWrapper,32,1,2048,32,4,256,0.6300373333333331,426.89493109403054,6.817004435715981
-BatchDecodeWithPagedKVCacheWrapper,32,1,4096,32,8,64,0.52494336,511.6104868913858,8.181772784019977
-BatchDecodeWithPagedKVCacheWrapper,32,1,4096,32,4,128,0.5080533333333332,528.8767583455806,8.453772496325847
-BatchDecodeWithPagedKVCacheWrapper,32,1,4096,32,4,256,1.2449493333333332,431.66029782202656,6.899826653186422
-BatchDecodeWithPagedKVCacheWrapper,32,1,8192,32,8,64,1.0273706666666667,522.6954607749491,8.361086091615102
-BatchDecodeWithPagedKVCacheWrapper,32,1,8192,32,4,128,1.0078719999999999,532.9377698755399,8.522842773685548
-BatchDecodeWithPagedKVCacheWrapper,32,1,8192,32,4,256,2.446784,439.05228741073995,7.021408176610604
-BatchDecodeWithPagedKVCacheWrapper,32,1,16384,32,8,64,2.0322986666666663,528.4030903594223,8.453417534430637
-BatchDecodeWithPagedKVCacheWrapper,32,1,16384,32,4,128,2.018026666666667,532.2050425498176,8.513202262276018
-BatchDecodeWithPagedKVCacheWrapper,32,1,16384,32,4,256,4.847957333333333,443.0748433429558,7.087467154826447
-BatchDecodeWithPagedKVCacheWrapper,64,1,512,32,8,64,0.13077504,515.1671756322919,8.210602145485865
-BatchDecodeWithPagedKVCacheWrapper,64,1,512,32,4,128,0.14377984000000002,470.39384659212305,7.4679581226408365
-BatchDecodeWithPagedKVCacheWrapper,64,1,512,32,4,256,0.27039743999999993,500.2499431947286,7.9419525865333656
-BatchDecodeWithPagedKVCacheWrapper,64,1,1024,32,8,64,0.231424,581.0973451327434,9.279433628318584
-BatchDecodeWithPagedKVCacheWrapper,64,1,1024,32,4,128,0.25729023999999995,523.6965692907746,8.34654143118682
-BatchDecodeWithPagedKVCacheWrapper,64,1,1024,32,4,256,0.502016,536.8036715961244,8.555439061703213
-BatchDecodeWithPagedKVCacheWrapper,64,1,2048,32,8,64,0.4335923199999999,619.7010131544766,9.905542828802874
-BatchDecodeWithPagedKVCacheWrapper,64,1,2048,32,4,128,0.47517866666666664,566.018137739068,9.038636616683128
-BatchDecodeWithPagedKVCacheWrapper,64,1,2048,32,4,256,0.9693866666666666,554.9070422535212,8.861205633802816
-BatchDecodeWithPagedKVCacheWrapper,64,1,4096,32,8,64,0.8388479999999999,640.3222705424583,10.240156252384224
-BatchDecodeWithPagedKVCacheWrapper,64,1,4096,32,4,128,0.9261013333333336,580.2768883462716,9.275372232844207
-BatchDecodeWithPagedKVCacheWrapper,64,1,4096,32,4,256,1.906474666666667,563.7580287805205,9.011328335161021
-BatchDecodeWithPagedKVCacheWrapper,64,1,8192,32,8,64,1.6543999999999999,649.1803481624759,10.384350328820116
-BatchDecodeWithPagedKVCacheWrapper,64,1,8192,32,4,128,1.8147413333333327,591.9665201137942,9.466841840453299
-BatchDecodeWithPagedKVCacheWrapper,64,1,8192,32,4,256,3.774634666666667,569.2026947596871,9.10279839037844
-BatchDecodeWithPagedKVCacheWrapper,64,1,16384,32,8,64,3.2680746666666667,657.1899393567508,10.513755612274872
-BatchDecodeWithPagedKVCacheWrapper,64,1,16384,32,4,128,3.5912106666666666,598.129192458031,9.567731207451676
-BatchDecodeWithPagedKVCacheWrapper,64,1,16384,32,4,256,7.526272,570.802632697835,9.130612969608327
-BatchDecodeWithPagedKVCacheWrapper,128,1,512,32,8,64,0.2176000000000001,619.2188235294115,9.86895058823529
-BatchDecodeWithPagedKVCacheWrapper,128,1,512,32,4,128,0.21536768,628.0715100798782,9.971243818942565
-BatchDecodeWithPagedKVCacheWrapper,128,1,512,32,4,256,0.45757866666666663,591.2264441232692,9.386292694298103
-BatchDecodeWithPagedKVCacheWrapper,128,1,1024,32,8,64,0.39856127999999985,674.826576229382,10.776177997019683
-BatchDecodeWithPagedKVCacheWrapper,128,1,1024,32,4,128,0.39381333333333335,684.2938244853738,10.906099241603465
-BatchDecodeWithPagedKVCacheWrapper,128,1,1024,32,4,256,0.8577493333333336,628.3514810853829,10.014504539010616
-BatchDecodeWithPagedKVCacheWrapper,128,1,2048,32,8,64,0.7606186666666664,706.5238121949853,11.293352330734283
-BatchDecodeWithPagedKVCacheWrapper,128,1,2048,32,4,128,0.7354026666666665,731.4625203063357,11.680586679043863
-BatchDecodeWithPagedKVCacheWrapper,128,1,2048,32,4,256,1.6673066666666665,645.2556074467406,10.30396478792144
-BatchDecodeWithPagedKVCacheWrapper,128,1,4096,32,8,64,1.4816639999999999,725.0403006349618,11.594983197270098
-BatchDecodeWithPagedKVCacheWrapper,128,1,4096,32,4,128,1.4354773333333333,748.7338009749138,11.968053263583403
-BatchDecodeWithPagedKVCacheWrapper,128,1,4096,32,4,256,3.2697173333333325,657.4209880731792,10.508473627893627
-BatchDecodeWithPagedKVCacheWrapper,128,1,8192,32,8,64,2.9226666666666676,734.9479708029195,11.75629734306569
-BatchDecodeWithPagedKVCacheWrapper,128,1,8192,32,4,128,2.825301333333333,760.4612643087983,12.161442024827087
-BatchDecodeWithPagedKVCacheWrapper,128,1,8192,32,4,256,6.484309333333334,662.6865294520187,10.5978097594357
-BatchDecodeWithPagedKVCacheWrapper,128,1,16384,32,8,64,5.794901333333332,741.2536188134122,11.858610316747416
-BatchDecodeWithPagedKVCacheWrapper,128,1,16384,32,4,128,5.61536,765.0472760428539,12.237768680191476
-BatchDecodeWithPagedKVCacheWrapper,128,1,16384,32,4,256,12.908458666666668,665.6125232199165,10.647200957222468
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashinfer_baseline/BatchMLAPagedAttentionWrapper_20260525_074200.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashinfer_baseline/BatchMLAPagedAttentionWrapper_20260525_074200.csv
@ -1,33 +0,0 @@
-api,batch_size,seq_len,num_heads,head_dim_ckv,head_dim_kpe,time_ms,bandwidth_GB_s,tflops
-BatchMLAPagedAttentionWrapper,1,1024,64,512,64,0.035975679999999996,34.83953604212624,3.963964989681919
-BatchMLAPagedAttentionWrapper,1,4096,64,512,64,0.05349631999999998,89.58223668469162,10.662889409963158
-BatchMLAPagedAttentionWrapper,1,8192,64,512,64,0.06174719999999999,154.0298507462687,18.47615257048093
-BatchMLAPagedAttentionWrapper,1,16384,64,512,64,0.08995584000000004,210.63775292410136,25.3646831156265
-BatchMLAPagedAttentionWrapper,4,1024,64,512,64,0.05086207999999998,98.5705657338434,11.215139923495071
-BatchMLAPagedAttentionWrapper,4,4096,64,512,64,0.08034559999999999,238.58531145451653,28.39858531145452
-BatchMLAPagedAttentionWrapper,4,8192,64,512,64,0.10866687999999997,350.0942329438373,41.99442140972485
-BatchMLAPagedAttentionWrapper,4,16384,64,512,64,0.16821760000000002,450.56155836250184,54.2559488662304
-BatchMLAPagedAttentionWrapper,16,1024,64,512,64,0.06735359999999997,297.7423033067276,33.87645762067656
-BatchMLAPagedAttentionWrapper,16,4096,64,512,64,0.14288383999999996,536.6395528003728,63.87570143691549
-BatchMLAPagedAttentionWrapper,16,8192,64,512,64,0.21618431999999987,703.9113289992544,84.43540682321462
-BatchMLAPagedAttentionWrapper,16,16384,64,512,64,0.39363328000000025,770.1826837405613,92.74424666532254
-BatchMLAPagedAttentionWrapper,64,1024,64,512,64,0.15278592,525.0226198853926,59.73590697362689
-BatchMLAPagedAttentionWrapper,64,4096,64,512,64,0.4850483199999999,632.3256206721838,75.26512413443676
-BatchMLAPagedAttentionWrapper,64,8192,64,512,64,0.9133465600000001,666.4484158127227,79.94166423750474
-BatchMLAPagedAttentionWrapper,64,16384,64,512,64,1.7720038399999998,684.3541287134007,82.40890045926764
-BatchMLAPagedAttentionWrapper,1,1024,128,512,64,0.04499968000000001,29.491409716691315,6.338104448742746
-BatchMLAPagedAttentionWrapper,1,4096,128,512,64,0.05375743999999999,90.51859612362495,21.222191532930147
-BatchMLAPagedAttentionWrapper,1,8192,128,512,64,0.08302080000000002,115.44865864939868,27.48349059512796
-BatchMLAPagedAttentionWrapper,1,16384,128,512,64,0.11321343999999998,168.01736613603475,40.30795947901592
-BatchMLAPagedAttentionWrapper,4,1024,128,512,64,0.05178880000000003,102.50123578843295,22.028907563025196
-BatchMLAPagedAttentionWrapper,4,4096,128,512,64,0.11032576,176.4247261926861,41.36298496380175
-BatchMLAPagedAttentionWrapper,4,8192,128,512,64,0.1688268800000001,227.08800873415404,54.06014435615937
-BatchMLAPagedAttentionWrapper,4,16384,128,512,64,0.30781695999999986,247.18357299091002,59.30021207408457
-BatchMLAPagedAttentionWrapper,16,1024,128,512,64,0.10527487999999995,201.69734698344004,43.34749896651511
-BatchMLAPagedAttentionWrapper,16,4096,128,512,64,0.2629478400000002,296.0920614521874,69.41913273750409
-BatchMLAPagedAttentionWrapper,16,8192,128,512,64,0.3962367999999998,387.02674764181444,92.13485980100793
-BatchMLAPagedAttentionWrapper,16,16384,128,512,64,0.7528985599999998,404.23663979381246,96.97779742333418
-BatchMLAPagedAttentionWrapper,64,1024,128,512,64,0.3242547199999998,261.9380714026308,56.29404872811108
-BatchMLAPagedAttentionWrapper,64,4096,128,512,64,1.1793126399999994,264.07507342582215,61.91271216426548
-BatchMLAPagedAttentionWrapper,64,8192,128,512,64,2.3186406399999986,264.55887532446616,62.98038839860932
-BatchMLAPagedAttentionWrapper,64,16384,128,512,64,4.6020608,264.53295358462015,63.462389746784744
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashinfer_baseline/BatchPrefillWithPagedKVCacheWrapper_20260526_150417.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashinfer_baseline/BatchPrefillWithPagedKVCacheWrapper_20260526_150417.csv
@ -1,33 +0,0 @@
-api,batch_size,seq_len,num_qo_heads,num_kv_heads,head_dim,time_ms,bandwidth_GB_s,tflops
-BatchPrefillWithPagedKVCacheWrapper,1,1024,32,4,128,0.3529011200000001,29.71302556364796,24.34091054174041
-BatchPrefillWithPagedKVCacheWrapper,1,4096,32,4,128,4.62532608,9.068126068205768,29.714435500296666
-BatchPrefillWithPagedKVCacheWrapper,1,8192,32,4,128,18.113853439999996,4.631045529757804,30.350019983820744
-BatchPrefillWithPagedKVCacheWrapper,1,16384,32,4,128,71.05519616000001,2.36115258372119,30.948099145350383
-BatchPrefillWithPagedKVCacheWrapper,4,1024,32,4,128,1.2374374399999997,33.89507917264893,27.766848858234006
-BatchPrefillWithPagedKVCacheWrapper,4,4096,32,4,128,17.896878079999997,9.374381344614939,30.71797279003423
-BatchPrefillWithPagedKVCacheWrapper,4,8192,32,4,128,71.25501952,4.709062214288198,30.861310127559136
-BatchPrefillWithPagedKVCacheWrapper,4,16384,32,4,128,283.27072767999994,2.3690716139159393,31.051895457919
-BatchPrefillWithPagedKVCacheWrapper,16,1024,32,4,128,4.752537600000002,35.301595509733566,28.919067041573737
-BatchPrefillWithPagedKVCacheWrapper,16,4096,32,4,128,70.51405312000001,9.517090711803915,31.185602844439067
-BatchPrefillWithPagedKVCacheWrapper,16,8192,32,4,128,284.16772266666663,4.723186952426669,30.953878011423416
-BatchPrefillWithPagedKVCacheWrapper,16,16384,32,4,128,1129.139136,2.377346134250013,31.160351250841774
-BatchPrefillWithPagedKVCacheWrapper,64,1024,32,4,128,18.757478399999997,35.77712449878125,29.3086203894016
-BatchPrefillWithPagedKVCacheWrapper,64,4096,32,4,128,281.4907093333333,9.536210151864244,31.248253425628754
-BatchPrefillWithPagedKVCacheWrapper,64,8192,32,4,128,1134.7048106666668,4.731370722616177,31.007511167737377
-BatchPrefillWithPagedKVCacheWrapper,64,16384,32,4,128,4514.139178666666,2.378619226173592,31.177037921302507
-BatchPrefillWithPagedKVCacheWrapper,1,1024,32,4,256,0.7928422399999997,26.4510629504301,21.668710768992337
-BatchPrefillWithPagedKVCacheWrapper,1,4096,32,4,256,12.533002240000002,6.69321511267838,21.932327281224513
-BatchPrefillWithPagedKVCacheWrapper,1,8192,32,4,256,49.81321727999999,3.368024977325858,22.072688491402744
-BatchPrefillWithPagedKVCacheWrapper,1,16384,32,4,256,190.01136128,1.765917141688929,23.14622915954513
-BatchPrefillWithPagedKVCacheWrapper,4,1024,32,4,256,3.111116800000001,26.963333552761494,22.088362846422218
-BatchPrefillWithPagedKVCacheWrapper,4,4096,32,4,256,47.738091520000026,7.02885912101079,23.032165567728153
-BatchPrefillWithPagedKVCacheWrapper,4,8192,32,4,256,190.14286336,3.529391680241077,23.130221315627924
-BatchPrefillWithPagedKVCacheWrapper,4,16384,32,4,256,759.6848640000004,1.76675532658763,23.157215416649382
-BatchPrefillWithPagedKVCacheWrapper,16,1024,32,4,256,12.28442624,27.31461066593534,22.376129057534232
-BatchPrefillWithPagedKVCacheWrapper,16,4096,32,4,256,191.34602666666663,7.014398487291994,22.984780963158407
-BatchPrefillWithPagedKVCacheWrapper,16,8192,32,4,256,759.7649706666668,3.5331380935403933,23.15477380982632
-BatchPrefillWithPagedKVCacheWrapper,16,16384,32,4,256,3028.668266666667,1.77263029400997,23.234219789647476
-BatchPrefillWithPagedKVCacheWrapper,64,1024,32,4,256,49.26948266666667,27.241554149868346,22.316281159572153
-BatchPrefillWithPagedKVCacheWrapper,64,4096,32,4,256,763.6229333333335,7.030576067909256,23.037791659325052
-BatchPrefillWithPagedKVCacheWrapper,64,8192,32,4,256,3037.7449386666663,3.534667477616765,23.16479678130923
-BatchPrefillWithPagedKVCacheWrapper,64,16384,32,4,256,12110.653866666667,1.7732185822854112,23.241930601731337
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashinfer_baseline/BatchPrefillWithRaggedKVCacheWrapper_20260526_145454.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashinfer_baseline/BatchPrefillWithRaggedKVCacheWrapper_20260526_145454.csv
@ -1,49 +0,0 @@
-api,batch_size,seq_len,num_qo_heads,num_kv_heads,head_dim_qk,head_dim_vo,time_ms,bandwidth_GB_s,tflops
-BatchPrefillWithRaggedKVCacheWrapper,1,1024,32,4,128,128,0.031580159999999996,66.66666666666667,272.00415045395596
-BatchPrefillWithRaggedKVCacheWrapper,1,4096,32,4,128,128,0.0424448,197.82870928829917,3238.0634016887816
-BatchPrefillWithRaggedKVCacheWrapper,1,8192,32,4,128,128,0.057313279999999994,292.871180989816,9592.119206717885
-BatchPrefillWithRaggedKVCacheWrapper,1,16384,32,4,128,128,0.06972416000000001,481.36290204141574,31538.89922161844
-BatchPrefillWithRaggedKVCacheWrapper,4,1024,32,4,128,128,0.04327423999999998,194.60482725982024,793.9998106956938
-BatchPrefillWithRaggedKVCacheWrapper,4,4096,32,4,128,128,0.06579199999999998,510.5058365758757,8355.967501945528
-BatchPrefillWithRaggedKVCacheWrapper,4,8192,32,4,128,128,0.09618432000000002,698.0517406579366,22862.596060896405
-BatchPrefillWithRaggedKVCacheWrapper,4,16384,32,4,128,128,0.15411199999999997,871.12292358804,57075.97735548174
-BatchPrefillWithRaggedKVCacheWrapper,16,1024,32,4,128,128,0.07452671999999999,451.99230557845567,1844.1567463588901
-BatchPrefillWithRaggedKVCacheWrapper,16,4096,32,4,128,128,0.1668906666666667,805.0108653969065,13176.43041083983
-BatchPrefillWithRaggedKVCacheWrapper,16,8192,32,4,128,128,0.2874026666666667,934.46080760095,30605.46766745843
-BatchPrefillWithRaggedKVCacheWrapper,16,16384,32,4,128,128,0.5342506666666667,1005.1498622369525,65857.42289917343
-BatchPrefillWithRaggedKVCacheWrapper,64,1024,32,4,128,128,0.15733333333333333,856.4111186440679,3494.2106814915255
-BatchPrefillWithRaggedKVCacheWrapper,64,4096,32,4,128,128,0.5614719999999999,957.1184315513509,15666.129428017784
-BatchPrefillWithRaggedKVCacheWrapper,64,8192,32,4,128,128,1.1031466666666667,973.8198414233224,31894.55505055115
-BatchPrefillWithRaggedKVCacheWrapper,64,16384,32,4,128,128,2.1813759999999998,984.7032038493136,64517.75776176505
-BatchPrefillWithRaggedKVCacheWrapper,1,1024,32,4,192,128,0.03564544000000001,73.88681413386956,301.22838264866414
-BatchPrefillWithRaggedKVCacheWrapper,1,4096,32,4,192,128,0.04922368,213.27231121281466,3490.1635115456625
-BatchPrefillWithRaggedKVCacheWrapper,1,8192,32,4,192,128,0.061327359999999984,342.16062781766584,11205.353815328106
-BatchPrefillWithRaggedKVCacheWrapper,1,16384,32,4,192,128,0.08377343999999999,500.8189707859675,32812.059161471705
-BatchPrefillWithRaggedKVCacheWrapper,4,1024,32,4,192,128,0.049623040000000056,212.29880313660726,865.5187783739157
-BatchPrefillWithRaggedKVCacheWrapper,4,4096,32,4,192,128,0.08634367999999998,486.3377609108161,7958.831119544594
-BatchPrefillWithRaggedKVCacheWrapper,4,8192,32,4,192,128,0.13644799999999999,615.1444652908068,20145.249981238278
-BatchPrefillWithRaggedKVCacheWrapper,4,16384,32,4,192,128,0.2321706666666666,722.8359827253516,47357.904577781876
-BatchPrefillWithRaggedKVCacheWrapper,16,1024,32,4,192,128,0.09042944,465.99479107688825,1899.8093081191257
-BatchPrefillWithRaggedKVCacheWrapper,16,4096,32,4,192,128,0.3087573333333334,544.0154770952807,8902.716705589717
-BatchPrefillWithRaggedKVCacheWrapper,16,8192,32,4,192,128,0.5995946666666665,559.9464882943145,18337.58185156195
-BatchPrefillWithRaggedKVCacheWrapper,16,16384,32,4,192,128,1.1809706666666668,568.4182232017052,37240.94624227753
-BatchPrefillWithRaggedKVCacheWrapper,64,1024,32,4,192,128,0.2555306666666667,659.6413424611787,2689.2849156787443
-BatchPrefillWithRaggedKVCacheWrapper,64,4096,32,4,192,128,0.9085866666666667,739.472740079831,12101.340115520075
-BatchPrefillWithRaggedKVCacheWrapper,64,8192,32,4,192,128,1.7810773333333334,754.017631276351,24693.1810808739
-BatchPrefillWithRaggedKVCacheWrapper,64,16384,32,4,192,128,3.5260586666666662,761.5134193267346,49891.92667360423
-BatchPrefillWithRaggedKVCacheWrapper,1,1024,32,4,256,256,0.044037119999999964,95.61678874549479,390.12245087780525
-BatchPrefillWithRaggedKVCacheWrapper,1,4096,32,4,256,256,0.08118271999999997,206.86175580222005,3385.916448032292
-BatchPrefillWithRaggedKVCacheWrapper,1,8192,32,4,256,256,0.11204607999999996,299.6161579235972,9813.030743922503
-BatchPrefillWithRaggedKVCacheWrapper,1,16384,32,4,256,256,0.14619648000000002,459.1440778875113,30083.12177628353
-BatchPrefillWithRaggedKVCacheWrapper,4,1024,32,4,256,256,0.07792639999999999,216.1366622864652,881.8510381077531
-BatchPrefillWithRaggedKVCacheWrapper,4,4096,32,4,256,256,0.13784064000000001,487.3337790654483,7976.686902904687
-BatchPrefillWithRaggedKVCacheWrapper,4,8192,32,4,256,256,0.22408533333333336,599.2505712109672,19626.65938766184
-BatchPrefillWithRaggedKVCacheWrapper,4,16384,32,4,256,256,0.3959893333333334,678.0510720827496,44425.908890852275
-BatchPrefillWithRaggedKVCacheWrapper,16,1024,32,4,256,256,0.15150079999999996,444.6907739101049,1814.366042581954
-BatchPrefillWithRaggedKVCacheWrapper,16,4096,32,4,256,256,0.4274346666666664,628.6284687562392,10289.400589339195
-BatchPrefillWithRaggedKVCacheWrapper,16,8192,32,4,256,256,0.7913173333333334,678.7833823093305,22231.518637802277
-BatchPrefillWithRaggedKVCacheWrapper,16,16384,32,4,256,256,1.5360853333333337,699.1824898616742,45810.43946625186
-BatchPrefillWithRaggedKVCacheWrapper,64,1024,32,4,256,256,0.43906133333333336,613.773091686507,2504.2324256352945
-BatchPrefillWithRaggedKVCacheWrapper,64,4096,32,4,256,256,1.6363946666666664,656.8039006075145,10750.576497692491
-BatchPrefillWithRaggedKVCacheWrapper,64,8192,32,4,256,256,3.234005333333333,664.3564257160574,21759.006842803803
-BatchPrefillWithRaggedKVCacheWrapper,64,16384,32,4,256,256,6.420821333333334,669.0757535484556,43837.84598543405
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashinfer_baseline/__MACOSX/._BatchDecodeWithPagedKVCacheWrapper_20260525_093143.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashinfer_baseline/__MACOSX/._BatchDecodeWithPagedKVCacheWrapper_20260525_093143.csv
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashinfer_baseline/__MACOSX/._BatchMLAPagedAttentionWrapper_20260525_074200.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashinfer_baseline/__MACOSX/._BatchMLAPagedAttentionWrapper_20260525_074200.csv
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashinfer_baseline/__MACOSX/._BatchPrefillWithPagedKVCacheWrapper_20260526_150417.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashinfer_baseline/__MACOSX/._BatchPrefillWithPagedKVCacheWrapper_20260526_150417.csv
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashinfer_baseline/__MACOSX/._BatchPrefillWithRaggedKVCacheWrapper_20260526_145454.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashinfer_baseline/__MACOSX/._BatchPrefillWithRaggedKVCacheWrapper_20260526_145454.csv
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashinfer_baseline/__MACOSX/._bench_batch_decode.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashinfer_baseline/__MACOSX/._bench_batch_decode.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashinfer_baseline/__MACOSX/._bench_batch_mla.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashinfer_baseline/__MACOSX/._bench_batch_mla.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashinfer_baseline/__MACOSX/._bench_batch_prefill_paged.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashinfer_baseline/__MACOSX/._bench_batch_prefill_paged.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashinfer_baseline/__MACOSX/._bench_batch_prefill_ragged.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashinfer_baseline/__MACOSX/._bench_batch_prefill_ragged.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashinfer_baseline/__MACOSX/._bench_common.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/flashinfer_baseline/__MACOSX/._bench_common.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/.vscode/settings.json
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/.vscode/settings.json
@ -1,3 +0,0 @@
-{
-    "cmake.sourceDirectory": "/root/Project/fusedmoe_v2/standalone/fused_moe_i8_tn"
-}
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/query/client-vscode/query.json
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/query/client-vscode/query.json
@ -1 +0,0 @@
-{"requests":[{"kind":"cache","version":2},{"kind":"codemodel","version":2},{"kind":"toolchains","version":1},{"kind":"cmakeFiles","version":1}]}
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/cache-v2-ea2ef11d05674d96d761.json
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/cache-v2-ea2ef11d05674d96d761.json
@ -1,471 +0,0 @@
-{
-	"entries" : 
-	[
-		{
-			"name" : "CMAKE_BUILD_TYPE",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "No help, variable specified on the command line."
-				}
-			],
-			"type" : "STRING",
-			"value" : "Debug"
-		},
-		{
-			"name" : "CMAKE_CACHEFILE_DIR",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "This is the directory where this CMakeCache.txt was created"
-				}
-			],
-			"type" : "INTERNAL",
-			"value" : "/root/Project/fusedmoe/build"
-		},
-		{
-			"name" : "CMAKE_CACHE_MAJOR_VERSION",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "Major version of cmake used to create the current loaded cache"
-				}
-			],
-			"type" : "INTERNAL",
-			"value" : "3"
-		},
-		{
-			"name" : "CMAKE_CACHE_MINOR_VERSION",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "Minor version of cmake used to create the current loaded cache"
-				}
-			],
-			"type" : "INTERNAL",
-			"value" : "28"
-		},
-		{
-			"name" : "CMAKE_CACHE_PATCH_VERSION",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "Patch version of cmake used to create the current loaded cache"
-				}
-			],
-			"type" : "INTERNAL",
-			"value" : "3"
-		},
-		{
-			"name" : "CMAKE_COMMAND",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "Path to CMake executable."
-				}
-			],
-			"type" : "INTERNAL",
-			"value" : "/usr/bin/cmake"
-		},
-		{
-			"name" : "CMAKE_CPACK_COMMAND",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "Path to cpack program executable."
-				}
-			],
-			"type" : "INTERNAL",
-			"value" : "/usr/bin/cpack"
-		},
-		{
-			"name" : "CMAKE_CTEST_COMMAND",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "Path to ctest program executable."
-				}
-			],
-			"type" : "INTERNAL",
-			"value" : "/usr/bin/ctest"
-		},
-		{
-			"name" : "CMAKE_CXX_COMPILER",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "No help, variable specified on the command line."
-				}
-			],
-			"type" : "FILEPATH",
-			"value" : "/usr/bin/g++"
-		},
-		{
-			"name" : "CMAKE_C_COMPILER",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "No help, variable specified on the command line."
-				}
-			],
-			"type" : "FILEPATH",
-			"value" : "/usr/bin/gcc"
-		},
-		{
-			"name" : "CMAKE_EXPORT_COMPILE_COMMANDS",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "No help, variable specified on the command line."
-				}
-			],
-			"type" : "BOOL",
-			"value" : "TRUE"
-		},
-		{
-			"name" : "CMAKE_EXTRA_GENERATOR",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "Name of external makefile project generator."
-				}
-			],
-			"type" : "INTERNAL",
-			"value" : ""
-		},
-		{
-			"name" : "CMAKE_FIND_PACKAGE_REDIRECTS_DIR",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "Value Computed by CMake."
-				}
-			],
-			"type" : "STATIC",
-			"value" : "/root/Project/fusedmoe/build/CMakeFiles/pkgRedirects"
-		},
-		{
-			"name" : "CMAKE_GENERATOR",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "Name of generator."
-				}
-			],
-			"type" : "INTERNAL",
-			"value" : "Ninja"
-		},
-		{
-			"name" : "CMAKE_GENERATOR_INSTANCE",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "Generator instance identifier."
-				}
-			],
-			"type" : "INTERNAL",
-			"value" : ""
-		},
-		{
-			"name" : "CMAKE_GENERATOR_PLATFORM",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "Name of generator platform."
-				}
-			],
-			"type" : "INTERNAL",
-			"value" : ""
-		},
-		{
-			"name" : "CMAKE_GENERATOR_TOOLSET",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "Name of generator toolset."
-				}
-			],
-			"type" : "INTERNAL",
-			"value" : ""
-		},
-		{
-			"name" : "CMAKE_HOME_DIRECTORY",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "Source directory with the top level CMakeLists.txt file for this project"
-				}
-			],
-			"type" : "INTERNAL",
-			"value" : "/root/Project/fusedmoe/standalone/fused_moe_i8_tn"
-		},
-		{
-			"name" : "CMAKE_INSTALL_PREFIX",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "Install path prefix, prepended onto install directories."
-				}
-			],
-			"type" : "PATH",
-			"value" : "/usr/local"
-		},
-		{
-			"name" : "CMAKE_INSTALL_SO_NO_EXE",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "Install .so files without execute permission."
-				}
-			],
-			"type" : "INTERNAL",
-			"value" : "1"
-		},
-		{
-			"name" : "CMAKE_MAKE_PROGRAM",
-			"properties" : 
-			[
-				{
-					"name" : "ADVANCED",
-					"value" : "1"
-				},
-				{
-					"name" : "HELPSTRING",
-					"value" : "Program used to build from build.ninja files."
-				}
-			],
-			"type" : "FILEPATH",
-			"value" : "/opt/conda/bin/ninja"
-		},
-		{
-			"name" : "CMAKE_NUMBER_OF_MAKEFILES",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "number of local generators"
-				}
-			],
-			"type" : "INTERNAL",
-			"value" : "1"
-		},
-		{
-			"name" : "CMAKE_PLATFORM_INFO_INITIALIZED",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "Platform information initialized"
-				}
-			],
-			"type" : "INTERNAL",
-			"value" : "1"
-		},
-		{
-			"name" : "CMAKE_PROJECT_DESCRIPTION",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "Value Computed by CMake"
-				}
-			],
-			"type" : "STATIC",
-			"value" : ""
-		},
-		{
-			"name" : "CMAKE_PROJECT_HOMEPAGE_URL",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "Value Computed by CMake"
-				}
-			],
-			"type" : "STATIC",
-			"value" : ""
-		},
-		{
-			"name" : "CMAKE_PROJECT_NAME",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "Value Computed by CMake"
-				}
-			],
-			"type" : "STATIC",
-			"value" : "fused_moe_i8_tn"
-		},
-		{
-			"name" : "CMAKE_ROOT",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "Path to CMake installation."
-				}
-			],
-			"type" : "INTERNAL",
-			"value" : "/usr/share/cmake-3.28"
-		},
-		{
-			"name" : "CMAKE_SKIP_INSTALL_RPATH",
-			"properties" : 
-			[
-				{
-					"name" : "ADVANCED",
-					"value" : "1"
-				},
-				{
-					"name" : "HELPSTRING",
-					"value" : "If set, runtime paths are not added when installing shared libraries, but are added when building."
-				}
-			],
-			"type" : "BOOL",
-			"value" : "NO"
-		},
-		{
-			"name" : "CMAKE_SKIP_RPATH",
-			"properties" : 
-			[
-				{
-					"name" : "ADVANCED",
-					"value" : "1"
-				},
-				{
-					"name" : "HELPSTRING",
-					"value" : "If set, runtime paths are not added when using shared libraries."
-				}
-			],
-			"type" : "BOOL",
-			"value" : "NO"
-		},
-		{
-			"name" : "CMAKE_UNAME",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "uname command"
-				}
-			],
-			"type" : "INTERNAL",
-			"value" : "/usr/bin/uname"
-		},
-		{
-			"name" : "CMAKE_VERBOSE_MAKEFILE",
-			"properties" : 
-			[
-				{
-					"name" : "ADVANCED",
-					"value" : "1"
-				},
-				{
-					"name" : "HELPSTRING",
-					"value" : "If this value is on, makefiles will be generated without the .SILENT directive, and all commands will be echoed to the console during the make.  This is useful for debugging only. With Visual Studio IDE projects all commands are done without /nologo."
-				}
-			],
-			"type" : "BOOL",
-			"value" : "FALSE"
-		},
-		{
-			"name" : "MACA_PATH",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "Path to MACA SDK"
-				}
-			],
-			"type" : "PATH",
-			"value" : "/opt/maca"
-		},
-		{
-			"name" : "MXCC",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "Path to a program."
-				}
-			],
-			"type" : "FILEPATH",
-			"value" : "/opt/maca/mxgpu_llvm/bin/mxcc"
-		},
-		{
-			"name" : "_CMAKE_LINKER_PUSHPOP_STATE_SUPPORTED",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "linker supports push/pop state"
-				}
-			],
-			"type" : "INTERNAL",
-			"value" : "FALSE"
-		},
-		{
-			"name" : "fused_moe_i8_tn_BINARY_DIR",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "Value Computed by CMake"
-				}
-			],
-			"type" : "STATIC",
-			"value" : "/root/Project/fusedmoe/build"
-		},
-		{
-			"name" : "fused_moe_i8_tn_IS_TOP_LEVEL",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "Value Computed by CMake"
-				}
-			],
-			"type" : "STATIC",
-			"value" : "ON"
-		},
-		{
-			"name" : "fused_moe_i8_tn_SOURCE_DIR",
-			"properties" : 
-			[
-				{
-					"name" : "HELPSTRING",
-					"value" : "Value Computed by CMake"
-				}
-			],
-			"type" : "STATIC",
-			"value" : "/root/Project/fusedmoe/standalone/fused_moe_i8_tn"
-		}
-	],
-	"kind" : "cache",
-	"version" : 
-	{
-		"major" : 2,
-		"minor" : 0
-	}
-}
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/cmakeFiles-v1-7899829d23c1c1ae3e98.json
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/cmakeFiles-v1-7899829d23c1c1ae3e98.json
@ -1,73 +0,0 @@
-{
-	"inputs" : 
-	[
-		{
-			"path" : "CMakeLists.txt"
-		},
-		{
-			"isCMake" : true,
-			"isExternal" : true,
-			"path" : "/usr/share/cmake-3.28/Modules/CMakeDetermineSystem.cmake"
-		},
-		{
-			"isCMake" : true,
-			"isExternal" : true,
-			"path" : "/usr/share/cmake-3.28/Modules/CMakeSystem.cmake.in"
-		},
-		{
-			"isGenerated" : true,
-			"path" : "/root/Project/fusedmoe/build/CMakeFiles/3.28.3/CMakeSystem.cmake"
-		},
-		{
-			"isCMake" : true,
-			"isExternal" : true,
-			"path" : "/usr/share/cmake-3.28/Modules/CMakeNinjaFindMake.cmake"
-		},
-		{
-			"isCMake" : true,
-			"isExternal" : true,
-			"path" : "/usr/share/cmake-3.28/Modules/CMakeSystemSpecificInitialize.cmake"
-		},
-		{
-			"isCMake" : true,
-			"isExternal" : true,
-			"path" : "/usr/share/cmake-3.28/Modules/Platform/Linux-Initialize.cmake"
-		},
-		{
-			"isCMake" : true,
-			"isExternal" : true,
-			"path" : "/usr/share/cmake-3.28/Modules/CMakeSystemSpecificInformation.cmake"
-		},
-		{
-			"isCMake" : true,
-			"isExternal" : true,
-			"path" : "/usr/share/cmake-3.28/Modules/CMakeGenericSystem.cmake"
-		},
-		{
-			"isCMake" : true,
-			"isExternal" : true,
-			"path" : "/usr/share/cmake-3.28/Modules/CMakeInitializeConfigs.cmake"
-		},
-		{
-			"isCMake" : true,
-			"isExternal" : true,
-			"path" : "/usr/share/cmake-3.28/Modules/Platform/Linux.cmake"
-		},
-		{
-			"isCMake" : true,
-			"isExternal" : true,
-			"path" : "/usr/share/cmake-3.28/Modules/Platform/UnixPaths.cmake"
-		}
-	],
-	"kind" : "cmakeFiles",
-	"paths" : 
-	{
-		"build" : "/root/Project/fusedmoe/build",
-		"source" : "/root/Project/fusedmoe/standalone/fused_moe_i8_tn"
-	},
-	"version" : 
-	{
-		"major" : 1,
-		"minor" : 0
-	}
-}
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/codemodel-v2-7cd0b1b00876e71f364f.json
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/codemodel-v2-7cd0b1b00876e71f364f.json
@ -1,69 +0,0 @@
-{
-	"configurations" : 
-	[
-		{
-			"directories" : 
-			[
-				{
-					"build" : ".",
-					"jsonFile" : "directory-.-Debug-f5ebdc15457944623624.json",
-					"minimumCMakeVersion" : 
-					{
-						"string" : "3.20"
-					},
-					"projectIndex" : 0,
-					"source" : ".",
-					"targetIndexes" : 
-					[
-						0,
-						1
-					]
-				}
-			],
-			"name" : "Debug",
-			"projects" : 
-			[
-				{
-					"directoryIndexes" : 
-					[
-						0
-					],
-					"name" : "fused_moe_i8_tn",
-					"targetIndexes" : 
-					[
-						0,
-						1
-					]
-				}
-			],
-			"targets" : 
-			[
-				{
-					"directoryIndex" : 0,
-					"id" : "build_fused_moe_i8_tn_example::@6890427a1f51a3e7e1df",
-					"jsonFile" : "target-build_fused_moe_i8_tn_example-Debug-a97a299baa6c6c6d83d0.json",
-					"name" : "build_fused_moe_i8_tn_example",
-					"projectIndex" : 0
-				},
-				{
-					"directoryIndex" : 0,
-					"id" : "run::@6890427a1f51a3e7e1df",
-					"jsonFile" : "target-run-Debug-0d66e135afa1376e0f20.json",
-					"name" : "run",
-					"projectIndex" : 0
-				}
-			]
-		}
-	],
-	"kind" : "codemodel",
-	"paths" : 
-	{
-		"build" : "/root/Project/fusedmoe/build",
-		"source" : "/root/Project/fusedmoe/standalone/fused_moe_i8_tn"
-	},
-	"version" : 
-	{
-		"major" : 2,
-		"minor" : 6
-	}
-}
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/directory-.-Debug-f5ebdc15457944623624.json
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/directory-.-Debug-f5ebdc15457944623624.json
@ -1,14 +0,0 @@
-{
-	"backtraceGraph" : 
-	{
-		"commands" : [],
-		"files" : [],
-		"nodes" : []
-	},
-	"installers" : [],
-	"paths" : 
-	{
-		"build" : ".",
-		"source" : "."
-	}
-}
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/index-2026-05-29T05-34-43-0302.json
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/index-2026-05-29T05-34-43-0302.json
@ -1,132 +0,0 @@
-{
-	"cmake" : 
-	{
-		"generator" : 
-		{
-			"multiConfig" : false,
-			"name" : "Ninja"
-		},
-		"paths" : 
-		{
-			"cmake" : "/usr/bin/cmake",
-			"cpack" : "/usr/bin/cpack",
-			"ctest" : "/usr/bin/ctest",
-			"root" : "/usr/share/cmake-3.28"
-		},
-		"version" : 
-		{
-			"isDirty" : false,
-			"major" : 3,
-			"minor" : 28,
-			"patch" : 3,
-			"string" : "3.28.3",
-			"suffix" : ""
-		}
-	},
-	"objects" : 
-	[
-		{
-			"jsonFile" : "codemodel-v2-7cd0b1b00876e71f364f.json",
-			"kind" : "codemodel",
-			"version" : 
-			{
-				"major" : 2,
-				"minor" : 6
-			}
-		},
-		{
-			"jsonFile" : "cache-v2-ea2ef11d05674d96d761.json",
-			"kind" : "cache",
-			"version" : 
-			{
-				"major" : 2,
-				"minor" : 0
-			}
-		},
-		{
-			"jsonFile" : "cmakeFiles-v1-7899829d23c1c1ae3e98.json",
-			"kind" : "cmakeFiles",
-			"version" : 
-			{
-				"major" : 1,
-				"minor" : 0
-			}
-		},
-		{
-			"jsonFile" : "toolchains-v1-8ae3cf416ede58af34e6.json",
-			"kind" : "toolchains",
-			"version" : 
-			{
-				"major" : 1,
-				"minor" : 0
-			}
-		}
-	],
-	"reply" : 
-	{
-		"client-vscode" : 
-		{
-			"query.json" : 
-			{
-				"requests" : 
-				[
-					{
-						"kind" : "cache",
-						"version" : 2
-					},
-					{
-						"kind" : "codemodel",
-						"version" : 2
-					},
-					{
-						"kind" : "toolchains",
-						"version" : 1
-					},
-					{
-						"kind" : "cmakeFiles",
-						"version" : 1
-					}
-				],
-				"responses" : 
-				[
-					{
-						"jsonFile" : "cache-v2-ea2ef11d05674d96d761.json",
-						"kind" : "cache",
-						"version" : 
-						{
-							"major" : 2,
-							"minor" : 0
-						}
-					},
-					{
-						"jsonFile" : "codemodel-v2-7cd0b1b00876e71f364f.json",
-						"kind" : "codemodel",
-						"version" : 
-						{
-							"major" : 2,
-							"minor" : 6
-						}
-					},
-					{
-						"jsonFile" : "toolchains-v1-8ae3cf416ede58af34e6.json",
-						"kind" : "toolchains",
-						"version" : 
-						{
-							"major" : 1,
-							"minor" : 0
-						}
-					},
-					{
-						"jsonFile" : "cmakeFiles-v1-7899829d23c1c1ae3e98.json",
-						"kind" : "cmakeFiles",
-						"version" : 
-						{
-							"major" : 1,
-							"minor" : 0
-						}
-					}
-				]
-			}
-		}
-	}
-}
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/target-build_fused_moe_i8_tn_example-Debug-a97a299baa6c6c6d83d0.json
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/target-build_fused_moe_i8_tn_example-Debug-a97a299baa6c6c6d83d0.json
@ -1,73 +0,0 @@
-{
-	"backtrace" : 1,
-	"backtraceGraph" : 
-	{
-		"commands" : 
-		[
-			"add_custom_target"
-		],
-		"files" : 
-		[
-			"CMakeLists.txt"
-		],
-		"nodes" : 
-		[
-			{
-				"file" : 0
-			},
-			{
-				"command" : 0,
-				"file" : 0,
-				"line" : 38,
-				"parent" : 0
-			}
-		]
-	},
-	"id" : "build_fused_moe_i8_tn_example::@6890427a1f51a3e7e1df",
-	"name" : "build_fused_moe_i8_tn_example",
-	"paths" : 
-	{
-		"build" : ".",
-		"source" : "."
-	},
-	"sourceGroups" : 
-	[
-		{
-			"name" : "",
-			"sourceIndexes" : 
-			[
-				0
-			]
-		},
-		{
-			"name" : "CMake Rules",
-			"sourceIndexes" : 
-			[
-				1,
-				2
-			]
-		}
-	],
-	"sources" : 
-	[
-		{
-			"backtrace" : 1,
-			"isGenerated" : true,
-			"path" : "/root/Project/fusedmoe/build/CMakeFiles/build_fused_moe_i8_tn_example",
-			"sourceGroupIndex" : 0
-		},
-		{
-			"backtrace" : 0,
-			"isGenerated" : true,
-			"path" : "/root/Project/fusedmoe/build/CMakeFiles/build_fused_moe_i8_tn_example.rule",
-			"sourceGroupIndex" : 1
-		},
-		{
-			"backtrace" : 0,
-			"isGenerated" : true,
-			"path" : "/root/Project/fusedmoe/build/fused_moe_i8_tn_example.rule",
-			"sourceGroupIndex" : 1
-		}
-	],
-	"type" : "UTILITY"
-}
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/target-run-Debug-0d66e135afa1376e0f20.json
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/target-run-Debug-0d66e135afa1376e0f20.json
@ -1,73 +0,0 @@
-{
-	"backtrace" : 1,
-	"backtraceGraph" : 
-	{
-		"commands" : 
-		[
-			"add_custom_target"
-		],
-		"files" : 
-		[
-			"CMakeLists.txt"
-		],
-		"nodes" : 
-		[
-			{
-				"file" : 0
-			},
-			{
-				"command" : 0,
-				"file" : 0,
-				"line" : 40,
-				"parent" : 0
-			}
-		]
-	},
-	"id" : "run::@6890427a1f51a3e7e1df",
-	"name" : "run",
-	"paths" : 
-	{
-		"build" : ".",
-		"source" : "."
-	},
-	"sourceGroups" : 
-	[
-		{
-			"name" : "",
-			"sourceIndexes" : 
-			[
-				0
-			]
-		},
-		{
-			"name" : "CMake Rules",
-			"sourceIndexes" : 
-			[
-				1,
-				2
-			]
-		}
-	],
-	"sources" : 
-	[
-		{
-			"backtrace" : 1,
-			"isGenerated" : true,
-			"path" : "/root/Project/fusedmoe/build/CMakeFiles/run",
-			"sourceGroupIndex" : 0
-		},
-		{
-			"backtrace" : 0,
-			"isGenerated" : true,
-			"path" : "/root/Project/fusedmoe/build/CMakeFiles/run.rule",
-			"sourceGroupIndex" : 1
-		},
-		{
-			"backtrace" : 0,
-			"isGenerated" : true,
-			"path" : "/root/Project/fusedmoe/build/fused_moe_i8_tn_example.rule",
-			"sourceGroupIndex" : 1
-		}
-	],
-	"type" : "UTILITY"
-}
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/toolchains-v1-8ae3cf416ede58af34e6.json
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/.cmake/api/v1/reply/toolchains-v1-8ae3cf416ede58af34e6.json
@ -1,18 +0,0 @@
-{
-	"kind" : "toolchains",
-	"toolchains" : 
-	[
-		{
-			"compiler" : 
-			{
-				"implicit" : {}
-			},
-			"language" : "NONE"
-		}
-	],
-	"version" : 
-	{
-		"major" : 1,
-		"minor" : 0
-	}
-}
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeCache.txt
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeCache.txt
@ -1,127 +0,0 @@
-# This is the CMakeCache file.
-# For build in directory: /root/Project/fusedmoe/build
-# It was generated by CMake: /usr/bin/cmake
-# You can edit this file to change values found and used by cmake.
-# If you do not want to change any of the values, simply exit the editor.
-# If you do want to change a value, simply edit, save, and exit the editor.
-# The syntax for the file is as follows:
-# KEY:TYPE=VALUE
-# KEY is the name of a variable in the cache.
-# TYPE is a hint to GUIs for the type of VALUE, DO NOT EDIT TYPE!.
-# VALUE is the current value for the KEY.
-
-########################
-# EXTERNAL cache entries
-########################
-
-//No help, variable specified on the command line.
-CMAKE_BUILD_TYPE:STRING=Debug
-
-//No help, variable specified on the command line.
-CMAKE_CXX_COMPILER:FILEPATH=/usr/bin/g++
-
-//No help, variable specified on the command line.
-CMAKE_C_COMPILER:FILEPATH=/usr/bin/gcc
-
-//No help, variable specified on the command line.
-CMAKE_EXPORT_COMPILE_COMMANDS:BOOL=TRUE
-
-//Value Computed by CMake.
-CMAKE_FIND_PACKAGE_REDIRECTS_DIR:STATIC=/root/Project/fusedmoe/build/CMakeFiles/pkgRedirects
-
-//Install path prefix, prepended onto install directories.
-CMAKE_INSTALL_PREFIX:PATH=/usr/local
-
-//Program used to build from build.ninja files.
-CMAKE_MAKE_PROGRAM:FILEPATH=/opt/conda/bin/ninja
-
-//Value Computed by CMake
-CMAKE_PROJECT_DESCRIPTION:STATIC=
-
-//Value Computed by CMake
-CMAKE_PROJECT_HOMEPAGE_URL:STATIC=
-
-//Value Computed by CMake
-CMAKE_PROJECT_NAME:STATIC=fused_moe_i8_tn
-
-//If set, runtime paths are not added when installing shared libraries,
-// but are added when building.
-CMAKE_SKIP_INSTALL_RPATH:BOOL=NO
-
-//If set, runtime paths are not added when using shared libraries.
-CMAKE_SKIP_RPATH:BOOL=NO
-
-//If this value is on, makefiles will be generated without the
-// .SILENT directive, and all commands will be echoed to the console
-// during the make.  This is useful for debugging only. With Visual
-// Studio IDE projects all commands are done without /nologo.
-CMAKE_VERBOSE_MAKEFILE:BOOL=FALSE
-
-//Path to MACA SDK
-MACA_PATH:PATH=/opt/maca
-
-//Path to a program.
-MXCC:FILEPATH=/opt/maca/mxgpu_llvm/bin/mxcc
-
-//Value Computed by CMake
-fused_moe_i8_tn_BINARY_DIR:STATIC=/root/Project/fusedmoe/build
-
-//Value Computed by CMake
-fused_moe_i8_tn_IS_TOP_LEVEL:STATIC=ON
-
-//Value Computed by CMake
-fused_moe_i8_tn_SOURCE_DIR:STATIC=/root/Project/fusedmoe/standalone/fused_moe_i8_tn
-
-
-########################
-# INTERNAL cache entries
-########################
-
-//This is the directory where this CMakeCache.txt was created
-CMAKE_CACHEFILE_DIR:INTERNAL=/root/Project/fusedmoe/build
-//Major version of cmake used to create the current loaded cache
-CMAKE_CACHE_MAJOR_VERSION:INTERNAL=3
-//Minor version of cmake used to create the current loaded cache
-CMAKE_CACHE_MINOR_VERSION:INTERNAL=28
-//Patch version of cmake used to create the current loaded cache
-CMAKE_CACHE_PATCH_VERSION:INTERNAL=3
-//Path to CMake executable.
-CMAKE_COMMAND:INTERNAL=/usr/bin/cmake
-//Path to cpack program executable.
-CMAKE_CPACK_COMMAND:INTERNAL=/usr/bin/cpack
-//Path to ctest program executable.
-CMAKE_CTEST_COMMAND:INTERNAL=/usr/bin/ctest
-//Name of external makefile project generator.
-CMAKE_EXTRA_GENERATOR:INTERNAL=
-//Name of generator.
-CMAKE_GENERATOR:INTERNAL=Ninja
-//Generator instance identifier.
-CMAKE_GENERATOR_INSTANCE:INTERNAL=
-//Name of generator platform.
-CMAKE_GENERATOR_PLATFORM:INTERNAL=
-//Name of generator toolset.
-CMAKE_GENERATOR_TOOLSET:INTERNAL=
-//Source directory with the top level CMakeLists.txt file for this
-// project
-CMAKE_HOME_DIRECTORY:INTERNAL=/root/Project/fusedmoe/standalone/fused_moe_i8_tn
-//Install .so files without execute permission.
-CMAKE_INSTALL_SO_NO_EXE:INTERNAL=1
-//ADVANCED property for variable: CMAKE_MAKE_PROGRAM
-CMAKE_MAKE_PROGRAM-ADVANCED:INTERNAL=1
-//number of local generators
-CMAKE_NUMBER_OF_MAKEFILES:INTERNAL=1
-//Platform information initialized
-CMAKE_PLATFORM_INFO_INITIALIZED:INTERNAL=1
-//Path to CMake installation.
-CMAKE_ROOT:INTERNAL=/usr/share/cmake-3.28
-//ADVANCED property for variable: CMAKE_SKIP_INSTALL_RPATH
-CMAKE_SKIP_INSTALL_RPATH-ADVANCED:INTERNAL=1
-//ADVANCED property for variable: CMAKE_SKIP_RPATH
-CMAKE_SKIP_RPATH-ADVANCED:INTERNAL=1
-//uname command
-CMAKE_UNAME:INTERNAL=/usr/bin/uname
-//ADVANCED property for variable: CMAKE_VERBOSE_MAKEFILE
-CMAKE_VERBOSE_MAKEFILE-ADVANCED:INTERNAL=1
-//linker supports push/pop state
-_CMAKE_LINKER_PUSHPOP_STATE_SUPPORTED:INTERNAL=FALSE
-
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeFiles/3.28.3/CMakeSystem.cmake
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeFiles/3.28.3/CMakeSystem.cmake
@ -1,15 +0,0 @@
-set(CMAKE_HOST_SYSTEM "Linux-5.15.0-58-generic")
-set(CMAKE_HOST_SYSTEM_NAME "Linux")
-set(CMAKE_HOST_SYSTEM_VERSION "5.15.0-58-generic")
-set(CMAKE_HOST_SYSTEM_PROCESSOR "x86_64")
-
-
-
-set(CMAKE_SYSTEM "Linux-5.15.0-58-generic")
-set(CMAKE_SYSTEM_NAME "Linux")
-set(CMAKE_SYSTEM_VERSION "5.15.0-58-generic")
-set(CMAKE_SYSTEM_PROCESSOR "x86_64")
-
-set(CMAKE_CROSSCOMPILING "FALSE")
-
-set(CMAKE_SYSTEM_LOADED 1)
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeFiles/CMakeConfigureLog.yaml
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeFiles/CMakeConfigureLog.yaml
@ -1,11 +0,0 @@
-
---
-events:
-  -
-    kind: "message-v1"
-    backtrace:
-      - "/usr/share/cmake-3.28/Modules/CMakeDetermineSystem.cmake:233 (message)"
-      - "CMakeLists.txt:3 (project)"
-    message: |
-      The system is: Linux - 5.15.0-58-generic - x86_64
-...
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeFiles/TargetDirectories.txt
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeFiles/TargetDirectories.txt
@ -1,4 +0,0 @@
-/root/Project/fusedmoe/build/CMakeFiles/build_fused_moe_i8_tn_example.dir
-/root/Project/fusedmoe/build/CMakeFiles/run.dir
-/root/Project/fusedmoe/build/CMakeFiles/edit_cache.dir
-/root/Project/fusedmoe/build/CMakeFiles/rebuild_cache.dir
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeFiles/cmake.check_cache
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeFiles/cmake.check_cache
@ -1 +0,0 @@
-# This file is generated by cmake for dependency checking of the CMakeCache.txt file
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeFiles/rules.ninja
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/CMakeFiles/rules.ninja
@ -1,45 +0,0 @@
-# CMAKE generated file: DO NOT EDIT!
-# Generated by "Ninja" Generator, CMake Version 3.28
-
-# This file contains all the rules used to get the outputs files
-# built from the input files.
-# It is included in the main 'build.ninja'.
-
-# =============================================================================
-# Project: fused_moe_i8_tn
-# Configurations: Debug
-# =============================================================================
-# =============================================================================
-
-#############################################
-# Rule for running custom commands.
-
-rule CUSTOM_COMMAND
-  command = $COMMAND
-  description = $DESC
-
-
-#############################################
-# Rule for re-running cmake.
-
-rule RERUN_CMAKE
-  command = /usr/bin/cmake --regenerate-during-build -S/root/Project/fusedmoe/standalone/fused_moe_i8_tn -B/root/Project/fusedmoe/build
-  description = Re-running CMake...
-  generator = 1
-
-
-#############################################
-# Rule for cleaning all built files.
-
-rule CLEAN
-  command = /opt/conda/bin/ninja $FILE_ARG -t clean $TARGETS
-  description = Cleaning all built files...
-
-
-#############################################
-# Rule for printing all primary targets available.
-
-rule HELP
-  command = /opt/conda/bin/ninja -t targets
-  description = All primary targets available:
-
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/build.ninja
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/build.ninja
@ -1,146 +0,0 @@
-# CMAKE generated file: DO NOT EDIT!
-# Generated by "Ninja" Generator, CMake Version 3.28
-
-# This file contains all the build statements describing the
-# compilation DAG.
-
-# =============================================================================
-# Write statements declared in CMakeLists.txt:
-# 
-# Which is the root file.
-# =============================================================================
-
-# =============================================================================
-# Project: fused_moe_i8_tn
-# Configurations: Debug
-# =============================================================================
-
-#############################################
-# Minimal version of Ninja required by this file
-
-ninja_required_version = 1.5
-
-
-#############################################
-# Set configuration variable for custom commands.
-
-CONFIGURATION = Debug
-# =============================================================================
-# Include auxiliary files.
-
-
-#############################################
-# Include rules file.
-
-include CMakeFiles/rules.ninja
-
-# =============================================================================
-
-#############################################
-# Logical path to working directory; prefix for absolute paths.
-
-cmake_ninja_workdir = /root/Project/fusedmoe/build/
-
-#############################################
-# Utility command for build_fused_moe_i8_tn_example
-
-build build_fused_moe_i8_tn_example: phony CMakeFiles/build_fused_moe_i8_tn_example fused_moe_i8_tn_example
-
-
-#############################################
-# Utility command for run
-
-build run: phony CMakeFiles/run fused_moe_i8_tn_example
-
-
-#############################################
-# Utility command for edit_cache
-
-build CMakeFiles/edit_cache.util: CUSTOM_COMMAND
-  COMMAND = cd /root/Project/fusedmoe/build && /usr/bin/cmake -E echo No\ interactive\ CMake\ dialog\ available.
-  DESC = No interactive CMake dialog available...
-  restat = 1
-
-build edit_cache: phony CMakeFiles/edit_cache.util
-
-
-#############################################
-# Utility command for rebuild_cache
-
-build CMakeFiles/rebuild_cache.util: CUSTOM_COMMAND
-  COMMAND = cd /root/Project/fusedmoe/build && /usr/bin/cmake --regenerate-during-build -S/root/Project/fusedmoe/standalone/fused_moe_i8_tn -B/root/Project/fusedmoe/build
-  DESC = Running CMake to regenerate build system...
-  pool = console
-  restat = 1
-
-build rebuild_cache: phony CMakeFiles/rebuild_cache.util
-
-
-#############################################
-# Phony custom command for CMakeFiles/build_fused_moe_i8_tn_example
-
-build CMakeFiles/build_fused_moe_i8_tn_example | ${cmake_ninja_workdir}CMakeFiles/build_fused_moe_i8_tn_example: phony fused_moe_i8_tn_example
-
-
-#############################################
-# Custom command for fused_moe_i8_tn_example
-
-build fused_moe_i8_tn_example | ${cmake_ninja_workdir}fused_moe_i8_tn_example: CUSTOM_COMMAND /root/Project/fusedmoe/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_example.cpp
-  COMMAND = cd /root/Project/fusedmoe/standalone/fused_moe_i8_tn && /opt/maca/mxgpu_llvm/bin/mxcc -std=c++17 -xmaca -I\"/root/Project/fusedmoe/standalone/fused_moe_i8_tn/src\" -I\"/opt/maca/include\" /root/Project/fusedmoe/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_example.cpp -L\"/opt/maca/lib\" -lmcruntime -o /root/Project/fusedmoe/build/fused_moe_i8_tn_example
-  DESC = Generating fused_moe_i8_tn_example
-  restat = 1
-
-
-#############################################
-# Custom command for CMakeFiles/run
-
-build CMakeFiles/run | ${cmake_ninja_workdir}CMakeFiles/run: CUSTOM_COMMAND fused_moe_i8_tn_example
-  COMMAND = cd /root/Project/fusedmoe/build && /root/Project/fusedmoe/build/fused_moe_i8_tn_example
-  pool = console
-
-# =============================================================================
-# Target aliases.
-
-# =============================================================================
-# Folder targets.
-
-# =============================================================================
-
-#############################################
-# Folder: /root/Project/fusedmoe/build
-
-build all: phony build_fused_moe_i8_tn_example
-
-# =============================================================================
-# Built-in targets
-
-
-#############################################
-# Re-run CMake if any of its inputs changed.
-
-build build.ninja: RERUN_CMAKE | /root/Project/fusedmoe/standalone/fused_moe_i8_tn/CMakeLists.txt /usr/share/cmake-3.28/Modules/CMakeDetermineSystem.cmake /usr/share/cmake-3.28/Modules/CMakeGenericSystem.cmake /usr/share/cmake-3.28/Modules/CMakeInitializeConfigs.cmake /usr/share/cmake-3.28/Modules/CMakeNinjaFindMake.cmake /usr/share/cmake-3.28/Modules/CMakeSystem.cmake.in /usr/share/cmake-3.28/Modules/CMakeSystemSpecificInformation.cmake /usr/share/cmake-3.28/Modules/CMakeSystemSpecificInitialize.cmake /usr/share/cmake-3.28/Modules/Platform/Linux-Initialize.cmake /usr/share/cmake-3.28/Modules/Platform/Linux.cmake /usr/share/cmake-3.28/Modules/Platform/UnixPaths.cmake CMakeCache.txt CMakeFiles/3.28.3/CMakeSystem.cmake
-  pool = console
-
-
-#############################################
-# A missing CMake input file is not an error.
-
-build /root/Project/fusedmoe/standalone/fused_moe_i8_tn/CMakeLists.txt /usr/share/cmake-3.28/Modules/CMakeDetermineSystem.cmake /usr/share/cmake-3.28/Modules/CMakeGenericSystem.cmake /usr/share/cmake-3.28/Modules/CMakeInitializeConfigs.cmake /usr/share/cmake-3.28/Modules/CMakeNinjaFindMake.cmake /usr/share/cmake-3.28/Modules/CMakeSystem.cmake.in /usr/share/cmake-3.28/Modules/CMakeSystemSpecificInformation.cmake /usr/share/cmake-3.28/Modules/CMakeSystemSpecificInitialize.cmake /usr/share/cmake-3.28/Modules/Platform/Linux-Initialize.cmake /usr/share/cmake-3.28/Modules/Platform/Linux.cmake /usr/share/cmake-3.28/Modules/Platform/UnixPaths.cmake CMakeCache.txt CMakeFiles/3.28.3/CMakeSystem.cmake: phony
-
-
-#############################################
-# Clean all the built files.
-
-build clean: CLEAN
-
-
-#############################################
-# Print all primary targets available.
-
-build help: HELP
-
-
-#############################################
-# Make the all target the default.
-
-default all
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/cmake_install.cmake
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/build/cmake_install.cmake
@ -1,49 +0,0 @@
-# Install script for directory: /root/Project/fusedmoe/standalone/fused_moe_i8_tn
-
-# Set the install prefix
-if(NOT DEFINED CMAKE_INSTALL_PREFIX)
-  set(CMAKE_INSTALL_PREFIX "/usr/local")
-endif()
-string(REGEX REPLACE "/$" "" CMAKE_INSTALL_PREFIX "${CMAKE_INSTALL_PREFIX}")
-
-# Set the install configuration name.
-if(NOT DEFINED CMAKE_INSTALL_CONFIG_NAME)
-  if(BUILD_TYPE)
-    string(REGEX REPLACE "^[^A-Za-z0-9_]+" ""
-           CMAKE_INSTALL_CONFIG_NAME "${BUILD_TYPE}")
-  else()
-    set(CMAKE_INSTALL_CONFIG_NAME "Debug")
-  endif()
-  message(STATUS "Install configuration: \"${CMAKE_INSTALL_CONFIG_NAME}\"")
-endif()
-
-# Set the component getting installed.
-if(NOT CMAKE_INSTALL_COMPONENT)
-  if(COMPONENT)
-    message(STATUS "Install component: \"${COMPONENT}\"")
-    set(CMAKE_INSTALL_COMPONENT "${COMPONENT}")
-  else()
-    set(CMAKE_INSTALL_COMPONENT)
-  endif()
-endif()
-
-# Install shared libraries without execute permission?
-if(NOT DEFINED CMAKE_INSTALL_SO_NO_EXE)
-  set(CMAKE_INSTALL_SO_NO_EXE "1")
-endif()
-
-# Is this installation the result of a crosscompile?
-if(NOT DEFINED CMAKE_CROSSCOMPILING)
-  set(CMAKE_CROSSCOMPILING "FALSE")
-endif()
-
-if(CMAKE_INSTALL_COMPONENT)
-  set(CMAKE_INSTALL_MANIFEST "install_manifest_${CMAKE_INSTALL_COMPONENT}.txt")
-else()
-  set(CMAKE_INSTALL_MANIFEST "install_manifest.txt")
-endif()
-
-string(REPLACE ";" "\n" CMAKE_INSTALL_MANIFEST_CONTENT
-       "${CMAKE_INSTALL_MANIFEST_FILES}")
-file(WRITE "/root/Project/fusedmoe/build/${CMAKE_INSTALL_MANIFEST}"
-     "${CMAKE_INSTALL_MANIFEST_CONTENT}")
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_bf16_tn/Makefile
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_bf16_tn/Makefile
@ -1,28 +0,0 @@
-MACA_PATH ?= /opt/maca
-MXCC := $(MACA_PATH)/mxgpu_llvm/bin/mxcc
-ROOT_DIR := $(abspath $(CURDIR)/../..)
-BUILD_DIR := $(CURDIR)/build
-SRC := $(CURDIR)/src/fused_moe_bf16_tn_example.cpp
-BIN := $(BUILD_DIR)/fused_moe_bf16_tn_example
-
-.PHONY: all build run clean
-
-all: build
-
-build: $(BIN)
-
-$(BIN): $(SRC)
-	mkdir -p $(BUILD_DIR)
-	$(MXCC) -std=c++17 -xmaca \
-		-I$(ROOT_DIR)/include \
-		-I$(MACA_PATH)/include \
-		$(SRC) \
-		-L$(MACA_PATH)/lib \
-		-lmcruntime \
-		-o $(BIN)
-
-run: $(BIN)
-	$(BIN)
-
-clean:
-	rm -rf $(BUILD_DIR)
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/Makefile
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/Makefile
@ -1,27 +0,0 @@
-MACA_PATH ?= /opt/maca
-MXCC := $(MACA_PATH)/mxgpu_llvm/bin/mxcc
-BUILD_DIR := $(CURDIR)/build
-SRC := $(CURDIR)/src/fused_moe_i8_tn_example.cpp
-BIN := $(BUILD_DIR)/fused_moe_i8_tn_example
-
-.PHONY: all build run clean
-
-all: build
-
-build: $(BIN)
-
-$(BIN): $(SRC)
-	mkdir -p $(BUILD_DIR)
-	$(MXCC) -std=c++17 -xmaca \
-		-I$(CURDIR)/src \
-		-I$(MACA_PATH)/include \
-		$(SRC) \
-		-L$(MACA_PATH)/lib \
-		-lmcruntime \
-		-o $(BIN)
-
-run: $(BIN)
-	$(BIN)
-
-clean:
-	rm -rf $(BUILD_DIR)
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/build/fused_moe_i8_tn_pybind.o
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/build/fused_moe_i8_tn_pybind.o
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/build/fused_moe_i8_tn_pybind.so
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/build/fused_moe_i8_tn_pybind.so
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/pycache/benchmark_fused_moe_i8_tn.cpython-312.pyc
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/pycache/benchmark_fused_moe_i8_tn.cpython-312.pyc
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/pycache/fused_moe_i8_tn_triton.cpython-310.pyc
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/pycache/fused_moe_i8_tn_triton.cpython-310.pyc
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/pycache/fused_moe_i8_tn_triton.cpython-312.pyc
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/pycache/fused_moe_i8_tn_triton.cpython-312.pyc
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/pycache/test_fused_moe_i8_tn_pybind.cpython-310.pyc
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/pycache/test_fused_moe_i8_tn_pybind.cpython-310.pyc
--- a/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/pycache/test_fused_moe_i8_tn_pybind.cpython-312.pyc
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/baselines/fused_moe/standalone/fused_moe_i8_tn/python/pycache/test_fused_moe_i8_tn_pybind.cpython-312.pyc
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/README.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/README.md
@ -0,0 +1,3 @@
+# README
+
+此目录下将持续更新baselines。
--- 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
@ -0,0 +1,146 @@
+from flash_attn.flash_attn_interface import flash_attn_with_kvcache
+import torch
+import math
+from einops import rearrange
+from datetime import datetime
+import csv
+
+
+def run_with_profiler(fn, warmup=10, reps=100, print_result=False, target_kernels=None):
+    """Run function with torch.profiler and return sum of specific kernel times in ms"""
+    for _ in range(warmup):
+        fn()
+    torch.cuda.synchronize()
+
+    with torch.profiler.profile(
+        activities=[torch.profiler.ProfilerActivity.CUDA],
+        record_shapes=False,
+        profile_memory=False,
+        with_stack=False,
+    ) as prof:
+        for _ in range(reps):
+            fn()
+    torch.cuda.synchronize()
+
+    if print_result:
+        print(prof.key_averages().table(sort_by="device_time", row_limit=20))
+
+    if target_kernels is None:
+        target_kernels = []
+
+    kernel_times_us = 0.0
+    for evt in prof.key_averages():
+        if any(k in evt.key for k in target_kernels):
+            kernel_times_us += evt.device_time
+
+    ms = kernel_times_us / 1e3
+    return ms
+
+
+def calc_bandwidth(batch_size, seqlen_q, seqlen_k, num_heads, num_heads_k, headdim, dtype, ms):
+    """Calculate bandwidth in GB/s"""
+    bytes_per_elem = 2 if dtype == torch.bfloat16 else 4
+    q_bytes = batch_size * seqlen_q * num_heads * headdim * bytes_per_elem
+    kv_bytes = batch_size * seqlen_k * num_heads_k * headdim * bytes_per_elem * 2
+    total_bytes = q_bytes + kv_bytes
+    bw_gb_s = (total_bytes / 1e9) / (ms / 1e3)
+    return bw_gb_s
+
+
+def benchmark_kvcache(batch_size, seqlen_k, seqlen_q, num_heads, num_heads_k, headdim, page_block_size, device, dtype=torch.bfloat16, causal=False):
+    num_blocks = math.ceil(seqlen_k / page_block_size) * batch_size * 3
+    num_blocks = max(1024, num_blocks)
+    paged_kv_block_size = page_block_size
+
+    nheads = num_heads
+    nheads_k = num_heads_k
+    d = headdim
+
+    torch.random.manual_seed(0)
+    window_size = (-1, -1)
+
+    q = torch.randn(batch_size, seqlen_q, nheads, d, device=device, dtype=dtype)
+
+    k_cache_paged = torch.randn(
+        num_blocks, paged_kv_block_size, nheads_k, d, device=device, dtype=dtype
+    )
+    v_cache_paged = torch.randn(
+        num_blocks, paged_kv_block_size, nheads_k, d, device=device, dtype=dtype
+    )
+    block_table = rearrange(
+        torch.randperm(num_blocks, dtype=torch.int32, device=device),
+        "(b nblocks) -> b nblocks",
+        b=batch_size,
+    )
+
+    cache_seqlens = torch.full((batch_size,), seqlen_k, dtype=torch.int32, device=device)
+
+    def run_fn():
+        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=causal,
+            window_size=window_size,
+            rotary_interleaved=False,
+            alibi_slopes=None,
+            num_splits=1,
+        )
+
+    return run_fn
+
+
+def main():
+    headdims = [256]
+    page_block_size = 16
+    batch_sizes = [1, 2, 4, 8, 16, 32, 64, 128]
+    seq_lens_kv = [512, 1024, 2048, 4096, 8192, 16384]
+
+    device = "cuda" if torch.cuda.is_available() else "cpu"
+    dtype = torch.bfloat16
+    causal = False
+    warmup = 10
+    repeat = 100
+
+    num_heads = 8
+    num_heads_k = 8
+    seqlen_q = 1
+
+    timestamp = datetime.now().strftime("%Y%m%d_%H%M%S")
+    csv_path = f"benchmark_kvcache_{timestamp}.csv"
+
+    with open(csv_path, "w", newline="") as f:
+        writer = csv.writer(f)
+        writer.writerow(["batch_size", "seq_len_kv", "heads", "headdim", "time_ms", "bandwidth_GB_s"])
+        print(f"{'batch_size':>10} {'seq_len_kv':>12} {'heads':>6} {'headdim':>8} {'time_ms':>10} {'bandwidth_GB_s':>15}")
+        print("-" * 75)
+        for headdim in headdims:
+            for seqlen_k in seq_lens_kv:
+                for batch_size in batch_sizes:
+                    try:
+                        run_fn = benchmark_kvcache(
+                            batch_size=batch_size,
+                            seqlen_k=seqlen_k,
+                            seqlen_q=seqlen_q,
+                            num_heads=num_heads,
+                            num_heads_k=num_heads_k,
+                            headdim=headdim,
+                            page_block_size=page_block_size,
+                            device=device,
+                            dtype=dtype,
+                            causal=causal,
+                        )
+                        ms = run_with_profiler(run_fn, warmup=warmup, reps=repeat, target_kernels=["flash"])
+                        bw = calc_bandwidth(batch_size, seqlen_q, seqlen_k, num_heads, num_heads_k, headdim, dtype, ms)
+                        writer.writerow([batch_size, seqlen_k, num_heads, headdim, f"{ms:.4f}", f"{bw:.2f}"])
+                        print(f"{batch_size:>10} {seqlen_k:>12} {num_heads:>6} {headdim:>8} {ms:>10.4f} {bw:>15.2f}")
+                    except Exception as e:
+                        writer.writerow([batch_size, seqlen_k, num_heads, headdim, "OOM", "OOM"])
+                        print(f"{batch_size:>10} {seqlen_k:>12} {num_heads:>6} {headdim:>8} {'OOM':>10} {'OOM':>15}  # {e}")
+
+    print(f"\nResults saved to {csv_path}")
+
+
+if __name__ == "__main__":
+    main()
--- 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/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/problem_20001/Agent
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/problem_20001/Agent
@ -0,0 +1,360 @@
+# Agent 推理算子库优化 - FlashInfer Ragged Prefill
+
+当前题目说明来源为 [*XPU-OJ 20001*](https://xpuoj.com/contest/2/problem/1)，内容以 XPU-OJ 线上为准。
+
+## 1. 题目描述
+你需要实现 FlashInfer ragged KV cache prefill 的CUDA C++前向算子。
+
+本题输入采用 FlashInfer BatchPrefillWithRaggedKVCacheWrapper 的 ragged NHD 布局。每个 batch 段的 query/KV 长度由 qo_indptr 和 kv_indptr 给出；seq_len 只是所有段长度的上界，真实总长度分别是 qo_indptr[batch_size] 和 kv_indptr[batch_size]。
+
+其中 query heads 采用 GQA 布局：num_qo_heads 个 query/output heads 共享 num_kv_heads 个 KV heads，G = num_qo_heads / num_kv_heads。
+
+评测程序会调用你提交代码中的 run_kernel 函数。你需要根据 qo_indptr 和 kv_indptr 读取 ragged Q/K/V，并将结果写入 output。
+
+baseline 使用 FlashInfer ragged prefill 的 Python API：
+
+``` python
+wrapper = flashinfer.BatchPrefillWithRaggedKVCacheWrapper(workspace, kv_layout="NHD", backend="auto")
+wrapper.plan(qo_indptr, kv_indptr, num_qo_heads, num_kv_heads,
+             head_dim_qk, head_dim_vo, causal=True,
+             q_data_type=torch.bfloat16, kv_data_type=torch.bfloat16)
+wrapper.run(q, k, v, out=output)
+
+```
+
+如何提交代码详见 [*评测指南*](https://xpuoj.com/d/2)。
+
+## 2. 接口约定
+
+### 2.1 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,
+    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
+);
+
+```
+
+**参数说明**
+
+- q：query tensor，shape (total_q, num_qo_heads, head_dim_qk)，连续 bf16，其中 total_q = qo_indptr[batch_size]
+- k：key tensor，shape (total_kv, num_kv_heads, head_dim_qk)，连续 bf16，其中 total_kv = kv_indptr[batch_size]
+- v：value tensor，shape (total_kv, num_kv_heads, head_dim_vo)，连续 bf16
+- output：输出缓冲区，shape (total_q, num_qo_heads, head_dim_vo)，连续 bf16
+- qo_indptr：query/output ragged indptr，shape (batch_size + 1)，连续 int32
+- kv_indptr：KV ragged indptr，shape (batch_size + 1)，连续 int32
+- seq_len：所有 query/KV 段长度的上界，可用于 launch grid；真实段长必须由 indptr 读取
+- causal：是否启用 causal mask，评测中固定为 1
+
+部分测试点是等长段，但也包含 q_len != kv_len 和不同 batch 段长度不相等的 ragged 测试点。实现不能假设 qo_indptr[b + 1] - qo_indptr[b] == seq_len 或 kv_indptr[b + 1] - kv_indptr[b] == seq_len。
+
+run_kernel 内部需要自行计算合适的 launch 配置并启动 CUDA kernel。为保证计时准确，不建议在 run_kernel 内部做 cudaDeviceSynchronize() 或显式同步。
+
+### 2.2 Triton
+
+你必须在提交的 Python 代码中提供 run_kernel 函数，函数名、参数顺序、类型必须完全一致：
+
+``` python
+import triton
+import triton.language as tl
+
+@triton.jit
+def your_kernel(...):
+    ...
+
+def run_kernel(
+    q,  # Tensor[bf16], shape (total_q, num_qo_heads, head_dim_qk)
+    k,  # Tensor[bf16], shape (total_kv, num_kv_heads, head_dim_qk)
+    v,  # Tensor[bf16], shape (total_kv, num_kv_heads, head_dim_vo)
+    output,  # Tensor[bf16], shape (total_q, num_qo_heads, head_dim_vo)
+    qo_indptr,  # Tensor[int32], shape (batch_size + 1)
+    kv_indptr,  # Tensor[int32], shape (batch_size + 1)
+    batch_size,  # int64
+    seq_len,  # int64, max segment length bound
+    num_qo_heads,  # int64
+    num_kv_heads,  # int64
+    head_dim_qk,  # int64
+    head_dim_vo,  # int64
+    causal,  # int64
+):
+    ...
+
+```
+
+**参数说明**
+
+- q/k/v：FlashInfer ragged prefill 输入 tensor，连续 bfloat16
+- output：输出缓冲区，连续 bfloat16，需要写入结果
+- qo_indptr/kv_indptr：ragged indptr，连续 int32；真实段长和 total_q/total_kv 以 indptr 为准
+- causal：是否启用 causal mask，评测中固定为 1
+
+run_kernel 内部需要自行计算合适的 grid/block，并 launch 你实现的 Triton kernel。
+
+### 2.3 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 (total_q, num_qo_heads, head_dim_qk)
+    k,  # Tensor[bf16], shape (total_kv, num_kv_heads, head_dim_qk)
+    v,  # Tensor[bf16], shape (total_kv, num_kv_heads, head_dim_vo)
+    output,  # Tensor[bf16], shape (total_q, num_qo_heads, head_dim_vo)
+    qo_indptr,  # Tensor[int32], shape (batch_size + 1)
+    kv_indptr,  # Tensor[int32], shape (batch_size + 1)
+    batch_size,  # int64
+    seq_len,  # int64, max segment length bound
+    num_qo_heads,  # int64
+    num_kv_heads,  # int64
+    head_dim_qk,  # int64
+    head_dim_vo,  # int64
+    causal,  # int64
+):
+    global real_kernel
+    if real_kernel is None:
+        real_kernel = build_kernel(...)
+    real_kernel(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)
+
+```
+
+**参数说明**
+
+- q/k/v：FlashInfer ragged prefill 输入 tensor，连续 bfloat16
+- output：输出缓冲区，连续 bfloat16，需要写入结果
+- qo_indptr/kv_indptr：ragged indptr，连续 int32；真实段长和 total_q/total_kv 以 indptr 为准
+- causal：是否启用 causal mask，评测中固定为 1
+
+run_kernel 内部需要自行计算合适的 grid/block，并 launch 你实现的 TileLang kernel。
+
+## 3. 输入格式
+
+本题输入由评测程序在 GPU 上构造，并按接口约定中的顺序传入 run_kernel。
+
+所有 q/k/v/output 均为连续 torch.bfloat16 CUDA tensor，qo_indptr/kv_indptr 为连续 torch.int32 CUDA tensor。
+
+张量布局固定为 FlashInfer ragged prefill 的 NHD 布局。
+
+## 4. 输出格式
+
+输出写入 output，shape 为 (total_q, num_qo_heads, head_dim_vo)，类型为 bfloat16，其中 total_q = qo_indptr[batch_size]。
+
+## 5. 样例
+
+若 batch_size = 1、seq_len = 4、num_qo_heads = 1、num_kv_heads = 1，则：
+
+```
+qo_indptr = [0, 4]
+kv_indptr = [0, 4]
+```
+
+第 t 个 query 会访问同一 batch 内的 KV token 前缀；启用 causal mask 时，只能看到位置不超过 t 的 token。例如 t = 2 时：
+
+```
+attention = softmax(q[2, 0, :] @ k[0:3, 0, :].T / sqrt(head_dim_qk))
+output[2, 0, :] = attention @ v[0:3, 0, :]
+```
+
+若某个 varlen case 中 q_len=2、kv_len=4，则 causal mask 采用 FlashInfer/sol-execbench 的 bottom-right 对齐：第 t 个 query 可见的 KV 上界为 t + 1 + (kv_len - q_len)。例如 t=0 时可见 k[0:3]，t=1 时可见 k[0:4]。
+
+## 6. 数据范围与提示
+
+- 数据类型：q/k/v/output 均为 bfloat16
+- KV layout：NHD
+- num_qo_heads = 32
+- num_kv_heads = 4
+- causal = 1
+- head_dim_qk, head_dim_vo 取值为 (128, 128)
+- batch_size 取值随测试点变化，覆盖 1, 2, 4, 15, 16, 27, 33
+- seq_len 参数表示所有 query/KV 段长度的上界，各测试点的段长上界覆盖 1, 65, 123, 873, 987, 1024, 1280, 2048, 4096, 16384（变长测试点内部还包含 512、640 等更短的真实段长）
+- total_q = qo_indptr[batch_size]
+- total_kv = kv_indptr[batch_size]
+
+注意：
+
+- G = num_qo_heads / num_kv_heads，同一个 KV head 服务连续的 G 个 query heads。
+- 对 query head h_q，对应的 KV head 为 h_q / G。
+- 真实段长必须从 qo_indptr 和 kv_indptr 读取，不能假设每段长度相同。
+- 启用 causal mask 后，采用 bottom-right 对齐。若当前段 q_len != kv_len，第 t 个 query 可访问的位置满足 kv_pos < t + 1 + (kv_len - q_len)。
+- 输出校验容差为 rtol=1.6e-2, atol=1.6e-2，且允许不超过 1% 的元素超差（匹配率需 ≥ 0.99）。
+- 被容忍的超差元素其绝对误差仍不得超过 8 × (atol + rtol · |ref|)，避免个别段被整段算错而蒙混通过。
+- 单 token 边界（用例 14）和非 2 的幂尾段（用例 15）为小规模确定性用例，要求逐元素通过（匹配率需 = 1.0）。
+- q/k/v 使用标准正态分布生成，避免均匀正输入导致长序列 softmax 退化成近似 prefix mean。
+
+## 7. 测试用例尺寸
+
+测试点顺序与 testcase_config.py 的 TESTCASES 一致。共 15 个测试点，全部 head_dim_qk = head_dim_vo = 128，覆盖等长长序列、变长 ragged、q_len < kv_len、短段和非 2 的幂长度。
+
+<table border="1" cellpadding="6" cellspacing="0" style="border-collapse:collapse; width:100%;">
+<thead>
+<tr style="text-align:center; vertical-align:middle;">
+<th style="padding:6px 10px;">测试用例ID</th>
+<th>类型</th>
+<th>batch</th>
+<th>total_q</th>
+<th>total_kv</th>
+<th>max_q</th>
+<th>max_kv</th>
+<th>heads</th>
+<th>head_dim</th>
+</tr>
+</thead>
+<tbody>
+<tr>
+<td>1</td>
+<td>混合 ragged 长序列</td>
+<td>33</td>
+<td colspan="2">16294</td>
+<td colspan="2">987</td>
+<td rowspan="15">32/4</td>
+<td rowspan="15">128/128</td>
+</tr>
+<tr>
+<td>2</td>
+<td rowspan="7">等长序列</td>
+<td rowspan="3">1</td>
+<td colspan="4">1024</td>
+</tr>
+<tr>
+<td>3</td>
+<td colspan="4">4096</td>
+</tr>
+<tr>
+<td>4</td>
+<td colspan="4">16384</td>
+</tr>
+<tr>
+<td>5</td>
+<td rowspan="2">4</td>
+<td colspan="2">4096</td>
+<td colspan="2">1024</td>
+</tr>
+<tr>
+<td>6</td>
+<td rowspan="2" colspan="2">16384</td>
+<td colspan="2">4096</td>
+</tr>
+<tr>
+<td>7</td>
+<td rowspan="2">16</td>
+<td colspan="2">1024</td>
+</tr>
+<tr>
+<td>8</td>
+<td colspan="2">32768</td>
+<td colspan="2">2048</td>
+</tr>
+<tr>
+<td>9</td>
+<td>变长 <code>q_len &lt; kv_len</code></td>
+<td rowspan="2">4</td>
+<td>2048</td>
+<td>4096</td>
+<td>512</td>
+<td>1024</td>
+</tr>
+<tr>
+<td>10</td>
+<td>混合变长 <code>q_len &lt; kv_len</code></td>
+<td>1536</td>
+<td>3584</td>
+<td>640</td>
+<td>1280</td>
+</tr>
+<tr>
+<td>11</td>
+<td>双段变长 <code>q_len &lt; kv_len</code></td>
+<td>2</td>
+<td>1024</td>
+<td>3072</td>
+<td>512</td>
+<td>2048</td>
+</tr>
+<tr>
+<td>12</td>
+<td>混合 ragged 中长序列</td>
+<td>27</td>
+<td colspan="2">12251</td>
+<td colspan="2">873</td>
+</tr>
+<tr>
+<td>13</td>
+<td>混合 ragged 短序列</td>
+<td>15</td>
+<td colspan="2">969</td>
+<td colspan="2">123</td>
+</tr>
+<tr>
+<td>14</td>
+<td>单 token 边界</td>
+<td colspan="5">1</td>
+</tr>
+<tr>
+<td>15</td>
+<td>非 2 的幂尾段</td>
+<td>2</td>
+<td colspan="2">98</td>
+<td colspan="2">65</td>
+</tr>
+</tbody>
+</table>
+
+说明：变长测试点的真实段长由 qo_indptr 和 kv_indptr 给出；参赛实现应始终以 indptr 为准，而不是从 seq_len、total_q 或 total_kv 反推出每段长度。
+
+## 8. PyTorch 参考实现
+
+``` python
+def baseline(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):
+    workspace_buffer = torch.empty(128 * 1024 * 1024, dtype=torch.uint8, device=q.device)
+    wrapper = flashinfer.BatchPrefillWithRaggedKVCacheWrapper(
+        workspace_buffer,
+        kv_layout="NHD",
+        backend="auto",
+    )
+    wrapper.plan(
+        qo_indptr,
+        kv_indptr,
+        num_qo_heads,
+        num_kv_heads,
+        head_dim_qk,
+        head_dim_vo,
+        causal=bool(causal),
+        q_data_type=torch.bfloat16,
+        kv_data_type=torch.bfloat16,
+    )
+    wrapper.run(q, k, v, out=output)
+
+```
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/problem_20002/Agent
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/problem_20002/Agent
@ -0,0 +1,189 @@
+# Agent 推理算子库优化 - FlashInfer Paged Prefill
+
+当前题目说明来源为 [*XPU-OJ 20002*](https://xpuoj.com/contest/2/problem/2)，内容以 XPU-OJ 线上为准。
+
+## 1. 题目描述
+你需要实现 FlashInfer paged KV cache prefill 的CUDA C++前向算子。
+
+本题输入采用 FlashInfer BatchPrefillWithPagedKVCacheWrapper 的 paged NHD 布局。每个 batch 中有 seq_len 个 query token，KV cache 也有 seq_len 个 token，并按 page 存储。
+
+评测程序会调用你提交代码中的 run_kernel 函数。你需要根据 qo_indptr、kv_indptr、kv_indices 和 last_page_len 读取 paged KV cache，并将结果写入 output。
+
+baseline 使用 FlashInfer paged prefill 的 Python API：
+
+``` python
+wrapper = flashinfer.BatchPrefillWithPagedKVCacheWrapper(workspace, kv_layout="NHD", backend="auto")
+wrapper.plan(qo_indptr, kv_indptr, kv_indices, last_page_len,
+             num_qo_heads, num_kv_heads, head_dim, page_block_size,
+             causal=bool(causal),
+             q_data_type=torch.bfloat16, kv_data_type=torch.bfloat16)
+wrapper.run(q, kv_data, out=output)
+
+```
+
+如何提交代码详见 [*评测指南*](https://xpuoj.com/d/2)。
+
+## 2. 接口约定
+
+### 2.1 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* kv_data,
+    __nv_bfloat16* output,
+    const int32_t* qo_indptr,
+    const int32_t* kv_indptr,
+    const int32_t* kv_indices,
+    const int32_t* last_page_len,
+    int64_t batch_size,
+    int64_t seq_len,
+    int64_t num_qo_heads,
+    int64_t num_kv_heads,
+    int64_t head_dim,
+    int64_t page_block_size,
+    int64_t causal
+);
+
+```
+
+**参数说明**
+
+- q：query tensor，shape (batch_size * seq_len, num_qo_heads, head_dim)，连续 bf16
+- kv_data：paged KV cache，shape (num_blocks, 2, page_block_size, num_kv_heads, head_dim)，连续 bf16，其中 kv_data[:, 0] 为 key，kv_data[:, 1] 为 value
+- output：输出缓冲区，shape (batch_size * seq_len, num_qo_heads, head_dim)，连续 bf16
+- qo_indptr：query/output indptr，shape (batch_size + 1)，连续 int32
+- kv_indptr：paged KV indptr，shape (batch_size + 1)，连续 int32
+- kv_indices：page index，shape (num_blocks)，连续 int32
+- last_page_len：每个 batch 最后一个 page 的有效 token 数，shape (batch_size)，连续 int32
+- page_block_size：page size，评测中固定为 16
+- causal：是否启用 causal mask，本题按 benchmark case 固定为 0
+
+run_kernel 内部需要自行计算合适的 launch 配置并启动 CUDA kernel。为保证计时准确，不建议在 run_kernel 内部做 cudaDeviceSynchronize() 或显式同步。
+
+### 2.2 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 * seq_len, num_qo_heads, head_dim)
+    kv_data,  # Tensor[bf16], shape (num_blocks, 2, page_block_size, num_kv_heads, head_dim)
+    output,  # Tensor[bf16], shape (batch_size * seq_len, num_qo_heads, head_dim)
+    qo_indptr,  # Tensor[int32], shape (batch_size + 1)
+    kv_indptr,  # Tensor[int32], shape (batch_size + 1)
+    kv_indices,  # Tensor[int32], shape (num_blocks)
+    last_page_len,  # Tensor[int32], shape (batch_size)
+    batch_size,  # int64
+    seq_len,  # int64
+    num_qo_heads,  # int64
+    num_kv_heads,  # int64
+    head_dim,  # int64
+    page_block_size,  # int64
+    causal,  # int64
+):
+    ...
+
+```
+
+**参数说明**
+
+- q：query tensor，连续 bfloat16
+- kv_data：paged KV cache，连续 bfloat16
+- output：输出缓冲区，连续 bfloat16，需要写入结果
+- qo_indptr/kv_indptr/kv_indices/last_page_len：paged KV metadata，连续 int32
+- page_block_size：评测中固定为 16
+- causal：评测中固定为 0
+
+run_kernel 内部需要自行计算合适的 grid/block，并 launch 你实现的 Triton kernel。
+
+### 2.3 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 * seq_len, num_qo_heads, head_dim)
+    kv_data,  # Tensor[bf16], shape (num_blocks, 2, page_block_size, num_kv_heads, head_dim)
+    output,  # Tensor[bf16], shape (batch_size * seq_len, num_qo_heads, head_dim)
+    qo_indptr,  # Tensor[int32], shape (batch_size + 1)
+    kv_indptr,  # Tensor[int32], shape (batch_size + 1)
+    kv_indices,  # Tensor[int32], shape (num_blocks)
+    last_page_len,  # Tensor[int32], shape (batch_size)
+    batch_size,  # int64
+    seq_len,  # int64
+    num_qo_heads,  # int64
+    num_kv_heads,  # int64
+    head_dim,  # int64
+    page_block_size,  # int64
+    causal,  # int64
+):
+    global real_kernel
+    if real_kernel is None:
+        real_kernel = build_kernel(...)
+    real_kernel(q, kv_data, output, qo_indptr, kv_indptr, kv_indices, last_page_len,
+                batch_size, seq_len, num_qo_heads, num_kv_heads,
+                head_dim, page_block_size, causal)
+
+```
+
+**参数说明**
+
+- q：query tensor，连续 bfloat16
+- kv_data：paged KV cache，连续 bfloat16
+- output：输出缓冲区，连续 bfloat16，需要写入结果
+- qo_indptr/kv_indptr/kv_indices/last_page_len：paged KV metadata，连续 int32
+- page_block_size：评测中固定为 16
+- causal：评测中固定为 0
+
+run_kernel 内部需要自行计算合适的 grid/block，并 launch 你实现的 TileLang kernel。
+
+## 3. 输入格式
+
+本题输入由评测程序在 GPU 上构造，并按接口约定中的顺序传入 run_kernel。
+
+q/kv_data/output 均为连续 torch.bfloat16 CUDA tensor，qo_indptr/kv_indptr/kv_indices/last_page_len 均为连续 torch.int32 CUDA tensor。
+
+KV layout 固定为 FlashInfer paged prefill 的 NHD 布局，page size 固定为 16。
+
+## 4. 输出格式
+
+输出写入 output，shape 为 (batch_size * seq_len, num_qo_heads, head_dim)，类型为 bfloat16。
+
+## 5. 样例
+
+若 batch_size = 1、seq_len = 32、page_block_size = 16，则：
+
+```
+qo_indptr = [0, 32]
+kv_indptr = [0, 2]
+kv_indices = [0, 1]
+last_page_len = [16]
+```
+
+第 0 个 batch 的 KV token 存放在 page 0 和 page 1 中，每个 page 有 16 个 token。
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/problem_20003/Agent
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/problem_20003/Agent
@ -0,0 +1,194 @@
+# Agent 推理算子库优化 - FlashInfer MLA Paged Attention
+
+当前题目说明来源为 [*XPU-OJ 20003*](https://xpuoj.com/contest/2/problem/3)，内容以 XPU-OJ 线上为准。
+
+## 1. 题目描述
+你需要实现 FlashInfer MLA paged attention 的 CUDA C++ 前向算子。
+
+本题使用 BatchMLAPagedAttentionWrapper 的 DeepSeek MLA decode 配置：q_nope 表示不含 RoPE 的 query，q_pe 表示 RoPE 部分的 query，ckv 与 kpe 分别表示 compressed KV cache 与 RoPE KV cache。MLA page size 固定为 1。
+
+评测程序会调用你提交代码中的 run_kernel 函数。你需要根据 q_indptr、kv_indptr、kv_indices 和 kv_lens 读取 cache，并将结果写入 output。
+
+baseline 使用 FlashInfer MLA 的 Python API：
+
+``` python
+wrapper = flashinfer.mla.BatchMLAPagedAttentionWrapper(workspace, backend="auto")
+wrapper.plan(q_indptr, kv_indptr, kv_indices, kv_lens,
+             num_heads, head_dim_ckv, head_dim_kpe,
+             page_size, False, sm_scale,
+             q_nope.dtype, ckv.dtype)
+wrapper.run(q_nope, q_pe, ckv, kpe, out=output, return_lse=False)
+
+```
+
+如何提交代码详见 [*评测指南*](https://xpuoj.com/d/2)。
+
+## 2. 接口约定
+
+### 2.1 CUDA
+
+你必须在提交的 CUDA 源码中提供如下 C 符号，函数名、参数类型、顺序必须完全一致，并使用 extern "C" 防止 name mangling：
+
+``` cpp
+#include <stdint.h>
+#include <cuda_bf16.h>
+
+extern "C" void run_kernel(
+    const __nv_bfloat16* q_nope,
+    const __nv_bfloat16* q_pe,
+    const __nv_bfloat16* ckv,
+    const __nv_bfloat16* kpe,
+    __nv_bfloat16* output,
+    const int32_t* q_indptr,
+    const int32_t* kv_indptr,
+    const int32_t* kv_indices,
+    const int32_t* kv_lens,
+    int64_t batch_size,
+    int64_t seq_len,
+    int64_t num_heads,
+    int64_t head_dim_ckv,
+    int64_t head_dim_kpe,
+    int64_t page_size,
+    int64_t causal
+);
+
+```
+
+**参数说明**
+
+- q_nope：query 的 compressed/nope 部分，shape (batch_size, num_heads, head_dim_ckv)，连续 bf16
+- q_pe：query 的 RoPE 部分，shape (batch_size, num_heads, head_dim_kpe)，连续 bf16
+- ckv：compressed KV cache，shape (batch_size * seq_len, 1, head_dim_ckv)，连续 bf16
+- kpe：RoPE KV cache，shape (batch_size * seq_len, 1, head_dim_kpe)，连续 bf16
+- output：输出缓冲区，shape (batch_size, num_heads, head_dim_ckv)，连续 bf16
+- q_indptr：decode query indptr，shape (batch_size + 1)，内容为 [0, 1, ..., batch_size]
+- kv_indptr：KV indptr，shape (batch_size + 1)，每段长度为 seq_len
+- kv_indices：page index，shape (batch_size * seq_len)，连续 int32
+- kv_lens：每个 batch 的 KV 长度，shape (batch_size)，连续 int32
+- page_size：评测中固定为 1
+- causal：评测中固定为 0
+
+run_kernel 内部需要自行计算合适的 launch 配置并启动 CUDA kernel。为保证计时准确，不建议在 run_kernel 内部做 cudaDeviceSynchronize() 或显式同步。
+
+### 2.2 Triton
+
+你必须在提交的 Python 代码中提供 run_kernel 函数，函数名、参数顺序、类型必须完全一致：
+
+``` python
+import triton
+import triton.language as tl
+
+@triton.jit
+def your_kernel(...):
+    ...
+
+def run_kernel(
+    q_nope,  # Tensor[bf16], shape (batch_size, num_heads, head_dim_ckv)
+    q_pe,  # Tensor[bf16], shape (batch_size, num_heads, head_dim_kpe)
+    ckv,  # Tensor[bf16], shape (batch_size * seq_len, 1, head_dim_ckv)
+    kpe,  # Tensor[bf16], shape (batch_size * seq_len, 1, head_dim_kpe)
+    output,  # Tensor[bf16], shape (batch_size, num_heads, head_dim_ckv)
+    q_indptr,  # Tensor[int32], shape (batch_size + 1)
+    kv_indptr,  # Tensor[int32], shape (batch_size + 1)
+    kv_indices,  # Tensor[int32], shape (batch_size * seq_len)
+    kv_lens,  # Tensor[int32], shape (batch_size)
+    batch_size,  # int64
+    seq_len,  # int64
+    num_heads,  # int64
+    head_dim_ckv,  # int64
+    head_dim_kpe,  # int64
+    page_size,  # int64
+    causal,  # int64
+):
+    ...
+
+```
+
+**参数说明**
+
+- q_nope/q_pe/ckv/kpe：MLA attention 输入 tensor，连续 bfloat16
+- output：输出缓冲区，连续 bfloat16，需要写入结果
+- q_indptr/kv_indptr/kv_indices/kv_lens：paged attention metadata，连续 int32
+- page_size：评测中固定为 1
+- causal：评测中固定为 0
+
+run_kernel 内部需要自行计算合适的 grid/block，并 launch 你实现的 Triton kernel。
+
+### 2.3 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_nope,  # Tensor[bf16], shape (batch_size, num_heads, head_dim_ckv)
+    q_pe,  # Tensor[bf16], shape (batch_size, num_heads, head_dim_kpe)
+    ckv,  # Tensor[bf16], shape (batch_size * seq_len, 1, head_dim_ckv)
+    kpe,  # Tensor[bf16], shape (batch_size * seq_len, 1, head_dim_kpe)
+    output,  # Tensor[bf16], shape (batch_size, num_heads, head_dim_ckv)
+    q_indptr,  # Tensor[int32], shape (batch_size + 1)
+    kv_indptr,  # Tensor[int32], shape (batch_size + 1)
+    kv_indices,  # Tensor[int32], shape (batch_size * seq_len)
+    kv_lens,  # Tensor[int32], shape (batch_size)
+    batch_size,  # int64
+    seq_len,  # int64
+    num_heads,  # int64
+    head_dim_ckv,  # int64
+    head_dim_kpe,  # int64
+    page_size,  # int64
+    causal,  # int64
+):
+    global real_kernel
+    if real_kernel is None:
+        real_kernel = build_kernel(...)
+    real_kernel(q_nope, q_pe, ckv, kpe, output,
+                q_indptr, kv_indptr, kv_indices, kv_lens,
+                batch_size, seq_len, num_heads,
+                head_dim_ckv, head_dim_kpe, page_size, causal)
+
+```
+
+**参数说明**
+
+- q_nope/q_pe/ckv/kpe：MLA attention 输入 tensor，连续 bfloat16
+- output：输出缓冲区，连续 bfloat16，需要写入结果
+- q_indptr/kv_indptr/kv_indices/kv_lens：paged attention metadata，连续 int32
+- page_size：评测中固定为 1
+- causal：评测中固定为 0
+
+run_kernel 内部需要自行计算合适的 grid/block，并 launch 你实现的 TileLang kernel。
+
+## 3. 输入格式
+
+本题输入由评测程序在 GPU 上构造，并按接口约定中的顺序传入 run_kernel。
+
+q_nope/q_pe/ckv/kpe/output 均为连续 torch.bfloat16 CUDA tensor，q_indptr/kv_indptr/kv_indices/kv_lens 均为连续 torch.int32 CUDA tensor。
+
+## 4. 输出格式
+
+输出写入 output，shape 为 (batch_size, num_heads, head_dim_ckv)，类型为 bfloat16。
+
+## 5. 样例
+
+若 batch_size = 2、seq_len = 4，则：
+
+```
+q_indptr = [0, 1, 2]
+kv_indptr = [0, 4, 8]
+kv_indices = [0, 1, 2, 3, 4, 5, 6, 7]
+kv_lens = [4, 4]
+```
+
+每个 batch 只有 1 个 decode query，会访问对应 batch 的全部 KV cache。
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/problem_20004/Agent
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/problem_20004/Agent
@ -0,0 +1,178 @@
+# Agent 推理算子库优化 - FlashInfer Paged Decode
+
+当前题目说明来源为 [*XPU-OJ 20004*](https://xpuoj.com/contest/2/problem/4)，内容以 XPU-OJ 线上为准。
+
+## 1. 题目描述
+你需要实现 FlashInfer paged KV cache decode 的 CUDA C++ 前向算子。
+
+本题输入采用 FlashInfer BatchDecodeWithPagedKVCacheWrapper 的 paged NHD 布局。每个 batch 只有 1 个 query token，KV cache 长度为 seq_len_kv。
+
+评测程序会调用你提交代码中的 run_kernel 函数。你需要根据 kv_indptr、kv_indices 和 last_page_len 读取 paged KV cache，并将结果写入 output。
+
+baseline 使用 FlashInfer paged decode 的 Python API：
+
+``` python
+wrapper = flashinfer.BatchDecodeWithPagedKVCacheWrapper(
+    workspace, kv_layout="NHD", use_tensor_cores=True)
+wrapper.plan(kv_indptr, kv_indices, last_page_len,
+             num_qo_heads, num_kv_heads, head_dim, page_block_size,
+             data_type=torch.bfloat16, q_data_type=torch.bfloat16)
+wrapper.run(q, kv_data, out=output)
+
+```
+
+如何提交代码详见 [*评测指南*](https://xpuoj.com/d/2)。
+
+## 2. 接口约定
+
+### 2.1 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* kv_data,
+    __nv_bfloat16* output,
+    const int32_t* kv_indptr,
+    const int32_t* kv_indices,
+    const int32_t* last_page_len,
+    int64_t batch_size,
+    int64_t seq_len_kv,
+    int64_t num_qo_heads,
+    int64_t num_kv_heads,
+    int64_t head_dim,
+    int64_t page_block_size
+);
+
+```
+
+**参数说明**
+
+- q：decode query tensor，shape (batch_size, num_qo_heads, head_dim)，连续 bf16
+- kv_data：paged KV cache，shape (num_blocks, 2, page_block_size, num_kv_heads, head_dim)，连续 bf16，其中 kv_data[:, 0] 为 key，kv_data[:, 1] 为 value
+- output：输出缓冲区，shape (batch_size, num_qo_heads, head_dim)，连续 bf16
+- kv_indptr：paged KV indptr，shape (batch_size + 1)，连续 int32
+- kv_indices：page index，shape (num_blocks)，连续 int32
+- last_page_len：每个 batch 最后一个 page 的有效 token 数，shape (batch_size)，连续 int32
+- page_block_size：page size，评测中固定为 16
+
+run_kernel 内部需要自行计算合适的 launch 配置并启动 CUDA kernel。为保证计时准确，不建议在 run_kernel 内部做 cudaDeviceSynchronize() 或显式同步。
+
+### 2.2 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, num_qo_heads, head_dim)
+    kv_data,  # Tensor[bf16], shape (num_blocks, 2, page_block_size, num_kv_heads, head_dim)
+    output,  # Tensor[bf16], shape (batch_size, num_qo_heads, head_dim)
+    kv_indptr,  # Tensor[int32], shape (batch_size + 1)
+    kv_indices,  # Tensor[int32], shape (num_blocks)
+    last_page_len,  # Tensor[int32], shape (batch_size)
+    batch_size,  # int64
+    seq_len_kv,  # int64
+    num_qo_heads,  # int64
+    num_kv_heads,  # int64
+    head_dim,  # int64
+    page_block_size,  # int64
+):
+    ...
+
+```
+
+**参数说明**
+
+- q：decode query tensor，连续 bfloat16
+- kv_data：paged KV cache，连续 bfloat16
+- output：输出缓冲区，连续 bfloat16，需要写入结果
+- kv_indptr/kv_indices/last_page_len：paged KV metadata，连续 int32
+- page_block_size：评测中固定为 16
+
+run_kernel 内部需要自行计算合适的 grid/block，并 launch 你实现的 Triton kernel。
+
+### 2.3 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, num_qo_heads, head_dim)
+    kv_data,  # Tensor[bf16], shape (num_blocks, 2, page_block_size, num_kv_heads, head_dim)
+    output,  # Tensor[bf16], shape (batch_size, num_qo_heads, head_dim)
+    kv_indptr,  # Tensor[int32], shape (batch_size + 1)
+    kv_indices,  # Tensor[int32], shape (num_blocks)
+    last_page_len,  # Tensor[int32], shape (batch_size)
+    batch_size,  # int64
+    seq_len_kv,  # int64
+    num_qo_heads,  # int64
+    num_kv_heads,  # int64
+    head_dim,  # int64
+    page_block_size,  # int64
+):
+    global real_kernel
+    if real_kernel is None:
+        real_kernel = build_kernel(...)
+    real_kernel(q, kv_data, output, kv_indptr, kv_indices, last_page_len,
+                batch_size, seq_len_kv, num_qo_heads,
+                num_kv_heads, head_dim, page_block_size)
+
+```
+
+**参数说明**
+
+- q：decode query tensor，连续 bfloat16
+- kv_data：paged KV cache，连续 bfloat16
+- output：输出缓冲区，连续 bfloat16，需要写入结果
+- kv_indptr/kv_indices/last_page_len：paged KV metadata，连续 int32
+- page_block_size：评测中固定为 16
+
+run_kernel 内部需要自行计算合适的 grid/block，并 launch 你实现的 TileLang kernel。
+
+## 3. 输入格式
+
+本题输入由评测程序在 GPU 上构造，并按接口约定中的顺序传入 run_kernel。
+
+q/kv_data/output 均为连续 torch.bfloat16 CUDA tensor，kv_indptr/kv_indices/last_page_len 均为连续 torch.int32 CUDA tensor。
+
+KV layout 固定为 FlashInfer paged decode 的 NHD 布局，page size 固定为 16。
+
+## 4. 输出格式
+
+输出写入 output，shape 为 (batch_size, num_qo_heads, head_dim)，类型为 bfloat16。
+
+## 5. 样例
+
+若 batch_size = 1、seq_len_kv = 32、page_block_size = 16，则：
+
+```
+kv_indptr = [0, 2]
+kv_indices = [0, 1]
+last_page_len = [16]
+```
+
+decode query 会访问该 batch 的全部 32 个 KV token。
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/scripts/build_fused_moe_i8_tn_pybind.sh
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/scripts/build_fused_moe_i8_tn_pybind.sh
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/scripts/run_fused_moe_i8_tn_benchmark.sh
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/scripts/run_fused_moe_i8_tn_benchmark.sh
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/scripts/run_fused_moe_i8_tn_pybind_test.sh
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/scripts/run_fused_moe_i8_tn_pybind_test.sh
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_bf16_tn/CMakeLists.txt
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_bf16_tn/CMakeLists.txt
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_bf16_tn/src/fused_moe_bf16_tn_example.cpp
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_bf16_tn/src/fused_moe_bf16_tn_example.cpp
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/CMakeLists.txt
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/CMakeLists.txt
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/benchmark_results.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/benchmark_results.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/fused-moe-i8-tn-pybind-triton-guide.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/fused-moe-i8-tn-pybind-triton-guide.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/python/benchmark_fused_moe_i8_tn.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/python/benchmark_fused_moe_i8_tn.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/python/fused_moe_i8_tn_triton.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/python/fused_moe_i8_tn_triton.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/python/test_fused_moe_i8_tn_pybind.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/python/test_fused_moe_i8_tn_pybind.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_example.cpp
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_example.cpp
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_kernel.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_kernel.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_macros.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_macros.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_pybind.cpp
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_pybind.cpp
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_runner.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_runner.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_types.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/fused_moe_i8_tn/src/fused_moe_i8_tn_types.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/i8_tn_256x256x128_raw_arrays/CMakeLists.txt
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/i8_tn_256x256x128_raw_arrays/CMakeLists.txt
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/i8_tn_256x256x128_raw_arrays/include/gemm_i8_tn_256x256x128_raw_arrays.hpp
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/i8_tn_256x256x128_raw_arrays/include/gemm_i8_tn_256x256x128_raw_arrays.hpp
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/i8_tn_256x256x128_raw_arrays/include/standalone_maca_kernel_utils.hpp
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/i8_tn_256x256x128_raw_arrays/include/standalone_maca_kernel_utils.hpp
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/i8_tn_256x256x128_raw_arrays/src/test.cpp
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/standalone/i8_tn_256x256x128_raw_arrays/src/test.cpp
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/attr.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/attr.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/buffer_info.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/buffer_info.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/cast.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/cast.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/chrono.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/chrono.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/common.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/common.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/complex.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/complex.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/conduit/README.txt
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/conduit/README.txt
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/conduit/pybind11_conduit_v1.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/conduit/pybind11_conduit_v1.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/conduit/pybind11_platform_abi_id.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/conduit/pybind11_platform_abi_id.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/conduit/wrap_include_python_h.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/conduit/wrap_include_python_h.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/critical_section.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/critical_section.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/argument_vector.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/argument_vector.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/class.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/class.h
--- a/Show More
+++ b/Show More
				`@ -1 +0,0 @@`
				`{"requests":[{"kind":"cache","version":2},{"kind":"codemodel","version":2},{"kind":"toolchains","version":1},{"kind":"cmakeFiles","version":1}]}`
				`@ -1 +0,0 @@`
				`# This file is generated by cmake for dependency checking of the CMakeCache.txt file`