docs: update flashinfer 文档链接, 公式, 图片及相关表述 #76
|
|
@ -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 账号进行注册。登录后显示页面:
|
||||
|
||||

|
||||
|
||||
点击页面左上角 “giteeAI - 模力方舟” 图标,进入用户控制台:
|
||||
|
||||

|
||||
|
||||
在左侧边栏选择 “费用中心”,点击右上角 “兑换” 使用兑换码兑换代金券:
|
||||
|
||||

|
||||
|
||||
|
|
@ -86,9 +94,9 @@
|
|||
|
||||

|
||||
|
||||
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**,点击下一步;
|
||||
|
||||

|
||||

|
||||
|
||||
3. 勾选同意服务条款,点击 “创建实例”;
|
||||
|
||||
|
|
@ -126,7 +134,7 @@
|
|||
|
||||

|
||||
|
||||
此部分内容可参考教程:[*模力方舟快速使用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
|
|||
|
||||
**预期结果:**
|
||||
|
||||

|
||||

|
||||
|
||||
#### 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/) 平台,使用组委会统一发放的账号和初始密码登录;
|
||||
|
||||

|
||||
|
||||
2. 登录后进入比赛 / 题目列表页面;
|
||||
2. 登录后进入 **比赛** 列表,找到 **沐曦 - 揭榜挂帅 - Agent 推理算子库优化 - FlashInfer 任务**;
|
||||
|
||||

|
||||

|
||||
|
||||
3. 找到对应题目,例如 **20001 FlashInfer Ragged Prefill**;
|
||||
|
||||

|
||||
3. 找到对应题目,例如 **Agent 推理算子库优化 - FlashInfer Ragged Prefill(OJ 题号 20001)** ;
|
||||
|
||||

|
||||
|
||||
4. 点击进入题目详情页,查看题目描述、接口约定、数据范围和提交入口。
|
||||
|
||||

|
||||

|
||||
|
||||
#### 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 评测流程**:
|
||||
|
||||
|
|
@ -690,45 +701,46 @@ FlashInfer 方向包含 **4 个可选算子题目**,每个对应独立的 benc
|
|||
|
||||
5. 查看结果
|
||||
|
||||
1. 单测试点分析 — 以 OJ 返回的一次评测为例
|
||||
1. 单测试点分析 — 以 OJ 返回的一次评测结果为例
|
||||
|
||||
提交后,OJ 平台对每个测试用例独立评测并返回结果。以下是一份优化后的真实评测输出(来自 20001 FlashInfer Ragged Prefill 第 1 个测试点):
|
||||
提交后,OJ 平台对每个测试用例独立评测并返回结果。以下是一份冒烟提交的典型评测输出(来自 20001 FlashInfer Ragged Prefill 第 1 个测试点):
|
||||
|
||||

|
||||

|
||||
|
||||
```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评分规则概要))。
|
||||
|
||||

|
||||

|
||||
|
||||
**榜单解读**:
|
||||
|
||||
- 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冒烟代码)
|
||||
|
||||
|
|
|
|||
Loading…
Reference in New Issue