新增沐曦通用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
247 changed files with 5720 additions and 6915 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 算力成本！

 ## 两大重磅赛题 直击推理成本核心痛点
@ -29,7 +46,7 @@

 ### 赛题二：基于 AI Agent 开发范式的国产 GPU 大模型推理算子库优化

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

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

@ -50,7 +67,7 @@
 - [模力方舟 Agent 部署准备教程](基于AI%20Agent开发范式的国产GPU大模型推理算子库优化/模力方舟Agent部署准备教程.md)
 - [赛题二说明及资料参考](基于AI%20Agent开发范式的国产GPU大模型推理算子库优化/赛题说明.md)

-###**两个赛题统一使用模力方舟上的镜像PyTorch-Agent / 2.8.0 / Python 3.12 / maca 3.7.2.1**
+### **两个赛题统一使用模力方舟上的镜像PyTorch-Agent / 2.8.0 / Python 3.12 / maca 3.7.1.5**

 ## 参赛对象

--- a/Agent开发范式的国产GPU大模型推理算子库优化/FlashAttention关键算子迁移与优化.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/FlashAttention关键算子迁移与优化.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/FlashInfer
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/FlashInfer
--- a/Agent开发范式的国产GPU大模型推理算子库优化/Fused
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/Fused
--- a/Agent开发范式的国产GPU大模型推理算子库优化/MCTLASS_Fused
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/MCTLASS_Fused
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/baseline结果实例/benchmark_kvcache_headdim128.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/baseline结果实例/benchmark_kvcache_headdim128.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/benchmark/baseline结果实例/benchmark_kvcache_headdim160.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/baseline结果实例/benchmark_kvcache_headdim160.csv
@ -1,49 +0,0 @@
-batch_size,seq_len_kv,heads,headdim,time_ms,bandwidth_GB_s
-1,512,8,160,0.0580,45.25
-2,512,8,160,0.0611,85.84
-4,512,8,160,0.0656,159.96
-8,512,8,160,0.0699,300.40
-16,512,8,160,0.1321,317.92
-32,512,8,160,0.2002,419.52
-64,512,8,160,0.3383,496.43
-128,512,8,160,0.6669,503.61
-1,1024,8,160,0.1129,46.45
-2,1024,8,160,0.1190,88.18
-4,1024,8,160,0.1224,171.49
-8,1024,8,160,0.1287,326.07
-16,1024,8,160,0.2479,338.49
-32,1024,8,160,0.3767,445.54
-64,1024,8,160,0.6419,523.01
-128,1024,8,160,1.2804,524.37
-1,2048,8,160,0.2270,46.20
-2,2048,8,160,0.2299,91.22
-4,2048,8,160,0.2349,178.63
-8,2048,8,160,0.2447,342.96
-16,2048,8,160,0.4773,351.60
-32,2048,8,160,0.7279,461.07
-64,2048,8,160,1.2559,534.49
-128,2048,8,160,2.5613,524.15
-1,4096,8,160,0.4460,47.02
-2,4096,8,160,0.4513,92.94
-4,4096,8,160,0.4593,182.64
-8,4096,8,160,0.4813,348.64
-16,4096,8,160,0.9363,358.43
-32,4096,8,160,1.4552,461.21
-64,4096,8,160,2.5615,524.05
-128,4096,8,160,5.1420,522.11
-1,8192,8,160,0.8847,47.41
-2,8192,8,160,0.8944,93.80
-4,8192,8,160,0.9094,184.51
-8,8192,8,160,0.9625,348.64
-16,8192,8,160,1.8550,361.80
-32,8192,8,160,2.9567,453.97
-64,8192,8,160,5.1398,522.30
-128,8192,8,160,10.2972,521.41
-1,16384,8,160,1.7608,47.64
-2,16384,8,160,1.7786,94.33
-4,16384,8,160,1.8143,184.95
-8,16384,8,160,1.9317,347.42
-16,16384,8,160,3.7301,359.83
-32,16384,8,160,5.9216,453.33
-64,16384,8,160,10.2668,522.94
-128,16384,8,160,20.6062,521.09
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/baseline结果实例/benchmark_kvcache_headdim192.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/baseline结果实例/benchmark_kvcache_headdim192.csv
@ -1,49 +0,0 @@
-batch_size,seq_len_kv,heads,headdim,time_ms,bandwidth_GB_s
-1,512,8,192,0.0458,68.82
-2,512,8,192,0.0515,122.32
-4,512,8,192,0.0574,219.28
-8,512,8,192,0.0607,414.80
-16,512,8,192,0.1147,439.27
-32,512,8,192,0.1763,571.40
-64,512,8,192,0.2978,676.79
-128,512,8,192,0.5874,686.11
-1,1024,8,192,0.0946,66.55
-2,1024,8,192,0.1033,121.85
-4,1024,8,192,0.1073,234.66
-8,1024,8,192,0.1131,445.23
-16,1024,8,192,0.2165,465.24
-32,1024,8,192,0.3347,601.80
-64,1024,8,192,0.5701,706.63
-128,1024,8,192,1.1302,712.88
-1,2048,8,192,0.1943,64.79
-2,2048,8,192,0.1992,126.38
-4,2048,8,192,0.2059,244.52
-8,2048,8,192,0.2174,463.13
-16,2048,8,192,0.4202,479.24
-32,2048,8,192,0.6503,619.36
-64,2048,8,192,1.1158,721.93
-128,2048,8,192,2.2250,724.05
-1,4096,8,192,0.3834,65.65
-2,4096,8,192,0.3904,128.95
-4,4096,8,192,0.4043,249.04
-8,4096,8,192,0.4267,471.92
-16,4096,8,192,0.8271,486.90
-32,4096,8,192,1.2840,627.28
-64,4096,8,192,2.2148,727.29
-128,4096,8,192,4.3819,735.21
-1,8192,8,192,0.7566,66.52
-2,8192,8,192,0.7712,130.54
-4,8192,8,192,0.7974,252.49
-8,8192,8,192,0.8433,477.47
-16,8192,8,192,1.6433,490.09
-32,8192,8,192,2.5573,629.84
-64,8192,8,192,4.3785,735.73
-128,8192,8,192,8.7303,737.99
-1,16384,8,192,1.5068,66.81
-2,16384,8,192,1.5350,131.16
-4,16384,8,192,1.5868,253.76
-8,16384,8,192,1.6778,479.99
-16,16384,8,192,3.2750,491.81
-32,16384,8,192,5.0659,635.88
-64,16384,8,192,8.7435,736.85
-128,16384,8,192,17.5040,736.13
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/baseline结果实例/benchmark_kvcache_headdim224.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/baseline结果实例/benchmark_kvcache_headdim224.csv
@ -1,49 +0,0 @@
-batch_size,seq_len_kv,heads,headdim,time_ms,bandwidth_GB_s
-1,512,8,224,0.1254,29.29
-2,512,8,224,0.1412,52.05
-4,512,8,224,0.1497,98.13
-8,512,8,224,0.1533,191.70
-16,512,8,224,0.1913,307.17
-32,512,8,224,0.3292,357.08
-64,512,8,224,0.5187,453.28
-128,512,8,224,0.9522,493.84
-1,1024,8,224,0.2727,26.93
-2,1024,8,224,0.2836,51.78
-4,1024,8,224,0.2890,101.63
-8,1024,8,224,0.2959,198.55
-16,1024,8,224,0.3696,317.93
-32,1024,8,224,0.6408,366.75
-64,1024,8,224,1.0081,466.21
-128,1024,8,224,1.8548,506.78
-1,2048,8,224,0.5515,26.63
-2,2048,8,224,0.5575,52.67
-4,2048,8,224,0.5666,103.65
-8,2048,8,224,0.5803,202.42
-16,2048,8,224,0.7250,324.05
-32,2048,8,224,1.2593,373.14
-64,2048,8,224,1.9890,472.48
-128,2048,8,224,3.6905,509.28
-1,4096,8,224,1.0939,26.84
-2,4096,8,224,1.1044,53.18
-4,4096,8,224,1.1219,104.69
-8,4096,8,224,1.1500,204.26
-16,4096,8,224,1.4390,326.48
-32,4096,8,224,2.4992,375.97
-64,4096,8,224,4.0082,468.86
-128,4096,8,224,7.3372,512.26
-1,8192,8,224,2.1775,26.97
-2,8192,8,224,2.1989,53.41
-4,8192,8,224,2.2338,105.15
-8,8192,8,224,2.3268,201.90
-16,8192,8,224,2.8806,326.18
-32,8192,8,224,5.0187,374.43
-64,8192,8,224,8.0323,467.90
-128,8192,8,224,14.6300,513.78
-1,16384,8,224,4.3360,27.09
-2,16384,8,224,4.3820,53.60
-4,16384,8,224,4.5006,104.38
-8,16384,8,224,4.6987,199.96
-16,16384,8,224,5.7361,327.59
-32,16384,8,224,10.1291,371.03
-64,16384,8,224,16.0745,467.60
-128,16384,8,224,OOM,OOM
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/baseline结果实例/benchmark_kvcache_headdim256.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/baseline结果实例/benchmark_kvcache_headdim256.csv
@ -1,49 +0,0 @@
-batch_size,seq_len_kv,heads,headdim,time_ms,bandwidth_GB_s
-1,512,8,256,0.0877,47.89
-2,512,8,256,0.0921,91.17
-4,512,8,256,0.0940,178.74
-8,512,8,256,0.0964,348.52
-16,512,8,256,0.1450,463.27
-32,512,8,256,0.2250,597.21
-64,512,8,256,0.3609,744.43
-128,512,8,256,0.6932,775.25
-1,1024,8,256,0.1747,48.04
-2,1024,8,256,0.1762,95.27
-4,1024,8,256,0.1784,188.22
-8,1024,8,256,0.1817,369.53
-16,1024,8,256,0.2796,480.25
-32,1024,8,256,0.4339,619.00
-64,1024,8,256,0.6960,771.73
-128,1024,8,256,1.3439,799.36
-1,2048,8,256,0.3410,49.21
-2,2048,8,256,0.3439,97.60
-4,2048,8,256,0.3469,193.52
-8,2048,8,256,0.3533,379.94
-16,2048,8,256,0.5461,491.67
-32,2048,8,256,0.8493,632.28
-64,2048,8,256,1.3667,785.82
-128,2048,8,256,2.6465,811.64
-1,4096,8,256,0.6742,49.77
-2,4096,8,256,0.6777,99.03
-4,4096,8,256,0.6836,196.36
-8,4096,8,256,0.6950,386.31
-16,4096,8,256,1.0803,497.02
-32,4096,8,256,1.6794,639.44
-64,4096,8,256,2.7101,792.50
-128,4096,8,256,5.2543,817.52
-1,8192,8,256,1.3375,50.18
-2,8192,8,256,1.3448,99.81
-4,8192,8,256,1.3564,197.91
-8,8192,8,256,1.3799,389.08
-16,8192,8,256,2.1465,500.25
-32,8192,8,256,3.3342,644.12
-64,8192,8,256,5.3983,795.67
-128,8192,8,256,10.4691,820.55
-1,16384,8,256,2.6697,50.28
-2,16384,8,256,2.6817,100.10
-4,16384,8,256,2.7049,198.49
-8,16384,8,256,2.7533,390.00
-16,16384,8,256,4.2789,501.89
-32,16384,8,256,6.6476,646.11
-64,16384,8,256,10.7723,797.43
-128,16384,8,256,OOM,OOM
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/baseline结果实例/benchmark_kvcache_headdim32.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/baseline结果实例/benchmark_kvcache_headdim32.csv
@ -1,49 +0,0 @@
-batch_size,seq_len_kv,heads,headdim,time_ms,bandwidth_GB_s
-1,512,8,32,0.0257,20.45
-2,512,8,32,0.0256,41.02
-4,512,8,32,0.0258,81.28
-8,512,8,32,0.0265,158.45
-16,512,8,32,0.0396,212.30
-32,512,8,32,0.0516,325.43
-64,512,8,32,0.0721,465.83
-128,512,8,32,0.1270,529.03
-1,1024,8,32,0.0461,22.75
-2,1024,8,32,0.0465,45.15
-4,1024,8,32,0.0477,88.04
-8,1024,8,32,0.0548,153.23
-16,1024,8,32,0.0734,228.71
-32,1024,8,32,0.0958,350.42
-64,1024,8,32,0.1334,503.15
-128,1024,8,32,0.2381,564.04
-1,2048,8,32,0.0872,24.06
-2,2048,8,32,0.0904,46.42
-4,2048,8,32,0.1028,81.59
-8,2048,8,32,0.1067,157.25
-16,2048,8,32,0.1428,235.10
-32,2048,8,32,0.1818,369.13
-64,2048,8,32,0.2554,525.57
-128,2048,8,32,0.4622,580.86
-1,4096,8,32,0.1730,24.25
-2,4096,8,32,0.1955,42.91
-4,4096,8,32,0.2020,83.05
-8,4096,8,32,0.2140,156.83
-16,4096,8,32,0.2777,241.65
-32,4096,8,32,0.3542,378.99
-64,4096,8,32,0.4990,538.05
-128,4096,8,32,0.9099,590.13
-1,8192,8,32,0.3820,21.96
-2,8192,8,32,0.3913,42.88
-4,8192,8,32,0.4127,81.31
-8,8192,8,32,0.4224,158.88
-16,8192,8,32,0.5490,244.51
-32,8192,8,32,0.6960,385.70
-64,8192,8,32,0.9870,543.98
-128,8192,8,32,1.8100,593.25
-1,16384,8,32,0.7655,21.92
-2,16384,8,32,0.8067,41.59
-4,16384,8,32,0.8228,81.56
-8,16384,8,32,0.8397,159.85
-16,16384,8,32,1.0910,246.04
-32,16384,8,32,1.3824,388.37
-64,16384,8,32,1.9663,546.08
-128,16384,8,32,3.6107,594.78
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/baseline结果实例/benchmark_kvcache_headdim512.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/baseline结果实例/benchmark_kvcache_headdim512.csv
@ -1,49 +0,0 @@
-batch_size,seq_len_kv,heads,headdim,time_ms,bandwidth_GB_s
-1,512,8,512,0.3588,23.40
-2,512,8,512,0.3651,46.00
-4,512,8,512,0.3736,89.89
-8,512,8,512,0.3856,174.22
-16,512,8,512,0.7472,179.80
-32,512,8,512,1.1447,234.72
-64,512,8,512,1.9549,274.89
-128,512,8,512,3.8962,275.85
-1,1024,8,512,0.7261,23.12
-2,1024,8,512,0.7354,45.65
-4,1024,8,512,0.7496,89.57
-8,1024,8,512,0.7746,173.35
-16,1024,8,512,1.5049,178.46
-32,1024,8,512,2.3111,232.42
-64,1024,8,512,3.9538,271.70
-128,1024,8,512,7.8811,272.62
-1,2048,8,512,1.4636,22.93
-2,2048,8,512,1.4826,45.27
-4,2048,8,512,1.5109,88.86
-8,2048,8,512,1.5549,172.68
-16,2048,8,512,3.0237,177.60
-32,2048,8,512,4.6439,231.27
-64,2048,8,512,7.9560,269.99
-128,2048,8,512,15.8741,270.63
-1,4096,8,512,2.9312,22.90
-2,4096,8,512,2.9675,45.24
-4,4096,8,512,3.0243,88.77
-8,4096,8,512,3.1127,172.50
-16,4096,8,512,6.0753,176.76
-32,4096,8,512,9.3182,230.49
-64,4096,8,512,15.9642,269.07
-128,4096,8,512,31.8313,269.89
-1,8192,8,512,5.8843,22.81
-2,8192,8,512,5.9344,45.24
-4,8192,8,512,6.0465,88.80
-8,8192,8,512,6.2334,172.27
-16,8192,8,512,12.1594,176.62
-32,8192,8,512,18.6826,229.90
-64,8192,8,512,32.0055,268.41
-128,8192,8,512,OOM,OOM
-1,16384,8,512,11.8153,22.72
-2,16384,8,512,11.9237,45.03
-4,16384,8,512,12.1671,88.25
-8,16384,8,512,12.4948,171.88
-16,16384,8,512,24.3414,176.45
-32,16384,8,512,37.3907,229.74
-64,16384,8,512,OOM,OOM
-128,16384,8,512,OOM,OOM
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/baseline结果实例/benchmark_kvcache_headdim64.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/baseline结果实例/benchmark_kvcache_headdim64.csv
@ -1,49 +0,0 @@
-batch_size,seq_len_kv,heads,headdim,time_ms,bandwidth_GB_s
-1,512,8,64,0.0404,25.99
-2,512,8,64,0.0399,52.60
-4,512,8,64,0.0413,101.69
-8,512,8,64,0.0482,174.25
-16,512,8,64,0.0540,310.86
-32,512,8,64,0.0629,533.75
-64,512,8,64,0.0833,806.14
-128,512,8,64,0.1104,1216.59
-1,1024,8,64,0.0747,28.08
-2,1024,8,64,0.0766,54.77
-4,1024,8,64,0.0891,94.17
-8,1024,8,64,0.0918,182.94
-16,1024,8,64,0.1044,321.41
-32,1024,8,64,0.1179,569.43
-64,1024,8,64,0.1566,857.28
-128,1024,8,64,0.2078,1292.17
-1,2048,8,64,0.1455,28.84
-2,2048,8,64,0.1684,49.82
-4,2048,8,64,0.1730,97.01
-8,2048,8,64,0.1850,181.39
-16,2048,8,64,0.2009,334.18
-32,2048,8,64,0.2268,592.01
-64,2048,8,64,0.3002,894.44
-128,2048,8,64,0.4027,1333.64
-1,4096,8,64,0.3265,25.69
-2,4096,8,64,0.3322,50.51
-4,4096,8,64,0.3522,95.27
-8,4096,8,64,0.3632,184.79
-16,4096,8,64,0.3942,340.56
-32,4096,8,64,0.4456,602.47
-64,4096,8,64,0.5927,905.94
-128,4096,8,64,0.7938,1352.87
-1,8192,8,64,0.6508,25.78
-2,8192,8,64,0.6879,48.78
-4,8192,8,64,0.7008,95.77
-8,8192,8,64,0.7199,186.44
-16,8192,8,64,0.7786,344.79
-32,8192,8,64,0.8798,610.25
-64,8192,8,64,1.1745,914.30
-128,8192,8,64,1.5728,1365.50
-1,16384,8,64,1.3524,24.81
-2,16384,8,64,1.3698,48.99
-4,16384,8,64,1.3923,96.40
-8,16384,8,64,1.4267,188.16
-16,16384,8,64,1.5451,347.47
-32,16384,8,64,1.7622,609.32
-64,16384,8,64,2.3392,918.09
-128,16384,8,64,3.1332,1370.84
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/baseline结果实例/benchmark_kvcache_headdim96.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/benchmark/baseline结果实例/benchmark_kvcache_headdim96.csv
@ -1,49 +0,0 @@
-batch_size,seq_len_kv,heads,headdim,time_ms,bandwidth_GB_s
-1,512,8,96,0.0407,38.67
-2,512,8,96,0.0398,79.02
-4,512,8,96,0.0431,146.08
-8,512,8,96,0.0495,254.61
-16,512,8,96,0.0698,360.64
-32,512,8,96,0.1117,450.87
-64,512,8,96,0.1780,566.16
-128,512,8,96,0.3329,605.28
-1,1024,8,96,0.0732,43.01
-2,1024,8,96,0.0794,79.29
-4,1024,8,96,0.0871,144.54
-8,1024,8,96,0.0934,269.52
-16,1024,8,96,0.1297,388.14
-32,1024,8,96,0.2114,476.36
-64,1024,8,96,0.3379,596.08
-128,1024,8,96,0.6327,636.68
-1,2048,8,96,0.1505,41.80
-2,2048,8,96,0.1619,77.76
-4,2048,8,96,0.1713,146.94
-8,2048,8,96,0.1780,282.84
-16,2048,8,96,0.2492,404.09
-32,2048,8,96,0.4088,492.55
-64,2048,8,96,0.6575,612.55
-128,2048,8,96,1.2457,646.61
-1,4096,8,96,0.3099,40.61
-2,4096,8,96,0.3259,77.23
-4,4096,8,96,0.3346,150.42
-8,4096,8,96,0.3467,290.41
-16,4096,8,96,0.4888,411.94
-32,4096,8,96,0.8055,499.94
-64,4096,8,96,1.3209,609.72
-128,4096,8,96,2.4810,649.25
-1,8192,8,96,0.6343,39.68
-2,8192,8,96,0.6437,78.20
-4,8192,8,96,0.6601,152.50
-8,8192,8,96,0.6826,294.97
-16,8192,8,96,0.9688,415.64
-32,8192,8,96,1.6057,501.55
-64,8192,8,96,2.6527,607.19
-128,8192,8,96,4.9464,651.27
-1,16384,8,96,1.2581,40.01
-2,16384,8,96,1.2812,78.57
-4,16384,8,96,1.3112,153.55
-8,16384,8,96,1.3653,294.92
-16,16384,8,96,1.9351,416.16
-32,16384,8,96,3.2277,499.01
-64,16384,8,96,5.3192,605.60
-128,16384,8,96,9.8747,652.44
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/starter/smoke_test-CUDAMaca.txt
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/starter/smoke_test-CUDAMaca.txt
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/starter/smoke_test-Tilelang.txt
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/starter/smoke_test-Tilelang.txt
@ -0,0 +1,187 @@
+"""FlashAttention KV Cache Decode in TileLang."""
+
+import tilelang
+import tilelang.language as T
+from tilelang import jit
+
+NUM_SPLITS = 4
+real_kernel = None
+
+@jit(
+    pass_configs={
+        tilelang.PassConfigKey.TL_ENABLE_FAST_MATH: False,
+    },
+)
+def build_kernel(
+    batch_size,
+    num_heads,
+    num_heads_k,
+    headdim,
+    page_block_size,
+    num_blocks,
+    causal,
+):
+    blocks_per_batch = num_blocks // batch_size
+    assert blocks_per_batch % NUM_SPLITS == 0, (
+        f"blocks_per_batch={blocks_per_batch} must be divisible by NUM_SPLITS={NUM_SPLITS}"
+    )
+    blocks_per_split = blocks_per_batch // NUM_SPLITS
+
+    BLOCK_M = 1
+    BLOCK_N = page_block_size
+    scale = (1.0 / headdim) ** 0.5 * 1.44269504  # log2(e)
+    dtype = "bfloat16"
+    accum_dtype = "float32"
+
+    # Use a large-negative-finite sentinel instead of -inf to avoid
+    # (-inf) - (-inf) = NaN when an entire split is masked out.
+    NEG_INF_SAFE = -1e30
+
+    @T.prim_func
+    def kernel(
+        Q: T.Tensor([batch_size, 1, num_heads, headdim], dtype),
+        K: T.Tensor([num_blocks, page_block_size, num_heads_k, headdim], dtype),
+        V: T.Tensor([num_blocks, page_block_size, num_heads_k, headdim], dtype),
+        Output: T.Tensor([batch_size, 1, num_heads, headdim], dtype),
+        cache_seqlens: T.Tensor([batch_size], "int32"),
+        block_table: T.Tensor([batch_size, blocks_per_batch], "int32"),
+    ):
+        # float32 workspace — avoids BF16StorageLegalize var-remap bug
+        glse = T.alloc_global([batch_size, num_heads, NUM_SPLITS], accum_dtype)
+        Output_partial = T.alloc_global(
+            [batch_size, 1, num_heads, NUM_SPLITS, headdim], accum_dtype
+        )
+
+        # ============= Stage 1: split kernel =============
+        with T.Kernel(NUM_SPLITS, num_heads, batch_size, threads=128) as (bs, bh, bz):
+            Q_shared = T.alloc_shared([BLOCK_M, headdim], dtype)
+            K_shared = T.alloc_shared([BLOCK_N, headdim], dtype)
+            V_shared = T.alloc_shared([BLOCK_N, headdim], dtype)
+            acc_s = T.alloc_fragment([BLOCK_M, BLOCK_N], accum_dtype)
+            acc_o = T.alloc_fragment([BLOCK_M, headdim], accum_dtype)
+            scores_max = T.alloc_fragment([BLOCK_M], accum_dtype)
+            scores_max_prev = T.alloc_fragment([BLOCK_M], accum_dtype)
+            scores_scale = T.alloc_fragment([BLOCK_M], accum_dtype)
+            scores_sum = T.alloc_fragment([BLOCK_M], accum_dtype)
+            logsum = T.alloc_fragment([BLOCK_M], accum_dtype)
+
+            T.copy(Q[bz, 0, bh, :], Q_shared)
+
+            kv_seqlen = cache_seqlens[bz]
+            split_k_start = bs * blocks_per_split
+
+            T.fill(acc_o, 0)
+            T.fill(logsum, 0)
+            # KEY FIX: use -1e30 instead of -inf to avoid (-inf)-(-inf)=NaN
+            T.fill(scores_max, NEG_INF_SAFE)
+
+            for k in T.Pipelined(blocks_per_split, num_stages=2):
+                global_k = split_k_start + k
+                physical_block = block_table[bz, global_k]
+                tok_offset = global_k * page_block_size
+
+                # ----- Q @ K^T (hand-written, M=1, masked) -----
+                T.copy(K[physical_block, 0:BLOCK_N, bh, :], K_shared)
+                T.fill(acc_s, 0)
+                for j in T.Parallel(BLOCK_N):
+                    if tok_offset + j < kv_seqlen:
+                        for d in T.serial(headdim):
+                            acc_s[0, j] = acc_s[0, j] + Q_shared[0, d] * K_shared[j, d]
+                    else:
+                        acc_s[0, j] = -T.infinity(accum_dtype)
+
+                # ----- online softmax -----
+                T.copy(scores_max, scores_max_prev)
+                # KEY FIX: use -1e30 instead of -inf here too
+                T.fill(scores_max, NEG_INF_SAFE)
+                T.reduce_max(acc_s, scores_max, dim=1, clear=False)
+                scores_max[0] = T.max(scores_max[0], scores_max_prev[0])
+                # (prev - cur) is now (finite - finite) = 0 when both are sentinel,
+                # never (-inf - (-inf)) = NaN
+                scores_scale[0] = T.exp2((scores_max_prev[0] - scores_max[0]) * scale)
+                for j in T.Parallel(BLOCK_N):
+                    acc_s[0, j] = T.exp2((acc_s[0, j] - scores_max[0]) * scale)
+                T.reduce_sum(acc_s, scores_sum, dim=1)
+                logsum[0] = logsum[0] * scores_scale[0] + scores_sum[0]
+                for d in T.Parallel(headdim):
+                    acc_o[0, d] = acc_o[0, d] * scores_scale[0]
+
+                # ----- P @ V (hand-written, fp32 accum) -----
+                T.copy(V[physical_block, 0:BLOCK_N, bh, :], V_shared)
+                for d in T.Parallel(headdim):
+                    for j in T.serial(BLOCK_N):
+                        acc_o[0, d] = acc_o[0, d] + acc_s[0, j] * V_shared[j, d]
+
+            # ----- final normalise & write partial state -----
+            # KEY FIX: add epsilon to avoid 0/0 = NaN when split is all-masked
+            safe_logsum = logsum[0] + 1e-30
+            for d in T.Parallel(headdim):
+                acc_o[0, d] = acc_o[0, d] / safe_logsum
+
+            lse_local = T.alloc_fragment([1], accum_dtype)
+            lse_local[0] = T.log2(safe_logsum) + scores_max[0] * scale
+            glse[bz, bh, bs] = lse_local[0]
+
+            for d in T.Parallel(headdim):
+                Output_partial[bz, 0, bh, bs, d] = acc_o[0, d]
+
+        # ============= Stage 2: combine kernel =============
+        with T.Kernel(num_heads, batch_size, threads=128) as (bh, bz):
+            lse_local = T.alloc_fragment([NUM_SPLITS], accum_dtype)
+            for s in T.serial(NUM_SPLITS):
+                lse_local[s] = glse[bz, bh, s]
+
+            lse_max = T.alloc_fragment([1], accum_dtype)
+            lse_max[0] = -T.infinity(accum_dtype)
+            for s in T.serial(NUM_SPLITS):
+                lse_max[0] = T.max(lse_max[0], lse_local[s])
+
+            lse_logsum = T.alloc_fragment([1], accum_dtype)
+            lse_logsum[0] = 0
+            for s in T.serial(NUM_SPLITS):
+                lse_logsum[0] = lse_logsum[0] + T.exp2(lse_local[s] - lse_max[0])
+            lse_logsum[0] = T.log2(lse_logsum[0]) + lse_max[0]
+
+            o_accum = T.alloc_fragment([headdim], accum_dtype)
+            T.fill(o_accum, 0)
+            for s in T.serial(NUM_SPLITS):
+                s_scale = T.exp2(lse_local[s] - lse_logsum[0])
+                for d in T.Parallel(headdim):
+                    o_accum[d] = o_accum[d] + Output_partial[bz, 0, bh, s, d] * s_scale
+
+            for d in T.Parallel(headdim):
+                Output[bz, 0, bh, d] = T.Cast(dtype, o_accum[d])
+
+    return kernel
+
+
+def run_kernel(
+    q,
+    k_cache_paged,
+    v_cache_paged,
+    output,
+    cache_seqlens,
+    block_table,
+    batch_size,
+    seqlen_k,
+    seqlen_q,
+    num_heads,
+    num_heads_k,
+    headdim,
+    page_block_size,
+    num_blocks,
+    causal,
+):
+    global real_kernel
+
+    B = int(batch_size)
+    H = int(num_heads)
+    HK = int(num_heads_k)
+    D = int(headdim)
+    PBS = int(page_block_size)
+    NB = int(num_blocks)
+
+    if real_kernel is None:
+        real_kernel = build_kernel(B, H, HK, D, PBS, NB, int(causal))
+
+    real_kernel(q, k_cache_paged, v_cache_paged, output, cache_seqlens, block_table)
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/starter/smoke_test-Triton.txt
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/starter/smoke_test-Triton.txt
@ -0,0 +1,106 @@
+import triton
+import triton.language as tl
+import torch
+
+@triton.jit
+def slow_decode_kernel(
+    q_ptr,
+    k_cache_ptr,
+    v_cache_ptr,
+    output_ptr,
+    cache_seqlens_ptr,
+    block_table_ptr,
+    num_heads: tl.constexpr,
+    num_heads_k: tl.constexpr,
+    headdim: tl.constexpr,
+    page_block_size: tl.constexpr,
+    blocks_per_batch,
+):
+    # 维度索引
+    pid_b = tl.program_id(0)  # Batch index
+    pid_h = tl.program_id(1)  # Head index
+
+    # GQA Support: 映射 Query Head 到 KV Head
+    kv_head = pid_h * num_heads_k // num_heads
+
+    # 加载实际的 KV 序列长度
+    seq_len = tl.load(cache_seqlens_ptr + pid_b).to(tl.int32)
+
+    # 维度偏移量 [0, 1, ..., headdim-1]
+    offs_d = tl.arange(0, headdim)
+
+    # Online Softmax 累加器
+    acc = tl.zeros([headdim], dtype=tl.float32)
+    l_i = 0.0
+    m_i = float('-inf')
+    scale = 1.0 / tl.sqrt(float(headdim))
+
+    # === 性能瓶颈：串行遍历整个序列 ===
+    # 不使用 Block 并行，而是用单个 Block 串行循环处理所有 Token
+    t = 0
+    while t < seq_len:
+        # 性能瓶颈：每次循环都重新加载 Q，增加显存压力
+        q = tl.load(q_ptr + pid_b * num_heads * headdim + pid_h * headdim + offs_d).to(tl.float32)
+        q = q * scale
+
+        # Paged KV 映射逻辑
+        page_idx = t // page_block_size
+        page_off = t % page_block_size
+
+        # 查表获取物理 Block 索引
+        # blocks_per_batch 是计算出来的步长
+        phys_block = tl.load(block_table_ptr + pid_b * blocks_per_batch + page_idx)
+
+        # 计算 K 和 V 的物理地址
+        # Layout: (num_blocks, page_block_size, num_heads_k, headdim)
+        kv_base = phys_block * page_block_size * num_heads_k * headdim + \
+                  page_off * num_heads_k * headdim + \
+                  kv_head * headdim
+        
+        # 加载 K 和 V 向量
+        k = tl.load(k_cache_ptr + kv_base + offs_d).to(tl.float32)
+        v = tl.load(v_cache_ptr + kv_base + offs_d).to(tl.float32)
+
+        # Attention 计算
+        s = tl.sum(q * k)  # 点积
+        
+        # Online Softmax 更新
+        m_new = tl.maximum(m_i, s)
+        p = tl.exp(s - m_new)
+        alpha = tl.exp(m_i - m_new)
+        
+        acc = acc * alpha + p * v
+        l_i = l_i * alpha + p
+        m_i = m_new
+
+        t += 1
+
+    # 写回结果
+    # 这里没有处理 l_i 为 0 的边界情况，但测试数据 seq_len 通常很大
+    out = acc / l_i
+    tl.store(output_ptr + pid_b * num_heads * headdim + pid_h * headdim + offs_d, out)
+
+def run_kernel(
+    q, k_cache_paged, v_cache_paged, output,
+    cache_seqlens, block_table,
+    batch_size, seqlen_k, seqlen_q, num_heads, num_heads_k, headdim,
+    page_block_size, num_blocks, causal,
+):
+    # 计算每个 batch 对应的 block_table 行宽
+    blocks_per_batch = num_blocks // batch_size
+    
+    # 启动配置：每个 Head 一个 Block
+    # 总 Block 数 = batch_size * num_heads (最大 128个)，并行度极低
+    grid = (batch_size, num_heads)
+    
+    slow_decode_kernel[grid](
+        q, k_cache_paged, v_cache_paged, output,
+        cache_seqlens, block_table,
+        num_heads=num_heads,
+        num_heads_k=num_heads_k,
+        headdim=headdim,
+        page_block_size=page_block_size,
+        blocks_per_batch=blocks_per_batch,
+        num_warps=1,    # 性能瓶颈：仅使用 1 个 warp，限制计算吞吐
+        num_stages=1,   # 性能瓶颈：禁用流水线并行
+    )
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/meta.json
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/meta.json
@ -1,16 +0,0 @@
-{
-  "id": 197,
-  "displayId": 20005,
-  "type": "Traditional",
-  "isPublic": false,
-  "locales": [
-    "zh_CN"
-  ],
-  "samples": [
-    {
-      "inputData": "1\n",
-      "outputData": ""
-    }
-  ],
-  "problemTagIds": []
-}
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/testcase_config.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/testcase_config.py
@ -1,298 +0,0 @@
-from __future__ import annotations
-
-
-HEAD_DIMS = [128]
-BATCH_SIZES = [1, 4, 16]
-SEQ_LENS_KV = [1024, 4096, 8192, 16384]
-SEQ_LEN_Q = 1
-NUM_HEADS = 8
-NUM_HEADS_K = 8
-PAGE_BLOCK_SIZE = 16
-CAUSAL = 0
-
-
-def _build_cases():
-    cases = []
-    for headdim in HEAD_DIMS:
-        for seqlen_k in SEQ_LENS_KV:
-            for batch_size in BATCH_SIZES:
-                cases.append(
-                    (
-                        batch_size,
-                        seqlen_k,
-                        SEQ_LEN_Q,
-                        NUM_HEADS,
-                        NUM_HEADS_K,
-                        headdim,
-                        PAGE_BLOCK_SIZE,
-                        CAUSAL,
-                    )
-                )
-    return cases
-
-
-TESTCASES = _build_cases()
-
-
-def getNumOfTestcases() -> int:
-    return len(TESTCASES)
-
-
-try:
-    from pathlib import Path
-    from typing import List, Tuple, Union
-    import math
-    import sys
-
-    import torch
-
-    KernelArg = Union[torch.Tensor, int, float]
-    CURRENT_CASE = None
-
-    def _ensure_flashattn_importable():
-        try:
-            from flash_attn.flash_attn_interface import flash_attn_with_kvcache  # noqa: F401
-
-            return
-        except ImportError:
-            pass
-
-        here = Path(__file__).resolve()
-        for parent in here.parents:
-            candidate = parent / "flashattn"
-            if (candidate / "flash_attn").is_dir():
-                sys.path.insert(0, str(candidate))
-                return
-
-    def _get_testcase_index() -> int:
-        try:
-            raw = input().strip()
-        except EOFError:
-            return 0
-        if raw == "":
-            return 0
-        try:
-            testcase_id = int(raw.split()[0])
-        except ValueError:
-            return 0
-        if 1 <= testcase_id <= len(TESTCASES):
-            return testcase_id - 1
-        if 0 <= testcase_id < len(TESTCASES):
-            return testcase_id
-        return 0
-
-    def _compute_reps(batch_size: int, seq_len: int, head_dim: int, base_reps: int = 100) -> int:
-        workload = batch_size * seq_len * head_dim
-        if workload < 1e5:
-            return base_reps
-        if workload < 1e6:
-            return base_reps // 2
-        if workload < 1e7:
-            return base_reps // 4
-        if workload < 1e8:
-            return base_reps // 8
-        if workload < 1e9:
-            return base_reps // 16
-        return base_reps // 32
-
-    def _get_num_blocks(batch_size: int, seqlen_k: int, page_block_size: int) -> int:
-        num_blocks = math.ceil(seqlen_k / page_block_size) * batch_size * 3
-        return max(1024, num_blocks)
-
-    def getTestCaseSize() -> Tuple[List[Tuple[int, ...]], Tuple[int, int]]:
-        testcase_id = _get_testcase_index()
-        global CURRENT_CASE
-        (
-            batch_size,
-            seqlen_k,
-            seqlen_q,
-            num_heads,
-            num_heads_k,
-            headdim,
-            page_block_size,
-            causal,
-        ) = TESTCASES[testcase_id]
-        num_blocks = _get_num_blocks(batch_size, seqlen_k, page_block_size)
-        blocks_per_batch = num_blocks // batch_size
-        CURRENT_CASE = (
-            batch_size,
-            seqlen_k,
-            seqlen_q,
-            num_heads,
-            num_heads_k,
-            headdim,
-            page_block_size,
-            num_blocks,
-            causal,
-            20260720 + testcase_id,
-        )
-        warmup = 3
-        iters = max(1, _compute_reps(batch_size, seqlen_k, headdim))
-        return [
-            (batch_size, seqlen_q, num_heads, headdim),
-            (num_blocks, page_block_size, num_heads_k, headdim),
-            (num_blocks, page_block_size, num_heads_k, headdim),
-            (batch_size, seqlen_q, num_heads, headdim),
-            (batch_size,),
-            (batch_size, blocks_per_batch),
-            (), (), (), (), (), (), (), (), (),
-        ], (warmup, iters)
-
-    def genTestCase(testcase_sizes, device: str = "cuda") -> List[KernelArg]:
-        del testcase_sizes
-        (
-            batch_size,
-            seqlen_k,
-            seqlen_q,
-            num_heads,
-            num_heads_k,
-            headdim,
-            page_block_size,
-            num_blocks,
-            causal,
-            seed,
-        ) = CURRENT_CASE
-        gen = torch.Generator(device=device)
-        gen.manual_seed(seed)
-        dtype = torch.bfloat16
-        blocks_per_batch = num_blocks // batch_size
-        q = torch.randn(
-            batch_size,
-            seqlen_q,
-            num_heads,
-            headdim,
-            dtype=dtype,
-            device=device,
-            generator=gen,
-        ).contiguous()
-        k_cache_paged = torch.randn(
-            num_blocks,
-            page_block_size,
-            num_heads_k,
-            headdim,
-            dtype=dtype,
-            device=device,
-            generator=gen,
-        ).contiguous()
-        v_cache_paged = torch.randn(
-            num_blocks,
-            page_block_size,
-            num_heads_k,
-            headdim,
-            dtype=dtype,
-            device=device,
-            generator=gen,
-        ).contiguous()
-        output = torch.empty(
-            batch_size,
-            seqlen_q,
-            num_heads,
-            headdim,
-            dtype=dtype,
-            device=device,
-        )
-        cache_seqlens = torch.full((batch_size,), seqlen_k, dtype=torch.int32, device=device)
-        block_table = torch.randperm(num_blocks, dtype=torch.int32, device=device, generator=gen).reshape(
-            batch_size,
-            blocks_per_batch,
-        )
-        return [
-            q,
-            k_cache_paged,
-            v_cache_paged,
-            output,
-            cache_seqlens,
-            block_table,
-            batch_size,
-            seqlen_k,
-            seqlen_q,
-            num_heads,
-            num_heads_k,
-            headdim,
-            page_block_size,
-            num_blocks,
-            causal,
-        ]
-
-    def baseline(
-        q,
-        k_cache_paged,
-        v_cache_paged,
-        output,
-        cache_seqlens,
-        block_table,
-        batch_size,
-        seqlen_k,
-        seqlen_q,
-        num_heads,
-        num_heads_k,
-        headdim,
-        page_block_size,
-        num_blocks,
-        causal,
-    ):
-        _ensure_flashattn_importable()
-        from flash_attn.flash_attn_interface import flash_attn_with_kvcache
-
-        out = flash_attn_with_kvcache(
-            q,
-            k_cache_paged,
-            v_cache_paged,
-            None,
-            None,
-            cache_seqlens=cache_seqlens,
-            cache_batch_idx=None,
-            block_table=block_table,
-            causal=bool(causal),
-            window_size=(-1, -1),
-            rotary_interleaved=False,
-            alibi_slopes=None,
-            num_splits=1,
-        )
-        output.copy_(out)
-        return [
-            q,
-            k_cache_paged,
-            v_cache_paged,
-            output,
-            cache_seqlens,
-            block_table,
-            batch_size,
-            seqlen_k,
-            seqlen_q,
-            num_heads,
-            num_heads_k,
-            headdim,
-            page_block_size,
-            num_blocks,
-            causal,
-        ]
-
-    def check(
-        testcase_sizes,
-        original_input_tensors,
-        target_kernel_input_tensors,
-        baseline_input_tensors,
-        rtol=1e-2,
-        atol=1e-2,
-    ) -> bool:
-        del testcase_sizes, original_input_tensors
-        output_t = target_kernel_input_tensors[3]
-        output_ref = baseline_input_tensors[3]
-        if output_t.shape != output_ref.shape:
-            print(f"[FAIL] shape mismatch: target {output_t.shape}, ref {output_ref.shape}", file=sys.stderr)
-            return False
-        if output_t.dtype != output_ref.dtype:
-            print(f"[FAIL] dtype mismatch: target {output_t.dtype}, ref {output_ref.dtype}", file=sys.stderr)
-            return False
-        if not torch.allclose(output_t.float(), output_ref.float(), rtol=rtol, atol=atol):
-            diff = (output_t.float() - output_ref.float()).abs()
-            print(
-                f"[FAIL] allclose failed: max_abs_diff={float(diff.max().item()):.6f}, "
-                f"mean_abs_diff={float(diff.mean().item()):.6f} (rtol={rtol}, atol={atol})",
-                file=sys.stderr,
-            )
-            return False
-        return True
-except Exception:
-    pass
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/00_棰樼洰鎻忚堪.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/00_棰樼洰鎻忚堪.md
@ -1,28 +0,0 @@
---
-sectionTitle: "题目描述"
-type: "Text"
---
-你需要实现 FlashAttention paged KV cache decode 的 CUDA C++ 前向算子。
-
-本题输入采用 `flash_attn_with_kvcache` 在 `flashattn/benchmarks/benchmark_kvcache.py` 中使用的 paged KV cache 配置。每个 batch 只有 1 个 query token，KV cache 长度为 `seqlen_k`，K/V cache 按 page 存储。
-
-评测程序会调用你提交代码中的 `run_kernel` 函数。你需要根据 `cache_seqlens` 和 `block_table` 读取 paged KV cache，并将结果写入 `output`。
-
-baseline 使用 benchmark 中的 FlashAttention Python API：
-
-```python
-out = flash_attn_with_kvcache(
-    q, k_cache_paged, v_cache_paged, None, None,
-    cache_seqlens=cache_seqlens,
-    cache_batch_idx=None,
-    block_table=block_table,
-    causal=False,
-    window_size=(-1, -1),
-    rotary_interleaved=False,
-    alibi_slopes=None,
-    num_splits=1,
-)
-output.copy_(out)
-```
-
-如何提交代码详见[评测指南](/d/2)。
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/01_鎺ュ彛绾﹀畾.cuda.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/01_鎺ュ彛绾﹀畾.cuda.md
@ -1,43 +0,0 @@
---
-sectionTitle: "接口约定"
-type: "codeSample"
-lang: "cuda"
---
-你必须在提交的 CUDA 源码中提供如下 **C 符号**，函数名、参数类型、顺序必须完全一致，并使用 `extern "C"` 防止 name mangling：
-
-```cpp
-#include <stdint.h>
-#include <cuda_bf16.h>
-
-extern "C" void run_kernel(
-    const __nv_bfloat16* q,
-    const __nv_bfloat16* k_cache_paged,
-    const __nv_bfloat16* v_cache_paged,
-    __nv_bfloat16* output,
-    const int32_t* cache_seqlens,
-    const int32_t* block_table,
-    int64_t batch_size,
-    int64_t seqlen_k,
-    int64_t seqlen_q,
-    int64_t num_heads,
-    int64_t num_heads_k,
-    int64_t headdim,
-    int64_t page_block_size,
-    int64_t num_blocks,
-    int64_t causal
-);
-```
-
-### 参数说明
-
-* `q`：decode query tensor，shape `(batch_size, seqlen_q, num_heads, headdim)`，连续 `bf16`
-* `k_cache_paged`：paged key cache，shape `(num_blocks, page_block_size, num_heads_k, headdim)`，连续 `bf16`
-* `v_cache_paged`：paged value cache，shape `(num_blocks, page_block_size, num_heads_k, headdim)`，连续 `bf16`
-* `output`：输出缓冲区，shape `(batch_size, seqlen_q, num_heads, headdim)`，连续 `bf16`
-* `cache_seqlens`：每个 batch 的 KV 长度，shape `(batch_size)`，连续 `int32`
-* `block_table`：每个 batch 的 page 映射表，shape `(batch_size, num_blocks / batch_size)`，连续 `int32`
-* `seqlen_q`：query 长度，评测中固定为 `1`
-* `page_block_size`：page size，评测中固定为 `16`
-* `causal`：是否启用 causal mask，评测中固定为 `0`
-
-`run_kernel` 内部需要自行计算合适的 launch 配置并启动 CUDA kernel。为保证计时准确，不建议在 `run_kernel` 内部做 `cudaDeviceSynchronize()` 或显式同步。
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/01_鎺ュ彛绾﹀畾.tilelang.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/01_鎺ュ彛绾﹀畾.tilelang.md
@ -1,57 +0,0 @@
---
-sectionTitle: "接口约定"
-type: "codeSample"
-lang: "tilelang"
---
-你必须在提交的 Python 代码中提供 `run_kernel` 函数，函数名、参数顺序、类型必须完全一致：
-
-```python
-import tilelang
-import tilelang.language as T
-from tilelang import jit
-
-real_kernel = None
-
-@jit
-def build_kernel(*args):
-    @T.prim_func
-    def kernel(*args):
-        ...
-    return kernel
-
-def run_kernel(
-    q,  # Tensor[bf16], shape (batch_size, seqlen_q, num_heads, headdim)
-    k_cache_paged,  # Tensor[bf16], shape (num_blocks, page_block_size, num_heads_k, headdim)
-    v_cache_paged,  # Tensor[bf16], shape (num_blocks, page_block_size, num_heads_k, headdim)
-    output,  # Tensor[bf16], shape (batch_size, seqlen_q, num_heads, headdim)
-    cache_seqlens,  # Tensor[int32], shape (batch_size)
-    block_table,  # Tensor[int32], shape (batch_size, num_blocks / batch_size)
-    batch_size,  # int64
-    seqlen_k,  # int64
-    seqlen_q,  # int64
-    num_heads,  # int64
-    num_heads_k,  # int64
-    headdim,  # int64
-    page_block_size,  # int64
-    num_blocks,  # int64
-    causal,  # int64
-):
-    global real_kernel
-    if real_kernel is None:
-        real_kernel = build_kernel(...)
-    real_kernel(q, k_cache_paged, v_cache_paged, output,
-                cache_seqlens, block_table,
-                batch_size, seqlen_k, seqlen_q, num_heads,
-                num_heads_k, headdim, page_block_size, num_blocks, causal)
-```
-
-### 参数说明
-
-* `q`：decode query tensor，连续 `bfloat16`
-* `k_cache_paged/v_cache_paged`：paged KV cache，连续 `bfloat16`
-* `output`：输出缓冲区，连续 `bfloat16`，需要写入结果
-* `cache_seqlens/block_table`：paged KV metadata，连续 `int32`
-* `page_block_size`：评测中固定为 `16`
-* `causal`：评测中固定为 `0`
-
-`run_kernel` 内部需要自行计算合适的 grid/block，并 launch 你实现的 TileLang kernel。
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/01_鎺ュ彛绾﹀畾.triton.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/01_鎺ュ彛绾﹀畾.triton.md
@ -1,45 +0,0 @@
---
-sectionTitle: "接口约定"
-type: "codeSample"
-lang: "triton"
---
-你必须在提交的 Python 代码中提供 `run_kernel` 函数，函数名、参数顺序、类型必须完全一致：
-
-```python
-import triton
-import triton.language as tl
-
-@triton.jit
-def your_kernel(...):
-    ...
-
-def run_kernel(
-    q,  # Tensor[bf16], shape (batch_size, seqlen_q, num_heads, headdim)
-    k_cache_paged,  # Tensor[bf16], shape (num_blocks, page_block_size, num_heads_k, headdim)
-    v_cache_paged,  # Tensor[bf16], shape (num_blocks, page_block_size, num_heads_k, headdim)
-    output,  # Tensor[bf16], shape (batch_size, seqlen_q, num_heads, headdim)
-    cache_seqlens,  # Tensor[int32], shape (batch_size)
-    block_table,  # Tensor[int32], shape (batch_size, num_blocks / batch_size)
-    batch_size,  # int64
-    seqlen_k,  # int64
-    seqlen_q,  # int64
-    num_heads,  # int64
-    num_heads_k,  # int64
-    headdim,  # int64
-    page_block_size,  # int64
-    num_blocks,  # int64
-    causal,  # int64
-):
-    ...
-```
-
-### 参数说明
-
-* `q`：decode query tensor，连续 `bfloat16`
-* `k_cache_paged/v_cache_paged`：paged KV cache，连续 `bfloat16`
-* `output`：输出缓冲区，连续 `bfloat16`，需要写入结果
-* `cache_seqlens/block_table`：paged KV metadata，连续 `int32`
-* `page_block_size`：评测中固定为 `16`
-* `causal`：评测中固定为 `0`
-
-`run_kernel` 内部需要自行计算合适的 grid/block，并 launch 你实现的 Triton kernel。
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/02_杈撳叆鏍煎紡.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/02_杈撳叆鏍煎紡.md
@ -1,9 +0,0 @@
---
-sectionTitle: "输入格式"
-type: "Text"
---
-本题输入由评测程序在 GPU 上构造，并按接口约定中的顺序传入 `run_kernel`。
-
-`q/k_cache_paged/v_cache_paged/output` 均为连续 `torch.bfloat16` CUDA tensor，`cache_seqlens/block_table` 均为连续 `torch.int32` CUDA tensor。
-
-KV cache layout 固定为 `flash_attn_with_kvcache` 的 paged cache 布局：`(num_blocks, page_block_size, num_heads_k, headdim)`。
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/03_杈撳嚭鏍煎紡.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/03_杈撳嚭鏍煎紡.md
@ -1,5 +0,0 @@
---
-sectionTitle: "输出格式"
-type: "Text"
---
-输出写入 `output`，shape 为 `(batch_size, 1, num_heads, headdim)`，类型为 `bfloat16`。
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/04_鏍蜂緥.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/04_鏍蜂緥.md
@ -1,12 +0,0 @@
---
-sectionTitle: "样例"
-type: "Text"
---
-若 `batch_size = 1`、`seqlen_k = 512`、`page_block_size = 16`，则每个序列需要访问 `32` 个有效 page：
-
-```text
-cache_seqlens = [512]
-block_table.shape = (1, num_blocks)
-```
-
-第 `t` 个 KV token 位于 `block_table[0, t / 16]` 指向的物理 page 中，page 内偏移为 `t % 16`。
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/_title.txt
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashattn_task_package/xpuoj_problem/zh_CN/_title.txt
@ -1 +0,0 @@
-FlashAttention KV Cache Decode
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/BatchDecodeWithPagedKVCacheWrapper_20260525_093143.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/BatchDecodeWithPagedKVCacheWrapper_20260525_093143.csv
@ -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/benchmark/BatchMLAPagedAttentionWrapper_20260525_074200.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/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/benchmark/BatchPrefillWithPagedKVCacheWrapper_20260526_150417.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/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/benchmark/BatchPrefillWithRaggedKVCacheWrapper_20260526_145454.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/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/benchmark/__MACOSX/._BatchDecodeWithPagedKVCacheWrapper_20260525_093143.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._BatchDecodeWithPagedKVCacheWrapper_20260525_093143.csv
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._BatchMLAPagedAttentionWrapper_20260525_074200.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._BatchMLAPagedAttentionWrapper_20260525_074200.csv
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._BatchPrefillWithPagedKVCacheWrapper_20260526_150417.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._BatchPrefillWithPagedKVCacheWrapper_20260526_150417.csv
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._BatchPrefillWithRaggedKVCacheWrapper_20260526_145454.csv
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._BatchPrefillWithRaggedKVCacheWrapper_20260526_145454.csv
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._bench_batch_decode.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._bench_batch_decode.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._bench_batch_mla.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._bench_batch_mla.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._bench_batch_prefill_paged.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._bench_batch_prefill_paged.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._bench_batch_prefill_ragged.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._bench_batch_prefill_ragged.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._bench_common.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/benchmark/__MACOSX/._bench_common.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/._problem_20001
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/._problem_20001
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/._problem_20002
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/._problem_20002
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/._problem_20003
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/._problem_20003
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/._problem_20004
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/._problem_20004
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/._problem_20005
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/._problem_20005
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/._meta.json
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/._meta.json
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/._testcase_config.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/._testcase_config.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/._zh_CN
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/._zh_CN
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/zh_CN/._00_棰樼洰鎻忚堪.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/zh_CN/._00_棰樼洰鎻忚堪.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/zh_CN/._01_鎺ュ彛绾﹀畾.cuda.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/zh_CN/._01_鎺ュ彛绾﹀畾.cuda.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/zh_CN/._02_杈撳叆鏍煎紡.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/zh_CN/._02_杈撳叆鏍煎紡.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/zh_CN/._03_杈撳嚭鏍煎紡.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/zh_CN/._03_杈撳嚭鏍煎紡.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/zh_CN/._04_鏍蜂緥.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/zh_CN/._04_鏍蜂緥.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/zh_CN/.__title.txt
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20001/zh_CN/.__title.txt
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/._meta.json
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/._meta.json
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/._testcase_config.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/._testcase_config.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/._zh_CN
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/._zh_CN
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/zh_CN/._00_棰樼洰鎻忚堪.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/zh_CN/._00_棰樼洰鎻忚堪.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/zh_CN/._01_鎺ュ彛绾﹀畾.cuda.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/zh_CN/._01_鎺ュ彛绾﹀畾.cuda.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/zh_CN/._02_杈撳叆鏍煎紡.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/zh_CN/._02_杈撳叆鏍煎紡.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/zh_CN/._03_杈撳嚭鏍煎紡.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/zh_CN/._03_杈撳嚭鏍煎紡.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/zh_CN/._04_鏍蜂緥.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/zh_CN/._04_鏍蜂緥.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/zh_CN/.__title.txt
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20002/zh_CN/.__title.txt
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/._meta.json
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/._meta.json
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/._testcase_config.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/._testcase_config.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/._zh_CN
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/._zh_CN
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/zh_CN/._00_棰樼洰鎻忚堪.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/zh_CN/._00_棰樼洰鎻忚堪.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/zh_CN/._01_鎺ュ彛绾﹀畾.cuda.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/zh_CN/._01_鎺ュ彛绾﹀畾.cuda.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/zh_CN/._02_杈撳叆鏍煎紡.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/zh_CN/._02_杈撳叆鏍煎紡.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/zh_CN/._03_杈撳嚭鏍煎紡.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/zh_CN/._03_杈撳嚭鏍煎紡.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/zh_CN/._04_鏍蜂緥.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/zh_CN/._04_鏍蜂緥.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/zh_CN/.__title.txt
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20003/zh_CN/.__title.txt
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/._meta.json
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/._meta.json
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/._testcase_config.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/._testcase_config.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/._zh_CN
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/._zh_CN
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/zh_CN/._00_棰樼洰鎻忚堪.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/zh_CN/._00_棰樼洰鎻忚堪.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/zh_CN/._01_鎺ュ彛绾﹀畾.cuda.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/zh_CN/._01_鎺ュ彛绾﹀畾.cuda.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/zh_CN/._02_杈撳叆鏍煎紡.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/zh_CN/._02_杈撳叆鏍煎紡.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/zh_CN/._03_杈撳嚭鏍煎紡.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/zh_CN/._03_杈撳嚭鏍煎紡.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/zh_CN/._04_鏍蜂緥.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/zh_CN/._04_鏍蜂緥.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/zh_CN/.__title.txt
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20004/zh_CN/.__title.txt
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/._meta.json
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/._meta.json
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/._testcase_config.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/._testcase_config.py
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/._zh_CN
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/._zh_CN
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/zh_CN/._00_棰樼洰鎻忚堪.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/zh_CN/._00_棰樼洰鎻忚堪.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/zh_CN/._01_鎺ュ彛绾﹀畾.cuda.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/zh_CN/._01_鎺ュ彛绾﹀畾.cuda.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/zh_CN/._02_杈撳叆鏍煎紡.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/zh_CN/._02_杈撳叆鏍煎紡.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/zh_CN/._03_杈撳嚭鏍煎紡.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/zh_CN/._03_杈撳嚭鏍煎紡.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/zh_CN/._04_鏍蜂緥.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/zh_CN/._04_鏍蜂緥.md
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/zh_CN/.__title.txt
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/__MACOSX/problem_20005/zh_CN/.__title.txt
--- 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_20001/meta.json
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/problem_20001/meta.json
@ -1,16 +0,0 @@
-{
-  "id": 193,
-  "displayId": 20001,
-  "type": "Traditional",
-  "isPublic": false,
-  "locales": [
-    "zh_CN"
-  ],
-  "samples": [
-    {
-      "inputData": "1\n",
-      "outputData": ""
-    }
-  ],
-  "problemTagIds": []
-}
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/problem_20001/testcase_config.py
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/problem_20001/testcase_config.py
@ -1,307 +0,0 @@
-from __future__ import annotations
-
-
-HEAD_DIM_CONFIGS = [(128, 128)]
-BATCH_SIZES = [1, 4, 16]
-SEQ_LENS = [1024, 4096, 8192, 16384]
-NUM_QO_HEADS = 32
-NUM_KV_HEADS = 4
-CAUSAL = 1
-
-
-def _build_cases():
-    cases = []
-    for head_dim_qk, head_dim_vo in HEAD_DIM_CONFIGS:
-        for batch_size in BATCH_SIZES:
-            for seq_len in SEQ_LENS:
-                cases.append(
-                    (
-                        batch_size,
-                        seq_len,
-                        NUM_QO_HEADS,
-                        NUM_KV_HEADS,
-                        head_dim_qk,
-                        head_dim_vo,
-                        CAUSAL,
-                    )
-                )
-    return cases
-
-
-TESTCASES = _build_cases()
-
-
-def getNumOfTestcases() -> int:
-    return len(TESTCASES)
-
-
-try:
-    from pathlib import Path
-    from typing import List, Tuple, Union
-    import sys
-
-    import torch
-
-    KernelArg = Union[torch.Tensor, int, float]
-    CURRENT_CASE = None
-
-    def _ensure_flashinfer_importable():
-        try:
-            import flashinfer  # noqa: F401
-
-            return
-        except ImportError:
-            pass
-
-        here = Path(__file__).resolve()
-        for parent in here.parents:
-            candidate = parent / "McFlashInfer"
-            if (candidate / "flashinfer").is_dir():
-                sys.path.insert(0, str(candidate))
-                return
-
-    def _get_testcase_index() -> int:
-        try:
-            raw = input().strip()
-        except EOFError:
-            return 0
-        if raw == "":
-            return 0
-        try:
-            testcase_id = int(raw.split()[0])
-        except ValueError:
-            return 0
-        if 1 <= testcase_id <= len(TESTCASES):
-            return testcase_id - 1
-        if 0 <= testcase_id < len(TESTCASES):
-            return testcase_id
-        return 0
-
-    def _compute_reps(batch_size: int, seq_len: int, head_dim: int, base_reps: int = 100) -> int:
-        workload = batch_size * seq_len * head_dim
-        if workload < 1e5:
-            return base_reps
-        if workload < 1e6:
-            return base_reps // 2
-        if workload < 1e7:
-            return base_reps // 4
-        if workload < 1e8:
-            return base_reps // 8
-        if workload < 1e9:
-            return base_reps // 16
-        return base_reps // 32
-
-    def getTestCaseSize() -> Tuple[List[Tuple[int, ...]], Tuple[int, int]]:
-        testcase_id = _get_testcase_index()
-        global CURRENT_CASE
-        batch_size, seq_len, num_qo_heads, num_kv_heads, head_dim_qk, head_dim_vo, causal = TESTCASES[testcase_id]
-        CURRENT_CASE = (
-            batch_size,
-            seq_len,
-            num_qo_heads,
-            num_kv_heads,
-            head_dim_qk,
-            head_dim_vo,
-            causal,
-            20260610 + testcase_id,
-        )
-        qo_len = batch_size * seq_len
-        kv_len = batch_size * seq_len
-        warmup = 3
-        iters = max(1, _compute_reps(batch_size, seq_len, head_dim_qk + head_dim_vo))
-        return [
-            (qo_len, num_qo_heads, head_dim_qk),
-            (kv_len, num_kv_heads, head_dim_qk),
-            (kv_len, num_kv_heads, head_dim_vo),
-            (qo_len, num_qo_heads, head_dim_vo),
-            (batch_size + 1,),
-            (batch_size + 1,),
-            (), (), (), (), (), (), (),
-        ], (warmup, iters)
-
-    def genTestCase(testcase_sizes, device: str = "cuda") -> List[KernelArg]:
-        del testcase_sizes
-        batch_size, seq_len, num_qo_heads, num_kv_heads, head_dim_qk, head_dim_vo, causal, seed = CURRENT_CASE
-        gen = torch.Generator(device=device)
-        gen.manual_seed(seed)
-        dtype = torch.bfloat16
-        qo_len = batch_size * seq_len
-        q = torch.rand(
-            qo_len,
-            num_qo_heads,
-            head_dim_qk,
-            dtype=dtype,
-            device=device,
-            generator=gen,
-        ).contiguous()
-        kv_len = batch_size * seq_len
-        k = torch.rand(
-            kv_len,
-            num_kv_heads,
-            head_dim_qk,
-            dtype=dtype,
-            device=device,
-            generator=gen,
-        ).contiguous()
-        v = torch.rand(
-            kv_len,
-            num_kv_heads,
-            head_dim_vo,
-            dtype=dtype,
-            device=device,
-            generator=gen,
-        ).contiguous()
-        output = torch.empty(
-            qo_len,
-            num_qo_heads,
-            head_dim_vo,
-            dtype=dtype,
-            device=device,
-        )
-        qo_indptr = torch.arange(0, batch_size + 1, dtype=torch.int32, device=device) * seq_len
-        kv_indptr = qo_indptr.clone()
-        return [
-            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,
-        ]
-
-    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,
-    ):
-        _ensure_flashinfer_importable()
-        import flashinfer
-
-        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,
-            int(num_qo_heads),
-            int(num_kv_heads),
-            int(head_dim_qk),
-            int(head_dim_vo),
-            causal=bool(causal),
-            q_data_type=torch.bfloat16,
-            kv_data_type=torch.bfloat16,
-        )
-        wrapper.run(q, k, v, out=output)
-        return [
-            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,
-        ]
-
-    def check(
-        testcase_sizes,
-        original_input_tensors,
-        target_kernel_input_tensors,
-        baseline_input_tensors,
-        rtol=1e-2,
-        atol=1e-2,
-    ) -> bool:
-        del testcase_sizes, original_input_tensors
-        output_t = target_kernel_input_tensors[3]
-        output_ref = baseline_input_tensors[3]
-        if output_t.shape != output_ref.shape:
-            print(f"[FAIL] shape mismatch: target {output_t.shape}, ref {output_ref.shape}", file=sys.stderr)
-            return False
-        if output_t.dtype != output_ref.dtype:
-            print(f"[FAIL] dtype mismatch: target {output_t.dtype}, ref {output_ref.dtype}", file=sys.stderr)
-            return False
-        if not torch.allclose(output_t.float(), output_ref.float(), rtol=rtol, atol=atol):
-            diff = (output_t.float() - output_ref.float()).abs()
-            print(
-                f"[FAIL] allclose failed: max_abs_diff={float(diff.max().item()):.6f}, "
-                f"mean_abs_diff={float(diff.mean().item()):.6f} (rtol={rtol}, atol={atol})",
-                file=sys.stderr,
-            )
-            return False
-        return True
-except Exception:
-    pass
-
-
-INPUT_CLASS = [
-    "INPUT",
-    "INPUT",
-    "INPUT",
-    "OUTPUT",
-    "INPUT",
-    "INPUT",
-    "INPUT",
-    "INPUT",
-    "INPUT",
-    "INPUT",
-    "INPUT",
-    "INPUT",
-    "INPUT",
-]
-
-
-def getWorkload(testcase_sizes) -> dict:
-    raw_sizes = testcase_sizes[0] if isinstance(testcase_sizes, tuple) and len(testcase_sizes) == 2 else testcase_sizes
-    q_shape, k_shape, v_shape, output_shape, qo_indptr_shape, kv_indptr_shape = raw_sizes[:6]
-    qo_len, num_qo_heads, head_dim_qk = q_shape
-    kv_len, num_kv_heads, k_dim = k_shape
-    v_len, v_heads, head_dim_vo = v_shape
-    assert k_dim == head_dim_qk
-    assert v_len == kv_len
-    assert v_heads == num_kv_heads
-    assert qo_len == kv_len
-    assert output_shape == (qo_len, num_qo_heads, head_dim_vo)
-    assert qo_indptr_shape == kv_indptr_shape
-    batch_size = qo_indptr_shape[0] - 1
-    seq_len = kv_len // batch_size
-    flops = batch_size * seq_len * seq_len * num_qo_heads * (head_dim_qk + head_dim_vo)
-    memory_bytes = (
-        qo_len * num_qo_heads * head_dim_qk * 2
-        + kv_len * num_kv_heads * head_dim_qk * 2
-        + kv_len * num_kv_heads * head_dim_vo * 2
-        + qo_len * num_qo_heads * head_dim_vo * 2
-        + (batch_size + 1) * 4 * 2
-    )
-    return {
-        "flops": flops,
-        "memory_bytes": memory_bytes,
-        "dtype": "bf16",
-    }
-
-
-DESIGNED_VRAM_SIZE = 48
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/problem_20001/zh_CN/00_题目描述.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/problem_20001/zh_CN/00_题目描述.md
@ -1,23 +0,0 @@
---
-sectionTitle: "题目描述"
-type: "Text"
---
-你需要实现 FlashInfer ragged KV cache prefill 的CUDA C++前向算子。
-
-本题输入采用 FlashInfer `BatchPrefillWithRaggedKVCacheWrapper` 的 ragged `NHD` 布局。每个 batch 中有 `seq_len` 个 query token，KV cache 中也有 `seq_len` 个 token：
-
-其中 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)
-```
-
-如何提交代码详见[评测指南](/d/2)。
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/problem_20001/zh_CN/01_接口约定cuda.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/problem_20001/zh_CN/01_接口约定cuda.md
@ -1,41 +0,0 @@
---
-sectionTitle: "接口约定"
-type: "codeSample"
-lang: "cuda"
---
-你必须在提交的 CUDA 源码中提供如下 **C 符号**，函数名、参数类型、顺序必须完全一致，并使用 `extern "C"` 防止 name mangling：
-
-```cpp
-#include <stdint.h>
-#include <cuda_bf16.h>
-
-extern "C" void run_kernel(
-    const __nv_bfloat16* q,
-    const __nv_bfloat16* k,
-    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 `(batch_size * seq_len, num_qo_heads, head_dim_qk)`，连续 `bf16`
-* `k`：key tensor，shape `(batch_size * seq_len, num_kv_heads, head_dim_qk)`，连续 `bf16`
-* `v`：value tensor，shape `(batch_size * seq_len, num_kv_heads, head_dim_vo)`，连续 `bf16`
-* `output`：输出缓冲区，shape `(batch_size * seq_len, 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`
-* `causal`：是否启用 causal mask，评测中固定为 `1`
-
-本题测试中 `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()` 或显式同步。
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/problem_20001/zh_CN/01_接口约定tilelang.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/problem_20001/zh_CN/01_接口约定tilelang.md
@ -1,52 +0,0 @@
---
-sectionTitle: "接口约定"
-type: "codeSample"
-lang: "tilelang"
---
-你必须在提交的 Python 代码中提供 `run_kernel` 函数，函数名、参数顺序、类型必须完全一致：
-
-```python
-import tilelang
-import tilelang.language as T
-from tilelang import jit
-
-real_kernel = None
-
-@jit
-def build_kernel(*args):
-    @T.prim_func
-    def kernel(*args):
-        ...
-    return kernel
-
-def run_kernel(
-    q,  # Tensor[bf16], shape (batch_size * seq_len, num_qo_heads, head_dim_qk)
-    k,  # Tensor[bf16], shape (batch_size * seq_len, num_kv_heads, head_dim_qk)
-    v,  # Tensor[bf16], shape (batch_size * seq_len, num_kv_heads, head_dim_vo)
-    output,  # Tensor[bf16], shape (batch_size * seq_len, 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
-    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`
-* `causal`：是否启用 causal mask，评测中固定为 `1`
-
-`run_kernel` 内部需要自行计算合适的 grid/block，并 launch 你实现的 TileLang kernel。
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/problem_20001/zh_CN/01_接口约定triton.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/problem_20001/zh_CN/01_接口约定triton.md
@ -1,41 +0,0 @@
---
-sectionTitle: "接口约定"
-type: "codeSample"
-lang: "triton"
---
-你必须在提交的 Python 代码中提供 `run_kernel` 函数，函数名、参数顺序、类型必须完全一致：
-
-```python
-import triton
-import triton.language as tl
-
-@triton.jit
-def your_kernel(...):
-    ...
-
-def run_kernel(
-    q,  # Tensor[bf16], shape (batch_size * seq_len, num_qo_heads, head_dim_qk)
-    k,  # Tensor[bf16], shape (batch_size * seq_len, num_kv_heads, head_dim_qk)
-    v,  # Tensor[bf16], shape (batch_size * seq_len, num_kv_heads, head_dim_vo)
-    output,  # Tensor[bf16], shape (batch_size * seq_len, 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
-    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`
-* `causal`：是否启用 causal mask，评测中固定为 `1`
-
-`run_kernel` 内部需要自行计算合适的 grid/block，并 launch 你实现的 Triton kernel。
--- a/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/problem_20001/zh_CN/02_数据范围.md
+++ b/Agent开发范式的国产GPU大模型推理算子库优化/operator_task_package/flashinfer_task_package/xpuoj_problem/problem_20001/zh_CN/02_数据范围.md
@ -1,26 +0,0 @@
---
-sectionTitle: "输入格式"
-type: "Text"
---
-## 数据范围
-
- 数据类型：`q/kv_data/output` 均为 `bfloat16`
- KV layout：`NHD`
- `num_qo_heads = 32`
- `num_kv_heads = 4`
- `page_block_size = 16`
- `causal = 0`
- `head_dim` 固定为 `128`
- `batch_size` 取值为 `1, 4, 16`
- `seq_len` 取值为 `1024, 4096, 8192, 16384`
-
-测试点顺序与 `McFlashInfer/benchmarks/bench_batch_prefill_paged.py` 中的 cases 一致，即：
-
-```python
-for head_dim in [128]:
-    for batch_size in [1, 4, 16]:
-        for seq_len in [1024, 4096, 8192, 16384]:
-            ...
-```
-
-输出参与 `torch.allclose` 校验，容差为 `rtol=1e-2, atol=1e-2`。
--- 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