diff --git a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/FlashInfer Benchmark实战：从性能基线到XPU-OJ冒烟提交.md b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/FlashInfer Benchmark实战：从性能基线到XPU-OJ冒烟提交.md index 730ee7f..236acc1 100644 --- a/基于AI Agent开发范式的国产GPU大模型推理算子库优化/FlashInfer Benchmark实战：从性能基线到XPU-OJ冒烟提交.md +++ b/基于AI Agent开发范式的国产GPU大模型推理算子库优化/FlashInfer Benchmark实战：从性能基线到XPU-OJ冒烟提交.md @@ -2,7 +2,7 @@ ## 1. 教程定位 -本教程是赛题二 FlashInfer 任务的 “benchmark 性能基线与 XPU-OJ 提交衔接” 模块，主要帮助学员跑通目标算子的 benchmark 脚本，理解原库 API、输入输出结构、性能指标和性能基线结果，并进一步读懂 XPU-OJ 题目中的接口约定、测试数据、参考输出和精度要求。 +本教程是 **沐曦 - 揭榜挂帅 - Agent 推理算子库优化 - FlashInfer 任务** 的 “benchmark 性能基线与 XPU-OJ 提交衔接” 模块，主要帮助学员跑通目标算子的 benchmark 脚本，理解原库 API、输入输出结构、性能指标和性能基线结果，并进一步读懂 XPU-OJ 题目中的接口约定、测试数据、参考输出和精度要求。 需要特别说明： @@ -75,7 +75,15 @@ 使用兑换码兑换 GPU 租用余额： -访问 [*模力方舟官网*](https://ai.gitee.com/)，在左侧边栏进入 “费用中心”，点击右上角 “兑换” 使用兑换码兑换代金券； +访问 [*模力方舟官网*](https://ai.gitee.com/)，首次登录需要使用手机号或者 Gitee 账号进行注册。登录后显示页面： + +![mainpage](https://origin.picgo.net/2026/07/08/Screenshot-2026-07-07-at-7.03.00PM4301a879ed958b6b.png) + +点击页面左上角 “giteeAI - 模力方舟” 图标，进入用户控制台： + +![dashboard](https://origin.picgo.net/2026/07/08/Screenshot-2026-07-07-at-7.08.25PMbdb17cde254ddf29.png) + +在左侧边栏选择 “费用中心”，点击右上角 “兑换” 使用兑换码兑换代金券： ![quote](https://origin.picgo.net/2026/06/23/-2026-06-23-1515027374ed9083ad0aea.png) @@ -86,9 +94,9 @@ ![rent](https://origin.picgo.net/2026/06/17/Weixin-Image_20260617180029_15_100c1e1fecb298bac54.png) -2. 进入 “创建实例” 页面，确认计费方式为 “按量收费”，预装镜像选择：基础镜像 - **PyTorch-Agent / 2.8.0 / Python 3.12 / maca 3.7.2.1**，点击下一步； +2. 进入 “创建实例” 页面，确认计费方式为 “按量收费”，预装镜像选择：基础镜像 - **PyTorch-Agent / 2.8.0 / Python 3.12 / maca 3.7.1.5**，点击下一步； - ![pytorch agent](https://origin.picgo.net/2026/06/18/pytorch-agent70c0e9dce299e26c.png) + ![pytorch agent](https://origin.picgo.net/2026/07/14/3177f86f0227834da1a.png) 3. 勾选同意服务条款，点击 “创建实例”； @@ -126,7 +134,7 @@ ![tools](https://origin.picgo.net/2026/06/23/-2026-06-23-1554030c167f9883fb1d79.png) - 此部分内容可参考教程：[*模力方舟快速使用SOP*](../模力方舟快速使用SOP.md) + 此部分内容可参考教程：[*模力方舟快速使用SOP*](https://gitlink.org.cn/metax-maca/op_optimization/tree/master/%E6%A8%A1%E5%8A%9B%E6%96%B9%E8%88%9F%E5%BF%AB%E9%80%9F%E4%BD%BF%E7%94%A8SOP.md) #### 4.1.4 深度学习环境配置 @@ -161,7 +169,7 @@ which mxcc && mxcc --version || echo "mxcc 未找到，请确认 MACA 工具链 | 问题 | 解决方法 | | --- | --- | -| `mxcc: command not found` | MACA 工具链未安装或 `PATH` 未配置，检查镜像是否预装或参考 [*模力方舟快速使用SOP*](../模力方舟快速使用SOP.md) | +| `mxcc: command not found` | MACA 工具链未安装或 `PATH` 未配置，检查镜像是否预装或参考 [*模力方舟快速使用SOP*](https://gitlink.org.cn/metax-maca/op_optimization/tree/master/%E6%A8%A1%E5%8A%9B%E6%96%B9%E8%88%9F%E5%BF%AB%E9%80%9F%E4%BD%BF%E7%94%A8SOP.md) | ### 4.2 工具准备 @@ -171,7 +179,7 @@ which mxcc && mxcc --version || echo "mxcc 未找到，请确认 MACA 工具链 - 已确认 Agent 可以正常调用模型。 -配置过程可参考 [*模力方舟Agent部署准备教程*](模力方舟Agent部署准备教程.md)。 +配置过程可参考 [*模力方舟Agent部署准备教程*](https://gitlink.org.cn/metax-maca/op_optimization/tree/master/%E5%9F%BA%E4%BA%8EAI%20Agent%E5%BC%80%E5%8F%91%E8%8C%83%E5%BC%8F%E7%9A%84%E5%9B%BD%E4%BA%A7GPU%E5%A4%A7%E6%A8%A1%E5%9E%8B%E6%8E%A8%E7%90%86%E7%AE%97%E5%AD%90%E5%BA%93%E4%BC%98%E5%8C%96%2F%E6%A8%A1%E5%8A%9B%E6%96%B9%E8%88%9FAgent%E9%83%A8%E7%BD%B2%E5%87%86%E5%A4%87%E6%95%99%E7%A8%8B.md)。 **以配置 OpenCode 为例** @@ -204,6 +212,8 @@ XPU-OJ 账号由组委会统一发放，参赛者无需自行注册。 ### 4.5 评分规则概要 +评分规则详情参考 [*基于AI Agent开发范式的国产GPU大模型推理算子库优化比赛方案*](https://gitlink.org.cn/metax-maca/op_optimization/tree/master/%E5%9F%BA%E4%BA%8EAI%20Agent%E5%BC%80%E5%8F%91%E8%8C%83%E5%BC%8F%E7%9A%84%E5%9B%BD%E4%BA%A7GPU%E5%A4%A7%E6%A8%A1%E5%9E%8B%E6%8E%A8%E7%90%86%E7%AE%97%E5%AD%90%E5%BA%93%E4%BC%98%E5%8C%96%2F%E5%9F%BA%E4%BA%8EAI%20Agent%E5%BC%80%E5%8F%91%E8%8C%83%E5%BC%8F%E7%9A%84%E5%9B%BD%E4%BA%A7GPU%E5%A4%A7%E6%A8%A1%E5%9E%8B%E7%AE%97%E5%AD%90%E6%8E%A8%E7%90%86%E5%BA%93%E4%BC%98%E5%8C%96%E6%96%B9%E6%A1%88.md)。 + 提交作品的最终评价采用 **100 分制**，由客观评测和专家评审共同组成： | 类别 | 评审维度 | 权重 | 说明 | @@ -214,7 +224,7 @@ XPU-OJ 账号由组委会统一发放，参赛者无需自行注册。 > - OJ 榜单分数与比赛最终得分属于不同分值体系。OJ 内部以各算子 baseline 为 50 分基准、硬件理论上限约 100 分；比赛性能部分满分 60 分，按排名映射； -> - FlashInfer 方向包含多个子任务，任选一种或多种提交，各子任务独立计分，取通过正确性和稳定性测试的最高成绩参与排名； +> - FlashInfer 方向包含多个子任务，任选一种或多种提交，各子任务独立计分，取通过正确性和稳定性测试的最高成绩参与排名，即 **FlashInfer 任务分数** = $\max$\{子题1分数, ..., 子题4分数\}； > - 组委会对提交内容严格审查，严禁抄袭，一经发现取消成绩。 ## 5. 知识预备 @@ -344,7 +354,7 @@ pip install pandas **预期结果：** - ![ls -l](https://origin.picgo.net/2026/06/24/-2026-06-24-1651399191c4b66abe6916.png) + ![ls -l](https://origin.picgo.net/2026/07/08/-2026-07-08-1444052f3264a5597c95d5.png) #### Step 3：验证项目脚本 @@ -479,14 +489,14 @@ else: **选择目标题目** -FlashInfer 方向包含 **4 个可选算子题目**，每个对应独立的 benchmark 脚本、题目说明、`run_kernel(...)` 接口和数据范围。 +FlashInfer 方向包含 **4 个可选算子题目**，均属于同一比赛通道，**FlashInfer 任务分** = $\max$\{各子题得分\}。每个对应独立的 benchmark 脚本、题目说明、`run_kernel(...)` 接口和数据范围。 -| OJ 题号 | 算子类型 | 核心特点 | Benchmark 脚本 | FlashInfer API | +| OJ 题号 | OJ 题目名称 | 核心特点 | Benchmark 脚本 | FlashInfer API | |---------|---------------|----------------|----------|----------| -| **20001** | Ragged Prefill | GQA布局，Q/K/V平坦存储，causal=1 | `bench_batch_prefill_ragged.py` | `BatchPrefillWithRaggedKVCacheWrapper` | -| **20002** | Paged Prefill | KV Cache分页存储，需解析page table | `bench_batch_prefill_paged.py` | `BatchPrefillWithPagedKVCacheWrapper` | -| **20003** | MLA Paged Attention | DeepSeek MLA特有，双路Q(nope+pe)/双路Cache(ckv+kpe) |`bench_batch_mla.py` | `BatchMLAPagedAttentionWrapper` | -| **20004** | Paged Decode | 每次只1个query token，memory-bound |`bench_batch_decode.py` | `BatchDecodeWithPagedKVCacheWrapper` | +| **20001** | FlashInfer Ragged Prefill | GQA布局，Q/K/V平坦存储，causal=1 | `bench_batch_prefill_ragged.py` | `BatchPrefillWithRaggedKVCacheWrapper` | +| **20002** | FlashInfer Paged Prefill | KV Cache分页存储，需解析page table | `bench_batch_prefill_paged.py` | `BatchPrefillWithPagedKVCacheWrapper` | +| **20003** | FlashInfer MLA Paged Attention | DeepSeek MLA特有，双路Q(nope+pe)/双路Cache(ckv+kpe) |`bench_batch_mla.py` | `BatchMLAPagedAttentionWrapper` | +| **20004** | FlashInfer Paged Decode | 每次只1个query token，memory-bound |`bench_batch_decode.py` | `BatchDecodeWithPagedKVCacheWrapper` | **每个子题的接口参数、数据范围和精度要求以对应题目说明为准。** 下文以题目 **20001 Flashinfer Ragged Prefill** 为例演示从 benchmark 到 XPU-OJ 提交的完整流程。 @@ -497,7 +507,7 @@ FlashInfer 方向包含 **4 个可选算子题目**，每个对应独立的 benc > 完成 benchmark 后，需要注意 benchmark 脚本主要用于建立性能基线，并不需要最终提交 > 最终评测以 XPU-OJ 为准，评测程序会调用选手提交代码中的 `run_kernel`，并将输出结果与 OJ 后台参考实现结果进行比较 -下面以题目 **20001 FlashInfer Ragged Prefill** 为例，从本地题目文档 [*Agent 推理算子库优化 - FlashInfer Ragged Prefill*](./operator_task_package/flashinfer_task_package/xpuoj_problem/problem_20001/Agent%20推理算子库优化%20-%20FlashInfer%20Ragged%20Prefill.md) 中逐节解读关键信息。 +下面以题目 **20001 FlashInfer Ragged Prefill** 为例，从本地题目文档 [*Agent 推理算子库优化 - FlashInfer Ragged Prefill*](https://gitlink.org.cn/metax-maca/op_optimization/tree/master/%E5%9F%BA%E4%BA%8EAI%20Agent%E5%BC%80%E5%8F%91%E8%8C%83%E5%BC%8F%E7%9A%84%E5%9B%BD%E4%BA%A7GPU%E5%A4%A7%E6%A8%A1%E5%9E%8B%E6%8E%A8%E7%90%86%E7%AE%97%E5%AD%90%E5%BA%93%E4%BC%98%E5%8C%96%2Foperator_task_package%2Fflashinfer_task_package%2Fxpuoj_problem%2Fproblem_20001%2FAgent%20%E6%8E%A8%E7%90%86%E7%AE%97%E5%AD%90%E5%BA%93%E4%BC%98%E5%8C%96%20-%20FlashInfer%20Ragged%20Prefill.md) 中逐节解读关键信息。 1. `## 1. 题目描述` — 我要实现什么？ @@ -585,23 +595,24 @@ FlashInfer 方向包含 **4 个可选算子题目**，每个对应独立的 benc #### Step 7：登录 XPU-OJ 并进入题目页面 -使用组委会统一发放的账号登录 XPU-OJ，并进入对应赛题页面。 +使用组委会统一发放的账号登录 XPU-OJ，并在 XPU-OJ 上找到 **比赛 -> 沐曦 - 揭榜挂帅 - Agent 推理算子库优化 - FlashInfer 任务**，进入后可见 4 个题目。 -1. 打开 [*XPU-OJ*](https://xpuoj.com/) 平台，使用组委会统一发放的账号和初始密码登录 **【后续发布】**； +1. 打开 [*XPU-OJ*](https://xpuoj.com/) 平台，使用组委会统一发放的账号和初始密码登录； ![OJ](https://origin.picgo.net/2026/06/16/image-20260616152804368aa0d44b72c4f9572.png) -2. 登录后进入比赛 / 题目列表页面； +2. 登录后进入 **比赛** 列表，找到 **沐曦 - 揭榜挂帅 - Agent 推理算子库优化 - FlashInfer 任务**； - ![OJ-2](https://origin.picgo.net/2026/06/16/image-20260616152953345986ce39fda69da55.png) + ![OJ-2](https://origin.picgo.net/2026/07/14/7c9be6f21b764bd4bb4d0d1ba03101aac3d2235116485c23.png) -3. 找到对应题目，例如 **20001 FlashInfer Ragged Prefill**； - ![OJ-3](https://origin.picgo.net/2026/06/16/image-202606161531423833b0f0428edf2e35e.png) +3. 找到对应题目，例如 **Agent 推理算子库优化 - FlashInfer Ragged Prefill（OJ 题号 20001）** ； + + ![OJ-3](https://origin.picgo.net/2026/07/14/screenshot-178399530855671a499a30efc93e3.png) 4. 点击进入题目详情页，查看题目描述、接口约定、数据范围和提交入口。 - ![OJ-4](https://origin.picgo.net/2026/06/17/Screenshot_17-6-2026_165526_xpuoj.com9a82c272c8ea5699.jpeg) + ![OJ-4](https://origin.picgo.net/2026/07/14/screenshot-17839970631786edace2ba4635016.png) #### Step 8：提交 OJ 冒烟代码 **目标**：完成一次最小提交，确认 OJ 提交链路、语言环境和 `run_kernel(...)` 接口可用。 @@ -648,7 +659,7 @@ FlashInfer 方向包含 **4 个可选算子题目**，每个对应独立的 benc - [*点击查看参考 Prompt*](#参考%20prompt) - - [*点击查看参考冒烟代码*](#20001%20flashinfer%20ragged%20prefill%20参考冒烟代码) + - [*点击查看参考冒烟代码*](#参考冒烟代码) 更多 OpenCode 使用教程和 Agent 使用技巧可见 @@ -660,11 +671,11 @@ FlashInfer 方向包含 **4 个可选算子题目**，每个对应独立的 benc 4. 点击提交，等待评测结果返回； - ![OJ-5](https://origin.picgo.net/2026/06/16/image-202606161550234433e54d703d581f858.png) + ![OJ-5](https://origin.picgo.net/2026/07/14/screenshot-178399559790299c95f05f67c4aea.png) - 评测时间与题目测试点数量、队列状态和平台负载有关，通常需要等待**数十秒到数分钟**，以平台实际返回为准。 + 评测时间与题目测试点数量、队列状态和平台负载有关，通常需要等待**数十秒到数分钟**，以平台实际运行状况为准。 - ![OJ-6](https://origin.picgo.net/2026/06/17/iwEcAqNwbmcDAQTRBt8F0QOuBrA3mrGf2xJ4fgoGfxH9GVsAB9MAAAABijZ-PQgACaJpbQoAC9IAAlPf771f660d13b07471.png) + ![OJ-6](https://origin.picgo.net/2026/07/14/screenshot-178399561798779de7c317dbbfb2d.png) **OJ 评测流程**： @@ -690,45 +701,46 @@ FlashInfer 方向包含 **4 个可选算子题目**，每个对应独立的 benc 5. 查看结果 - 1. 单测试点分析 — 以 OJ 返回的一次评测为例 + 1. 单测试点分析 — 以 OJ 返回的一次评测结果为例 - 提交后，OJ 平台对每个测试用例独立评测并返回结果。以下是一份优化后的真实评测输出（来自 20001 FlashInfer Ragged Prefill 第 1 个测试点）： + 提交后，OJ 平台对每个测试用例独立评测并返回结果。以下是一份冒烟提交的典型评测输出（来自 20001 FlashInfer Ragged Prefill 第 1 个测试点）： - ![OJ-7](https://origin.picgo.net/2026/06/24/Screenshot-2026-06-24-at-5.57.34PMbe24088f85c638ca.png) + ![OJ-7](https://origin.picgo.net/2026/07/08/-2026-07-08-1451020acc658c43662f1a.png) ```plaintext - Testcase #1 + 测试点 #1 Accepted - 112 pts - 125 ms + 1 pts + 119 ms 22.0 G + 输入文件 1 - Your output + 你的输出 - OJCHAL v1 1bHB14u2jFxvzumsIYHDPA== - OJRESULT v1 f84db85c... eyJ0aW1lX21zIjoxMjQuNjY1LCJzcGVlZHVwIjoyLjgwMjY1NSwidGtfdGltZV9tcyI6MTI0LjY2NSwidGJfdGltZV9tcyI6MzQ5LjM5MywidGhfdGltZV9tcyI6MTQ2LjYwMTU1LCJzY29yZV9yYXRpbyI6MS4xMjEyOTQsInBhc3MiOnRydWV9 + OJCHAL v1 dQ0CZtDCX0JLxT1SF4h/Pw== + OJRESULT v1 c6583f61291371b771de67188893f7d47a745beb2929edce17a67f44a0d80def eyJzY2hlbWFfdmVyc2lvbiI6MiwidGltZV9tcyI6MTE5LjI0Mywic3BlZWR1cCI6MC4wMTM0NTIsInRrX3RpbWVfbXMiOjExOS4yNDMsInRiX3RpbWVfbXMiOjEuNjA0LCJ0aF90aW1lX21zIjowLjM3MzMzNCwic2NvcmVfcmF0aW8iOjAuMDEwMjQ3LCJwYXNzIjp0cnVlfQ== + 你的标准错误输出 - Checker message + {"schema_version":2,"time_ms":119.243,"speedup":0.013452,"tk_time_ms":119.243,"tb_time_ms":1.604,"th_time_ms":0.373334,"score_ratio":0.010247,"pass":true} + 检查器信息 === SPJ Report - FlashInfer Batch Prefill === ---------------------------------------------------------------- Testcase #1 - Config: batch=16, seq_len=16384, q_heads=32, kv_heads=4, - head_dim_qk=128, head_dim_vo=128, causal=1 + Config: sol016_trace_synthetic_b33_total16294: batch=33, total_q=16294, total_kv=16294, max_q=987, max_kv=987, q_heads=32, kv_heads=4, head_dim_qk=128, head_dim_vo=128, causal=1, source=sol-execbench 016 axes - Baseline: 349.393000 ms - User kernel: 124.665000 ms - Hardware bound: 146.601550 ms - Speedup vs base: 2.803 x + Baseline: 1.604000 ms + User kernel: 119.243000 ms + Speedup vs base: 0.013 x - Score ratio: 1.121294 (112.13%) - Display score: 112 / 100 + Score ratio: 0.010247 (1.02%) + Display score: 1 / 100 Pass: OK ---------------------------------------------------------------- ``` - > 以上是**优化后**的结果（得分 112 分，加速比 2.8x），非冒烟测试预期。冒烟测试通常 `speedup < 1`，这是正常的起始点。另外 `OJCHAL` 和 `OJRESULT` 为平台元信息，参赛者无需关注，SPJ Report 中已包含所有评测指标的可读版本 + > 以上是一份冒烟提交的评测输出示例：`speedup = 0.013x`（远慢于 baseline），`Display score = 1`。 `OJCHAL` 和 `OJRESULT` 为平台元信息，参赛者无需关注，SPJ Report 中已包含所有评测指标的可读版本。 **OJ 输出中各项指标的含义：** @@ -737,16 +749,16 @@ FlashInfer 方向包含 **4 个可选算子题目**，每个对应独立的 benc ```plaintext Testcase #1 Accepted ← 状态 - 112 pts ← 显示得分 - 125 ms ← kernel 耗时 + 1 pts ← 显示得分 + 119 ms ← kernel 耗时 22.0 G ← 内存占用 ``` | 项目 | 含义 | 说明 | |------|------|------| - | `Accepted` | 评测状态 | 通过正确性校验；若为 `Wrong Answer` 则未通过，不参与排名 | - | `112 pts` | 显示得分 | 基于评分公式 + 对数压缩后的分数（满分参考值 100） | - | `125 ms` | kernel 耗时 | 该测试点 OJ 测速阶段你的 kernel 平均执行时间 | + | `Accepted` | 评测状态 | 通过正确性校验；若为 `Wrong Answer`、`Time Limit Exceeded` 则未通过，不参与排名 | + | `1 pts` | 显示得分 | 基于评分公式 + 对数压缩后的分数。冒烟阶段通常极低，正常现象 | + | `119 ms` | kernel 耗时 | 该测试点 OJ 测速阶段你的 kernel 平均执行时间 | | `22.0 G` | 内存占用 | CPU 侧 RSS，仅供平台监控进程占用、防止 OOM | 其次，`Checker message` 中的 SPJ Report 给出了详细的性能对比： @@ -755,16 +767,15 @@ FlashInfer 方向包含 **4 个可选算子题目**，每个对应独立的 benc === SPJ Report - FlashInfer Batch Prefill === ---------------------------------------------------------------- Testcase #1 - Config: batch=16, seq_len=16384, q_heads=32, kv_heads=4, - head_dim_qk=128, head_dim_vo=128, causal=1 + Config: sol016_trace_synthetic_b33_total16294: batch=33, total_q=16294, total_kv=16294, max_q=987, max_kv=987, + q_heads=32, kv_heads=4, head_dim_qk=128, head_dim_vo=128, causal=1 - Baseline: 349.393000 ms - User kernel: 124.665000 ms - Hardware bound: 146.601550 ms - Speedup vs base: 2.803 x + Baseline: 1.604000 ms + User kernel: 119.243000 ms + Speedup vs base: 0.013 x - Score ratio: 1.121294 (112.13%) - Display score: 112 / 100 + Score ratio: 0.010247 (1.02%) + Display score: 1 / 100 Pass: OK ---------------------------------------------------------------- ``` @@ -773,44 +784,41 @@ FlashInfer 方向包含 **4 个可选算子题目**，每个对应独立的 benc | 字段 | 含义 | 本例值 | 解读 | |------|------|--------|------| - | `Pass` | 正确性校验 | `OK` | 通过了 `allclose(rtol=1e-2, atol=1e-2)`，可参与排名 | - | `User kernel` | 你的 kernel 耗时 (ms) | `124.67 ms` | **核心性能指标** | - | `Baseline` | OJ 参考实现耗时 (ms) | `349.39 ms` | 你的基准对比对象 | - | `Hardware bound` | 硬件理论下限 (ms) | `146.60 ms` | 基于 FLOPs / 带宽估算的理论最快时间 | - | `Speedup vs base` | 加速比 | `2.80x` | `Baseline / User kernel`，`> 1` 表示优于 OJ 参考实现 | - | `Score ratio` | 归一化得分（原始值） | `1.121` | 综合评分原始值 | - | `Display score` | 显示得分 | `112` | 经过对数压缩后的得分，超过 100 表示超越硬件理论估算 | - | `Config` | 测试配置 | — | 该测试点的 `batch`、`seq_len`、`q_heads`、`kv_heads` 等参数 | + | `Config` | 测试配置 | — | 该测试点的 `batch`、`total_q`、`max_q`、`q_heads`、`kv_heads` 等参数 | + | `Baseline` | OJ 参考实现耗时 (ms) | `1.60 ms` | 你的基准对比对象 | + | `User kernel` | 你的 kernel 耗时 (ms) | `119.24 ms` | **核心性能指标** | + | `Speedup vs base` | 加速比 | `0.013x` | `Baseline / User kernel`。`< 1` 表示比 baseline 慢 | + | `Score ratio` | 归一化得分（原始值） | `0.0102` | 综合评分原始值 | + | `Display score` | 显示得分 | `1` | OJ 内部评分，冒烟阶段通常极低 | + | `Pass` | 正确性校验 | `OK` | 通过了正确性校验，可参与排名 | **从 SPJ Report 分析优化方向：** - 以上述结果为例： - - `Speedup vs base = 2.80x`，已超越 OJ 参考实现 2.8 倍 + 以上述冒烟结果为例： + - `Speedup vs base = 0.013x`：**冒烟代码远慢于 OJ baseline**，这是正常的起始点——冒烟阶段的目的是验证接口和提交链路，不追求性能； - - `User kernel = 124.67ms` 对比 `Hardware bound = 146.60ms`：**你的 kernel 已超越硬件理论下限**，说明该测试点的实现已非常接近硬件极限，继续优化的收益空间有限 - - - `Score ratio = 1.121` 对应 `Display score = 112`，已超过 100 分的满分线 + - `Baseline = 1.60ms` 对比 `User kernel = 119.24ms`：两者之间有约 75 倍的性能差距，说明**优化空间巨大**。后续可以通过引入 tiling、shared memory、向量化访存等优化手段逐步缩小差距。 **优化优先级判断法：** | 情况 | 优化空间 | 建议方向 | |------|----------|----------| - | `User kernel ≈ Baseline` | 大 | 先保证正确性，再分析 compute / memory 瓶颈 | - | `User kernel` 在 `Baseline` 和 `Hardware bound` 之间 | 中 | 使用 profiler 分析，优化瓶颈阶段 | - | `User kernel ≈ Hardware bound` | 小 | 该测试点已接近最优，转而优化其他测试点 | + | `User kernel ≫ Baseline`（speedup < 1） | 巨大 | 冒烟阶段正常状态。优先保证正确性，再引入 tiling / shared memory 等基础优化 | + | `User kernel ≈ Baseline` | 大 | 已追平 baseline，进一步分析 compute / memory 瓶颈 | + | `User kernel ≪ Baseline`（speedup > 1） | 中 | 已优于 baseline，使用 profiler 精调瓶颈阶段 | **查看所有测试用例结果：** - OJ 平台每次提交会评测所有测试用例，每个测试用例返回一组独立的 SPJ Report。建议将各组结果整理成表格，方便跟踪优化进展： + OJ 平台每次提交会评测所有测试用例，每个测试用例返回一组独立的 SPJ Report。每个子题独立计分，FlashInfer 任务总分取各子题得分的最高值。建议将各组结果整理成表格，方便跟踪优化进展： - | 测试点 | batch_size | seq_len | Baseline | User kernel | Speedup | Score | Pass | + | 测试点 | batch | total_q | Baseline | User kernel | Speedup | Score | Pass | |:---:|:---:|:---:|:---:|:---:|:---:|:---:|:---:| - | 1 | 1 | 1024 | — | — | — | — | — | - | 2 | 1 | 4096 | — | — | — | — | — | + | 1 | 33 | 16294 | 1.60 | 119.24 | 0.013 | 1 | ✓ | + | 2 | 1 | 1024 | — | — | — | — | — | | ... | ... | ... | ... | ... | ... | ... | ... | - | 12 | 16 | 16384 | 349.39 | 124.67 | 2.80 | 112.13 | ✓ | + | 15 | 2 | 98 | — | — | — | — | — | - 2. XPU-OJ 平台评分机制（仅供参考） + 2. XPU-OJ 平台评分机制（**仅供参考**） 以下为 XPU-OJ 平台的内部评分算法，用于生成榜单中每个题目的单题分数。请注意： @@ -822,7 +830,7 @@ FlashInfer 方向包含 **4 个可选算子题目**，每个对应独立的 benc **单测试点评分公式**（参考 Sol-ExecBench）： - $$S(T_k) = \frac{100}{1 + \left(\frac{1}{0.5} - 1\right) \cdot \frac{T_k - T_h}{T_b - T_h}}$$ + $S(T_k) = \frac{100}{1 + \left(\frac{1}{0.5} - 1\right) \cdot \frac{T_k - T_h}{T_b - T_h}}$ 其中： - $T_k$：你的 kernel 平均时间 @@ -848,17 +856,16 @@ FlashInfer 方向包含 **4 个可选算子题目**，每个对应独立的 benc 3. 榜单查看 — 如何阅读排行榜 - 排行榜位于 XPU-OJ 比赛页面，展示各参赛者在每个题目的 OJ 内部得分。请注意：**OJ 榜单中各题的分数与比赛最终得分不是同一体系。** 比赛最终成绩中，性能提升效果仅取你得分最高的一个任务，按该任务的 OJ 排名换算（详见 [*4.5 评分规则概要*](#45%20评分规则概要)）。 + 排行榜位于 **比赛 -> 沐曦 - 揭榜挂帅 - Agent 推理算子库优化 - FlashInfer 任务 -> 排行榜** 页面，展示各参赛者在每个题目的 OJ 内部得分。请注意：**OJ 榜单中各题的分数与比赛最终得分不是同一体系。** 比赛最终成绩中，性能提升效果仅取你得分最高的一个任务，按该任务的 OJ 排名换算（详见 [*4.5 评分规则概要*](#45%20评分规则概要)）。 - ![ranklist](https://origin.picgo.net/2026/07/07/-2026-07-07-173642c2d1ab65ebb06d55.png) + ![ranklist](https://origin.picgo.net/2026/07/14/screenshot-1783995971130bb450447de868379.png) **榜单解读**： - - 6 列对应 6 个 OJ 子题：Ragged Prefill (20001)、Paged Prefill (20002)、MLA Attention (20003)、Paged Decode (20004)、FlashAttention KV Cache Decode (20005)、Fused MoE i8 tn (20006) + - FlashInfer 任务有 4 个子题，各子题独立计分，FlashInfer 任务得分取各子题的最高值； + - 各题分数为 OJ 内部评分（OJ baseline = 50 分），排名按 Total（各题 OJ 得分之和）降序。此为 OJ 平台汇总逻辑，比赛最终只取单一最高分任务，不做多任务累加 - - 各题分数为 OJ 内部评分（OJ baseline = 50 分），排名按 Total（各题 OJ 得分之和）降序。**比赛最终只取单一最高分任务**，而非多任务累加 - - - 每列下方的括号数值（如 `(48)`）表示该题提交次数 + - 每列下方的括号数值（如 `(24)`）表示该题提交次数 - `pass = false` 的提交不参与排名（对应 0 分） @@ -930,7 +937,18 @@ FlashInfer 方向包含 **4 个可选算子题目**，每个对应独立的 benc - online softmax 的 m/l 更新逻辑是否正确 ``` -### 7.7 问题排查 +### 7.7 调试 TLE 超时 +``` plaintext +我的 run_kernel 提交到 OJ 后显示 Time Limit Exceeded，请帮我排查： +1. 这是我的 submit 代码：[粘贴你的 run_kernel 实现] +请帮我检查： +- run_kernel 中是否调用了 cudaDeviceSynchronize()（应删除） +- kernel 中 for 循环的终止条件是否有死循环风险 +- __syncthreads() 是否在条件分支内（应移到分支外） +- grid/block 配置是否过大 +``` + +### 7.8 问题排查 ``` plaintext 运行 bench_batch_prefill_ragged.py 时报错 out of memory，请帮我分析原因并给出解决方案。 @@ -1244,7 +1262,7 @@ Before output, confirm ALL items. Any failure = 0 points. [*回退到 Step 8*](#step%208提交%20oj%20冒烟代码) -### 20001 FlashInfer Ragged Prefill 参考冒烟代码 +### 参考冒烟代码 [*回退到 Step 8*](#step%208提交%20oj%20冒烟代码)