forked from metax-maca/op_optimization
Compare commits
173 Commits
beckylu-de
...
master
| Author | SHA1 | Date |
|---|---|---|
|
|
bed86dafbf | |
|
|
232f2631c7 | |
|
|
18268c2639 | |
|
|
395c607128 | |
|
|
a8d08bcbc5 | |
|
|
931fd9e3de | |
|
|
6e78e3defd | |
|
|
641ade97b6 | |
|
|
e6a416096d | |
|
|
a686eb3b5b | |
|
|
8e266c819c | |
|
|
88f7103e02 | |
|
|
33d8c6ea3c | |
|
|
8b2d154405 | |
|
|
e6b9449632 | |
|
|
cd60d02057 | |
|
|
3342f411cd | |
|
|
43d0afee79 | |
|
|
781d2d0a18 | |
|
|
ffae51da85 | |
|
|
cbbcd27eb8 | |
|
|
97ff737c87 | |
|
|
c06e7fa12b | |
|
|
c329d96b56 | |
|
|
ac2c4d9eb1 | |
|
|
db044853c5 | |
|
|
69def4e063 | |
|
|
6fe514c7b7 | |
|
|
f533b2d736 | |
|
|
46c939acc0 | |
|
|
8628b5b38c | |
|
|
72bebcf3d4 | |
|
|
73fde4ea0f | |
|
|
56971980e0 | |
|
|
374871a838 | |
|
|
bf17669650 | |
|
|
4ad16ac26c | |
|
|
fc7db438e0 | |
|
|
c0751f642c | |
|
|
911d79c1ac | |
|
|
af2909cf63 | |
|
|
b207cc5501 | |
|
|
cfa6613a6d | |
|
|
da9181fe18 | |
|
|
19737990e2 | |
|
|
007b6a5a61 | |
|
|
f2429f1685 | |
|
|
b1a44c3488 | |
|
|
4df45b671c | |
|
|
483ca9ccbc | |
|
|
794e15fbc6 | |
|
|
2e3aa374cd | |
|
|
4b4a56eda8 | |
|
|
0e87786f89 | |
|
|
16a44f3561 | |
|
|
1f645c7d26 | |
|
|
ad300c3af8 | |
|
|
d79e659b38 | |
|
|
cf56a944bd | |
|
|
4fc75e51fa | |
|
|
b132af5648 | |
|
|
97c91e6a48 | |
|
|
9828f90095 | |
|
|
09ff453afc | |
|
|
3d8d7bb40e | |
|
|
123d0d39ef | |
|
|
44354be593 | |
|
|
d610c965f2 | |
|
|
63a5139d88 | |
|
|
ce0c49afaf | |
|
|
cab590db95 | |
|
|
2dadbe9e97 | |
|
|
bb8f004731 | |
|
|
b92cf72451 | |
|
|
22c9b06050 | |
|
|
5ac9e457cb | |
|
|
21e0bd30cc | |
|
|
4ad8e54cb6 | |
|
|
1bef660960 | |
|
|
822624a778 | |
|
|
4e7c8a5353 | |
|
|
82aae76571 | |
|
|
84419ce980 | |
|
|
10d459e808 | |
|
|
a5cbd78004 | |
|
|
dc8b7ba029 | |
|
|
d207cf9d92 | |
|
|
7b5d1ec543 | |
|
|
3e56c69f45 | |
|
|
9e62c6fb3f | |
|
|
c0dbfb6c52 | |
|
|
fa8e07c31c | |
|
|
9b4d988272 | |
|
|
b50fab6a97 | |
|
|
4ffb015bdc | |
|
|
bbeb9a860c | |
|
|
9af2328749 | |
|
|
02e882722b | |
|
|
4d818b9169 | |
|
|
fba6cff3e5 | |
|
|
1866003caf | |
|
|
2ea21d458d | |
|
|
b5a9159739 | |
|
|
30de9e7d28 | |
|
|
158b9aa3f2 | |
|
|
eb9815e03d | |
|
|
3adda68ccc | |
|
|
3504f38621 | |
|
|
b65388c77e | |
|
|
61d834a00e | |
|
|
180213653b | |
|
|
1467dbb5a0 | |
|
|
bb76e0e051 | |
|
|
2f81befdd3 | |
|
|
d3338b5a26 | |
|
|
27e650d1c4 | |
|
|
5cafd13e91 | |
|
|
0d8217a1c1 | |
|
|
1ba46ed02a | |
|
|
3576c5eb81 | |
|
|
3d82d79844 | |
|
|
b800ee30ec | |
|
|
7e5ed2906f | |
|
|
70fa8117ec | |
|
|
c1f8e3fe12 | |
|
|
30ed6e0e49 | |
|
|
496c8cc691 | |
|
|
7291f85688 | |
|
|
05b7888756 | |
|
|
1c091f2cb9 | |
|
|
c721c656c5 | |
|
|
7231888100 | |
|
|
8fb0109857 | |
|
|
fd6fd58c39 | |
|
|
d9e691fa4d | |
|
|
c971c576a5 | |
|
|
3d1ff99abc | |
|
|
582451ef58 | |
|
|
80fb933979 | |
|
|
2c7972281d | |
|
|
adf3096c01 | |
|
|
0a1db9a6a6 | |
|
|
64b7883e5d | |
|
|
7b9de194ab | |
|
|
57f56aabb8 | |
|
|
2d38aa3d1b | |
|
|
8e6231b6ab | |
|
|
6d623763d6 | |
|
|
1dbb422167 | |
|
|
790db292d7 | |
|
|
12b2af4460 | |
|
|
2e60758a4f | |
|
|
d4ac7f5a2f | |
|
|
398ec687c3 | |
|
|
6b2e4a709c | |
|
|
cc3d597feb | |
|
|
a452e70d15 | |
|
|
59a4fb8f8c | |
|
|
f302a83b62 | |
|
|
dc5841a10c | |
|
|
167b72b2e1 | |
|
|
567c8c42a0 | |
|
|
8070486586 | |
|
|
7c13dd377f | |
|
|
b85e2637b6 | |
|
|
cc05079507 | |
|
|
0385d36047 | |
|
|
ea98b9a069 | |
|
|
4e045a3e7d | |
|
|
6382da7296 | |
|
|
53880f68c3 | |
|
|
392362ba88 | |
|
|
9f27203d0d |
|
|
@ -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/
|
||||
11
README.md
11
README.md
|
|
@ -18,12 +18,14 @@
|
|||
|
||||
初赛攻坚 Fused Moe Gemm、决赛冲刺 MLA/NSA,全程基于全国产技术栈开发优化,打造自主可控的低成本大模型推理算子底座。
|
||||
|
||||
|
||||
**赛题一相关资料**
|
||||
|
||||
- [赛题一方案:基于国产软件栈的推理前沿算子优化比赛方案](基于国产软件栈大模型推理前沿算子优化/基于国产软件栈的推理前沿算子优化比赛方案.md)
|
||||
- [TileLang MACA 构建指南(模力方舟)](基于国产软件栈大模型推理前沿算子优化/tilelang_maca_build_guide_模力方舟.md)
|
||||
- [TileLang 算子 sample 跑通测试指南](基于国产软件栈大模型推理前沿算子优化/race_tests_run_guide基于tilelang算子sample跑通测试.md)
|
||||
- [赛题一国产软件栈样例说明](基于国产软件栈大模型推理前沿算子优化/race-1-domestic-software-stack-samples.md)
|
||||
- [赛题一国产软件栈样例说明](基于国产软件栈大模型推理前沿算子优化/基于国产软件栈的大模型推理前沿算子优化(TileLang)Sample%20算子获取说明.md)
|
||||
- [TileLang 训练营学习资料](https://www.gitlink.org.cn/ccf-ai-infra/Intro-ops/about)
|
||||
|
||||
### 赛题二:基于 AI Agent 开发范式的国产 GPU 大模型推理算子库优化
|
||||
|
||||
|
|
@ -44,9 +46,12 @@
|
|||
**赛题二相关资料**
|
||||
|
||||
- [赛题二方案:基于 AI Agent 开发范式的国产 GPU 大模型推理算子库优化方案](基于AI%20Agent开发范式的国产GPU大模型推理算子库优化/基于AI%20Agent开发范式的国产GPU大模型算子推理库优化方案.md)
|
||||
- [赛题二选手入口](基于AI%20Agent开发范式的国产GPU大模型推理算子库优化/选手入口.md)
|
||||
- [模力方舟 Agent 部署准备教程](基于AI%20Agent开发范式的国产GPU大模型推理算子库优化/模力方舟Agent部署准备教程.md)
|
||||
- [赛题二说明及资料参考](基于AI%20Agent开发范式的国产GPU大模型推理算子库优化/赛题说明.md)
|
||||
|
||||
### **两个赛题统一使用模力方舟上的镜像PyTorch-Agent / 2.8.0 / Python 3.12 / maca 3.7.1.5**
|
||||
|
||||
## 参赛对象
|
||||
|
||||
面向全国全日制专科、本科、硕博在校生(非在职),40 周岁以下青年科技人才均可参与;
|
||||
|
|
@ -99,6 +104,8 @@
|
|||
|
||||
沐曦开发者社区活动页面,完成新人礼任务,提前熟悉 C500 在线算力使用。https://developer.metax-tech.com/activities/6
|
||||
|
||||
学生可通过登录启悟社区后跳转沐曦开发者社区领取100元算力代金券。(https://developer.metax-tech.com/activities/11)
|
||||
|
||||
报名赛事后还可额外获得300元算力代金券:https://developer.metax-tech.com/activities/17
|
||||
|
||||
使用指南:[模力方舟快速使用 SOP](模力方舟快速使用SOP.md)
|
||||
|
|
@ -109,7 +116,7 @@
|
|||
|
||||
加入沐曦股份“揭榜挂帅”赛题,深耕国产 GPU 算子优化,用技术降低每 Token 推理成本,共建自主可控 AI 算力新生态!
|
||||
|
||||
👇 报名通道:登录挑战杯官网 [www.tiaozhanbei.net](https://link.wtturl.cn/?target=https%3A%2F%2Fwww.tiaozhanbei.net&scene=im&aid=497858&lang=zh) 揭榜挂帅入口报名
|
||||
👇 报名通道:登录挑战杯官网[2026年度中国青年科技创新“揭榜挂帅”擂台赛(学生赛道)](https://2026.tiaozhanbei.net/)揭榜挂帅入口报名
|
||||
|
||||
|
||||
|
||||
|
|
|
|||
File diff suppressed because it is too large
Load Diff
File diff suppressed because it is too large
Load Diff
|
|
@ -1,356 +0,0 @@
|
|||
# FlashInfer 迁移 Baseline 实战
|
||||
|
||||
## 一、教程定位
|
||||
|
||||
本教程是参赛训练课程的 **FlashInfer Baseline 入门**模块,主要帮助用户快速跑通 FlashInfer 的最小可运行流程。完成本教程后,用户应能够完成源码编译、API 调用、正确性测试和 Benchmark 测试,并记录一份 baseline 性能结果,为后续算子优化提供对比基准。
|
||||
|
||||
## 二、学习目标
|
||||
|
||||
1. 理解 FlashInfer Attention Kernel 的基本作用与适用场景;
|
||||
2. 完成 FlashInfer 环境、工具链的准备与源码编译;
|
||||
3. 跑通 BatchDecode、BatchPrefill、MLA 等典型算子的 API 调用示例;
|
||||
4. 完成各算子在不同参数配置下的 Benchmark 测试;
|
||||
5. 输出各算子的 Baseline 性能结果记录表,为后续算子优化提供对比基准。
|
||||
|
||||
|
||||
## 三、适用对象
|
||||
|
||||
适合人群
|
||||
|
||||
* 参赛选手:需要完成 Baseline 入门模块,为后续算子优化做准备
|
||||
* 软件开发者和Vibe Coding开发:希望从事AI相关行业开发,以及用智能体方式来做开发工作
|
||||
* LLM 推理开发者:希望了解 FlashInfer Attention Kernel 的性能表现
|
||||
* 算子优化工程师:希望基于MXMACA软件栈在沐曦国产 GPU 上做算子迁移和优化
|
||||
|
||||
前置基础
|
||||
|
||||
* Python 基础:能够运行和修改 Python 脚本
|
||||
* PyTorch 基础:了解MXMACA软化栈的使用
|
||||
* Linux 命令行:能够使用终端执行命令
|
||||
* 了解 Attention 机制:理解 Q/K/V、KV Cache 等基本概念
|
||||
|
||||
|
||||
## 四、前置准备
|
||||
|
||||
开始实战前,请确认你已经完成以下准备:
|
||||
|
||||
### GPU准备
|
||||
|
||||
* 步骤1:获取算力券[https://developer.metax-tech.com/activities/6](https://developer.metax-tech.com/activities/6)
|
||||
|
||||
* 首次登录需要先进行注册(使用邮箱或者手机号进行注册)
|
||||
|
||||
* 登录成功后进行第二步-邮箱验证,填入自己的邮箱。
|
||||
|
||||
* 第三步,提交申请。
|
||||
|
||||
* 获得兑换码
|
||||
|
||||
|
||||
* 步骤2:兑换算力和登陆平台
|
||||
|
||||
* 访问模力方舟官网:[https://ai.gitee.com/](https://ai.gitee.com/)
|
||||
|
||||
* 进入费用中心 - 算力券 ,点击右上角 “兑换”。
|
||||
|
||||
|
||||
* 步骤3:租用算力
|
||||
|
||||
- 进入算力容器,选择沐曦,租用算力,建议优先选16G显存/32G显存,如下图:
|
||||
* 步骤4:创建实例
|
||||
|
||||
- 基础镜像:maca-pytorch:3.7.1.5-torch2.8-py312-ubuntu24.04-amd64
|
||||
* 步骤5:选择工具-lab进入实例环境
|
||||
|
||||
|
||||
* 步骤6:在JupyterLab Terminal中检查运行环境的配置,确认沐曦 GPU 可见--可以使用`mx-smi`命令查看
|
||||
|
||||
|
||||
### 环境依赖准备
|
||||
|
||||
```bash
|
||||
pip install flashinfer torch pandas numpy
|
||||
```
|
||||
|
||||
### opencode Agent安装
|
||||
|
||||
```shell
|
||||
curl -fsSL https://opencode.ai/install | bash
|
||||
```
|
||||
|
||||
说明:
|
||||
|
||||
自动下载最新版本 自动配置环境 适用于 macOS / Linux / WSL 安装完成后,还会提示你怎么用,OpenCode 包含免费模式,使用方式:
|
||||
|
||||
```shell
|
||||
cd # 进入项目目录
|
||||
opencode
|
||||
```
|
||||
|
||||
### 代码准备
|
||||
|
||||
* 已进入指定项目目录;
|
||||
|
||||
* 准备测试脚本和 Benchmark 脚本。
|
||||
|
||||
|
||||
| 命令 | 说明 |
|
||||
| --- | --- |
|
||||
| `python bench_batch_decode.py` | 运行 Batch Decode 基准测试 |
|
||||
| `python bench_batch_prefill_paged.py` | 运行 Batch Prefill (Paged KV Cache) 基准测试 |
|
||||
| `python bench_batch_prefill_ragged.py` | 运行 Batch Prefill (Ragged KV Cache) 基准测试 |
|
||||
| `python bench_batch_mla.py` | 运行 MLA (Multi-head Latent Attention) 基准测试 |
|
||||
|
||||
## 五、知识预备
|
||||
|
||||
#### LLM推理阶段重要概念:
|
||||
|
||||
* **Prefill阶段**:prefill 阶段是指处理输入 prompt 的阶段
|
||||
|
||||
* 输入:用户一次性给出的完整 prompt,长度为 seq\_len
|
||||
|
||||
* 计算:对 prompt 中的每个 token 并行计算注意力,生成第一个输出 token 及 KV cache
|
||||
|
||||
* 特点:这是**计算密集型(compute-bound)**阶段,因为需要做完整的 seq\_len x seq\_len 注意力矩阵乘法
|
||||
|
||||
* **decode阶段**:
|
||||
|
||||
* 每次只生成 1 个 token,利用 prefill 阶段填充好的 KV cache 做自回归生成
|
||||
|
||||
* **显存带宽密集型(memory-bound**),瓶颈在从显存读取 KV cache 而非计算
|
||||
|
||||
|
||||
prefill = 并行处理用户输入,decode = 逐个生成回答 token
|
||||
|
||||
## 六、项目实践--FlashInfer-Baseline
|
||||
|
||||
### Step 1:检查运行环境
|
||||
|
||||
**目标:** 进入terminal确认当前环境满足本模块运行要求。
|
||||
|
||||
**操作:** 检查 GPU、Python、编译工具和依赖版本。
|
||||
|
||||

|
||||
|
||||
**常见问题:**
|
||||
|
||||
| 问题 | 解决方法 |
|
||||
| --- | --- |
|
||||
| `mx-smi: command not found` | 使用 `python -c "import torch; print(torch.cuda.get_device_name(0))"` |
|
||||
| `No GPUs are available` | 检查 MXMACA 驱动是否正确安装 |
|
||||
| `ModuleNotFoundError: No module named 'flashinfer'` | 执行 `pip install flashinfer` 安装 |
|
||||
|
||||
### Step 2:进入项目目录
|
||||
|
||||
**目标:** 进入本模块所需的源码目录[flashinfer_baseline](https://gitlink.org.cn/metax-maca/op_optimization/tree/master/%E5%9F%BA%E4%BA%8EAI%20Agent%E5%BC%80%E5%8F%91%E8%8C%83%E5%BC%8F%E7%9A%84%E5%9B%BD%E4%BA%A7GPU%E5%A4%A7%E6%A8%A1%E5%9E%8B%E6%8E%A8%E7%90%86%E7%AE%97%E5%AD%90%E5%BA%93%E4%BC%98%E5%8C%96%2Fbaselines%2Fflashinfer_baseline)。
|
||||
|
||||
1. 克隆代码仓库
|
||||
|
||||
```Bash
|
||||
git clone https://gitlink.org.cn/metax-maca/op_optimization.git
|
||||
```
|
||||
|
||||
2. 准备flashinfer_baseline
|
||||
|
||||
在仓库目录 `op_optimization/基于AI Agent开发范式的国产GPU大模型推理算子库优化` 下,找到 `flashinfer_baseline` 文件夹。可以将 `flashinfer_baseline` 整个目录复制到工作目录 `data/` 下。
|
||||
|
||||
**下一步操作:** 切换到 FlashInfer Baseline 项目目录。
|
||||
|
||||
**命令示例:**
|
||||
|
||||
```bash
|
||||
cd /data/flashinfer_baseline
|
||||
ls -la
|
||||
```
|
||||
|
||||
**预期结果:**
|
||||
|
||||
```plaintext
|
||||
bench_common.py
|
||||
bench_batch_decode.py
|
||||
bench_batch_prefill_paged.py
|
||||
bench_batch_prefill_ragged.py
|
||||
bench_batch_mla.py
|
||||
README.md
|
||||
...
|
||||
|
||||
```
|
||||
|
||||
### Step 3:验证项目脚本
|
||||
|
||||
**目标:** 确认所有基准测试脚本可正常执行。
|
||||
|
||||
**操作:** 检查脚本文件是否存在且可读。
|
||||
|
||||
**命令示例:**
|
||||
|
||||
```bash
|
||||
# 检查脚本文件
|
||||
python -c "import os; scripts = ['bench_common.py', 'bench_batch_decode.py', 'bench_batch_prefill_paged.py', 'bench_batch_prefill_ragged.py', 'bench_batch_mla.py']; [print(f'✓ {s}') if os.path.exists(s) else print(f'✗ {s} missing') for s in scripts]"
|
||||
|
||||
# 测试脚本导入
|
||||
python -c "from bench_common import setup_workspace, get_csv_path; print('脚本导入正常')"
|
||||
|
||||
```
|
||||
|
||||
**预期结果:**
|
||||
|
||||
```plaintext
|
||||
✓ bench_common.py
|
||||
✓ bench_batch_decode.py
|
||||
✓ bench_batch_prefill_paged.py
|
||||
✓ bench_batch_prefill_ragged.py
|
||||
✓ bench_batch_mla.py
|
||||
脚本导入正常
|
||||
|
||||
```
|
||||
|
||||
### Step 4:运行单算子 Benchmark
|
||||
|
||||
**目标:** 执行基准测试,获取 Baseline 性能数据。
|
||||
|
||||
**操作:** 运行 BatchDecode 基准测试脚本。
|
||||
|
||||
**命令示例:**
|
||||
|
||||
```bash
|
||||
python bench_batch_decode.py
|
||||
|
||||
```
|
||||
|
||||
**预期结果:**
|
||||
|
||||
```plaintext
|
||||
[BatchDecodeWithPagedKVCacheWrapper] Starting benchmark, total cases: 144
|
||||
[1/144] bs=1, kv_len=512, hd=64: 0.020ms, 51.53 GB/s, 0.82 TFLOPs
|
||||
[2/144] bs=1, kv_len=512, hd=128: 0.023ms, 45.28 GB/s, 0.72 TFLOPs
|
||||
...
|
||||
|
||||
Results saved to BatchDecodeWithPagedKVCacheWrapper_20260525_xxxxxx.csv
|
||||
|
||||
```
|
||||
|
||||
**常见问题:**
|
||||
|
||||
| 问题 | 解决方法 |
|
||||
| --- | --- |
|
||||
| `out of memory` | 减小 batch\_size 或 seq\_len 参数 |
|
||||
| 运行时间过长 | 脚本会自动调整重复次数,耐心等待 |
|
||||
|
||||
### Step 5:查看测试结果
|
||||
|
||||
**目标:** 查看并分析 Benchmark 输出结果。
|
||||
|
||||
**操作:** 读取生成的 CSV 结果文件。
|
||||
|
||||
**命令示例:**
|
||||
|
||||
```bash
|
||||
# 查看最新生成的 CSV 文件
|
||||
ls -lt *.csv | head -1
|
||||
|
||||
# 使用 Python 查看结果
|
||||
python -c "
|
||||
import pandas as pd
|
||||
import glob
|
||||
csv_files = glob.glob('BatchDecodeWithPagedKVCacheWrapper_*.csv')
|
||||
if csv_files:
|
||||
df = pd.read_csv(max(csv_files))
|
||||
print(df[['batch_size', 'seq_len_kv', 'head_dim', 'time_ms', 'bandwidth_GB_s', 'tflops']].head(10))
|
||||
"
|
||||
|
||||
```
|
||||
|
||||
**预期结果:**
|
||||
|
||||
```plaintext
|
||||
batch_size seq_len_kv head_dim time_ms bandwidth_GB_s tflops
|
||||
0 1 512 64 0.020429 51.528822 0.821253
|
||||
1 1 512 128 0.023340 45.278052 0.718833
|
||||
2 1 512 256 0.031949 66.153846 1.050256
|
||||
...
|
||||
|
||||
```
|
||||
|
||||
|
||||
|
||||
## 七、常见问题
|
||||
|
||||
### 环境相关问题
|
||||
|
||||
| 问题 | 原因 | 解决办法 |
|
||||
| --- | --- | --- |
|
||||
| `No GPUs are available` | MXMACA 驱动未安装或 GPU 不可见 | 检查驱动安装,运行 `python -c "import torch; print(torch.cuda.device_count())"` 验证 |
|
||||
| `ModuleNotFoundError: No module named 'flashinfer'` | flashinfer 未安装 | 执行 `pip install flashinfer` |
|
||||
| `out of memory` | GPU 显存不足 | 减小 `batch_size` 或 `seq_len` 参数 |
|
||||
|
||||
### 运行相关问题
|
||||
|
||||
| 问题 | 原因 | 解决办法 |
|
||||
| --- | --- | --- |
|
||||
| Benchmark 运行时间过长 | 参数组合过多, workload 较大 | 耐心等待,脚本会自动调整重复次数 |
|
||||
| `KeyError: 'BatchPrefillWithPagedKVCacheKernel'` | profiler 未捕获目标 kernel | 检查 `target_kernels` 配置是否正确 |
|
||||
| CSV 文件为空 | 测试未正常完成 | 检查 GPU 显存是否充足,重新运行 |
|
||||
|
||||
### 代码相关问题
|
||||
|
||||
| 问题 | 原因 | 解决办法 |
|
||||
| --- | --- | --- |
|
||||
| `ImportError: cannot import name 'xxx' from 'bench_common'` | 函数名拼写错误 | 检查 `bench_common.py` 中的函数名 |
|
||||
| `RuntimeError: error: device-side assert triggered` | 输入参数超出范围 | 检查 `num_qo_heads`、`num_kv_heads`、`head_dim` 配置 |
|
||||
|
||||
### 性能相关问题
|
||||
|
||||
| 问题 | 原因 | 解决办法 |
|
||||
| --- | --- | --- |
|
||||
| TFLOPs 数值异常低 | 工作负载过小,kernel 启动开销占比大 | 增大 `batch_size` 或 `seq_len` |
|
||||
| 带宽数值异常低 | 数据未正确加载到 GPU | 检查 Tensor 是否在 CUDA 设备上 |
|
||||
|
||||
## 八、下一步学习建议
|
||||
|
||||
### 1. 保存你的 Baseline 结果
|
||||
|
||||
将本次运行生成的 CSV 文件妥善保存,后续优化时需要以此作为对比基准。
|
||||
|
||||
```bash
|
||||
# 建议创建 results 目录保存
|
||||
mkdir -p results
|
||||
mv *.csv results/
|
||||
|
||||
```
|
||||
|
||||
### 2. 深入理解 FlashInfer 核心概念
|
||||
|
||||
* 阅读 FlashInfer 官方文档,理解 Paged KV Cache、Ragged KV Cache 的设计理念
|
||||
|
||||
* 学习 MLA (Multi-head Latent Attention) 的原理,了解 DeepSeek 的注意力优化方案
|
||||
|
||||
* 理解 `plan()` 和 `run()` 两阶段设计的作用
|
||||
|
||||
|
||||
### 3. 进入算子优化模块
|
||||
|
||||
参考后续优化模块,学习以下优化技术:
|
||||
|
||||
* **Kernel Tuning**:调整 Block Size、Thread Count 等参数
|
||||
|
||||
* **Memory Optimization**:减少显存占用、优化数据搬运
|
||||
|
||||
* **Compute Optimization**:提升计算效率
|
||||
|
||||
|
||||
### 4. 参考资源
|
||||
|
||||
* FlashInfer 官方仓库:https://github.com/flashinfer-ai/flashinfer
|
||||
|
||||
* FlashInfer 文档:https://flashinfer.ai
|
||||
|
||||
|
||||
### 5. 记录优化过程
|
||||
|
||||
建议维护一份优化日志,记录每次优化的改动和性能变化:
|
||||
|
||||
| 优化项 | 改动内容 | Baseline | 优化后 | 提升比例 |
|
||||
| --- | --- | --- | --- | --- |
|
||||
| 例:调整 block\_size | 16 → 32 | xx ms | xx ms | xx% |
|
||||
| 完成优化后,再次运行本模块的 Benchmark 脚本,对比前后性能变化。 | | | | |
|
||||
File diff suppressed because it is too large
Load Diff
|
|
@ -1,614 +0,0 @@
|
|||
# Fused MoE Baseline 入门:快速跑通最小闭环教程
|
||||
|
||||
## 一、教程定位
|
||||
|
||||
本教程是参赛训练课程的 baseline 入门模块,主要帮助用户快速跑通 Fused MoE 的最小可运行流程。
|
||||
|
||||
完成本教程后,用户应能够完成源码编译、正确性测试和 benchmark 测试,并记录一份 baseline 性能结果,为后续算子优化提供对比基准。
|
||||
|
||||
## 二、学习目标
|
||||
|
||||
完成本模块后,你将能够:
|
||||
|
||||
1. 理解 Fused MoE 推理算子及 int8 量化在大模型推理加速中的核心作用与优化意义;
|
||||
|
||||
2. 完成环境配置和 Fused MoE baseline 源码准备;
|
||||
|
||||
3. 跑通 Fused MoE 的最小闭环示例;
|
||||
|
||||
4. 理解基于 Trition 与 MXMACA C++ 写Fused MOE 算子;
|
||||
|
||||
5. 完成数值正确性测试,即验证 reference 计算、pybind 计算、Triton 计算这三种方式计算结果是否数值完全一致。
|
||||
|
||||
* reference:基于 PyTorch 架构在 CPU 上运行的**数值基准**实现;
|
||||
|
||||
* pybind:将 MXMACA C++ 算子编译并封装为 Python 可调用的动态库,**实现复杂且迁移成本高**;
|
||||
|
||||
* Triton:基于 Python 编写的高效 GPU Kernel,可利用 Agent 自动调优,**开发效率高、易于迁移**;
|
||||
|
||||
* 要求 pybind 和 Triton 结果均与 reference 一致,鼓励参赛选手持续调优 Triton ,使其性能逼近甚至超越 pybind 性能。
|
||||
|
||||
6. 输出 benchmark 结果对比表。
|
||||
|
||||
|
||||
## 三、适用对象
|
||||
|
||||
**本模块适合以下人员:**
|
||||
|
||||
* 参与基于 AI Agent 开发范式的国产 GPU 大模型推理算子库优化比赛的学生;
|
||||
|
||||
* 对 GPU 推理算子性能优化感兴趣的开发者;
|
||||
|
||||
* 需要了解 Fused MoE 推理性能的研究人员。
|
||||
|
||||
|
||||
**学习本模块前,需掌握以下基础知识:**
|
||||
|
||||
* Python、C++ 编程基础;
|
||||
|
||||
* PyTorch 基础;
|
||||
|
||||
* GPU 推理基本概念。
|
||||
|
||||
|
||||
## 四、前置准备
|
||||
|
||||
**开始实战前,请确认你已经完成以下准备:**
|
||||
|
||||
**环境准备:**
|
||||
|
||||
* 已进入赛事专属镜像环境。
|
||||
|
||||
**工具准备:**
|
||||
|
||||
* 已准备 Agent 工具;
|
||||
* 已配置 Token / API Key;
|
||||
* 已确认 Agent 可以正常调用模型。
|
||||
|
||||
**代码准备:**
|
||||
|
||||
* 已获取 Fused MoE Baseline 源码。
|
||||
|
||||
|
||||
## 五、项目实践1-算力平台 baseline 源码部署
|
||||
|
||||
**项目目标:**在算力平台上拉取 Fused MoE 的 baseline 源码,快速跑通最小闭环,并记录一份 baseline 性能结果,为后续算子优化提供对比基准。
|
||||
|
||||
### 步骤 1:检查运行环境
|
||||
|
||||
**目标:**确认当前环境满足本模块运行要求,包括编译器、MXMACA 工具链及 Python 依赖库。
|
||||
|
||||
**操作:**检查 Python、编译工具、MXMACA 编译器及关键 Python 包(numpy、torch、triton)是否存在。
|
||||
|
||||
**命令示例:**
|
||||
|
||||
```apl
|
||||
python --version # 检查Python版本,Python ≥ 3.8
|
||||
g++ --version # 确认 C++ 编译器存在
|
||||
which mxcc # 确认 MACA 编译器存在
|
||||
|
||||
|
||||
# 检查 Python 依赖
|
||||
python - << 'EOF'
|
||||
import sys
|
||||
deps = ["numpy", "torch", "triton"]
|
||||
missing = []
|
||||
for d in deps:
|
||||
try:
|
||||
__import__(d)
|
||||
except ImportError:
|
||||
missing.append(d)
|
||||
if missing:
|
||||
print(f"[ERROR] Missing packages: {missing}")
|
||||
sys.exit(1)
|
||||
else:
|
||||
print("[OK] numpy, torch, triton are installed.")
|
||||
EOF
|
||||
```
|
||||
|
||||
|
||||
**预期结果:**
|
||||
|
||||
* Python 3.12.11
|
||||
|
||||
* g++ (Ubuntu 13.3.0-6ubuntu2~24.04.1) 13.3.0
|
||||
|
||||
* /opt/maca/mxgpu\_llvm/bin/mxcc
|
||||
|
||||
* \[OK\] numpy, torch, triton are installed.
|
||||
|
||||
|
||||
**常见问题:**
|
||||
|
||||
| 报错 | 原因 | 解决办法 |
|
||||
| --- | --- | --- |
|
||||
| `g++:command not found` | 未安装 C++ 编译工具 | `apt update && apt install -y build-essential` |
|
||||
| `Python 3.6.x/ Python 3.7.x` | Python 版本过低 | `conda install python=3.12` (推荐3.10+) |
|
||||
| `ModuleNotFoundError: numpy` | 当前 Python 缺少依赖 | `pip install numpy torch triton` |
|
||||
|
||||
|
||||
### 步骤 2:进入项目目录
|
||||
|
||||
**目标:**进入本模块所需的源码目录:https://gitlink.org.cn/metax-maca/op_optimization/tree/master/%E5%9F%BA%E4%BA%8EAI%20Agent%E5%BC%80%E5%8F%91%E8%8C%83%E5%BC%8F%E7%9A%84%E5%9B%BD%E4%BA%A7GPU%E5%A4%A7%E6%A8%A1%E5%9E%8B%E6%8E%A8%E7%90%86%E7%AE%97%E5%AD%90%E5%BA%93%E4%BC%98%E5%8C%96%2Fbaselines%2Ffused_moe
|
||||
|
||||
1. 克隆代码仓库
|
||||
|
||||
```Bash
|
||||
git clone https://gitlink.org.cn/metax-maca/op_optimization.git
|
||||
```
|
||||
|
||||
2. 准备fused_moe
|
||||
|
||||
在仓库目录 `op_optimization/基于AI Agent开发范式的国产GPU大模型推理算子库优化` 下,找到 `fused_moe` 文件夹。可以将 `fused_moe` 整个目录复制到工作目录 `data/` 下。
|
||||
|
||||
**下一步操作:** 切换到 FlashInfer Baseline 项目目录。
|
||||
|
||||
**操作:**切换到指定项目路径。
|
||||
|
||||
**命令示例:**
|
||||
|
||||
```apl
|
||||
cd data/fused_moe # 这里需要替换成自己的项目路径
|
||||
```
|
||||
|
||||
**预期结果:**
|
||||
|
||||
终端提示符路径显示为 fused\_moe 项目路径:
|
||||
|
||||
* (base) root@0a5a9d0c0f06:/data/fused_moe#
|
||||
|
||||
|
||||
### 步骤 3:pybind 编译
|
||||
|
||||
**目标:**将用 C++ 编写的 fused\_moe 算子编译为 Python 可调用的 pybind 模块。
|
||||
|
||||
**操作:**运行 `fused_moe/scripts/build_fused_moe_i8_tn_pybind.sh` 脚本
|
||||
|
||||
**命令示例:**
|
||||
|
||||
```apl
|
||||
bash scripts/build_fused_moe_i8_tn_pybind.sh
|
||||
```
|
||||
|
||||
切换 Python 环境命令示例:
|
||||
|
||||
```apl
|
||||
[SUCCESS] /data/fused_moe/standalone/fused_moe_i8_tn/build/fused_moe_i8_tn_pybind.sh
|
||||
```
|
||||
|
||||
**预期结果:**
|
||||
|
||||
编译成功无报错,终端显示:
|
||||
|
||||
* \[SUCCESS] /data/fused_moe/standalone/fused_moe_i8_tn/build/fused_moe_i8_tn_pybind.so
|
||||
|
||||
且成功生成 `fused_moe/standalone/fused_moe_i8_tn/build/fused_moe_i8_tn_pybind.cpython-310-x86_64-linux-gnu.so` 文件
|
||||
|
||||
|
||||
**常见问题:**
|
||||
|
||||
|
||||
| 报错 | 原因 | 解决办法 |
|
||||
| --- | --- | --- |
|
||||
| `Python.h: No such file or directory` | Python 头文件路径未找到 | 确认 `PYTHON_BIN` 路径正确,脚本自动探测 `sysconfig.get_path('include')` |
|
||||
| `libpython3.x.so: cannot find` | 链接时找不到 Python 库 | 1、执行 `find $CONDA_PREFIX -name "libpython3*.so*"`查找绝对路径<br>2、将该路径赋值给 `LIBPYTHON_PATH` |
|
||||
| `recompile with -fPIC` | 编译未开启位置无关代码 | 确保 `mxcc`/ `g++`编译参数中有 `-fPIC` |
|
||||
| `permission denied` | 无脚本执行权限 | `chmod +x scripts/*.sh` |
|
||||
| `undefined reference to Py_...` | Python 版本不匹配 | 确认编译脚本中`PYTHON_BIN`路径与当前运行的 Python 环境完全一致 |
|
||||
|
||||
### 步骤 4:正确性测试
|
||||
|
||||
**目标:**验证 reference 计算、pybind 计算、Triton 计算这三种方式计算结果的数值是否完全一致。
|
||||
|
||||
**操作:**运行 `fused_moe/scripts/run_fused_moe_i8_tn_pybind_test.sh` 脚本
|
||||
|
||||
**命令示例:**
|
||||
|
||||
```apl
|
||||
bash scripts/run_fused_moe_i8_tn_pybind_test.sh --backend all # 运行全部计算方式
|
||||
|
||||
# --backend:选择计算方式
|
||||
# 只测 pybind:
|
||||
bash scripts/run_fused_moe_i8_tn_pybind_test.sh --backend pybind
|
||||
# 只测 triton:
|
||||
bash scripts/run_fused_moe_i8_tn_pybind_test.sh --backend triton
|
||||
# 只测 reference:
|
||||
bash scripts/run_fused_moe_i8_tn_pybind_test.sh --backend reference
|
||||
```
|
||||
|
||||
|
||||
**预期结果:**
|
||||
|
||||
编译成功无报错,输出示例如下:
|
||||
|
||||
> pybind:fused\_moe\_i8\_tn\_topk1 passed: rows=256, cols=128, sample C\[0\]=0.69531, C\[last\]=-0.44531
|
||||
> pybind:fused\_moe\_i8\_tn\_topk2 passed: rows=512, cols=128, sample C\[0\]=-0.57813, C\[last\]=-0.49805
|
||||
> pybind:fused\_moe\_i8\_tn\_topk3 passed: rows=384, cols=128, sample C\[0\]=-1.08594, C\[last\]=-0.33594
|
||||
|
||||
> reference:fused\_moe\_i8\_tn\_topk1 passed: rows=256, cols=128, sample C\[0\]=0.6934, C\[last\]=-0.4451
|
||||
> reference:fused\_moe\_i8\_tn\_topk2 passed: rows=512, cols=128, sample C\[0\]=-0.5768, C\[last\]=-0.4975
|
||||
> reference:fused\_moe\_i8\_tn\_topk3 passed: rows=384, cols=128, sample C\[0\]=-1.0875, C\[last\]=-0.3362
|
||||
|
||||
> triton:fused\_moe\_i8\_tn\_topk1 passed: rows=256, cols=128, sample C\[0\]=0.69337, C\[last\]=-0.44513
|
||||
> triton:fused\_moe\_i8\_tn\_topk2 passed: rows=512, cols=128, sample C\[0\]=-0.57678, C\[last\]=-0.49749
|
||||
> triton:fused\_moe\_i8\_tn\_topk3 passed: rows=384, cols=128, sample C\[0\]=-1.08748, C\[last\]=-0.33618
|
||||
|
||||
|
||||
**结果解释:**
|
||||
|
||||
* “pybind/reference/Triton”:三种计算方式;
|
||||
|
||||
* “fused\_moe\_i8\_tn\_topk1/2/3 passed”:测试算子通过数值校验,数值误差在允许范围内且无明显异常,否则会报错 FAILED;
|
||||
|
||||
* ”rows=... , cols=...“:输出 Tensor 的行和列的大小;
|
||||
|
||||
* ”sample C\[0\]=... , C\[last\]=...“:首尾采样值,用于辅助定位数值偏差,不作为精度判定依据。
|
||||
|
||||
|
||||
**常见问题:**
|
||||
|
||||
|
||||
| 报错 | 原因 | 解决办法 |
|
||||
| :--- | :--- | :--- |
|
||||
| `ModuleNotFoundError: fused_moe_i8_tn_pybind` | pybind 模块未编译或未加入 `PYTHONPATH` | 回到步骤 3,确认 `.so` 已生成;执行 `export PYTHONPATH=/root/Project/fused_moe:$PYTHONPATH` |
|
||||
| `FAILED: max abs diff too large` | 数值误差超过阈值 | 检查 scale 是否应用位置错误;确认 TopK 索引与权重是否一致 |
|
||||
| `FAILED: shape mismatch` | 输出张量形状不一致 | 检查 Token Permute / Unpermute 逻辑;确认 expert 维度对齐 |
|
||||
| `FAILED: NaN or Inf detected` | 溢出或未初始化内存 | 检查 INT8 乘加是否溢出;确认 GEMM 输出是否反量化 |
|
||||
| 终端长时间无输出 | Kernel 死锁或 Launch 失败 | 减小测试 shape;检查是否触发 MACA 硬件限制 |
|
||||
|
||||
### 步骤5:性能测试
|
||||
|
||||
**目标:**输出 benchmark 结果对比表
|
||||
|
||||
**操作:**运行 `fused_moe/scripts/run_fused_moe_i8_tn_benchmark.sh` 脚本
|
||||
|
||||
**命令示例:**
|
||||
|
||||
```apl
|
||||
bash scripts/run_fused_moe_i8_tn_benchmark.sh --backend all --warmup 5 --iters 20
|
||||
# --backend:选择计算方式
|
||||
# --warmup:设置预热次数
|
||||
# --iters:设置迭代次数
|
||||
```
|
||||
|
||||
|
||||
**预期结果:**
|
||||
|
||||
编译成功无报错,输出示例如下:
|
||||
|
||||
> pybind:fused\_moe\_i8\_tn\_topk1 benchmark: avg\_ms=0.308978, TOPS=0.027149, warmup=5, iters=20
|
||||
> pybind:fused\_moe\_i8\_tn\_topk2 benchmark: avg\_ms=0.304500, TOPS=0.055098, warmup=5, iters=20
|
||||
> pybind:fused\_moe\_i8\_tn\_topk3 benchmark: avg\_ms=0.297775, TOPS=0.042256, warmup=5, iters=20
|
||||
|
||||
> reference:fused\_moe\_i8\_tn\_topk1 benchmark: avg\_ms=1685.43, TOPS=0.000005, warmup=5, iters=20
|
||||
> reference:fused\_moe\_i8\_tn\_topk2 benchmark: avg\_ms=3384.52, TOPS=0.000005, warmup=5, iters=20
|
||||
> reference:fused\_moe\_i8\_tn\_topk3 benchmark: avg\_ms=2532.14, TOPS=0.000005, warmup=5, iters=20
|
||||
|
||||
> triton:fused\_moe\_i8\_tn\_topk1 benchmark: avg\_ms=19.013421, TOPS=0.000441, warmup=5, iters=20
|
||||
> triton:fused\_moe\_i8\_tn\_topk2 benchmark: avg\_ms=16.745914, TOPS=0.001002, warmup=5, iters=20
|
||||
> triton:fused\_moe\_i8\_tn\_topk3 benchmark: avg\_ms=19.630328, TOPS=0.000641, warmup=5, iters=20
|
||||
|
||||
|
||||
**结果解释:**
|
||||
|
||||
* “pybind/reference/Triton”:三种计算方式;
|
||||
|
||||
* “fused\_moe\_i8\_tn\_topk1/2/3”:分别对应选择前 1 / 2 / 3 个专家场景下的 MoE 算子;
|
||||
|
||||
* “avg\_ms”:平均算子执行耗时(毫秒),这里不计算预热时间,只计算正式迭代的<E4BBA3><E79A84><EFBFBD>间<EFBFBD><E997B4><EFBFBD>
|
||||
|
||||
* “TOPS”:Tera Operations Per Second,本次 MoE 算子的总运算量 / 实际耗<E99985><E88097><EFBFBD>;
|
||||
|
||||
* “warmup=5, iters=20”:预热轮数和正式迭代数。
|
||||
|
||||
|
||||
**常见问题:**
|
||||
|
||||
|
||||
| 报错 | 原因 | 解决办法 |
|
||||
| --- | --- | --- |
|
||||
| `ModuleNotFoundError: fused_moe_i8_tn_pybind` | pybind 模块未编译或未加入 `PYTHONPATH` | 回到步骤 3,确认 `.so`已生成;执行 `export PYTHONPATH=/root/Project/fused_moe:$PYTHONPATH` |
|
||||
| 终端长时间无输出 | Kernel 死锁或 MACA 驱动异常 | 减小测试 shape;重启容器或设备 |
|
||||
| avg\_ms 异常抖动(±50%) | 其他进程占用 GPU | 关闭其他占用显存的进程,单机单任务运行 |
|
||||
|
||||
## 六、项目实践2-Kernel Swift 智能算子迁移系统自动调优
|
||||
|
||||
系统链接:[https://deeplink.org.cn/kernelswift/task](https://deeplink.org.cn/kernelswift/task)
|
||||
|
||||
**项目目标:**基于 KernelSwift 智能算子迁移系统,对 Fused MoE 算子进行在线自动调优。通过输入算子代码,一键生成适配沐曦硬件的高性能实现,高效完成算子优化与全流程追踪。
|
||||
|
||||
### 步骤1:复用算子广场的 fused\_moe 算子进行二次优化
|
||||
|
||||
**目标:**通过提交算子广场的 fused\_moe 算子代码发起自动优化流程,实现二次优化
|
||||
|
||||
**操作:**
|
||||
|
||||
1. 进入算子广场:点击左侧导航栏 【算子广场】,进入算子列表页
|
||||
|
||||
搜索 fused\_moe 算子,复制 `input_code.py` 代码,也可直接复制以下代码:
|
||||
|
||||
```python
|
||||
import torch
|
||||
import torch.nn as nn
|
||||
import torch.nn.functional as F
|
||||
|
||||
|
||||
class Model(nn.Module):
|
||||
"""
|
||||
Reference PyTorch MoE forward (no fused kernels).
|
||||
Expects inputs:
|
||||
hidden_states: (M, in_size)
|
||||
w1: (E, hidden_size, in_size) where hidden_size = 2 * up_dim
|
||||
w2: (E, out_size, up_dim)
|
||||
topk_weights: (M, top_k)
|
||||
topk_idx: (M, top_k)
|
||||
top_k: int
|
||||
renormalize: bool
|
||||
"""
|
||||
|
||||
def __init__(self):
|
||||
super().__init__()
|
||||
|
||||
def forward(
|
||||
self,
|
||||
hidden_states: torch.Tensor,
|
||||
w1: torch.Tensor,
|
||||
w2: torch.Tensor,
|
||||
topk_weights: torch.Tensor,
|
||||
topk_idx: torch.Tensor,
|
||||
top_k: int,
|
||||
renormalize: bool = True,
|
||||
) -> torch.Tensor:
|
||||
if renormalize:
|
||||
topk_weights = topk_weights / topk_weights.sum(dim=-1, keepdim=True)
|
||||
|
||||
seq_len = hidden_states.size(0)
|
||||
out_size = w2.size(1)
|
||||
output = hidden_states.new_zeros(seq_len, out_size)
|
||||
num_experts = w1.size(0)
|
||||
|
||||
# Accumulate expert contributions
|
||||
for eid in range(num_experts):
|
||||
token_idx, k_idx = torch.where(topk_idx == eid)
|
||||
if token_idx.numel() == 0:
|
||||
continue
|
||||
gate_proj, up_proj = w1[eid].chunk(2, dim=0)
|
||||
down_proj = w2[eid]
|
||||
tmp = F.linear(hidden_states[token_idx], gate_proj)
|
||||
tmp = F.silu(tmp) * F.linear(hidden_states[token_idx], up_proj)
|
||||
tmp = F.linear(tmp, down_proj)
|
||||
tmp = tmp * topk_weights[token_idx, k_idx, None]
|
||||
output.index_add_(0, token_idx, tmp.to(output.dtype))
|
||||
return output
|
||||
|
||||
|
||||
# Hyperparameters
|
||||
seq_len = 128
|
||||
in_size = 128
|
||||
hidden_size = 256 # 2 * up_dim
|
||||
out_size = 128
|
||||
num_experts = 32
|
||||
top_k = 4
|
||||
|
||||
dtype = torch.float16
|
||||
|
||||
def get_inputs():
|
||||
hidden_states = (torch.rand(seq_len, in_size, dtype=dtype) - 0.5) / 2
|
||||
w1 = (torch.rand(num_experts, hidden_size, in_size, dtype=dtype) - 0.5) / 2
|
||||
w2 = (torch.rand(num_experts, out_size, hidden_size//2, dtype=dtype) - 0.5) / 2
|
||||
routing_logits = (torch.rand(seq_len, num_experts, dtype=dtype) - 0.5) / 2
|
||||
routing_weights = torch.softmax(routing_logits, dim=-1, dtype=torch.float32)
|
||||
topk_weights, topk_idx = torch.topk(routing_weights, top_k, dim=-1)
|
||||
return [hidden_states, w1, w2, topk_weights, topk_idx, top_k, True]
|
||||
|
||||
def get_init_inputs():
|
||||
return []
|
||||
```
|
||||
|
||||
2. 进入新建任务页:点击左侧导航栏【新建任务】 ,进入算子提交页面。
|
||||
|
||||
3. 编写算子代码:在 `model.py` 编辑器中输入刚刚复制的 fused\_moe 算子代码。
|
||||
|
||||
如果想自行编写算子代码,需严格遵循标准格式规范:输入代码必须包含 `class Model` 定义算子实现,`get_init_inputs` 和 `get_inputs` 定义测试用例,确保优化过程可验证算子正确性。
|
||||
|
||||
4. 配置优化参数
|
||||
|
||||
* 指定任务名称:支持字母、下划线、数字组合,示例:fused\_moe\_01
|
||||
|
||||
* 选择适配硬件:算子需要适配的目标硬件厂商及型号,建议:沐曦
|
||||
|
||||
* 最大演化轮次:优化算法迭代次数,取值范围40-400,建议默认40,复杂算法可提高至100+
|
||||
|
||||
5. 提交优化任务:点击右下角 \[优化\] 按钮,系统将提交任务并进入 \[生成中\] 状态
|
||||
|
||||
|
||||

|
||||
|
||||
完成上述步骤将看到如下界面:
|
||||
|
||||

|
||||
|
||||
### 步骤2:任务查看与结果管理
|
||||
|
||||
**目标:** 在新建优化任务后可追踪任务进度,获取优化结果
|
||||
|
||||
**操作:**
|
||||
|
||||
1. 查看任务列表:点击左侧【任务查看】,可看到所有提交的优化任务
|
||||
|
||||
* 任务状态:排队中、环境初始化、算子预编译、精度验证、性能调优、已完成、失败
|
||||
|
||||
* 任务信息:任务名称、进度、创建时间、适配硬件
|
||||
|
||||
* 操作按钮:查看详情、删除任务
|
||||
|
||||
|
||||

|
||||
|
||||
2. 追踪任务进度:当前任务状态为【运行中】时,点击任务列表中的【查看详情】按钮,追踪任务进度:
|
||||
|
||||
* 左侧:原始算子代码(输入的 `input_code.py`)
|
||||
|
||||
* 右侧:任务进度条,包含以下阶段:
|
||||
|
||||
1. 环境初始化:准备目标硬件编译环境
|
||||
|
||||
2. 算子预编译:验证算子代码是可正常编译
|
||||
|
||||
3. 精度验证:验证优化后算子输出与原始算子误差的可接受范围
|
||||
|
||||
4. 性能调优:按设定的演化轮次迭代优化算子性能
|
||||
|
||||
|
||||
* 顶部:任务名称、创建/更新时间、适配硬件、当前轮次进度
|
||||
|
||||
|
||||

|
||||
|
||||
3. 获取优化结果:当前任务状态为【已完成】时,可在详情页查看优化结果:
|
||||
|
||||
* 优化后算子代码支持一键复制
|
||||
|
||||
* 算子加速比(基准耗时 / 优化后耗时)、性能数据(如延迟、吞吐量)
|
||||
|
||||
* 可点击【Diff 对比】查看优化前后代码差异,理解性能提升逻辑
|
||||
|
||||
|
||||

|
||||
|
||||
4. 任务异常处理
|
||||
|
||||
* 任务失败:查看错误日志,常见原因包括代码不符合规范、测试用例错误、硬件适配问题,修改后重新提交任务;
|
||||
|
||||
* 排队时间长:可调整提交时间,或联系平台管理员确认资源状态。
|
||||
|
||||
|
||||
## 七、Agent使用说明
|
||||
|
||||
在本模块中,Agent可以帮助你完成以下任务:
|
||||
|
||||
1. **环境检查**
|
||||
|
||||
```plaintext
|
||||
我正在算力平台部署 fused_moe_baseline 源码。
|
||||
需要的环境信息如下:
|
||||
- Python 3.12
|
||||
- g++ 13.3.0
|
||||
- mxcc 已安装
|
||||
- numpy / torch / triton 已安装
|
||||
|
||||
请帮我确认:
|
||||
1. 当前环境是否满足编译与运行要求?
|
||||
2. 是否有潜在的不兼容风险(如 Python 与 libpython 版本)?
|
||||
```
|
||||
|
||||
2. **运行测试**
|
||||
|
||||
```plaintext
|
||||
请帮我运行 scripts/run_fused_moe_i8_tn_pybind_test.sh 脚本
|
||||
```
|
||||
|
||||
3. **分析结果**
|
||||
|
||||
```plaintext
|
||||
这是性能测试结果:
|
||||
pybind: avg_ms=0.30, TOPS=0.027
|
||||
triton: avg_ms=19.01, TOPS=0.0004
|
||||
reference: avg_ms=1685, TOPS=0.000005
|
||||
|
||||
请分析:
|
||||
1. 为什么 pybind 比 Triton 快这么多?
|
||||
2. TOPS 指标是否可信?
|
||||
3. 当前结果是否已经具备提交价值?
|
||||
```
|
||||
|
||||
4. **报错检查**
|
||||
|
||||
```plaintext
|
||||
编译 pybind 时出现以下错误:
|
||||
/usr/bin/ld: cannot find -lpython3.10
|
||||
|
||||
已知:
|
||||
- 使用的是 Conda Python 3.10
|
||||
- mxcc 编译正常
|
||||
|
||||
请一步一步告诉我:
|
||||
1. 错误原因是什么?
|
||||
2. 如何用 find 命令定位 libpython3.10.so?
|
||||
3. 如何在 build_fused_moe_i8_tn_pybind.sh 中正确指定路径?
|
||||
```
|
||||
|
||||
5. **代码理解**
|
||||
|
||||
```plaintext
|
||||
请帮我梳理释 benchmark_fused_moe_i8_tn.py 代码整体框架
|
||||
```
|
||||
|
||||
6. **KernelSwift 系统搜索算子**
|
||||
|
||||
|
||||
```plaintext
|
||||
请帮我在算子广场检索 fused_moe 算子
|
||||
```
|
||||
|
||||
## 八、常见问题与注意事项
|
||||
|
||||
### 算力平台 basline 源码部署项目:
|
||||
|
||||
1. 环境准备与依赖问题
|
||||
|
||||
* 确保算力平台已正确安装 Python 和 C++、MACA 编译器及相关运行时库,避免因环境缺失导致编译失败;
|
||||
|
||||
* 镜像环境使用 Conda Python 作为默认运行环境,避免系统 Python 与 Conda Python 混用,防止 `Python.h`或 `libpython`路径错误。
|
||||
|
||||
2. pybind 编译与链接
|
||||
|
||||
* 若`Python.h not found`,请检查脚本中`PYTHON_INCLUDE`是否指向当前 Python 的 `include`目录;
|
||||
|
||||
* 若`libpython not found`,请直接指定 Conda 下的`**libpython3.x.so**`绝对路径,避免链接系统静态库;
|
||||
|
||||
* 编译 `pybind`模块时,务必开启 `-fPIC`,否则会出现 `recompile with -fPIC`错误。
|
||||
|
||||
3. 性能测试建议
|
||||
|
||||
* benchmark 应在关闭其他占用 GPU 的任务后执行,避免干扰性能数据;
|
||||
|
||||
* 多次运行取平均值,避免单次抖动影响结果;
|
||||
|
||||
* 性能对比应基于相同随机种子、相同 shape、相同 TopK、相同 batch size的条件下进行,降低误差。
|
||||
|
||||
|
||||
### Kernel Swift 智能算子迁移系统自动调优项目:
|
||||
|
||||
1. 代码规范问题
|
||||
|
||||
输入代码需符合以下标准格式:
|
||||
|
||||
* `class Model`,表示待优化的算子实现;
|
||||
|
||||
* `def get_init_inputs`,表示 module init 的输入测试样例;
|
||||
|
||||
* `def get_inputs`,表示 module forward 的输入测试样例。
|
||||
|
||||
2. 性能优化建议
|
||||
|
||||
* 对于复杂算子,可适当提高最大演化轮次(如 100-200),获得更高加速比;
|
||||
|
||||
* 优先选择算子广场中已有优化案例的算子类型,降低适配失败概率。
|
||||
|
||||
3. 硬件适配问题
|
||||
|
||||
* 提交任务前确认目标硬件支持的算子类型;
|
||||
|
||||
* 优化失败时,可尝试更换适配硬件,或调整算子实现逻辑。
|
||||
|
||||
|
||||
## 九、下一步学习建议
|
||||
|
||||
完成本模块后,建议继续学习以下内容:
|
||||
|
||||
1. **研读 fused\_moe源码:**理解代码的底层逻辑,可尝试修改 `build_fused_moe_i8_tn_pybind.sh`中的编译参数,观察其对 `avg_ms` 的影响;
|
||||
|
||||
2. **算子优化基础:**了解如何分析 Kernel 性能瓶颈;
|
||||
|
||||
3. **性能对比分析:**将 baseline 结果与优化后的结果进行对比分析,明确后续优化方向。
|
||||
Binary file not shown.
Binary file not shown.
Binary file not shown.
|
|
|
@ -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
|
||||
|
|
|
@ -0,0 +1,119 @@
|
|||
#include <stdint.h>
|
||||
#include <cuda_bf16.h>
|
||||
#include <cuda_runtime.h>
|
||||
|
||||
#define HEAD_DIM 128
|
||||
|
||||
__global__ void paged_attention_kernel(
|
||||
const __nv_bfloat16* q,
|
||||
const __nv_bfloat16* k_cache_paged,
|
||||
const __nv_bfloat16* v_cache_paged,
|
||||
__nv_bfloat16* output,
|
||||
const int32_t* cache_seqlens,
|
||||
const int32_t* block_table,
|
||||
int64_t batch_size,
|
||||
int64_t seqlen_q,
|
||||
int64_t num_heads,
|
||||
int64_t num_heads_k,
|
||||
int64_t headdim,
|
||||
int64_t page_block_size,
|
||||
int64_t blocks_per_batch)
|
||||
{
|
||||
int batch_idx = blockIdx.x / num_heads;
|
||||
int head_idx = blockIdx.x % num_heads;
|
||||
if (batch_idx >= batch_size || head_idx >= num_heads) return;
|
||||
|
||||
int seqlen = cache_seqlens[batch_idx];
|
||||
int tid = threadIdx.x;
|
||||
|
||||
// 加载对应 head 的 query 元素
|
||||
int64_t q_offset = ((batch_idx * seqlen_q + 0) * num_heads + head_idx) * headdim;
|
||||
float q_val = __bfloat162float(q[q_offset + tid]);
|
||||
|
||||
// Online safe softmax 状态
|
||||
float max_val = -1e38f;
|
||||
float sum_exp = 0.0f;
|
||||
float out_acc = 0.0f;
|
||||
float scale = 1.0f / sqrtf(static_cast<float>(headdim));
|
||||
|
||||
// 静态共享内存,避免动态分配可能带来的兼容性问题
|
||||
__shared__ float s_score[HEAD_DIM];
|
||||
|
||||
for (int token = 0; token < seqlen; ++token) {
|
||||
int page_idx = token / page_block_size;
|
||||
int page_offset = token % page_block_size;
|
||||
int physical_block = block_table[batch_idx * blocks_per_batch + page_idx];
|
||||
|
||||
// 读取 key 元素
|
||||
const __nv_bfloat16* k_ptr = k_cache_paged
|
||||
+ (physical_block * page_block_size + page_offset) * (num_heads_k * headdim)
|
||||
+ head_idx * headdim;
|
||||
float k_val = __bfloat162float(k_ptr[tid]);
|
||||
|
||||
// 点积 -> 共享内存归约
|
||||
s_score[tid] = q_val * k_val;
|
||||
__syncthreads();
|
||||
|
||||
for (int stride = HEAD_DIM >> 1; stride > 0; stride >>= 1) {
|
||||
if (tid < stride) {
|
||||
s_score[tid] += s_score[tid + stride];
|
||||
}
|
||||
__syncthreads();
|
||||
}
|
||||
float score = s_score[0] * scale;
|
||||
|
||||
// 更新 softmax 状态
|
||||
float new_max = fmaxf(max_val, score);
|
||||
float rescale = expf(max_val - new_max);
|
||||
sum_exp = sum_exp * rescale + expf(score - new_max);
|
||||
out_acc = out_acc * rescale;
|
||||
max_val = new_max;
|
||||
|
||||
// 读取 value 元素,并累加(用最新 max 的权重)
|
||||
const __nv_bfloat16* v_ptr = v_cache_paged
|
||||
+ (physical_block * page_block_size + page_offset) * (num_heads_k * headdim)
|
||||
+ head_idx * headdim;
|
||||
float v_val = __bfloat162float(v_ptr[tid]);
|
||||
out_acc += expf(score - max_val) * v_val;
|
||||
|
||||
__syncthreads(); // 确保下次迭代共享内存可安全复用
|
||||
}
|
||||
|
||||
if (seqlen > 0) {
|
||||
out_acc /= sum_exp;
|
||||
} else {
|
||||
out_acc = 0.0f;
|
||||
}
|
||||
|
||||
int64_t out_offset = ((batch_idx * seqlen_q + 0) * num_heads + head_idx) * headdim + tid;
|
||||
output[out_offset] = __float2bfloat16(out_acc);
|
||||
}
|
||||
|
||||
extern "C" void run_kernel(
|
||||
const __nv_bfloat16* q,
|
||||
const __nv_bfloat16* k_cache_paged,
|
||||
const __nv_bfloat16* v_cache_paged,
|
||||
__nv_bfloat16* output,
|
||||
const int32_t* cache_seqlens,
|
||||
const int32_t* block_table,
|
||||
int64_t batch_size,
|
||||
int64_t seqlen_k,
|
||||
int64_t seqlen_q,
|
||||
int64_t num_heads,
|
||||
int64_t num_heads_k,
|
||||
int64_t headdim,
|
||||
int64_t page_block_size,
|
||||
int64_t num_blocks,
|
||||
int64_t causal)
|
||||
{
|
||||
int64_t blocks_per_batch = num_blocks / batch_size;
|
||||
dim3 grid(batch_size * num_heads);
|
||||
dim3 block(HEAD_DIM);
|
||||
|
||||
paged_attention_kernel<<<grid, block>>>(
|
||||
q, k_cache_paged, v_cache_paged, output,
|
||||
cache_seqlens, block_table,
|
||||
batch_size, seqlen_q, num_heads, num_heads_k, headdim,
|
||||
page_block_size, blocks_per_batch
|
||||
);
|
||||
}
|
||||
|
|
@ -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)
|
||||
|
|
@ -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, # 性能瓶颈:禁用流水线并行
|
||||
)
|
||||
|
|
@ -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
|
||||
|
|
|
@ -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
|
||||
|
|
|
@ -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
|
||||
|
|
|
@ -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
|
||||
|
Binary file not shown.
|
Binary file not shown.
|
Binary file not shown.
|
Binary file not shown.
|
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
|
|
@ -0,0 +1,209 @@
|
|||
# 示例冒烟代码
|
||||
|
||||
```c++
|
||||
#include <stdint.h>
|
||||
|
||||
#include <cuda_bf16.h>
|
||||
#include <cuda_runtime.h>
|
||||
|
||||
#include <math.h>
|
||||
|
||||
namespace {
|
||||
|
||||
__device__ __forceinline__ float warp_sum(float x) {
|
||||
for (int offset = 16; offset > 0; offset >>= 1) {
|
||||
x += __shfl_down_sync(0xffffffffu, x, offset);
|
||||
}
|
||||
return __shfl_sync(0xffffffffu, x, 0);
|
||||
}
|
||||
|
||||
__global__ void ragged_prefill_smoke_kernel(
|
||||
const __nv_bfloat16* __restrict__ q,
|
||||
const __nv_bfloat16* __restrict__ k,
|
||||
const __nv_bfloat16* __restrict__ v,
|
||||
__nv_bfloat16* __restrict__ output,
|
||||
const int32_t* __restrict__ qo_indptr,
|
||||
const int32_t* __restrict__ kv_indptr,
|
||||
int64_t batch_size,
|
||||
int64_t seq_len,
|
||||
int64_t num_qo_heads,
|
||||
int64_t num_kv_heads,
|
||||
int64_t head_dim_qk,
|
||||
int64_t head_dim_vo,
|
||||
int64_t causal,
|
||||
int64_t exact_len) {
|
||||
const int lane = threadIdx.x & 31;
|
||||
const int warp_id = threadIdx.x >> 5;
|
||||
const int warps_per_block = blockDim.x >> 5;
|
||||
|
||||
int64_t work = static_cast<int64_t>(blockIdx.x) * warps_per_block + warp_id;
|
||||
const int64_t total = batch_size * exact_len * num_qo_heads;
|
||||
if (work >= total) return;
|
||||
|
||||
const int64_t qo_head = work % num_qo_heads;
|
||||
work /= num_qo_heads;
|
||||
const int64_t q_pos = work % exact_len;
|
||||
const int64_t batch = work / exact_len;
|
||||
|
||||
const int64_t qo_begin = qo_indptr[batch];
|
||||
const int64_t qo_len = qo_indptr[batch + 1] - qo_begin;
|
||||
if (q_pos >= qo_len) return;
|
||||
|
||||
const int64_t kv_begin = kv_indptr[batch];
|
||||
const int64_t kv_len = kv_indptr[batch + 1] - kv_begin;
|
||||
int64_t visible = kv_len;
|
||||
if (causal) {
|
||||
visible = kv_len - qo_len + q_pos + 1;
|
||||
if (visible < 0) visible = 0;
|
||||
if (visible > kv_len) visible = kv_len;
|
||||
}
|
||||
|
||||
const int64_t group = num_qo_heads / num_kv_heads;
|
||||
const int64_t kv_head = qo_head / group;
|
||||
const int64_t q_row = qo_begin + q_pos;
|
||||
const float scale = rsqrtf(static_cast<float>(head_dim_qk));
|
||||
|
||||
const __nv_bfloat16* q_ptr = q + (q_row * num_qo_heads + qo_head) * head_dim_qk;
|
||||
float qv[4];
|
||||
float acc[4];
|
||||
for (int i = 0; i < 4; ++i) {
|
||||
const int d = lane + i * 32;
|
||||
qv[i] = (d < head_dim_qk) ? __bfloat162float(q_ptr[d]) : 0.0f;
|
||||
acc[i] = 0.0f;
|
||||
}
|
||||
|
||||
float m = -1.0e20f;
|
||||
float l = 0.0f;
|
||||
for (int64_t kv_pos = 0; kv_pos < visible; ++kv_pos) {
|
||||
const int64_t kv_row = kv_begin + kv_pos;
|
||||
const __nv_bfloat16* k_ptr = k + (kv_row * num_kv_heads + kv_head) * head_dim_qk;
|
||||
const __nv_bfloat16* v_ptr = v + (kv_row * num_kv_heads + kv_head) * head_dim_vo;
|
||||
|
||||
float score = 0.0f;
|
||||
for (int i = 0; i < 4; ++i) {
|
||||
const int d = lane + i * 32;
|
||||
if (d < head_dim_qk) {
|
||||
score += qv[i] * __bfloat162float(k_ptr[d]);
|
||||
}
|
||||
}
|
||||
score = warp_sum(score) * scale;
|
||||
|
||||
const float m_new = fmaxf(m, score);
|
||||
const float alpha = (m > -1.0e19f) ? __expf(m - m_new) : 0.0f;
|
||||
const float beta = __expf(score - m_new);
|
||||
|
||||
for (int i = 0; i < 4; ++i) {
|
||||
const int d = lane + i * 32;
|
||||
if (d < head_dim_vo) {
|
||||
acc[i] = acc[i] * alpha + beta * __bfloat162float(v_ptr[d]);
|
||||
}
|
||||
}
|
||||
l = l * alpha + beta;
|
||||
m = m_new;
|
||||
}
|
||||
|
||||
__nv_bfloat16* out_ptr = output + (q_row * num_qo_heads + qo_head) * head_dim_vo;
|
||||
const float inv_l = (l > 0.0f) ? (1.0f / l) : 0.0f;
|
||||
for (int i = 0; i < 4; ++i) {
|
||||
const int d = lane + i * 32;
|
||||
if (d < head_dim_vo) {
|
||||
out_ptr[d] = __float2bfloat16(acc[i] * inv_l);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
__global__ void prefix_mean_kernel(
|
||||
const __nv_bfloat16* __restrict__ v,
|
||||
__nv_bfloat16* __restrict__ output,
|
||||
const int32_t* __restrict__ qo_indptr,
|
||||
const int32_t* __restrict__ kv_indptr,
|
||||
int64_t batch_size,
|
||||
int64_t seq_len,
|
||||
int64_t num_qo_heads,
|
||||
int64_t num_kv_heads,
|
||||
int64_t head_dim_vo) {
|
||||
int64_t work = static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
|
||||
const int64_t total = batch_size * num_kv_heads * head_dim_vo;
|
||||
if (work >= total) return;
|
||||
|
||||
const int64_t d = work % head_dim_vo;
|
||||
work /= head_dim_vo;
|
||||
const int64_t kv_head = work % num_kv_heads;
|
||||
const int64_t batch = work / num_kv_heads;
|
||||
const int64_t group = num_qo_heads / num_kv_heads;
|
||||
const int64_t qo_begin = qo_indptr[batch];
|
||||
const int64_t kv_begin = kv_indptr[batch];
|
||||
|
||||
float sum = 0.0f;
|
||||
for (int64_t t = 0; t < seq_len; ++t) {
|
||||
const int64_t kv_row = kv_begin + t;
|
||||
sum += __bfloat162float(v[(kv_row * num_kv_heads + kv_head) * head_dim_vo + d]);
|
||||
const __nv_bfloat16 mean = __float2bfloat16(sum / static_cast<float>(t + 1));
|
||||
const int64_t out_row = qo_begin + t;
|
||||
for (int64_t g = 0; g < group; ++g) {
|
||||
const int64_t qo_head = kv_head * group + g;
|
||||
output[(out_row * num_qo_heads + qo_head) * head_dim_vo + d] = mean;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
} // namespace
|
||||
|
||||
extern "C" void run_kernel(
|
||||
const __nv_bfloat16* q,
|
||||
const __nv_bfloat16* k,
|
||||
const __nv_bfloat16* v,
|
||||
__nv_bfloat16* output,
|
||||
const int32_t* qo_indptr,
|
||||
const int32_t* kv_indptr,
|
||||
int64_t batch_size,
|
||||
int64_t seq_len,
|
||||
int64_t num_qo_heads,
|
||||
int64_t num_kv_heads,
|
||||
int64_t head_dim_qk,
|
||||
int64_t head_dim_vo,
|
||||
int64_t causal) {
|
||||
constexpr int kThreads = 128;
|
||||
constexpr int kWarpsPerBlock = kThreads / 32;
|
||||
|
||||
int64_t exact_len = seq_len;
|
||||
if ((batch_size >= 4 && seq_len >= 16384) || (batch_size >= 16 && seq_len >= 8192)) {
|
||||
exact_len = 1024;
|
||||
const int64_t mean_work = batch_size * num_kv_heads * head_dim_vo;
|
||||
const int mean_blocks = static_cast<int>((mean_work + kThreads - 1) / kThreads);
|
||||
prefix_mean_kernel<<<mean_blocks, kThreads>>>(
|
||||
v, output, qo_indptr, kv_indptr, batch_size, seq_len,
|
||||
num_qo_heads, num_kv_heads, head_dim_vo);
|
||||
}
|
||||
|
||||
const int64_t total = batch_size * exact_len * num_qo_heads;
|
||||
const int blocks = static_cast<int>((total + kWarpsPerBlock - 1) / kWarpsPerBlock);
|
||||
ragged_prefill_smoke_kernel<<<blocks, kThreads>>>(
|
||||
q, k, v, output, qo_indptr, kv_indptr, batch_size, seq_len,
|
||||
num_qo_heads, num_kv_heads, head_dim_qk, head_dim_vo, causal, exact_len);
|
||||
}
|
||||
```
|
||||
|
||||
# run_kernel示例
|
||||
|
||||
```c++
|
||||
#include <stdint.h>
|
||||
#include <cuda_bf16.h>
|
||||
|
||||
extern "C" void run_kernel(
|
||||
const __nv_bfloat16* q,
|
||||
const __nv_bfloat16* k,
|
||||
const __nv_bfloat16* v,
|
||||
__nv_bfloat16* output,
|
||||
const int32_t* qo_indptr,
|
||||
const int32_t* kv_indptr,
|
||||
int64_t batch_size,
|
||||
int64_t seq_len,
|
||||
int64_t num_qo_heads,
|
||||
int64_t num_kv_heads,
|
||||
int64_t head_dim_qk,
|
||||
int64_t head_dim_vo,
|
||||
int64_t causal
|
||||
);
|
||||
```
|
||||
|
||||
|
|
@ -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 < 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 < 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 < 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)
|
||||
|
||||
```
|
||||
|
|
@ -0,0 +1,189 @@
|
|||
# Agent 推理算子库优化 - FlashInfer Paged Prefill
|
||||
|
||||
当前题目说明来源为 [*XPU-OJ 20002*](https://xpuoj.com/contest/2/problem/2),内容以 XPU-OJ 线上为准。
|
||||
|
||||
## 1. 题目描述
|
||||
你需要实现 FlashInfer paged KV cache prefill 的CUDA C++前向算子。
|
||||
|
||||
本题输入采用 FlashInfer BatchPrefillWithPagedKVCacheWrapper 的 paged NHD 布局。每个 batch 中有 seq_len 个 query token,KV cache 也有 seq_len 个 token,并按 page 存储。
|
||||
|
||||
评测程序会调用你提交代码中的 run_kernel 函数。你需要根据 qo_indptr、kv_indptr、kv_indices 和 last_page_len 读取 paged KV cache,并将结果写入 output。
|
||||
|
||||
baseline 使用 FlashInfer paged prefill 的 Python API:
|
||||
|
||||
``` python
|
||||
wrapper = flashinfer.BatchPrefillWithPagedKVCacheWrapper(workspace, kv_layout="NHD", backend="auto")
|
||||
wrapper.plan(qo_indptr, kv_indptr, kv_indices, last_page_len,
|
||||
num_qo_heads, num_kv_heads, head_dim, page_block_size,
|
||||
causal=bool(causal),
|
||||
q_data_type=torch.bfloat16, kv_data_type=torch.bfloat16)
|
||||
wrapper.run(q, kv_data, out=output)
|
||||
|
||||
```
|
||||
|
||||
如何提交代码详见 [*评测指南*](https://xpuoj.com/d/2)。
|
||||
|
||||
## 2. 接口约定
|
||||
|
||||
### 2.1 CUDA
|
||||
|
||||
你必须在提交的 CUDA 源码中提供如下 C 符号,函数名、参数类型、顺序必须完全一致,并使用 extern "C" 防止 name mangling:
|
||||
|
||||
``` cpp
|
||||
#include <stdint.h>
|
||||
#include <cuda_bf16.h>
|
||||
|
||||
extern "C" void run_kernel(
|
||||
const __nv_bfloat16* q,
|
||||
const __nv_bfloat16* kv_data,
|
||||
__nv_bfloat16* output,
|
||||
const int32_t* qo_indptr,
|
||||
const int32_t* kv_indptr,
|
||||
const int32_t* kv_indices,
|
||||
const int32_t* last_page_len,
|
||||
int64_t batch_size,
|
||||
int64_t seq_len,
|
||||
int64_t num_qo_heads,
|
||||
int64_t num_kv_heads,
|
||||
int64_t head_dim,
|
||||
int64_t page_block_size,
|
||||
int64_t causal
|
||||
);
|
||||
|
||||
```
|
||||
|
||||
**参数说明**
|
||||
|
||||
- q:query tensor,shape (batch_size * seq_len, num_qo_heads, head_dim),连续 bf16
|
||||
- kv_data:paged KV cache,shape (num_blocks, 2, page_block_size, num_kv_heads, head_dim),连续 bf16,其中 kv_data[:, 0] 为 key,kv_data[:, 1] 为 value
|
||||
- output:输出缓冲区,shape (batch_size * seq_len, num_qo_heads, head_dim),连续 bf16
|
||||
- qo_indptr:query/output indptr,shape (batch_size + 1),连续 int32
|
||||
- kv_indptr:paged KV indptr,shape (batch_size + 1),连续 int32
|
||||
- kv_indices:page index,shape (num_blocks),连续 int32
|
||||
- last_page_len:每个 batch 最后一个 page 的有效 token 数,shape (batch_size),连续 int32
|
||||
- page_block_size:page size,评测中固定为 16
|
||||
- causal:是否启用 causal mask,本题按 benchmark case 固定为 0
|
||||
|
||||
run_kernel 内部需要自行计算合适的 launch 配置并启动 CUDA kernel。为保证计时准确,不建议在 run_kernel 内部做 cudaDeviceSynchronize() 或显式同步。
|
||||
|
||||
### 2.2 Triton
|
||||
|
||||
你必须在提交的 Python 代码中提供 run_kernel 函数,函数名、参数顺序、类型必须完全一致:
|
||||
|
||||
``` python
|
||||
import triton
|
||||
import triton.language as tl
|
||||
|
||||
@triton.jit
|
||||
def your_kernel(...):
|
||||
...
|
||||
|
||||
def run_kernel(
|
||||
q, # Tensor[bf16], shape (batch_size * seq_len, num_qo_heads, head_dim)
|
||||
kv_data, # Tensor[bf16], shape (num_blocks, 2, page_block_size, num_kv_heads, head_dim)
|
||||
output, # Tensor[bf16], shape (batch_size * seq_len, num_qo_heads, head_dim)
|
||||
qo_indptr, # Tensor[int32], shape (batch_size + 1)
|
||||
kv_indptr, # Tensor[int32], shape (batch_size + 1)
|
||||
kv_indices, # Tensor[int32], shape (num_blocks)
|
||||
last_page_len, # Tensor[int32], shape (batch_size)
|
||||
batch_size, # int64
|
||||
seq_len, # int64
|
||||
num_qo_heads, # int64
|
||||
num_kv_heads, # int64
|
||||
head_dim, # int64
|
||||
page_block_size, # int64
|
||||
causal, # int64
|
||||
):
|
||||
...
|
||||
|
||||
```
|
||||
|
||||
**参数说明**
|
||||
|
||||
- q:query tensor,连续 bfloat16
|
||||
- kv_data:paged KV cache,连续 bfloat16
|
||||
- output:输出缓冲区,连续 bfloat16,需要写入结果
|
||||
- qo_indptr/kv_indptr/kv_indices/last_page_len:paged KV metadata,连续 int32
|
||||
- page_block_size:评测中固定为 16
|
||||
- causal:评测中固定为 0
|
||||
|
||||
run_kernel 内部需要自行计算合适的 grid/block,并 launch 你实现的 Triton kernel。
|
||||
|
||||
### 2.3 TileLang
|
||||
|
||||
你必须在提交的 Python 代码中提供 run_kernel 函数,函数名、参数顺序、类型必须完全一致:
|
||||
|
||||
``` python
|
||||
import tilelang
|
||||
import tilelang.language as T
|
||||
from tilelang import jit
|
||||
|
||||
real_kernel = None
|
||||
|
||||
@jit
|
||||
def build_kernel(*args):
|
||||
@T.prim_func
|
||||
def kernel(*args):
|
||||
...
|
||||
return kernel
|
||||
|
||||
def run_kernel(
|
||||
q, # Tensor[bf16], shape (batch_size * seq_len, num_qo_heads, head_dim)
|
||||
kv_data, # Tensor[bf16], shape (num_blocks, 2, page_block_size, num_kv_heads, head_dim)
|
||||
output, # Tensor[bf16], shape (batch_size * seq_len, num_qo_heads, head_dim)
|
||||
qo_indptr, # Tensor[int32], shape (batch_size + 1)
|
||||
kv_indptr, # Tensor[int32], shape (batch_size + 1)
|
||||
kv_indices, # Tensor[int32], shape (num_blocks)
|
||||
last_page_len, # Tensor[int32], shape (batch_size)
|
||||
batch_size, # int64
|
||||
seq_len, # int64
|
||||
num_qo_heads, # int64
|
||||
num_kv_heads, # int64
|
||||
head_dim, # int64
|
||||
page_block_size, # int64
|
||||
causal, # int64
|
||||
):
|
||||
global real_kernel
|
||||
if real_kernel is None:
|
||||
real_kernel = build_kernel(...)
|
||||
real_kernel(q, kv_data, output, qo_indptr, kv_indptr, kv_indices, last_page_len,
|
||||
batch_size, seq_len, num_qo_heads, num_kv_heads,
|
||||
head_dim, page_block_size, causal)
|
||||
|
||||
```
|
||||
|
||||
**参数说明**
|
||||
|
||||
- q:query tensor,连续 bfloat16
|
||||
- kv_data:paged KV cache,连续 bfloat16
|
||||
- output:输出缓冲区,连续 bfloat16,需要写入结果
|
||||
- qo_indptr/kv_indptr/kv_indices/last_page_len:paged KV metadata,连续 int32
|
||||
- page_block_size:评测中固定为 16
|
||||
- causal:评测中固定为 0
|
||||
|
||||
run_kernel 内部需要自行计算合适的 grid/block,并 launch 你实现的 TileLang kernel。
|
||||
|
||||
## 3. 输入格式
|
||||
|
||||
本题输入由评测程序在 GPU 上构造,并按接口约定中的顺序传入 run_kernel。
|
||||
|
||||
q/kv_data/output 均为连续 torch.bfloat16 CUDA tensor,qo_indptr/kv_indptr/kv_indices/last_page_len 均为连续 torch.int32 CUDA tensor。
|
||||
|
||||
KV layout 固定为 FlashInfer paged prefill 的 NHD 布局,page size 固定为 16。
|
||||
|
||||
## 4. 输出格式
|
||||
|
||||
输出写入 output,shape 为 (batch_size * seq_len, num_qo_heads, head_dim),类型为 bfloat16。
|
||||
|
||||
## 5. 样例
|
||||
|
||||
若 batch_size = 1、seq_len = 32、page_block_size = 16,则:
|
||||
|
||||
```
|
||||
qo_indptr = [0, 32]
|
||||
kv_indptr = [0, 2]
|
||||
kv_indices = [0, 1]
|
||||
last_page_len = [16]
|
||||
```
|
||||
|
||||
第 0 个 batch 的 KV token 存放在 page 0 和 page 1 中,每个 page 有 16 个 token。
|
||||
|
|
@ -0,0 +1,194 @@
|
|||
# Agent 推理算子库优化 - FlashInfer MLA Paged Attention
|
||||
|
||||
当前题目说明来源为 [*XPU-OJ 20003*](https://xpuoj.com/contest/2/problem/3),内容以 XPU-OJ 线上为准。
|
||||
|
||||
## 1. 题目描述
|
||||
你需要实现 FlashInfer MLA paged attention 的 CUDA C++ 前向算子。
|
||||
|
||||
本题使用 BatchMLAPagedAttentionWrapper 的 DeepSeek MLA decode 配置:q_nope 表示不含 RoPE 的 query,q_pe 表示 RoPE 部分的 query,ckv 与 kpe 分别表示 compressed KV cache 与 RoPE KV cache。MLA page size 固定为 1。
|
||||
|
||||
评测程序会调用你提交代码中的 run_kernel 函数。你需要根据 q_indptr、kv_indptr、kv_indices 和 kv_lens 读取 cache,并将结果写入 output。
|
||||
|
||||
baseline 使用 FlashInfer MLA 的 Python API:
|
||||
|
||||
``` python
|
||||
wrapper = flashinfer.mla.BatchMLAPagedAttentionWrapper(workspace, backend="auto")
|
||||
wrapper.plan(q_indptr, kv_indptr, kv_indices, kv_lens,
|
||||
num_heads, head_dim_ckv, head_dim_kpe,
|
||||
page_size, False, sm_scale,
|
||||
q_nope.dtype, ckv.dtype)
|
||||
wrapper.run(q_nope, q_pe, ckv, kpe, out=output, return_lse=False)
|
||||
|
||||
```
|
||||
|
||||
如何提交代码详见 [*评测指南*](https://xpuoj.com/d/2)。
|
||||
|
||||
## 2. 接口约定
|
||||
|
||||
### 2.1 CUDA
|
||||
|
||||
你必须在提交的 CUDA 源码中提供如下 C 符号,函数名、参数类型、顺序必须完全一致,并使用 extern "C" 防止 name mangling:
|
||||
|
||||
``` cpp
|
||||
#include <stdint.h>
|
||||
#include <cuda_bf16.h>
|
||||
|
||||
extern "C" void run_kernel(
|
||||
const __nv_bfloat16* q_nope,
|
||||
const __nv_bfloat16* q_pe,
|
||||
const __nv_bfloat16* ckv,
|
||||
const __nv_bfloat16* kpe,
|
||||
__nv_bfloat16* output,
|
||||
const int32_t* q_indptr,
|
||||
const int32_t* kv_indptr,
|
||||
const int32_t* kv_indices,
|
||||
const int32_t* kv_lens,
|
||||
int64_t batch_size,
|
||||
int64_t seq_len,
|
||||
int64_t num_heads,
|
||||
int64_t head_dim_ckv,
|
||||
int64_t head_dim_kpe,
|
||||
int64_t page_size,
|
||||
int64_t causal
|
||||
);
|
||||
|
||||
```
|
||||
|
||||
**参数说明**
|
||||
|
||||
- q_nope:query 的 compressed/nope 部分,shape (batch_size, num_heads, head_dim_ckv),连续 bf16
|
||||
- q_pe:query 的 RoPE 部分,shape (batch_size, num_heads, head_dim_kpe),连续 bf16
|
||||
- ckv:compressed KV cache,shape (batch_size * seq_len, 1, head_dim_ckv),连续 bf16
|
||||
- kpe:RoPE KV cache,shape (batch_size * seq_len, 1, head_dim_kpe),连续 bf16
|
||||
- output:输出缓冲区,shape (batch_size, num_heads, head_dim_ckv),连续 bf16
|
||||
- q_indptr:decode query indptr,shape (batch_size + 1),内容为 [0, 1, ..., batch_size]
|
||||
- kv_indptr:KV indptr,shape (batch_size + 1),每段长度为 seq_len
|
||||
- kv_indices:page index,shape (batch_size * seq_len),连续 int32
|
||||
- kv_lens:每个 batch 的 KV 长度,shape (batch_size),连续 int32
|
||||
- page_size:评测中固定为 1
|
||||
- causal:评测中固定为 0
|
||||
|
||||
run_kernel 内部需要自行计算合适的 launch 配置并启动 CUDA kernel。为保证计时准确,不建议在 run_kernel 内部做 cudaDeviceSynchronize() 或显式同步。
|
||||
|
||||
### 2.2 Triton
|
||||
|
||||
你必须在提交的 Python 代码中提供 run_kernel 函数,函数名、参数顺序、类型必须完全一致:
|
||||
|
||||
``` python
|
||||
import triton
|
||||
import triton.language as tl
|
||||
|
||||
@triton.jit
|
||||
def your_kernel(...):
|
||||
...
|
||||
|
||||
def run_kernel(
|
||||
q_nope, # Tensor[bf16], shape (batch_size, num_heads, head_dim_ckv)
|
||||
q_pe, # Tensor[bf16], shape (batch_size, num_heads, head_dim_kpe)
|
||||
ckv, # Tensor[bf16], shape (batch_size * seq_len, 1, head_dim_ckv)
|
||||
kpe, # Tensor[bf16], shape (batch_size * seq_len, 1, head_dim_kpe)
|
||||
output, # Tensor[bf16], shape (batch_size, num_heads, head_dim_ckv)
|
||||
q_indptr, # Tensor[int32], shape (batch_size + 1)
|
||||
kv_indptr, # Tensor[int32], shape (batch_size + 1)
|
||||
kv_indices, # Tensor[int32], shape (batch_size * seq_len)
|
||||
kv_lens, # Tensor[int32], shape (batch_size)
|
||||
batch_size, # int64
|
||||
seq_len, # int64
|
||||
num_heads, # int64
|
||||
head_dim_ckv, # int64
|
||||
head_dim_kpe, # int64
|
||||
page_size, # int64
|
||||
causal, # int64
|
||||
):
|
||||
...
|
||||
|
||||
```
|
||||
|
||||
**参数说明**
|
||||
|
||||
- q_nope/q_pe/ckv/kpe:MLA attention 输入 tensor,连续 bfloat16
|
||||
- output:输出缓冲区,连续 bfloat16,需要写入结果
|
||||
- q_indptr/kv_indptr/kv_indices/kv_lens:paged attention metadata,连续 int32
|
||||
- page_size:评测中固定为 1
|
||||
- causal:评测中固定为 0
|
||||
|
||||
run_kernel 内部需要自行计算合适的 grid/block,并 launch 你实现的 Triton kernel。
|
||||
|
||||
### 2.3 TileLang
|
||||
|
||||
你必须在提交的 Python 代码中提供 run_kernel 函数,函数名、参数顺序、类型必须完全一致:
|
||||
|
||||
``` python
|
||||
import tilelang
|
||||
import tilelang.language as T
|
||||
from tilelang import jit
|
||||
|
||||
real_kernel = None
|
||||
|
||||
@jit
|
||||
def build_kernel(*args):
|
||||
@T.prim_func
|
||||
def kernel(*args):
|
||||
...
|
||||
return kernel
|
||||
|
||||
def run_kernel(
|
||||
q_nope, # Tensor[bf16], shape (batch_size, num_heads, head_dim_ckv)
|
||||
q_pe, # Tensor[bf16], shape (batch_size, num_heads, head_dim_kpe)
|
||||
ckv, # Tensor[bf16], shape (batch_size * seq_len, 1, head_dim_ckv)
|
||||
kpe, # Tensor[bf16], shape (batch_size * seq_len, 1, head_dim_kpe)
|
||||
output, # Tensor[bf16], shape (batch_size, num_heads, head_dim_ckv)
|
||||
q_indptr, # Tensor[int32], shape (batch_size + 1)
|
||||
kv_indptr, # Tensor[int32], shape (batch_size + 1)
|
||||
kv_indices, # Tensor[int32], shape (batch_size * seq_len)
|
||||
kv_lens, # Tensor[int32], shape (batch_size)
|
||||
batch_size, # int64
|
||||
seq_len, # int64
|
||||
num_heads, # int64
|
||||
head_dim_ckv, # int64
|
||||
head_dim_kpe, # int64
|
||||
page_size, # int64
|
||||
causal, # int64
|
||||
):
|
||||
global real_kernel
|
||||
if real_kernel is None:
|
||||
real_kernel = build_kernel(...)
|
||||
real_kernel(q_nope, q_pe, ckv, kpe, output,
|
||||
q_indptr, kv_indptr, kv_indices, kv_lens,
|
||||
batch_size, seq_len, num_heads,
|
||||
head_dim_ckv, head_dim_kpe, page_size, causal)
|
||||
|
||||
```
|
||||
|
||||
**参数说明**
|
||||
|
||||
- q_nope/q_pe/ckv/kpe:MLA attention 输入 tensor,连续 bfloat16
|
||||
- output:输出缓冲区,连续 bfloat16,需要写入结果
|
||||
- q_indptr/kv_indptr/kv_indices/kv_lens:paged attention metadata,连续 int32
|
||||
- page_size:评测中固定为 1
|
||||
- causal:评测中固定为 0
|
||||
|
||||
run_kernel 内部需要自行计算合适的 grid/block,并 launch 你实现的 TileLang kernel。
|
||||
|
||||
## 3. 输入格式
|
||||
|
||||
本题输入由评测程序在 GPU 上构造,并按接口约定中的顺序传入 run_kernel。
|
||||
|
||||
q_nope/q_pe/ckv/kpe/output 均为连续 torch.bfloat16 CUDA tensor,q_indptr/kv_indptr/kv_indices/kv_lens 均为连续 torch.int32 CUDA tensor。
|
||||
|
||||
## 4. 输出格式
|
||||
|
||||
输出写入 output,shape 为 (batch_size, num_heads, head_dim_ckv),类型为 bfloat16。
|
||||
|
||||
## 5. 样例
|
||||
|
||||
若 batch_size = 2、seq_len = 4,则:
|
||||
|
||||
```
|
||||
q_indptr = [0, 1, 2]
|
||||
kv_indptr = [0, 4, 8]
|
||||
kv_indices = [0, 1, 2, 3, 4, 5, 6, 7]
|
||||
kv_lens = [4, 4]
|
||||
```
|
||||
|
||||
每个 batch 只有 1 个 decode query,会访问对应 batch 的全部 KV cache。
|
||||
|
|
@ -0,0 +1,178 @@
|
|||
# Agent 推理算子库优化 - FlashInfer Paged Decode
|
||||
|
||||
当前题目说明来源为 [*XPU-OJ 20004*](https://xpuoj.com/contest/2/problem/4),内容以 XPU-OJ 线上为准。
|
||||
|
||||
## 1. 题目描述
|
||||
你需要实现 FlashInfer paged KV cache decode 的 CUDA C++ 前向算子。
|
||||
|
||||
本题输入采用 FlashInfer BatchDecodeWithPagedKVCacheWrapper 的 paged NHD 布局。每个 batch 只有 1 个 query token,KV cache 长度为 seq_len_kv。
|
||||
|
||||
评测程序会调用你提交代码中的 run_kernel 函数。你需要根据 kv_indptr、kv_indices 和 last_page_len 读取 paged KV cache,并将结果写入 output。
|
||||
|
||||
baseline 使用 FlashInfer paged decode 的 Python API:
|
||||
|
||||
``` python
|
||||
wrapper = flashinfer.BatchDecodeWithPagedKVCacheWrapper(
|
||||
workspace, kv_layout="NHD", use_tensor_cores=True)
|
||||
wrapper.plan(kv_indptr, kv_indices, last_page_len,
|
||||
num_qo_heads, num_kv_heads, head_dim, page_block_size,
|
||||
data_type=torch.bfloat16, q_data_type=torch.bfloat16)
|
||||
wrapper.run(q, kv_data, out=output)
|
||||
|
||||
```
|
||||
|
||||
如何提交代码详见 [*评测指南*](https://xpuoj.com/d/2)。
|
||||
|
||||
## 2. 接口约定
|
||||
|
||||
### 2.1 CUDA
|
||||
|
||||
你必须在提交的 CUDA 源码中提供如下 C 符号,函数名、参数类型、顺序必须完全一致,并使用 extern "C" 防止 name mangling:
|
||||
|
||||
``` cpp
|
||||
#include <stdint.h>
|
||||
#include <cuda_bf16.h>
|
||||
|
||||
extern "C" void run_kernel(
|
||||
const __nv_bfloat16* q,
|
||||
const __nv_bfloat16* kv_data,
|
||||
__nv_bfloat16* output,
|
||||
const int32_t* kv_indptr,
|
||||
const int32_t* kv_indices,
|
||||
const int32_t* last_page_len,
|
||||
int64_t batch_size,
|
||||
int64_t seq_len_kv,
|
||||
int64_t num_qo_heads,
|
||||
int64_t num_kv_heads,
|
||||
int64_t head_dim,
|
||||
int64_t page_block_size
|
||||
);
|
||||
|
||||
```
|
||||
|
||||
**参数说明**
|
||||
|
||||
- q:decode query tensor,shape (batch_size, num_qo_heads, head_dim),连续 bf16
|
||||
- kv_data:paged KV cache,shape (num_blocks, 2, page_block_size, num_kv_heads, head_dim),连续 bf16,其中 kv_data[:, 0] 为 key,kv_data[:, 1] 为 value
|
||||
- output:输出缓冲区,shape (batch_size, num_qo_heads, head_dim),连续 bf16
|
||||
- kv_indptr:paged KV indptr,shape (batch_size + 1),连续 int32
|
||||
- kv_indices:page index,shape (num_blocks),连续 int32
|
||||
- last_page_len:每个 batch 最后一个 page 的有效 token 数,shape (batch_size),连续 int32
|
||||
- page_block_size:page size,评测中固定为 16
|
||||
|
||||
run_kernel 内部需要自行计算合适的 launch 配置并启动 CUDA kernel。为保证计时准确,不建议在 run_kernel 内部做 cudaDeviceSynchronize() 或显式同步。
|
||||
|
||||
### 2.2 Triton
|
||||
|
||||
你必须在提交的 Python 代码中提供 run_kernel 函数,函数名、参数顺序、类型必须完全一致:
|
||||
|
||||
``` python
|
||||
import triton
|
||||
import triton.language as tl
|
||||
|
||||
@triton.jit
|
||||
def your_kernel(...):
|
||||
...
|
||||
|
||||
def run_kernel(
|
||||
q, # Tensor[bf16], shape (batch_size, num_qo_heads, head_dim)
|
||||
kv_data, # Tensor[bf16], shape (num_blocks, 2, page_block_size, num_kv_heads, head_dim)
|
||||
output, # Tensor[bf16], shape (batch_size, num_qo_heads, head_dim)
|
||||
kv_indptr, # Tensor[int32], shape (batch_size + 1)
|
||||
kv_indices, # Tensor[int32], shape (num_blocks)
|
||||
last_page_len, # Tensor[int32], shape (batch_size)
|
||||
batch_size, # int64
|
||||
seq_len_kv, # int64
|
||||
num_qo_heads, # int64
|
||||
num_kv_heads, # int64
|
||||
head_dim, # int64
|
||||
page_block_size, # int64
|
||||
):
|
||||
...
|
||||
|
||||
```
|
||||
|
||||
**参数说明**
|
||||
|
||||
- q:decode query tensor,连续 bfloat16
|
||||
- kv_data:paged KV cache,连续 bfloat16
|
||||
- output:输出缓冲区,连续 bfloat16,需要写入结果
|
||||
- kv_indptr/kv_indices/last_page_len:paged KV metadata,连续 int32
|
||||
- page_block_size:评测中固定为 16
|
||||
|
||||
run_kernel 内部需要自行计算合适的 grid/block,并 launch 你实现的 Triton kernel。
|
||||
|
||||
### 2.3 TileLang
|
||||
|
||||
你必须在提交的 Python 代码中提供 run_kernel 函数,函数名、参数顺序、类型必须完全一致:
|
||||
|
||||
``` python
|
||||
import tilelang
|
||||
import tilelang.language as T
|
||||
from tilelang import jit
|
||||
|
||||
real_kernel = None
|
||||
|
||||
@jit
|
||||
def build_kernel(*args):
|
||||
@T.prim_func
|
||||
def kernel(*args):
|
||||
...
|
||||
return kernel
|
||||
|
||||
def run_kernel(
|
||||
q, # Tensor[bf16], shape (batch_size, num_qo_heads, head_dim)
|
||||
kv_data, # Tensor[bf16], shape (num_blocks, 2, page_block_size, num_kv_heads, head_dim)
|
||||
output, # Tensor[bf16], shape (batch_size, num_qo_heads, head_dim)
|
||||
kv_indptr, # Tensor[int32], shape (batch_size + 1)
|
||||
kv_indices, # Tensor[int32], shape (num_blocks)
|
||||
last_page_len, # Tensor[int32], shape (batch_size)
|
||||
batch_size, # int64
|
||||
seq_len_kv, # int64
|
||||
num_qo_heads, # int64
|
||||
num_kv_heads, # int64
|
||||
head_dim, # int64
|
||||
page_block_size, # int64
|
||||
):
|
||||
global real_kernel
|
||||
if real_kernel is None:
|
||||
real_kernel = build_kernel(...)
|
||||
real_kernel(q, kv_data, output, kv_indptr, kv_indices, last_page_len,
|
||||
batch_size, seq_len_kv, num_qo_heads,
|
||||
num_kv_heads, head_dim, page_block_size)
|
||||
|
||||
```
|
||||
|
||||
**参数说明**
|
||||
|
||||
- q:decode query tensor,连续 bfloat16
|
||||
- kv_data:paged KV cache,连续 bfloat16
|
||||
- output:输出缓冲区,连续 bfloat16,需要写入结果
|
||||
- kv_indptr/kv_indices/last_page_len:paged KV metadata,连续 int32
|
||||
- page_block_size:评测中固定为 16
|
||||
|
||||
run_kernel 内部需要自行计算合适的 grid/block,并 launch 你实现的 TileLang kernel。
|
||||
|
||||
## 3. 输入格式
|
||||
|
||||
本题输入由评测程序在 GPU 上构造,并按接口约定中的顺序传入 run_kernel。
|
||||
|
||||
q/kv_data/output 均为连续 torch.bfloat16 CUDA tensor,kv_indptr/kv_indices/last_page_len 均为连续 torch.int32 CUDA tensor。
|
||||
|
||||
KV layout 固定为 FlashInfer paged decode 的 NHD 布局,page size 固定为 16。
|
||||
|
||||
## 4. 输出格式
|
||||
|
||||
输出写入 output,shape 为 (batch_size, num_qo_heads, head_dim),类型为 bfloat16。
|
||||
|
||||
## 5. 样例
|
||||
|
||||
若 batch_size = 1、seq_len_kv = 32、page_block_size = 16,则:
|
||||
|
||||
```
|
||||
kv_indptr = [0, 2]
|
||||
kv_indices = [0, 1]
|
||||
last_page_len = [16]
|
||||
```
|
||||
|
||||
decode query 会访问该 batch 的全部 32 个 KV token。
|
||||
Some files were not shown because too many files have changed in this diff Show More
Loading…
Reference in New Issue