新增沐曦通用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 7067 additions and 3092 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)
+- [赛题二选手入口](基于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大模型推理算子库优化/FlashAttention关键算子迁移与优化.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/FlashAttention关键算子迁移与优化.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/FlashInfer
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/FlashInfer
--- a/Agent开发范式的国产GPU大模型推理算子库优化/FlashInfer关键算子迁移与优化.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/FlashInfer关键算子迁移与优化.md
@ -1,356 +0,0 @@
-# FlashInfer 迁移 Baseline 实战
-
-## 一、教程定位
-
-本教程是参赛训练课程的 **FlashInfer Baseline 入门**模块，主要帮助用户快速跑通 FlashInfer 的最小可运行流程。完成本教程后，用户应能够完成源码编译、API 调用、正确性测试和 Benchmark 测试，并记录一份 baseline 性能结果，为后续算子优化提供对比基准。
-
-## 二、学习目标
-
-1.  理解 FlashInfer Attention Kernel 的基本作用与适用场景；
-2.  完成 FlashInfer 环境、工具链的准备与源码编译；
-3.  跑通 BatchDecode、BatchPrefill、MLA 等典型算子的 API 调用示例；
-4.  完成各算子在不同参数配置下的 Benchmark 测试；
-5.  输出各算子的 Baseline 性能结果记录表，为后续算子优化提供对比基准。
-
-
-## 三、适用对象
-
-适合人群
-
-*   参赛选手：需要完成 Baseline 入门模块，为后续算子优化做准备
-*   软件开发者和Vibe Coding开发：希望从事AI相关行业开发，以及用智能体方式来做开发工作
-*   LLM 推理开发者：希望了解 FlashInfer Attention Kernel 的性能表现
-*   算子优化工程师：希望基于MXMACA软件栈在沐曦国产 GPU 上做算子迁移和优化
-
-前置基础
-
-*   Python 基础：能够运行和修改 Python 脚本
-*   PyTorch 基础：了解MXMACA软化栈的使用
-*   Linux 命令行：能够使用终端执行命令
-*   了解 Attention 机制：理解 Q/K/V、KV Cache 等基本概念
-    
-
-## 四、前置准备
-
-开始实战前，请确认你已经完成以下准备：
-
-### GPU准备
-
-* 步骤1：获取算力券[https://developer.metax-tech.com/activities/6](https://developer.metax-tech.com/activities/6)
-
-  *   首次登录需要先进行注册（使用邮箱或者手机号进行注册）
-      
-*   登录成功后进行第二步-邮箱验证，填入自己的邮箱。
-      
-  *   第三步，提交申请。
-    
-  *   获得兑换码
-      
-
-* 步骤2：兑换算力和登陆平台
-
-  *   访问模力方舟官网：[https://ai.gitee.com/](https://ai.gitee.com/)
-      
-  *   进入费用中心 - 算力券 ，点击右上角 “兑换”。
-      
-
-*   步骤3：租用算力
-    
-    - 进入算力容器，选择沐曦，租用算力，建议优先选16G显存/32G显存，如下图：
-*   步骤4：创建实例
-
-    - 基础镜像：maca-pytorch:3.7.1.5-torch2.8-py312-ubuntu24.04-amd64
-*   步骤5：选择工具-lab进入实例环境
-
-
-*   步骤6：在JupyterLab Terminal中检查运行环境的配置，确认沐曦 GPU 可见--可以使用`mx-smi`命令查看
-
-
-### 环境依赖准备
-
-```bash
-pip install flashinfer torch pandas numpy
-```
-
-### opencode Agent安装
-
-```shell
-curl -fsSL https://opencode.ai/install | bash
-```
-
-说明：
-
-自动下载最新版本 自动配置环境 适用于 macOS / Linux / WSL 安装完成后，还会提示你怎么用，OpenCode 包含免费模式，使用方式：
-
-```shell
-cd   # 进入项目目录
-opencode
-```
-
-### 代码准备
-
-*   已进入指定项目目录；
-    
-*   准备测试脚本和 Benchmark 脚本。
-    
-
-| 命令 | 说明 |
-| --- | --- |
-| `python bench_batch_decode.py` | 运行 Batch Decode 基准测试 |
-| `python bench_batch_prefill_paged.py` | 运行 Batch Prefill (Paged KV Cache) 基准测试 |
-| `python bench_batch_prefill_ragged.py` | 运行 Batch Prefill (Ragged KV Cache) 基准测试 |
-| `python bench_batch_mla.py` | 运行 MLA (Multi-head Latent Attention) 基准测试 |
-
-## 五、知识预备
-
-#### LLM推理阶段重要概念：
-
-*   **Prefill阶段**：prefill 阶段是指处理输入 prompt 的阶段
-    
-    *   输入：用户一次性给出的完整 prompt，长度为 seq\_len
-        
-    *   计算：对 prompt 中的每个 token 并行计算注意力，生成第一个输出 token 及 KV cache
-        
-    *   特点：这是**计算密集型（compute-bound）**阶段，因为需要做完整的 seq\_len x seq\_len 注意力矩阵乘法
-    
-*   **decode阶段**：
-    
-    *   每次只生成 1 个 token，利用 prefill 阶段填充好的 KV cache 做自回归生成
-        
-    *   **显存带宽密集型（memory-bound**），瓶颈在从显存读取 KV cache 而非计算
-        
-
-prefill = 并行处理用户输入，decode = 逐个生成回答 token
-
-## 六、项目实践--FlashInfer-Baseline
-
-### Step 1：检查运行环境
-
-**目标：** 进入terminal确认当前环境满足本模块运行要求。
-
-**操作：** 检查 GPU、Python、编译工具和依赖版本。
-
-![giteeai 实例 12](https://origin.picgo.net/2026/06/04/giteeai--12c1772b12867f6be0.png)
-
-**常见问题：**
-
-| 问题 | 解决方法 |
-| --- | --- |
-| `mx-smi: command not found` | 使用 `python -c "import torch; print(torch.cuda.get_device_name(0))"` |
-| `No GPUs are available` | 检查 MXMACA 驱动是否正确安装 |
-| `ModuleNotFoundError: No module named 'flashinfer'` | 执行 `pip install flashinfer` 安装 |
-
-### Step 2：进入项目目录
-
-**目标：** 进入本模块所需的源码目录[flashinfer_baseline](https://gitlink.org.cn/metax-maca/op_optimization/tree/master/%E5%9F%BA%E4%BA%8EAI%20Agent%E5%BC%80%E5%8F%91%E8%8C%83%E5%BC%8F%E7%9A%84%E5%9B%BD%E4%BA%A7GPU%E5%A4%A7%E6%A8%A1%E5%9E%8B%E6%8E%A8%E7%90%86%E7%AE%97%E5%AD%90%E5%BA%93%E4%BC%98%E5%8C%96%2Fbaselines%2Fflashinfer_baseline)。
-
-1. 克隆代码仓库
-
-   ```Bash
-   git clone https://gitlink.org.cn/metax-maca/op_optimization.git
-   ```
-
-2. 准备flashinfer_baseline
-
-   在仓库目录 `op_optimization/基于AI Agent开发范式的国产GPU大模型推理算子库优化` 下，找到 `flashinfer_baseline` 文件夹。可以将 `flashinfer_baseline` 整个目录复制到工作目录 `data/` 下。
-
-**下一步操作：** 切换到 FlashInfer Baseline 项目目录。
-
-**命令示例：**
-
-```bash
-cd /data/flashinfer_baseline
-ls -la
-```
-
-**预期结果：**
-
-```plaintext
-bench_common.py
-bench_batch_decode.py
-bench_batch_prefill_paged.py
-bench_batch_prefill_ragged.py
-bench_batch_mla.py
-README.md
-...
-
-```
-
-### Step 3：验证项目脚本
-
-**目标：** 确认所有基准测试脚本可正常执行。
-
-**操作：** 检查脚本文件是否存在且可读。
-
-**命令示例：**
-
-```bash
-# 检查脚本文件
-python -c "import os; scripts = ['bench_common.py', 'bench_batch_decode.py', 'bench_batch_prefill_paged.py', 'bench_batch_prefill_ragged.py', 'bench_batch_mla.py']; [print(f'✓ {s}') if os.path.exists(s) else print(f'✗ {s} missing') for s in scripts]"
-
-# 测试脚本导入
-python -c "from bench_common import setup_workspace, get_csv_path; print('脚本导入正常')"
-
-```
-
-**预期结果：**
-
-```plaintext
-✓ bench_common.py
-✓ bench_batch_decode.py
-✓ bench_batch_prefill_paged.py
-✓ bench_batch_prefill_ragged.py
-✓ bench_batch_mla.py
-脚本导入正常
-
-```
-
-### Step 4：运行单算子 Benchmark
-
-**目标：** 执行基准测试，获取 Baseline 性能数据。
-
-**操作：** 运行 BatchDecode 基准测试脚本。
-
-**命令示例：**
-
-```bash
-python bench_batch_decode.py
-
-```
-
-**预期结果：**
-
-```plaintext
-[BatchDecodeWithPagedKVCacheWrapper] Starting benchmark, total cases: 144
-  [1/144] bs=1, kv_len=512, hd=64: 0.020ms, 51.53 GB/s, 0.82 TFLOPs
-  [2/144] bs=1, kv_len=512, hd=128: 0.023ms, 45.28 GB/s, 0.72 TFLOPs
-  ...
-
-Results saved to BatchDecodeWithPagedKVCacheWrapper_20260525_xxxxxx.csv
-
-```
-
-**常见问题：**
-
-| 问题 | 解决方法 |
-| --- | --- |
-| `out of memory` | 减小 batch\_size 或 seq\_len 参数 |
-| 运行时间过长 | 脚本会自动调整重复次数，耐心等待 |
-
-### Step 5：查看测试结果
-
-**目标：** 查看并分析 Benchmark 输出结果。
-
-**操作：** 读取生成的 CSV 结果文件。
-
-**命令示例：**
-
-```bash
-# 查看最新生成的 CSV 文件
-ls -lt *.csv | head -1
-
-# 使用 Python 查看结果
-python -c "
-import pandas as pd
-import glob
-csv_files = glob.glob('BatchDecodeWithPagedKVCacheWrapper_*.csv')
-if csv_files:
-    df = pd.read_csv(max(csv_files))
-    print(df[['batch_size', 'seq_len_kv', 'head_dim', 'time_ms', 'bandwidth_GB_s', 'tflops']].head(10))
-"
-
-```
-
-**预期结果：**
-
-```plaintext
-   batch_size  seq_len_kv  head_dim   time_ms  bandwidth_GB_s    tflops
-0           1         512        64  0.020429       51.528822  0.821253
-1           1         512       128  0.023340       45.278052  0.718833
-2           1         512       256  0.031949       66.153846  1.050256
-...
-
-```
-
-
-
-## 七、常见问题
-
-### 环境相关问题
-
-| 问题 | 原因 | 解决办法 |
-| --- | --- | --- |
-| `No GPUs are available` | MXMACA 驱动未安装或 GPU 不可见 | 检查驱动安装，运行 `python -c "import torch; print(torch.cuda.device_count())"` 验证 |
-| `ModuleNotFoundError: No module named 'flashinfer'` | flashinfer 未安装 | 执行 `pip install flashinfer` |
-| `out of memory` | GPU 显存不足 | 减小 `batch_size` 或 `seq_len` 参数 |
-
-### 运行相关问题
-
-| 问题 | 原因 | 解决办法 |
-| --- | --- | --- |
-| Benchmark 运行时间过长 | 参数组合过多， workload 较大 | 耐心等待，脚本会自动调整重复次数 |
-| `KeyError: 'BatchPrefillWithPagedKVCacheKernel'` | profiler 未捕获目标 kernel | 检查 `target_kernels` 配置是否正确 |
-| CSV 文件为空 | 测试未正常完成 | 检查 GPU 显存是否充足，重新运行 |
-
-### 代码相关问题
-
-| 问题 | 原因 | 解决办法 |
-| --- | --- | --- |
-| `ImportError: cannot import name 'xxx' from 'bench_common'` | 函数名拼写错误 | 检查 `bench_common.py` 中的函数名 |
-| `RuntimeError: error: device-side assert triggered` | 输入参数超出范围 | 检查 `num_qo_heads`、`num_kv_heads`、`head_dim` 配置 |
-
-### 性能相关问题
-
-| 问题 | 原因 | 解决办法 |
-| --- | --- | --- |
-| TFLOPs 数值异常低 | 工作负载过小，kernel 启动开销占比大 | 增大 `batch_size` 或 `seq_len` |
-| 带宽数值异常低 | 数据未正确加载到 GPU | 检查 Tensor 是否在 CUDA 设备上 |
-
-## 八、下一步学习建议
-
-### 1. 保存你的 Baseline 结果
-
-将本次运行生成的 CSV 文件妥善保存，后续优化时需要以此作为对比基准。
-
-```bash
-# 建议创建 results 目录保存
-mkdir -p results
-mv *.csv results/
-
-```
-
-### 2. 深入理解 FlashInfer 核心概念
-
-*   阅读 FlashInfer 官方文档，理解 Paged KV Cache、Ragged KV Cache 的设计理念
-    
-*   学习 MLA (Multi-head Latent Attention) 的原理，了解 DeepSeek 的注意力优化方案
-    
-*   理解 `plan()` 和 `run()` 两阶段设计的作用
-    
-
-### 3. 进入算子优化模块
-
-参考后续优化模块，学习以下优化技术：
-
-*   **Kernel Tuning**：调整 Block Size、Thread Count 等参数
-    
-*   **Memory Optimization**：减少显存占用、优化数据搬运
-    
-*   **Compute Optimization**：提升计算效率
-    
-
-### 4. 参考资源
-
-*   FlashInfer 官方仓库：https://github.com/flashinfer-ai/flashinfer
-    
-*   FlashInfer 文档：https://flashinfer.ai
-    
-
-### 5. 记录优化过程
-
-建议维护一份优化日志，记录每次优化的改动和性能变化：
-
-| 优化项 | 改动内容 | Baseline | 优化后 | 提升比例 |
-| --- | --- | --- | --- | --- |
-| 例：调整 block\_size | 16 → 32 | xx ms | xx ms | xx% |
-| 完成优化后，再次运行本模块的 Benchmark 脚本，对比前后性能变化。 |  |  |  |  |
--- a/Agent开发范式的国产GPU大模型推理算子库优化/Fused
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/Fused
--- a/Agent开发范式的国产GPU大模型推理算子库优化/MCTLASS_Fused
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/MCTLASS_Fused
@ -1,614 +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：进入项目目录
-
-**目标：**进入本模块所需的源码目录：https://gitlink.org.cn/metax-maca/op_optimization/tree/master/%E5%9F%BA%E4%BA%8EAI%20Agent%E5%BC%80%E5%8F%91%E8%8C%83%E5%BC%8F%E7%9A%84%E5%9B%BD%E4%BA%A7GPU%E5%A4%A7%E6%A8%A1%E5%9E%8B%E6%8E%A8%E7%90%86%E7%AE%97%E5%AD%90%E5%BA%93%E4%BC%98%E5%8C%96%2Fbaselines%2Ffused_moe
-
-1. 克隆代码仓库
-
-   ```Bash
-   git clone https://gitlink.org.cn/metax-maca/op_optimization.git
-   ```
-
-2. 准备fused_moe
-
-   在仓库目录 `op_optimization/基于AI Agent开发范式的国产GPU大模型推理算子库优化` 下，找到 `fused_moe` 文件夹。可以将 `fused_moe` 整个目录复制到工作目录 `data/` 下。
-
-**下一步操作：** 切换到 FlashInfer Baseline 项目目录。
-
-**操作：**切换到指定项目路径。
-
-**命令示例：**
-
-```apl
-cd data/fused_moe # 这里需要替换成自己的项目路径
-```
-
-**预期结果：**
-
-终端提示符路径显示为 fused\_moe 项目路径：
-
-*   (base) root@0a5a9d0c0f06:/data/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
-[SUCCESS] /data/fused_moe/standalone/fused_moe_i8_tn/build/fused_moe_i8_tn_pybind.sh
-```
-
-**预期结果：**
-
-编译成功无报错，终端显示：
-
-*   \[SUCCESS] /data/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 算子的总运算量 / 实际耗<E99985><E88097><EFBFBD>；
-    
-*   “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大模型推理算子库优化/operator_task_package/flashattn_task_package/__MACOSX/._benchmark_kvcache.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/__MACOSX/._benchmark_kvcache.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/__MACOSX/._benchmark_kvcache_1.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/__MACOSX/._benchmark_kvcache_1.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/__MACOSX/._benchmark_kvcache_20260526_150953.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/__MACOSX/._benchmark_kvcache_20260526_150953.csv
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/benchmark_kvcache.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/benchmark_kvcache.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark_kvcache_20260526_150953.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark_kvcache_20260526_150953.csv
@ -1,49 +0,0 @@
-batch_size,seq_len_kv,heads,headdim,time_ms,bandwidth_GB_s
-1,512,8,128,0.0322,65.27
-2,512,8,128,0.0324,129.61
-4,512,8,128,0.0332,253.01
-8,512,8,128,0.0355,472.66
-16,512,8,128,0.0546,614.98
-32,512,8,128,0.0817,822.60
-64,512,8,128,0.1297,1035.62
-128,512,8,128,0.2453,1095.50
-1,1024,8,128,0.0578,72.55
-2,1024,8,128,0.0586,143.14
-4,1024,8,128,0.0597,281.24
-8,1024,8,128,0.0625,536.80
-16,1024,8,128,0.0982,683.59
-32,1024,8,128,0.1493,899.64
-64,1024,8,128,0.2403,1117.62
-128,1024,8,128,0.4594,1169.27
-1,2048,8,128,0.1101,76.24
-2,2048,8,128,0.1107,151.64
-4,2048,8,128,0.1119,299.88
-8,2048,8,128,0.1159,578.98
-16,2048,8,128,0.1849,726.23
-32,2048,8,128,0.2843,944.47
-64,2048,8,128,0.4607,1165.56
-128,2048,8,128,0.8868,1211.07
-1,4096,8,128,0.2139,78.46
-2,4096,8,128,0.2151,156.01
-4,4096,8,128,0.2163,310.36
-8,4096,8,128,0.2227,602.81
-16,4096,8,128,0.3574,751.13
-32,4096,8,128,0.5540,969.23
-64,4096,8,128,0.9016,1191.07
-128,4096,8,128,1.7414,1233.34
-1,8192,8,128,0.4215,79.61
-2,8192,8,128,0.4226,158.81
-4,8192,8,128,0.4242,316.39
-8,8192,8,128,0.4362,615.46
-16,8192,8,128,0.7035,763.14
-32,8192,8,128,1.0934,982.11
-64,8192,8,128,1.7814,1205.57
-128,8192,8,128,3.4505,1244.82
-1,16384,8,128,0.8356,80.32
-2,16384,8,128,0.8377,160.23
-4,16384,8,128,0.8407,319.30
-8,16384,8,128,0.8625,622.51
-16,16384,8,128,1.3934,770.60
-32,16384,8,128,2.1695,989.88
-64,16384,8,128,3.5397,1213.41
-128,16384,8,128,6.8668,1250.98
--- 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/BatchDecodeWithPagedKVCacheWrapper_20260525_093143.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/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大模型推理算子库优化/operator_task_package/flashinfer_task_package/BatchMLAPagedAttentionWrapper_20260525_074200.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/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大模型推理算子库优化/operator_task_package/flashinfer_task_package/BatchPrefillWithPagedKVCacheWrapper_20260526_150417.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/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大模型推理算子库优化/operator_task_package/flashinfer_task_package/BatchPrefillWithRaggedKVCacheWrapper_20260526_145454.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/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大模型推理算子库优化/operator_task_package/flashinfer_task_package/__MACOSX/._BatchDecodeWithPagedKVCacheWrapper_20260525_093143.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/__MACOSX/._BatchDecodeWithPagedKVCacheWrapper_20260525_093143.csv
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/__MACOSX/._BatchMLAPagedAttentionWrapper_20260525_074200.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/__MACOSX/._BatchMLAPagedAttentionWrapper_20260525_074200.csv
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/__MACOSX/._BatchPrefillWithPagedKVCacheWrapper_20260526_150417.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/__MACOSX/._BatchPrefillWithPagedKVCacheWrapper_20260526_150417.csv
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/__MACOSX/._BatchPrefillWithRaggedKVCacheWrapper_20260526_145454.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/__MACOSX/._BatchPrefillWithRaggedKVCacheWrapper_20260526_145454.csv
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/__MACOSX/._bench_batch_decode.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/__MACOSX/._bench_batch_decode.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/__MACOSX/._bench_batch_mla.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/__MACOSX/._bench_batch_mla.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/__MACOSX/._bench_batch_prefill_paged.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/__MACOSX/._bench_batch_prefill_paged.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/__MACOSX/._bench_batch_prefill_ragged.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/__MACOSX/._bench_batch_prefill_ragged.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/__MACOSX/._bench_common.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/__MACOSX/._bench_common.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/bench_batch_decode.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/bench_batch_decode.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/bench_batch_mla.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/bench_batch_mla.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/bench_batch_prefill_paged.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/bench_batch_prefill_paged.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/bench_batch_prefill_ragged.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/bench_batch_prefill_ragged.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/bench_common.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/bench_common.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/starter/示例冒烟代码.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/starter/示例冒烟代码.md
@ -0,0 +1,209 @@
+# 示例冒烟代码
+
+```c++
+#include <stdint.h>
+
+#include <cuda_bf16.h>
+#include <cuda_runtime.h>
+
+#include <math.h>
+
+namespace {
+
+__device__ __forceinline__ float warp_sum(float x) {
+    for (int offset = 16; offset > 0; offset >>= 1) {
+        x += __shfl_down_sync(0xffffffffu, x, offset);
+    }
+    return __shfl_sync(0xffffffffu, x, 0);
+}
+
+__global__ void ragged_prefill_smoke_kernel(
+    const __nv_bfloat16* __restrict__ q,
+    const __nv_bfloat16* __restrict__ k,
+    const __nv_bfloat16* __restrict__ v,
+    __nv_bfloat16* __restrict__ output,
+    const int32_t* __restrict__ qo_indptr,
+    const int32_t* __restrict__ kv_indptr,
+    int64_t batch_size,
+    int64_t seq_len,
+    int64_t num_qo_heads,
+    int64_t num_kv_heads,
+    int64_t head_dim_qk,
+    int64_t head_dim_vo,
+    int64_t causal,
+    int64_t exact_len) {
+    const int lane = threadIdx.x & 31;
+    const int warp_id = threadIdx.x >> 5;
+    const int warps_per_block = blockDim.x >> 5;
+
+    int64_t work = static_cast<int64_t>(blockIdx.x) * warps_per_block + warp_id;
+    const int64_t total = batch_size * exact_len * num_qo_heads;
+    if (work >= total) return;
+
+    const int64_t qo_head = work % num_qo_heads;
+    work /= num_qo_heads;
+    const int64_t q_pos = work % exact_len;
+    const int64_t batch = work / exact_len;
+
+    const int64_t qo_begin = qo_indptr[batch];
+    const int64_t qo_len = qo_indptr[batch + 1] - qo_begin;
+    if (q_pos >= qo_len) return;
+
+    const int64_t kv_begin = kv_indptr[batch];
+    const int64_t kv_len = kv_indptr[batch + 1] - kv_begin;
+    int64_t visible = kv_len;
+    if (causal) {
+        visible = kv_len - qo_len + q_pos + 1;
+        if (visible < 0) visible = 0;
+        if (visible > kv_len) visible = kv_len;
+    }
+
+    const int64_t group = num_qo_heads / num_kv_heads;
+    const int64_t kv_head = qo_head / group;
+    const int64_t q_row = qo_begin + q_pos;
+    const float scale = rsqrtf(static_cast<float>(head_dim_qk));
+
+    const __nv_bfloat16* q_ptr = q + (q_row * num_qo_heads + qo_head) * head_dim_qk;
+    float qv[4];
+    float acc[4];
+    for (int i = 0; i < 4; ++i) {
+        const int d = lane + i * 32;
+        qv[i] = (d < head_dim_qk) ? __bfloat162float(q_ptr[d]) : 0.0f;
+        acc[i] = 0.0f;
+    }
+
+    float m = -1.0e20f;
+    float l = 0.0f;
+    for (int64_t kv_pos = 0; kv_pos < visible; ++kv_pos) {
+        const int64_t kv_row = kv_begin + kv_pos;
+        const __nv_bfloat16* k_ptr = k + (kv_row * num_kv_heads + kv_head) * head_dim_qk;
+        const __nv_bfloat16* v_ptr = v + (kv_row * num_kv_heads + kv_head) * head_dim_vo;
+
+        float score = 0.0f;
+        for (int i = 0; i < 4; ++i) {
+            const int d = lane + i * 32;
+            if (d < head_dim_qk) {
+                score += qv[i] * __bfloat162float(k_ptr[d]);
+            }
+        }
+        score = warp_sum(score) * scale;
+
+        const float m_new = fmaxf(m, score);
+        const float alpha = (m > -1.0e19f) ? __expf(m - m_new) : 0.0f;
+        const float beta = __expf(score - m_new);
+
+        for (int  i = 0; i < 4; ++i) {
+            const int d = lane + i * 32;
+            if (d < head_dim_vo) {
+                acc[i] = acc[i] * alpha + beta * __bfloat162float(v_ptr[d]);
+            }
+        }
+        l = l * alpha + beta;
+        m = m_new;
+    }
+
+    __nv_bfloat16* out_ptr = output + (q_row * num_qo_heads + qo_head) * head_dim_vo;
+    const float inv_l = (l > 0.0f) ? (1.0f / l) : 0.0f;
+    for (int i = 0; i < 4; ++i) {
+        const int d = lane + i * 32;
+        if (d < head_dim_vo) {
+            out_ptr[d] = __float2bfloat16(acc[i] * inv_l);
+        }
+    }
+}
+
+__global__ void prefix_mean_kernel(
+    const __nv_bfloat16* __restrict__ v,
+    __nv_bfloat16* __restrict__ output,
+    const int32_t* __restrict__ qo_indptr,
+    const int32_t* __restrict__ kv_indptr,
+    int64_t batch_size,
+    int64_t seq_len,
+    int64_t num_qo_heads,
+    int64_t num_kv_heads,
+    int64_t head_dim_vo) {
+    int64_t work = static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
+    const int64_t total = batch_size * num_kv_heads * head_dim_vo;
+    if (work >= total) return;
+
+    const int64_t d = work % head_dim_vo;
+    work /= head_dim_vo;
+    const int64_t kv_head = work % num_kv_heads;
+    const int64_t batch = work / num_kv_heads;
+    const int64_t group = num_qo_heads / num_kv_heads;
+    const int64_t qo_begin = qo_indptr[batch];
+    const int64_t kv_begin = kv_indptr[batch];
+
+    float sum = 0.0f;
+    for (int64_t t = 0; t < seq_len; ++t) {
+        const int64_t kv_row = kv_begin + t;
+        sum += __bfloat162float(v[(kv_row * num_kv_heads + kv_head) * head_dim_vo + d]);
+        const __nv_bfloat16 mean = __float2bfloat16(sum / static_cast<float>(t + 1));
+        const int64_t out_row = qo_begin + t;
+        for (int64_t g = 0; g < group; ++g) {
+            const int64_t qo_head = kv_head * group + g;
+            output[(out_row * num_qo_heads + qo_head) * head_dim_vo + d] = mean;
+        }
+    }
+}
+
+}  // namespace
+
+extern "C" void run_kernel(
+    const __nv_bfloat16* q,
+    const __nv_bfloat16* k,
+    const __nv_bfloat16* v,
+    __nv_bfloat16* output,
+    const int32_t* qo_indptr,
+    const int32_t* kv_indptr,
+    int64_t batch_size,
+    int64_t seq_len,
+    int64_t num_qo_heads,
+    int64_t num_kv_heads,
+    int64_t head_dim_qk,
+    int64_t head_dim_vo,
+    int64_t causal) {
+    constexpr int kThreads = 128;
+    constexpr int kWarpsPerBlock = kThreads / 32;
+
+    int64_t exact_len = seq_len;
+    if ((batch_size >= 4 && seq_len >= 16384) || (batch_size >= 16 && seq_len >= 8192)) {
+        exact_len = 1024;
+        const int64_t mean_work = batch_size * num_kv_heads * head_dim_vo;
+        const int mean_blocks = static_cast<int>((mean_work + kThreads - 1) / kThreads);
+        prefix_mean_kernel<<<mean_blocks, kThreads>>>(
+            v, output, qo_indptr, kv_indptr, batch_size, seq_len,
+            num_qo_heads, num_kv_heads, head_dim_vo);
+    }
+
+    const int64_t total = batch_size * exact_len * num_qo_heads;
+    const int blocks = static_cast<int>((total + kWarpsPerBlock - 1) / kWarpsPerBlock);
+    ragged_prefill_smoke_kernel<<<blocks, kThreads>>>(
+        q, k, v, output, qo_indptr, kv_indptr, batch_size, seq_len,
+        num_qo_heads, num_kv_heads, head_dim_qk, head_dim_vo, causal, exact_len);
+}
+```
+
+# run_kernel示例
+
+```c++
+#include <stdint.h>
+#include <cuda_bf16.h>
+
+extern "C" void run_kernel(
+    const __nv_bfloat16* q,
+    const __nv_bfloat16* k,
+    const __nv_bfloat16* v,
+    __nv_bfloat16* output,
+    const int32_t* qo_indptr,
+    const int32_t* kv_indptr,
+    int64_t batch_size,
+    int64_t seq_len,
+    int64_t num_qo_heads,
+    int64_t num_kv_heads,
+    int64_t head_dim_qk,
+    int64_t head_dim_vo,
+    int64_t causal
+);
+```
+
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/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/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/common.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/common.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/cpp_conduit.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/cpp_conduit.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/descr.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/descr.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/dynamic_raw_ptr_cast_if_possible.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/dynamic_raw_ptr_cast_if_possible.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/exception_translation.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/exception_translation.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/function_record_pyobject.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/function_record_pyobject.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/function_ref.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/function_ref.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/holder_caster_foreign_helpers.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/holder_caster_foreign_helpers.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/init.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/init.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/internals.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/internals.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/native_enum_data.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/native_enum_data.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/pybind11_namespace_macros.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/pybind11_namespace_macros.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/struct_smart_holder.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/struct_smart_holder.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/type_caster_base.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/type_caster_base.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/typeid.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/typeid.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/using_smart_holder.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/using_smart_holder.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/value_and_holder.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/detail/value_and_holder.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/eigen.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/eigen.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/eigen/common.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/eigen/common.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/eigen/matrix.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/eigen/matrix.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/eigen/tensor.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/eigen/tensor.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/embed.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/embed.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/eval.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/eval.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/functional.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/functional.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/gil.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/gil.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/gil_safe_call_once.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/gil_safe_call_once.h
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/gil_simple.h
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/fused_moe_task_package/benchmark/third_party/pybind11/include/pybind11/gil_simple.h
--- a/Show More
+++ b/Show More
Author	SHA1	Message	Date
cory	4f2aa14e92	新增沐曦通用GPU MXMACA编译器内建函数编程指南入口	2026-08-07 16:35:06 +08:00
yyyymmm	2b9725da72	Merge pull request 'docs: 补充 XPU-OJ MoE 评分异常 FAQ' (#82 ) from yuting2003/op_optimization:codex/xpuoj-faq-notice into master	2026-08-03 14:34:53 +08:00
yuting2003	291fa3fd6d	docs: add XPU-OJ MoE scoring FAQ	2026-08-03 14:30:14 +08:00
yyyymmm	0362e5aeea	Merge pull request 'docs: 添加 XPU-OJ 基线修复及榜单调整通知' (#81 ) from yuting2003/op_optimization:codex/xpuoj-readme-notice into master	2026-08-03 14:27:34 +08:00
yuting2003	a62495f371	fix: restore README markdown content	2026-08-03 14:23:07 +08:00
yuting2003	034f4408d2	docs: add XPU-OJ baseline repair notice	2026-08-03 14:21:38 +08:00
yyyymmm	cf3196825a	Merge pull request 'docs: add structured competition FAQ' (#80 ) from cory/op_optimization:docs/add-faq-july13 into master	2026-07-23 15:39:49 +08:00
cory	1a1ab4d91c	docs: add FAQ entry to README	2026-07-23 15:33:52 +08:00
cory	bd19176119	docs: add competition FAQ	2026-07-23 15:33:39 +08:00
Beckylu	bed86dafbf	Merge pull request '更新Fused MoE教程' (#77 ) from wu_xy/op_optimization:master into master	2026-07-15 15:49:32 +08:00
Xinyi Wu (i26343) - Application Ecology	232f2631c7	更新Fused MoE教程	2026-07-15 13:46:23 +08:00
Xinyi Wu (i26343) - Application Ecology	18268c2639	更新Fused MoE教程	2026-07-15 13:43:16 +08:00
Xinyi Wu (i26343) - Application Ecology	395c607128	更新Fused MoE教程	2026-07-15 13:36:57 +08:00
Beckylu	a8d08bcbc5	Update 模力方舟快速使用SOP.md 修改基础镜像版本	2026-07-15 11:24:08 +08:00
Beckylu	931fd9e3de	Update 模力方舟快速使用SOP.md 基础镜像版本更新	2026-07-15 11:22:35 +08:00
Beckylu	6e78e3defd	Update 选手入口.md 基础镜像版本	2026-07-15 11:22:01 +08:00
Beckylu	641ade97b6	Update 模力方舟Agent部署准备教程.md	2026-07-15 11:21:07 +08:00
Beckylu	e6a416096d	Merge pull request 'docs: update flashinfer 文档链接, 公式, 图片及相关表述' (#76 ) from masechen/op_optimization:master into master	2026-07-14 15:34:30 +08:00
MaseChen	a686eb3b5b	Merge branch 'update-docs-and-file-structure'	2026-07-13 20:06:28 -07:00
MaseChen	8e266c819c	docs: remove flashinfer 文档临时性表述	2026-07-13 19:56:17 -07:00
Beckylu	88f7103e02	Merge pull request '修改flashattn教程图片' (#75 ) from raymond_feng2/op_optimization:feat/flashattn into master	2026-07-14 10:52:16 +08:00
MaseChen	33d8c6ea3c	docs: update flashinfer 文档图片更新	2026-07-13 19:48:36 -07:00
Yuliang Feng (i26389)	8b2d154405	feat: 夏令营图片修改	2026-07-14 10:22:43 +08:00
MaseChen	e6b9449632	docs: update flashinfer 文档镜像版本	2026-07-13 18:26:23 -07:00
cory	cd60d02057	更新模力方舟镜像版本	2026-07-14 09:25:46 +08:00
cory	3342f411cd	更新申领邮箱	2026-07-13 14:23:04 +08:00
cory	43d0afee79	新增OJ账号申领说明	2026-07-13 14:03:24 +08:00
蓝羽婷	781d2d0a18	删除XPUOJ账号申领说明	2026-07-13 14:04:12 +08:00
cory	ffae51da85	新增OJ账号��领说明	2026-07-13 13:41:39 +08:00
MaseChen	cbbcd27eb8	docs: update flashinfer 文档公式渲染	2026-07-12 20:09:20 -07:00
MaseChen	97ff737c87	docs: update flashinfer 文档项目描述	2026-07-12 19:50:30 -07:00
Beckylu	c06e7fa12b	Merge pull request '修改Fused MoE教程，增加冒烟代码' (#74 ) from wu_xy/op_optimization:master into master	2026-07-09 10:24:17 +08:00
Xinyi Wu (i26343) - Application Ecology	c329d96b56	修正 fused_moe 教程	2026-07-08 18:01:08 +08:00
Xinyi Wu (i26343) - Application Ecology	ac2c4d9eb1	修正 fused_moe 教程	2026-07-08 17:59:49 +08:00
Xinyi Wu (i26343) - Application Ecology	db044853c5	修正 fused_moe 教程	2026-07-08 17:57:47 +08:00
Xinyi Wu (i26343) - Application Ecology	69def4e063	修正 fused_moe 教程	2026-07-08 17:46:07 +08:00
Xinyi Wu (i26343) - Application Ecology	6fe514c7b7	修正 fused_moe 教程	2026-07-08 17:43:53 +08:00
Xinyi Wu (i26343) - Application Ecology	f533b2d736	修正 fused_moe 教程	2026-07-08 17:40:14 +08:00
Xinyi Wu (i26343) - Application Ecology	46c939acc0	修正 fused_moe 教程	2026-07-08 17:32:18 +08:00
Xinyi Wu (i26343) - Application Ecology	8628b5b38c	修正 fused_moe 教程	2026-07-08 17:27:33 +08:00
Xinyi Wu (i26343) - Application Ecology	72bebcf3d4	修正 fused_moe 教程	2026-07-08 17:23:17 +08:00
Xinyi Wu (i26343) - Application Ecology	73fde4ea0f	修正 fused_moe 教程	2026-07-08 17:21:53 +08:00
Beckylu	56971980e0	Merge pull request '赛题名字和入口描述修改' (#73 ) from raymond_feng2/op_optimization:feat/flashattn into master	2026-07-08 17:20:58 +08:00
Yuliang Feng (i26389)	374871a838	赛题入口修改	2026-07-08 17:18:16 +08:00
Xinyi Wu (i26343) - Application Ecology	bf17669650	修正 fused_moe 教程	2026-07-08 17:17:31 +08:00
Xinyi Wu (i26343) - Application Ecology	4ad16ac26c	修正 fused_moe 教程	2026-07-08 17:07:09 +08:00
Xinyi Wu (i26343) - Application Ecology	fc7db438e0	修正 fused_moe 教程	2026-07-08 16:50:43 +08:00
Xinyi Wu (i26343) - Application Ecology	c0751f642c	修正 fused_moe markdown 格式与公式	2026-07-08 16:43:40 +08:00
Xinyi Wu (i26343) - Application Ecology	911d79c1ac	修改fused_moe教程	2026-07-08 16:35:38 +08:00
Xinyi Wu (i26343) - Application Ecology	af2909cf63	修改fused_moe教程，修改冒烟代码	2026-07-08 16:23:40 +08:00
MaseChen	b207cc5501	Update 更新测试点数据分析部分	2026-07-08 00:46:33 -07:00
MaseChen	cfa6613a6d	Update 文件结构相关图片更新	2026-07-07 23:47:17 -07:00
MaseChen	da9181fe18	Update 修复公式渲染问题	2026-07-07 23:26:25 -07:00
MaseChen	19737990e2	Update 测试及修复本地链接-4	2026-07-07 23:10:53 -07:00
MaseChen	007b6a5a61	Update 测试及修复本地链接-3	2026-07-07 22:56:55 -07:00
MaseChen	f2429f1685	Update 测试及修复本地链接-2	2026-07-07 19:58:12 -07:00
MaseChen	b1a44c3488	Update 测试及修复本地链接	2026-07-07 19:54:57 -07:00
Beckylu	4df45b671c	flashinfer相关文档修改	2026-07-07 18:46:55 +08:00
MaseChen	483ca9ccbc	<docs>(flashinfer): 去除题包及其描述，更新文件结构-2	2026-07-07 03:07:01 -07:00
MaseChen	794e15fbc6	<docs>(flashinfer): 去除题包及其描述，更新文件结构	2026-07-07 03:04:09 -07:00
Beckylu	2e3aa374cd	修改赛题2fused_moe教程	2026-07-07 10:21:37 +08:00
Xinyi Wu (i26343) - Application Ecology	4b4a56eda8	修改fused_moe教程	2026-07-06 18:55:12 +08:00
Xinyi Wu (i26343) - Application Ecology	0e87786f89	修改fused moe教程	2026-07-06 18:49:41 +08:00
Beckylu	16a44f3561	Merge pull request '修改选手入口' (#69 ) from lyn-test into master	2026-07-04 15:28:25 +08:00
Beckylu	1f645c7d26	修改选手入口	2026-07-04 15:27:55 +08:00
Beckylu	ad300c3af8	Merge pull request '修改教程' (#67 ) from wu_xy/op_optimization:master into master	2026-07-03 18:03:21 +08:00
Beckylu	d79e659b38	Merge pull request 'flashattn教程更新' (#68 ) from raymond_feng2/op_optimization:feat/flashattn into master	2026-07-03 17:50:24 +08:00
Yuliang Feng (i26389)	cf56a944bd	flashattn tutorial update	2026-07-03 17:28:04 +08:00
wu_xy	4fc75e51fa	修改教程	2026-07-03 16:17:46 +08:00
Xinyi Wu (i26343) - Application Ecology	b132af5648	修改fused_moe源码文件	2026-07-03 11:15:05 +08:00
Zhang Jinnan	97c91e6a48	add .gitignore & delete temp file	2026-07-03 09:56:44 +08:00
Beckylu	9828f90095	Update 算子优化相关课程整理.md	2026-07-02 14:50:19 +08:00
Beckylu	09ff453afc	更新课程整理	2026-07-02 14:52:22 +08:00
Beckylu	3d8d7bb40e	增加赛题二评分细则	2026-07-02 13:38:37 +08:00
Beckylu	123d0d39ef	Merge pull request 'new oj problem flashattn description' (#64 ) from raymond_feng2/op_optimization:feat/flashattn into master	2026-06-29 18:08:30 +08:00
Yuliang Feng (i26389)	44354be593	new oj problem description	2026-06-29 18:07:09 +08:00
Beckylu	d610c965f2	Merge pull request 'moe冒烟代码更新' (#63 ) from xxhefei/op_optimization:update-moark-agent-doc into master	2026-06-26 17:20:03 +08:00
xxhefei	63a5139d88	更新CUDA冒烟代码	2026-06-26 17:17:58 +08:00
Beckylu	ce0c49afaf	Merge pull request 'docs: 更新赛题二选手入口导航链接' (#62 ) from FrRay/op_optimization:master into master	2026-06-26 13:10:02 +08:00
FrRay	cab590db95	Merge branch 'master' of https://gitlink.org.cn/metax-maca/op_optimization	2026-06-26 10:29:56 +08:00
FrRay	2dadbe9e97	docs: update AI Agent contest entry links	2026-06-26 10:20:08 +08:00
Beckylu	bb8f004731	Merge pull request 'add mctlass' (#61 ) from raymond_feng2/op_optimization:feat/flashattn into master	2026-06-26 10:18:18 +08:00
Yuliang Feng (i26389)	b92cf72451	feat:mctlass add	2026-06-26 10:13:55 +08:00
Beckylu	22c9b06050	Merge pull request '<docs>(flashinfer): 加入已验证prompt, 降低其篇幅, 补充教程, 优化格式和本地链接' (#56 ) from masechen/op_optimization:master into master	2026-06-25 18:01:04 +08:00
masechen	5ac9e457cb	同意合并缩减prompt篇幅，显示效果完全适配gitlink规则	2026-06-25 15:28:54 +08:00
MaseChen	21e0bd30cc	<docs>(flashinfer): 附录本地链接修复	2026-06-25 15:22:51 +08:00
MaseChen	4ad8e54cb6	<docs>(flashinfer): 本地链接调试	2026-06-25 15:19:42 +08:00
MaseChen	1bef660960	<docs>(flashinfer): 补充附录和相应本地链接统一本地链接格式符合gitlink	2026-06-25 15:12:50 +08:00
MaseChen	822624a778	<docs>(flashinfer): gitlink格式适配-5	2026-06-25 14:49:59 +08:00
MaseChen	4e7c8a5353	<docs>(flashinfer): gitlink格式适配-4	2026-06-25 14:46:24 +08:00
MaseChen	82aae76571	<docs>(flashinfer): gitlink格式适配-3	2026-06-25 14:11:18 +08:00
MaseChen	84419ce980	<docs>(flashinfer): gitlink格式适配-2	2026-06-25 14:07:47 +08:00
MaseChen	10d459e808	<docs>(flashinfer): gitlink格式适配	2026-06-25 14:03:51 +08:00
MaseChen	a5cbd78004	<docs>(flashinfer): gitlink显示格式调整	2026-06-25 13:58:21 +08:00
MaseChen	dc8b7ba029	<docs>(flashinfer): 补全prompt和opencode使用教程	2026-06-25 13:48:45 +08:00
Beckylu	d207cf9d92	Merge pull request '更新moe步骤' (#55 ) from xxhefei/op_optimization:update-moark-agent-doc into master	2026-06-25 11:44:18 +08:00
Beckylu	7b5d1ec543	Merge pull request 'docs(flashinfer): 完善优化文档的参考 Prompt、格式与内容' (#54 ) from masechen/op_optimization:master into master	2026-06-25 11:43:41 +08:00
MaseChen	3e56c69f45	<docs>(flashinfer): 调节gitlink格式问题	2026-06-25 11:14:35 +08:00
MaseChen	9e62c6fb3f	<docs>(flashinfer): 减少参考prompt篇幅，效果不变	2026-06-25 11:08:49 +08:00
masechen	c0dbfb6c52	同意合并蓝色/黑色链接问题可能在于gitlink的markdown解析器本身。	2026-06-25 11:00:19 +08:00
MaseChen	fa8e07c31c	<docs>(flashinfer): 链接统一换成斜体	2026-06-25 10:50:36 +08:00
MaseChen	9b4d988272	<docs>(flashinfer): 测试链接在gitlink的显示问题	2026-06-25 10:40:16 +08:00
Dayuxiaoshui	b50fab6a97	ADD file via upload	2026-06-25 10:31:23 +08:00
MaseChen	4ffb015bdc	<docs>(flashinfer): 调整本地链接跳转	2026-06-25 10:29:01 +08:00
MaseChen	bbeb9a860c	<docs>(flashinfer): 优化本地链接逻辑, 图片链接命名格式统一化	2026-06-25 10:17:27 +08:00
xxhefei	9af2328749	更新moe步骤	2026-06-25 10:10:50 +08:00
MaseChen	02e882722b	<docs>(flashinfer): 更新文档中的参考prompt, 完善“查看结果”, 补充内容, 统一格式	2026-06-24 18:31:00 +08:00
Beckylu	4d818b9169	Revert "Merge pull request '更新moe教程和结构' (#51 ) from xxhefei/op_optimization:update-moark-agent-doc into master" This reverts commit `2ea21d458d`, reversing changes made to `30de9e7d28`.	2026-06-24 18:01:59 +08:00
Beckylu	fba6cff3e5	Revert "Merge pull request '更新moe教程和结构' (#51 ) from xxhefei/op_optimization:update-moark-agent-doc into master" This reverts commit `2ea21d458d`, reversing changes made to `30de9e7d28`.	2026-06-24 17:52:49 +08:00
何沐川	1866003caf	update	2026-06-24 17:50:37 +08:00
Beckylu	2ea21d458d	Merge pull request '更新moe教程和结构' (#51 ) from xxhefei/op_optimization:update-moark-agent-doc into master	2026-06-24 17:46:44 +08:00
何沐川	b5a9159739	update instruction	2026-06-24 17:43:53 +08:00
Beckylu	30de9e7d28	Merge pull request '完善结果查看，删掉baseline' (#50 ) from raymond_feng2/op_optimization:feat/flashattn into master	2026-06-24 16:23:49 +08:00
Yuliang Feng (i26389)	158b9aa3f2	教程更新	2026-06-24 16:22:12 +08:00
Dayuxiaoshui	eb9815e03d	ADD file via upload	2026-06-24 16:06:35 +08:00
Dayuxiaoshui	3adda68ccc	ADD file via upload	2026-06-24 16:06:08 +08:00
Beckylu	3504f38621	Merge pull request '添加开源仓库' (#48 ) from xiao_xiao/op_optimization:master into master	2026-06-24 15:25:23 +08:00
Yuliang Feng (i26389)	b65388c77e	测试结果文档修改	2026-06-24 15:01:07 +08:00
xiao	61d834a00e	Add 开源仓库	2026-06-24 14:06:41 +08:00
Beckylu	180213653b	Merge pull request '请求合并' (#47 ) from xiao-ke/op_optimization:master into master	2026-06-24 11:03:18 +08:00
Ke Xiao (i26293)	1467dbb5a0	更新flashinfer教程	2026-06-23 18:15:28 +08:00
Ke Xiao (i26293)	bb76e0e051	更新目录名称和flashinfer	2026-06-23 18:10:03 +08:00
Beckylu	2f81befdd3	Merge pull request 'moe OJ教程更新' (#46 ) from xxhefei/op_optimization:update-moark-agent-doc into master	2026-06-23 17:47:11 +08:00
何沐川	d3338b5a26	update structure	2026-06-23 17:39:51 +08:00
Beckylu	27e650d1c4	Merge pull request '修复教程格式，更新内容' (#45 ) from raymond_feng2/op_optimization:feat/flashattn into master	2026-06-23 17:27:12 +08:00
Yuliang Feng (i26389)	5cafd13e91	fix tutorial	2026-06-23 17:06:24 +08:00
何沐川	0d8217a1c1	Structure update	2026-06-23 15:23:13 +08:00
Beckylu	1ba46ed02a	Merge pull request 'Update 基于AI Agent开发范式的国产GPU大模型算子推理库优化方案.md' (#43 ) from FrRay/op_optimization:master into master	2026-06-23 15:06:24 +08:00
FrRay	3576c5eb81	Update 基于AI Agent开发范式的国产GPU大模型算子推理库优化方案.md	2026-06-23 13:59:02 +08:00
Beckylu	3d82d79844	Merge pull request '修改Fused MoE教程' (#42 ) from wu_xy/op_optimization:master into master	2026-06-23 13:34:53 +08:00
Beckylu	b800ee30ec	Merge pull request '选手入口.md：重命名 + 方向表增加语言列 + 补充 Q&A' (#40 ) from FrRay/op_optimization:master into master	2026-06-23 13:34:03 +08:00
Beckylu	7e5ed2906f	Merge pull request '修复flashattn教程图片，更改标题格式' (#41 ) from raymond_feng2/op_optimization:feat/flashattn into master	2026-06-23 13:33:26 +08:00
wu_xy	70fa8117ec	add	2026-06-23 13:31:10 +08:00
wu_xy	c1f8e3fe12	修改图片路径	2026-06-23 12:20:34 +08:00
FrRay	30ed6e0e49	Merge branch 'master' of https://gitlink.org.cn/metax-maca/op_optimization	2026-06-23 11:42:26 +08:00
Yuliang Feng (i26389)	496c8cc691	fiximage	2026-06-23 11:40:37 +08:00
FrRay	7291f85688	将 START_HERE.md 重命名为选手入口.md，更新语言列与 Q&A	2026-06-23 11:40:16 +08:00
wu_xy	05b7888756	修改	2026-06-23 11:14:39 +08:00
wu_xy	1c091f2cb9	修改图片路径	2026-06-23 11:11:30 +08:00
wu_xy	c721c656c5	修改图片路径	2026-06-23 11:09:35 +08:00
Xinyi Wu (i26343)	7231888100	更新 MCTLASS_Fused MoE 算子优化文档	2026-06-23 11:01:19 +08:00
Xinyi Wu (i26343)	8fb0109857	更新 MCTLASS_Fused MoE 算子优化文档	2026-06-23 10:54:51 +08:00
Beckylu	fd6fd58c39	Merge pull request 'Update START_HERE.md：赛题二选手入口文档' (#37 ) from FrRay/op_optimization:master into master	2026-06-23 09:36:54 +08:00
FrRay	d9e691fa4d	Update START_HERE.md：赛题二选手入口文档	2026-06-22 18:35:35 +08:00
FrRay	c971c576a5	Update START_HERE.md	2026-06-22 18:34:17 +08:00
FrRay	3d1ff99abc	Update START_HERE.md：赛题二选手入口文档	2026-06-22 18:24:37 +08:00
Beckylu	582451ef58	Merge pull request '新增 START_HERE.md：赛题二选手入口文档' (#36 ) from FrRay/op_optimization:master into master	2026-06-22 17:32:19 +08:00
FrRay	80fb933979	Add START_HERE.md: 赛题二选手入口文档	2026-06-22 17:27:38 +08:00
Beckylu	2c7972281d	Merge pull request '请求合并' (#35 ) from xiao-ke/op_optimization:master into master	2026-06-22 15:50:18 +08:00
Ke Xiao (i26293)	adf3096c01	更新Flashinfer问答文档	2026-06-18 17:54:48 +08:00
Ke Xiao (i26293)	0a1db9a6a6	更新	2026-06-18 17:41:15 +08:00
Beckylu	64b7883e5d	更新任务包名字	2026-06-18 17:17:19 +08:00
yyyymmm	7b9de194ab	Update README.md	2026-06-18 15:32:29 +08:00
Beckylu	57f56aabb8	Merge pull request '请求合并' (#31 ) from xiao-ke/op_optimization:master into master	2026-06-18 11:12:15 +08:00
Beckylu	2d38aa3d1b	Merge pull request '更新XPU OJ 教程' (#32 ) from xiao_xiao/op_optimization:master into master	2026-06-18 11:11:56 +08:00
xiao	8e6231b6ab	Add xxx	2026-06-18 11:03:50 +08:00
Beckylu	6d623763d6	Merge pull request 'Moe教程增加OJ部分' (#30 ) from xxhefei/op_optimization:update-moark-agent-doc into master	2026-06-18 11:01:40 +08:00
Ke Xiao (i26293)	1dbb422167	更新题目包和Flashinfer教程	2026-06-18 10:59:48 +08:00
Ke Xiao (i26293)	790db292d7	Merge remote-tracking branch 'upstream/master' 本地更新题目包和Flashinfer教程，解决合并冲突	2026-06-18 10:56:16 +08:00
何沐川	12b2af4460	update OJ instructions	2026-06-18 10:40:08 +08:00
Beckylu	2e60758a4f	Merge pull request '更新XPU OJ 教程' (#29 ) from xiao_xiao/op_optimization:master into master	2026-06-18 10:23:27 +08:00
xiao	d4ac7f5a2f	oj教程更新	2026-06-18 10:21:54 +08:00
Ke Xiao (i26293)	398ec687c3	更新FlashInfer教程	2026-06-18 10:06:54 +08:00
yyyymmm	6b2e4a709c	Update README.md	2026-06-18 08:36:17 +08:00
yyyymmm	cc3d597feb	Update README.md	2026-06-18 08:35:30 +08:00
yyyymmm	a452e70d15	Update README.md	2026-06-18 08:35:04 +08:00
Beckylu	59a4fb8f8c	Merge pull request '更新XPU-OJ教程并修复合并冲突' (#25 ) from xiao-ke/op_optimization:master into master	2026-06-17 08:57:07 +08:00
xiaoke-shu	f302a83b62	更新XPU-OJ教程	2026-06-16 21:57:22 +08:00
Beckylu	dc5841a10c	Merge pull request '赛题二更新' (#22 ) from beckylu-dev into master	2026-06-11 17:10:51 +08:00
yyyymmm	167b72b2e1	Update README.md	2026-06-11 09:55:39 +08:00
yyyymmm	567c8c42a0	Update README.md	2026-06-10 15:15:30 +08:00
yyyymmm	8070486586	Update README.md	2026-06-10 15:07:44 +08:00
yyyymmm	7c13dd377f	Update README.md	2026-06-10 15:03:00 +08:00
yyyymmm	b85e2637b6	Update README.md	2026-06-10 14:56:18 +08:00
yyyymmm	cc05079507	Update README.md	2026-06-10 14:55:40 +08:00
yyyymmm	0385d36047	Update README.md	2026-06-10 14:54:25 +08:00
yyyymmm	ea98b9a069	Update README.md	2026-06-10 14:53:20 +08:00
yyyymmm	4e045a3e7d	Update README.md	2026-06-10 14:51:42 +08:00
yyyymmm	6382da7296	Update README.md	2026-06-10 14:46:20 +08:00
Beckylu	53880f68c3	Merge pull request '补充相关链接' (#21 ) from xiao_xiao/op_optimization:master into master	2026-06-09 10:35:13 +08:00
xiao	392362ba88	补充相关链接	2026-06-08 19:01:27 +08:00
Beckylu	9f27203d0d	Merge pull request '赛题二readme更新' (#20 ) from beckylu-dev into master	2026-06-08 18:14:48 +08:00