AddingZero
/
op_optimization
forked from metax-maca/op_optimization
0
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Update 更新测试点数据分析部分

This commit is contained in:
MaseChen 2026-07-08 00:46:33 -07:00
parent cfa6613a6d
commit b207cc5501
1 changed files with 62 additions and 54 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

116
基于AI Agent开发范式的国产GPU大模型推理算子库优化/FlashInfer Benchmark实战：从性能基线到XPU-OJ冒烟提交.md
View File

@ -700,45 +700,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 输出中各项指标的含义：**
@ -747,16 +748,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 给出了详细的性能对比：
@ -765,16 +766,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
----------------------------------------------------------------
```
@ -783,44 +783,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。建议将各组结果整理成表格，方便跟踪优化进展：
| 测试点 | 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 平台的内部评分算法，用于生成榜单中每个题目的单题分数。请注意：
@ -940,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，请帮我分析原因并给出解决方案。