docs(perf): profile CUDA graph coverage and PLE costs on Spark
This commit is contained in:
@@ -8,7 +8,7 @@
|
||||
PLE 查找表通过磁盘映射按需读取;原始模型权重未重新量化。
|
||||
|
||||
> 这是针对固定 nightly 的社区适配,非 NVIDIA/vLLM 官方支持方案。
|
||||
> 已验证短问答和工具调用;未压测完整 262K 输入。仅适用于此模型、单 GPU / ETP=1。
|
||||
> 已验证短问答、工具调用及有限的 128K 单请求检索;接近 260K 已触发主机 OOM。仅适用于此模型、单 GPU / ETP=1。
|
||||
|
||||
## 快速开始
|
||||
|
||||
@@ -109,3 +109,8 @@ python3 scripts/summarize-benchmark.py benchmark.jsonl
|
||||
```
|
||||
|
||||
测试只使用合成输入。方法、首字时间口径与小样本限制见 docs/benchmark-method.md。
|
||||
|
||||
## CUDA Graph 性能剖析
|
||||
|
||||
Nsight 已确认图重放生效,短回答约提升 4.5%;长输入预填充未被当前小尺寸图覆盖,
|
||||
GPU 活动本已接近连续。详见 [性能剖析报告](docs/cuda-graph-profile.md),包括 PLE 等待/查表区分及原始证据摘要。
|
||||
|
||||
@@ -0,0 +1,127 @@
|
||||
# CUDA Graph 性能剖析:图已生效,主要收益受 GPU 工作占比限制
|
||||
|
||||
2026-09-18,Asia/Shanghai。方案与脚本见 [profile128k](../experiments/profile128k/plan.md)。
|
||||
|
||||
CUDA Graph 确实被重放,并非参数失效。本轮短回答生成速度中位数由 **29.954 → 31.316 tokens/s(约 +4.5%)**,
|
||||
与此前约 +4.4% 的复测一致。32K 首次输入首 token 则为 **15.416 → 15.376 秒**,基本不变。
|
||||
|
||||
原因有两层:当前小尺寸图没有覆盖实际的长输入预填充分块;同时,GPU 时间线上本来就没有很多空隙可消除。
|
||||
不能据此认定 SSD、PLE mmap 或 CPU 调度是主要瓶颈,也不能把 CUDA Graph 的小收益解释为 CUDA/GPU 没有发挥作用。
|
||||
**eager 和 graph 两组都在使用 GPU 执行 CUDA 内核;本次比较的是是否通过 CUDA Graph 提交这些工作。**
|
||||
|
||||
## 正常计时与正确性
|
||||
|
||||
- 固定 nightly `0bfc7a15`、相同模型 revision、BF16 KV、MTP=2、前缀缓存、两处 Mamba 修正、GPU memory=0.80。
|
||||
- 实验上下文上限 131072;max_num_seqs=4,但测试串行;max_num_batched_tokens=2048。
|
||||
- 小尺寸图为 breakable PIECEWISE,capture sizes `[1,2,4,8,12]`。
|
||||
- 短回答预算 512,长检索预算 2048,思考与正文共用;temperature=0、seed=42、reasoning_effort=low。
|
||||
- 两组各 19 次请求:1 次算术预热、三轮各 5 次计时请求、3 次 trace 请求。38 次均通过各自判据。
|
||||
长检索要求精确答案和正常 stop,短说明仅检查非空及正常 stop,不代表进行了语言质量评估。
|
||||
- 19 个对应请求的输入 SHA256 全部一致;服务返回 prompt_tokens 与预检逐条一致。
|
||||
- 每轮首次输入使用不同前缀,避免跨轮缓存命中;重复输入紧随首次输入。
|
||||
- 下表为三轮中位数,不含预热及 trace 请求。短回答每次均输出 214 tokens;长检索输出约 62–74 tokens,有少量差异。
|
||||
|
||||
| 指标 | eager | 小尺寸 graph |
|
||||
|---|---:|---:|
|
||||
| 短回答解码速度(近似 tokens/s) | 29.954 | 31.316 |
|
||||
| 短回答首 token(秒,含思考) | 0.2627 | 0.2624 |
|
||||
| 短回答完整耗时(秒) | 7.374 | 7.064 |
|
||||
| 8K 首次首 token(秒) | 3.788 | 3.770 |
|
||||
| 8K 复用首 token(秒) | 0.880 | 0.874 |
|
||||
| 32K 首次首 token(秒) | 15.416 | 15.376 |
|
||||
| 32K 复用首 token(秒) | 1.368 | 1.327 |
|
||||
|
||||
原始请求记录和自动汇总见 [results/profile128k](results/profile128k/benchmark-summary.json)。
|
||||
计时阶段没有开启 capture,但 Nsight launcher 和 NVTX 包装仍存在;不是完全卸载 profiler 的独立测速。
|
||||
执行顺序 eager → graph,未随机化、未清 OS 页缓存,小样本不代表统计显著性或长期稳定性。
|
||||
|
||||
## 时间线证据
|
||||
|
||||
Nsight Systems 2025.3.2,CUDA profiler API 控制三个短采样窗口,node 级 graph trace。
|
||||
NVTX 临时标记 PLE 同步、CPU 去重、CPU gather、lookup,以及图调度描述符。
|
||||
计时阶段和 trace 阶段分开;下表不用于跨模式的端到端性能排名。
|
||||
|
||||
“GPU 活动占比”是 kernel、memcpy、memset 时间区间并集,除以首次至末次 GPU 活动的跨度。
|
||||
它不是 SM 利用率、内存带宽利用率,也不包含完整 HTTP 请求前后的所有工作。
|
||||
|
||||
| 采样请求 | eager 活动占比 | graph 活动占比 | graph 中来自图节点的 kernel 数量占比 |
|
||||
|---|---:|---:|---:|
|
||||
| 短回答 | 93.01% | 95.22% | 82.43% |
|
||||
| 32K 首次输入及回答 | 97.38% | 97.50% | 44.25% |
|
||||
| 32K 复用及回答 | 95.22% | 95.29% | 77.15% |
|
||||
|
||||
### 图生效,但覆盖范围有限且分成多段
|
||||
|
||||
- eager 的三个窗口都没有 graph-node kernel。
|
||||
- graph 短回答窗口记录 **4,664 次 `cudaGraphLaunch`,160,160 / 194,303 个 kernel 来自图节点**。
|
||||
- 调度记录显示解码 `actual=3` 被 padding 到 4,模式为 PIECEWISE;MTP 路径也存在图重放。
|
||||
同时保留部分 NONE 路径,不能说整次请求被捕获成一张图。
|
||||
- 该窗口有 88 次 `actual=3` 的图调度,累计 graph launch 约为其 53 倍。
|
||||
这是目标模型与辅助路径合计的实测比值,说明提交仍分成多段;不是每生成一个 token 就固定调用 53 次。
|
||||
- 32K 首次请求中的 **1600 和 742-token 预填充调度均为 NONE**;后续解码才命中图。
|
||||
- 日志确认 attention/KV block size 为 1600。虽然批处理上限是 2048,实际主要分块是 1600,
|
||||
因此不能仅凭配置里的 2048 推断捕获或预填充形状。
|
||||
|
||||
### PLE 的“等待时间”不是磁盘读表时间
|
||||
|
||||
32K 首次输入 trace:
|
||||
|
||||
| 时间项 | eager | graph |
|
||||
|---|---:|---:|
|
||||
| 首次至末次 GPU 活动跨度 | 17.458 s | 17.614 s |
|
||||
| GPU 活动区间并集 | 17.000 s | 17.173 s |
|
||||
| PLE CPU gather | 0.275 s | 0.292 s |
|
||||
| PLE CPU 去重(标记中的 np.unique) | 0.019 s | 0.021 s |
|
||||
| PLE 等待 GPU | 6.389 s | 7.292 s |
|
||||
| 其中与 GPU 活动重叠 | 6.341 s | 7.233 s |
|
||||
|
||||
CPU wait 与 GPU 工作几乎完全重叠,不是额外叠加在 GPU 计算之后的同等开销。
|
||||
不能把所有 API/NVTX/GPU 时长相加,也不能依据同步函数耗时最长就优先优化同步本身。
|
||||
图版本 wait 更长不证明 PLE 退化:提交时序和输出长度都不同。
|
||||
|
||||
本轮 gather 包括页缓存或实际存储读取,未测物理 NVMe I/O;没有证据说明每次都从 SSD 读取,
|
||||
也没有证据支持将整张 47.7 GiB 表强制驻留能带来显著收益。
|
||||
短回答中 CPU gather 约 0.20 秒,相比约 7 秒 GPU 活动跨度亦非主导项。
|
||||
|
||||
## 决策与后续方向
|
||||
|
||||
本轮**不扩大图捕获尺寸**。长输入 GPU 活动已接近连续,消除提交间隙的空间有限;
|
||||
扩大图还可能增加 padding、捕获内存和启动成本。当前证据不足以让它成为稳定基线的优先改动。
|
||||
这不等于证明大图一定没有收益;该候选未实测,不报告虚构加速或稳定性结果。
|
||||
按方案,只有新候选通过才追加 128K 验证,因此本轮未重跑 128K,也未再次触碰 260K。
|
||||
|
||||
稳定基线建议仍是:**eager + 前缀缓存 + MTP=2,实际使用限制在 128K 以内**。
|
||||
愿意维护实验补丁时,小尺寸图可作为约 4–5% 短回答收益的可选项;目前证据不足以要求默认切换。
|
||||
下一个有依据的性能工作应聚焦耗时 GPU 内核及其 GB10 后端:32K trace 的主要内核包含
|
||||
CUTLASS 分组量化 GEMM、稀疏注意力和归一化组合;短回答还有显著 BF16 GEMM 时间。
|
||||
要判定算力受限还是带宽受限,需要另做针对性内核分析,不能用本轮活动占比替代结论。
|
||||
|
||||
## 资源与实验修正
|
||||
|
||||
实验采样最低 MemAvailable **15.748 GiB**,主机 `oom_kill` 从 25 到 25,没有新增 OOM。
|
||||
监控每秒采样,低于 1 GiB 立即停止,持续低于 2 GiB 或主机 OOM 增加则停止。
|
||||
监控覆盖两组实验,恢复阶段另做最终系统检查;没有为了验证保护阈值主动制造内存压力。
|
||||
|
||||
首次诊断启动因临时包装函数缺少 PyTorch schema 所需类型注解而失败,尚未加载完成或执行请求。
|
||||
自动恢复触发后,修复注解并验证 `torch.library.infer_schema`,随后才运行完整两组。
|
||||
该失败属于诊断脚本问题,不是模型/vLLM 的新 bug;失败日志保留在 Spark 的 attempt1 目录。
|
||||
|
||||
原始 trace 和 SQLite 只保留在 Spark `optimization-results/profile128k/{eager,graph}`,父目录权限 0700。
|
||||
它们可能包含进程参数,不提交 Git。仓库的 [证据清单](results/profile128k/evidence-manifest.json)
|
||||
记录文件大小、SHA256、Nsight 版本、配置要点;仅提交 CUDA/NVTX 摘要、请求结果及资源记录。
|
||||
|
||||
## 恢复状态
|
||||
|
||||
2026-09-18 01:21(Asia/Shanghai)已恢复原 eager 镜像,容器 healthy、重启策略 `unless-stopped`。
|
||||
算术、中文、工具调用 smoke test 全部通过;原 Compose 与测试前逐字一致,临时 profiler 参数与挂载均已移除。
|
||||
可用内存约 18.85 GiB,主机 OOM 计数仍为 25,系统及用户 failed units 均为 0。
|
||||
恢复证据见 [recovery.json](results/profile128k/recovery.json) 和 [smoke log](results/profile128k/recovery-smoke.log)。
|
||||
|
||||
默认配置仍保留原来的 262144 上限,本轮未修改服务容量;该配置值不代表已验证的稳定容量。
|
||||
实际使用仍建议限制在 128K 以内,接近 260K 的既有 OOM 结论不变。
|
||||
|
||||
## 参考
|
||||
|
||||
- [vLLM profiling 文档](https://docs.vllm.ai/en/latest/contributing/profiling/)
|
||||
- [NVIDIA Nsight Systems User Guide](https://docs.nvidia.com/nsight-systems/UserGuide/)
|
||||
- [NVIDIA CUDA Graph 性能问题说明](https://docs.nvidia.com/dl-cuda-graph/latest/troubleshooting/performance-issues/)
|
||||
@@ -0,0 +1,394 @@
|
||||
{
|
||||
"eager": {
|
||||
"total_requests": 19,
|
||||
"correct": 19,
|
||||
"groups": {
|
||||
"short": {
|
||||
"n": 3,
|
||||
"correct": 3,
|
||||
"ttft_s": {
|
||||
"median": 0.2627224419993581,
|
||||
"values": [
|
||||
0.33354272600263357,
|
||||
0.2627224419993581,
|
||||
0.26176839799882146
|
||||
]
|
||||
},
|
||||
"elapsed_s": {
|
||||
"median": 7.373869219998596,
|
||||
"values": [
|
||||
7.68168595500174,
|
||||
7.373869219998596,
|
||||
7.35648784199293
|
||||
]
|
||||
},
|
||||
"decode_tps_approx": {
|
||||
"median": 29.953694756255302,
|
||||
"values": [
|
||||
28.987645227728436,
|
||||
29.953694756255302,
|
||||
30.02293954286349
|
||||
]
|
||||
},
|
||||
"output_tokens": [
|
||||
214,
|
||||
214,
|
||||
214
|
||||
],
|
||||
"reasoning_tokens": [
|
||||
111,
|
||||
111,
|
||||
111
|
||||
]
|
||||
},
|
||||
"prefill-8192": {
|
||||
"n": 3,
|
||||
"correct": 3,
|
||||
"ttft_s": {
|
||||
"median": 3.7875278329956927,
|
||||
"values": [
|
||||
4.006232126004761,
|
||||
3.7754420729979756,
|
||||
3.7875278329956927
|
||||
]
|
||||
},
|
||||
"elapsed_s": {
|
||||
"median": 5.739436170995759,
|
||||
"values": [
|
||||
5.835497587999271,
|
||||
5.626387860997056,
|
||||
5.739436170995759
|
||||
]
|
||||
},
|
||||
"decode_tps_approx": {
|
||||
"median": 35.041074313706375,
|
||||
"values": [
|
||||
35.584837627199924,
|
||||
35.041074313706375,
|
||||
32.64277985613848
|
||||
]
|
||||
},
|
||||
"output_tokens": [
|
||||
63,
|
||||
63,
|
||||
62
|
||||
],
|
||||
"reasoning_tokens": [
|
||||
55,
|
||||
55,
|
||||
54
|
||||
]
|
||||
},
|
||||
"reuse-8192": {
|
||||
"n": 3,
|
||||
"correct": 3,
|
||||
"ttft_s": {
|
||||
"median": 0.8803304789980757,
|
||||
"values": [
|
||||
0.8856951920024585,
|
||||
0.8737719170021592,
|
||||
0.8803304789980757
|
||||
]
|
||||
},
|
||||
"elapsed_s": {
|
||||
"median": 2.811309409997193,
|
||||
"values": [
|
||||
2.7496982970042154,
|
||||
2.835690722997242,
|
||||
2.811309409997193
|
||||
]
|
||||
},
|
||||
"decode_tps_approx": {
|
||||
"median": 36.047200027797416,
|
||||
"values": [
|
||||
34.74998178667343,
|
||||
36.047200027797416,
|
||||
36.74879746387076
|
||||
]
|
||||
},
|
||||
"output_tokens": [
|
||||
63,
|
||||
69,
|
||||
69
|
||||
],
|
||||
"reasoning_tokens": [
|
||||
55,
|
||||
61,
|
||||
61
|
||||
]
|
||||
},
|
||||
"prefill-32768": {
|
||||
"n": 3,
|
||||
"correct": 3,
|
||||
"ttft_s": {
|
||||
"median": 15.41581593700539,
|
||||
"values": [
|
||||
15.595506220997777,
|
||||
15.403379977004079,
|
||||
15.41581593700539
|
||||
]
|
||||
},
|
||||
"elapsed_s": {
|
||||
"median": 17.35202403000585,
|
||||
"values": [
|
||||
17.837196191001567,
|
||||
17.295454844999767,
|
||||
17.35202403000585
|
||||
]
|
||||
},
|
||||
"decode_tps_approx": {
|
||||
"median": 36.438073002765556,
|
||||
"values": [
|
||||
33.795316059955795,
|
||||
36.438073002765556,
|
||||
36.713832650404036
|
||||
]
|
||||
},
|
||||
"output_tokens": [
|
||||
74,
|
||||
67,
|
||||
69
|
||||
],
|
||||
"reasoning_tokens": [
|
||||
66,
|
||||
59,
|
||||
61
|
||||
]
|
||||
},
|
||||
"reuse-32768": {
|
||||
"n": 3,
|
||||
"correct": 3,
|
||||
"ttft_s": {
|
||||
"median": 1.367758443004277,
|
||||
"values": [
|
||||
1.368702627994935,
|
||||
1.367758443004277,
|
||||
1.357998088002205
|
||||
]
|
||||
},
|
||||
"elapsed_s": {
|
||||
"median": 3.5468969909998123,
|
||||
"values": [
|
||||
3.579059300995141,
|
||||
3.5468969909998123,
|
||||
3.3747847780032316
|
||||
]
|
||||
},
|
||||
"decode_tps_approx": {
|
||||
"median": 32.877578996597116,
|
||||
"values": [
|
||||
32.419700549347965,
|
||||
32.877578996597116,
|
||||
33.71865583479496
|
||||
]
|
||||
},
|
||||
"output_tokens": [
|
||||
70,
|
||||
70,
|
||||
69
|
||||
],
|
||||
"reasoning_tokens": [
|
||||
62,
|
||||
62,
|
||||
61
|
||||
]
|
||||
}
|
||||
}
|
||||
},
|
||||
"graph": {
|
||||
"total_requests": 19,
|
||||
"correct": 19,
|
||||
"groups": {
|
||||
"short": {
|
||||
"n": 3,
|
||||
"correct": 3,
|
||||
"ttft_s": {
|
||||
"median": 0.2623967160034226,
|
||||
"values": [
|
||||
0.32573997599683935,
|
||||
0.2600282569983392,
|
||||
0.2623967160034226
|
||||
]
|
||||
},
|
||||
"elapsed_s": {
|
||||
"median": 7.0642776130043785,
|
||||
"values": [
|
||||
7.360314821999054,
|
||||
7.060739258995454,
|
||||
7.0642776130043785
|
||||
]
|
||||
},
|
||||
"decode_tps_approx": {
|
||||
"median": 31.315723764990278,
|
||||
"values": [
|
||||
30.279974628892248,
|
||||
31.32107217692348,
|
||||
31.315723764990278
|
||||
]
|
||||
},
|
||||
"output_tokens": [
|
||||
214,
|
||||
214,
|
||||
214
|
||||
],
|
||||
"reasoning_tokens": [
|
||||
111,
|
||||
111,
|
||||
111
|
||||
]
|
||||
},
|
||||
"prefill-8192": {
|
||||
"n": 3,
|
||||
"correct": 3,
|
||||
"ttft_s": {
|
||||
"median": 3.7701075279983343,
|
||||
"values": [
|
||||
3.921214998998039,
|
||||
3.7701075279983343,
|
||||
3.768740819999948
|
||||
]
|
||||
},
|
||||
"elapsed_s": {
|
||||
"median": 5.670499198000471,
|
||||
"values": [
|
||||
6.191493427999376,
|
||||
5.596617848997994,
|
||||
5.670499198000471
|
||||
]
|
||||
},
|
||||
"decode_tps_approx": {
|
||||
"median": 35.516432222041566,
|
||||
"values": [
|
||||
33.338899955373805,
|
||||
35.516432222041566,
|
||||
37.32072215199369
|
||||
]
|
||||
},
|
||||
"output_tokens": [
|
||||
74,
|
||||
63,
|
||||
69
|
||||
],
|
||||
"reasoning_tokens": [
|
||||
66,
|
||||
55,
|
||||
61
|
||||
]
|
||||
},
|
||||
"reuse-8192": {
|
||||
"n": 3,
|
||||
"correct": 3,
|
||||
"ttft_s": {
|
||||
"median": 0.8737782960015466,
|
||||
"values": [
|
||||
0.8737782960015466,
|
||||
0.8690901810041396,
|
||||
0.8816422219970264
|
||||
]
|
||||
},
|
||||
"elapsed_s": {
|
||||
"median": 2.8319971849996364,
|
||||
"values": [
|
||||
2.8319971849996364,
|
||||
3.0074559670028975,
|
||||
2.711807350999152
|
||||
]
|
||||
},
|
||||
"decode_tps_approx": {
|
||||
"median": 35.510793077788634,
|
||||
"values": [
|
||||
34.055335354489465,
|
||||
35.510793077788634,
|
||||
35.522341967132114
|
||||
]
|
||||
},
|
||||
"output_tokens": [
|
||||
65,
|
||||
74,
|
||||
63
|
||||
],
|
||||
"reasoning_tokens": [
|
||||
57,
|
||||
66,
|
||||
55
|
||||
]
|
||||
},
|
||||
"prefill-32768": {
|
||||
"n": 3,
|
||||
"correct": 3,
|
||||
"ttft_s": {
|
||||
"median": 15.375977513001999,
|
||||
"values": [
|
||||
15.511759164000978,
|
||||
15.357147006005107,
|
||||
15.375977513001999
|
||||
]
|
||||
},
|
||||
"elapsed_s": {
|
||||
"median": 17.268514374001825,
|
||||
"values": [
|
||||
17.691939716001798,
|
||||
17.239325528003974,
|
||||
17.268514374001825
|
||||
]
|
||||
},
|
||||
"decode_tps_approx": {
|
||||
"median": 37.54862656190945,
|
||||
"values": [
|
||||
32.87688689931954,
|
||||
37.728914829613814,
|
||||
37.54862656190945
|
||||
]
|
||||
},
|
||||
"output_tokens": [
|
||||
70,
|
||||
69,
|
||||
69
|
||||
],
|
||||
"reasoning_tokens": [
|
||||
62,
|
||||
61,
|
||||
61
|
||||
]
|
||||
},
|
||||
"reuse-32768": {
|
||||
"n": 3,
|
||||
"correct": 3,
|
||||
"ttft_s": {
|
||||
"median": 1.3270910260034725,
|
||||
"values": [
|
||||
1.325252010996337,
|
||||
1.3272808809997514,
|
||||
1.3270910260034725
|
||||
]
|
||||
},
|
||||
"elapsed_s": {
|
||||
"median": 3.293290553992847,
|
||||
"values": [
|
||||
3.293290553992847,
|
||||
3.4540864679947845,
|
||||
3.2060768180017476
|
||||
]
|
||||
},
|
||||
"decode_tps_approx": {
|
||||
"median": 36.2184980346298,
|
||||
"values": [
|
||||
36.2184980346298,
|
||||
33.71980234098872,
|
||||
37.82214005937634
|
||||
]
|
||||
},
|
||||
"output_tokens": [
|
||||
69,
|
||||
70,
|
||||
69
|
||||
],
|
||||
"reasoning_tokens": [
|
||||
61,
|
||||
62,
|
||||
61
|
||||
]
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,39 @@
|
||||
{"label": "eager", "time": 1789663943.2948034, "event": "REQUEST_START", "name": "warmup", "profile": false, "prompt_tokens_preflight": 51}
|
||||
{"label": "eager", "time": 1789663945.101216, "event": "RESULT", "name": "warmup", "profile": false, "correct": true, "content": "\n\n323", "finish_reason": "stop", "usage": {"prompt_tokens": 51, "total_tokens": 83, "completion_tokens": 32, "completion_tokens_details": {"reasoning_tokens": 26}}, "max_tokens": 2048, "prompt_sha256": "4181e1a44a0bed10d00c74e822377bd75a6c6c9a1ea458f5170b8193e9fce959", "output_sha256": "4477e82dcd9f3f8da1dd56be073d72e4ba20121d6ce6fbc97ee54258fb4b2959", "ttft_s": 0.8800418169994373, "first_content_s": 1.7162362779999967, "elapsed_s": 1.800662048997765, "decode_tps_approx": 37.07271627093361, "metrics_delta": {"vllm:prefix_cache_queries_total": 51.0, "vllm:prefix_cache_hits_total": 0.0}}
|
||||
{"label": "eager", "time": 1789663945.1120813, "event": "REQUEST_START", "name": "short-0", "profile": false, "prompt_tokens_preflight": 65}
|
||||
{"label": "eager", "time": 1789663952.799862, "event": "RESULT", "name": "short-0", "profile": false, "correct": true, "content": "\n\n数据库索引类似于书籍的目录,它通过B+树等数据结构将数据按特定字段有序组织,使查询时能快速定位目标记录,避免全表扫描,从而大幅提升读取效率。然而,索引对写入操作有负面影响:每次插入、更新或删除数据时,数据库都需要同步维护索引结构,这会增加额外的I/O和计算开销,导致写入性能下降。因此,索引是\"以空间换时间\"的典型策略,需根据实际读写比例合理设计。", "finish_reason": "stop", "usage": {"prompt_tokens": 65, "total_tokens": 279, "completion_tokens": 214, "completion_tokens_details": {"reasoning_tokens": 111}}, "max_tokens": 512, "prompt_sha256": "40e7d1df3e96ad760d661cfaff312f5d732cb2edd5eab555b4b5c43cea4512ca", "output_sha256": "126232af3bc1889efc9c0dad67f2d47fc60d3c4ff56e699a83f1576c3763c274", "ttft_s": 0.33354272600263357, "first_content_s": 3.9721969300007913, "elapsed_s": 7.68168595500174, "decode_tps_approx": 28.987645227728436, "metrics_delta": {"vllm:prefix_cache_queries_total": 65.0, "vllm:prefix_cache_hits_total": 0.0}}
|
||||
{"label": "eager", "time": 1789663952.8174412, "event": "REQUEST_START", "name": "prefill-8192-0", "profile": false, "prompt_tokens_preflight": 8173}
|
||||
{"label": "eager", "time": 1789663958.6588385, "event": "RESULT", "name": "prefill-8192-0", "profile": false, "correct": true, "content": "\n\n21573", "finish_reason": "stop", "usage": {"prompt_tokens": 8173, "total_tokens": 8236, "completion_tokens": 63, "completion_tokens_details": {"reasoning_tokens": 55}}, "max_tokens": 2048, "prompt_sha256": "f59f163e1c7bb8a79d08df830d94d6d2e98fec457e9de75e3c09f6fca7ab0813", "output_sha256": "09d17d110b06ec13d532382c41ea29955b7edfa8dd4bd34163a08fa746e56610", "ttft_s": 4.006232126004761, "first_content_s": 5.657178637004108, "elapsed_s": 5.835497587999271, "decode_tps_approx": 35.584837627199924, "metrics_delta": {"vllm:prefix_cache_queries_total": 8173.0, "vllm:prefix_cache_hits_total": 0.0}}
|
||||
{"label": "eager", "time": 1789663958.6799748, "event": "REQUEST_START", "name": "reuse-8192-0", "profile": false, "prompt_tokens_preflight": 8173}
|
||||
{"label": "eager", "time": 1789663961.4358463, "event": "RESULT", "name": "reuse-8192-0", "profile": false, "correct": true, "content": "\n\n21573", "finish_reason": "stop", "usage": {"prompt_tokens": 8173, "total_tokens": 8236, "completion_tokens": 63, "completion_tokens_details": {"reasoning_tokens": 55}}, "max_tokens": 2048, "prompt_sha256": "f59f163e1c7bb8a79d08df830d94d6d2e98fec457e9de75e3c09f6fca7ab0813", "output_sha256": "09d17d110b06ec13d532382c41ea29955b7edfa8dd4bd34163a08fa746e56610", "ttft_s": 0.8856951920024585, "first_content_s": 2.5918558570047026, "elapsed_s": 2.7496982970042154, "decode_tps_approx": 34.74998178667343, "metrics_delta": {"vllm:prefix_cache_queries_total": 8173.0, "vllm:prefix_cache_hits_total": 6400.0}}
|
||||
{"label": "eager", "time": 1789663961.4779031, "event": "REQUEST_START", "name": "prefill-32768-0", "profile": false, "prompt_tokens_preflight": 32745}
|
||||
{"label": "eager", "time": 1789663979.3212829, "event": "RESULT", "name": "prefill-32768-0", "profile": false, "correct": true, "content": "\n\n90812", "finish_reason": "stop", "usage": {"prompt_tokens": 32745, "total_tokens": 32819, "completion_tokens": 74, "completion_tokens_details": {"reasoning_tokens": 66}}, "max_tokens": 2048, "prompt_sha256": "605f7b492252f50eaf60e8d128dde3f2d2e1b50dd26f0b45db88dac00e36a5ca", "output_sha256": "997501e1996b1c37af707622b1c4caf8dba93113e868f28210bd47ba700b5804", "ttft_s": 15.595506220997777, "first_content_s": 17.658189941001183, "elapsed_s": 17.837196191001567, "decode_tps_approx": 33.795316059955795, "metrics_delta": {"vllm:prefix_cache_queries_total": 32745.0, "vllm:prefix_cache_hits_total": 0.0}}
|
||||
{"label": "eager", "time": 1789663979.3644226, "event": "REQUEST_START", "name": "reuse-32768-0", "profile": false, "prompt_tokens_preflight": 32745}
|
||||
{"label": "eager", "time": 1789663982.9494953, "event": "RESULT", "name": "reuse-32768-0", "profile": false, "correct": true, "content": "\n\n90812", "finish_reason": "stop", "usage": {"prompt_tokens": 32745, "total_tokens": 32815, "completion_tokens": 70, "completion_tokens_details": {"reasoning_tokens": 62}}, "max_tokens": 2048, "prompt_sha256": "605f7b492252f50eaf60e8d128dde3f2d2e1b50dd26f0b45db88dac00e36a5ca", "output_sha256": "7eec57291da7ee95ae6d9214f8ed940a405e64beb9f49a4693387362aac3e14c", "ttft_s": 1.368702627994935, "first_content_s": 3.42049944599421, "elapsed_s": 3.579059300995141, "decode_tps_approx": 32.419700549347965, "metrics_delta": {"vllm:prefix_cache_queries_total": 32745.0, "vllm:prefix_cache_hits_total": 30400.0}}
|
||||
{"label": "eager", "time": 1789663982.957122, "event": "REQUEST_START", "name": "short-1", "profile": false, "prompt_tokens_preflight": 65}
|
||||
{"label": "eager", "time": 1789663990.3364515, "event": "RESULT", "name": "short-1", "profile": false, "correct": true, "content": "\n\n数据库索引类似于书籍的目录,它通过B+树等数据结构将数据按特定字段有序组织,使查询时能快速定位目标记录,避免全表扫描,从而大幅提升读取效率。然而,索引对写入操作有负面影响:每次插入、更新或删除数据时,数据库都需要同步维护索引结构,这会增加额外的I/O和计算开销,导致写入性能下降。因此,索引是\"以空间换时间\"的典型策略,需根据实际读写比例合理设计。", "finish_reason": "stop", "usage": {"prompt_tokens": 65, "total_tokens": 279, "completion_tokens": 214, "completion_tokens_details": {"reasoning_tokens": 111}}, "max_tokens": 512, "prompt_sha256": "40e7d1df3e96ad760d661cfaff312f5d732cb2edd5eab555b4b5c43cea4512ca", "output_sha256": "126232af3bc1889efc9c0dad67f2d47fc60d3c4ff56e699a83f1576c3763c274", "ttft_s": 0.2627224419993581, "first_content_s": 3.7929709769960027, "elapsed_s": 7.373869219998596, "decode_tps_approx": 29.953694756255302, "metrics_delta": {"vllm:prefix_cache_queries_total": 65.0, "vllm:prefix_cache_hits_total": 0.0}}
|
||||
{"label": "eager", "time": 1789663990.3527281, "event": "REQUEST_START", "name": "prefill-8192-1", "profile": false, "prompt_tokens_preflight": 8173}
|
||||
{"label": "eager", "time": 1789663995.9852123, "event": "RESULT", "name": "prefill-8192-1", "profile": false, "correct": true, "content": "\n\n21573", "finish_reason": "stop", "usage": {"prompt_tokens": 8173, "total_tokens": 8236, "completion_tokens": 63, "completion_tokens_details": {"reasoning_tokens": 55}}, "max_tokens": 2048, "prompt_sha256": "6c13deb2d4ddee54db95d7ebc8a1b87d24ed8a4112c33644c18b6b5d61ad7a16", "output_sha256": "09d17d110b06ec13d532382c41ea29955b7edfa8dd4bd34163a08fa746e56610", "ttft_s": 3.7754420729979756, "first_content_s": 5.466695101997175, "elapsed_s": 5.626387860997056, "decode_tps_approx": 35.041074313706375, "metrics_delta": {"vllm:prefix_cache_queries_total": 8173.0, "vllm:prefix_cache_hits_total": 0.0}}
|
||||
{"label": "eager", "time": 1789663996.0020921, "event": "REQUEST_START", "name": "reuse-8192-1", "profile": false, "prompt_tokens_preflight": 8173}
|
||||
{"label": "eager", "time": 1789663998.84421, "event": "RESULT", "name": "reuse-8192-1", "profile": false, "correct": true, "content": "\n\n21573", "finish_reason": "stop", "usage": {"prompt_tokens": 8173, "total_tokens": 8242, "completion_tokens": 69, "completion_tokens_details": {"reasoning_tokens": 61}}, "max_tokens": 2048, "prompt_sha256": "6c13deb2d4ddee54db95d7ebc8a1b87d24ed8a4112c33644c18b6b5d61ad7a16", "output_sha256": "f72eba21a095aa0a43356b25cb120346b064d7f0a253a63cee7579d2c3422b82", "ttft_s": 0.8737719170021592, "first_content_s": 2.675332922997768, "elapsed_s": 2.835690722997242, "decode_tps_approx": 36.047200027797416, "metrics_delta": {"vllm:prefix_cache_queries_total": 8173.0, "vllm:prefix_cache_hits_total": 6400.0}}
|
||||
{"label": "eager", "time": 1789663998.8886485, "event": "REQUEST_START", "name": "prefill-32768-1", "profile": false, "prompt_tokens_preflight": 32745}
|
||||
{"label": "eager", "time": 1789664016.1901195, "event": "RESULT", "name": "prefill-32768-1", "profile": false, "correct": true, "content": "\n\n90812", "finish_reason": "stop", "usage": {"prompt_tokens": 32745, "total_tokens": 32812, "completion_tokens": 67, "completion_tokens_details": {"reasoning_tokens": 59}}, "max_tokens": 2048, "prompt_sha256": "0a5e4ba33ff264d180b8a56d89184c0fefad59a96fdd348d655bd54fe5b5ed45", "output_sha256": "ca2e8f89b165bee277a12cac7bb1cbc8ce34eef2ee29ecd1a8a46aa71c502cfa", "ttft_s": 15.403379977004079, "first_content_s": 17.130364165001083, "elapsed_s": 17.295454844999767, "decode_tps_approx": 36.438073002765556, "metrics_delta": {"vllm:prefix_cache_queries_total": 32745.0, "vllm:prefix_cache_hits_total": 0.0}}
|
||||
{"label": "eager", "time": 1789664016.2312467, "event": "REQUEST_START", "name": "reuse-32768-1", "profile": false, "prompt_tokens_preflight": 32745}
|
||||
{"label": "eager", "time": 1789664019.7841623, "event": "RESULT", "name": "reuse-32768-1", "profile": false, "correct": true, "content": "\n\n90812", "finish_reason": "stop", "usage": {"prompt_tokens": 32745, "total_tokens": 32815, "completion_tokens": 70, "completion_tokens_details": {"reasoning_tokens": 62}}, "max_tokens": 2048, "prompt_sha256": "0a5e4ba33ff264d180b8a56d89184c0fefad59a96fdd348d655bd54fe5b5ed45", "output_sha256": "7eec57291da7ee95ae6d9214f8ed940a405e64beb9f49a4693387362aac3e14c", "ttft_s": 1.367758443004277, "first_content_s": 3.3833834170000046, "elapsed_s": 3.5468969909998123, "decode_tps_approx": 32.877578996597116, "metrics_delta": {"vllm:prefix_cache_queries_total": 32745.0, "vllm:prefix_cache_hits_total": 30400.0}}
|
||||
{"label": "eager", "time": 1789664019.7920012, "event": "REQUEST_START", "name": "short-2", "profile": false, "prompt_tokens_preflight": 65}
|
||||
{"label": "eager", "time": 1789664027.1539514, "event": "RESULT", "name": "short-2", "profile": false, "correct": true, "content": "\n\n数据库索引类似于书籍的目录,它通过B+树等数据结构将数据按特定字段有序组织,使查询时能快速定位目标记录,避免全表扫描,从而大幅提升读取效率。然而,索引对写入操作有负面影响:每次插入、更新或删除数据时,数据库都需要同步维护索引结构,这会增加额外的I/O和计算开销,导致写入性能下降。因此,索引是\"以空间换时间\"的典型策略,需根据实际读写比例合理设计。", "finish_reason": "stop", "usage": {"prompt_tokens": 65, "total_tokens": 279, "completion_tokens": 214, "completion_tokens_details": {"reasoning_tokens": 111}}, "max_tokens": 512, "prompt_sha256": "40e7d1df3e96ad760d661cfaff312f5d732cb2edd5eab555b4b5c43cea4512ca", "output_sha256": "126232af3bc1889efc9c0dad67f2d47fc60d3c4ff56e699a83f1576c3763c274", "ttft_s": 0.26176839799882146, "first_content_s": 3.792491918997257, "elapsed_s": 7.35648784199293, "decode_tps_approx": 30.02293954286349, "metrics_delta": {"vllm:prefix_cache_queries_total": 65.0, "vllm:prefix_cache_hits_total": 0.0}}
|
||||
{"label": "eager", "time": 1789664027.17096, "event": "REQUEST_START", "name": "prefill-8192-2", "profile": false, "prompt_tokens_preflight": 8173}
|
||||
{"label": "eager", "time": 1789664032.9165242, "event": "RESULT", "name": "prefill-8192-2", "profile": false, "correct": true, "content": "\n\n21573", "finish_reason": "stop", "usage": {"prompt_tokens": 8173, "total_tokens": 8235, "completion_tokens": 62, "completion_tokens_details": {"reasoning_tokens": 54}}, "max_tokens": 2048, "prompt_sha256": "b49f6dbbf80f49ad860a60b0a7c58dfc2122e38bb8493ae155790f22305d8f7a", "output_sha256": "0eb9ce304c0547f0814d0590d36cb2144dcc242a2c550f3ba9f8a70b2e22b203", "ttft_s": 3.7875278329956927, "first_content_s": 5.574596539998311, "elapsed_s": 5.739436170995759, "decode_tps_approx": 32.64277985613848, "metrics_delta": {"vllm:prefix_cache_queries_total": 8173.0, "vllm:prefix_cache_hits_total": 0.0}}
|
||||
{"label": "eager", "time": 1789664032.9353032, "event": "REQUEST_START", "name": "reuse-8192-2", "profile": false, "prompt_tokens_preflight": 8173}
|
||||
{"label": "eager", "time": 1789664035.7520719, "event": "RESULT", "name": "reuse-8192-2", "profile": false, "correct": true, "content": "\n\n21573", "finish_reason": "stop", "usage": {"prompt_tokens": 8173, "total_tokens": 8242, "completion_tokens": 69, "completion_tokens_details": {"reasoning_tokens": 61}}, "max_tokens": 2048, "prompt_sha256": "b49f6dbbf80f49ad860a60b0a7c58dfc2122e38bb8493ae155790f22305d8f7a", "output_sha256": "f72eba21a095aa0a43356b25cb120346b064d7f0a253a63cee7579d2c3422b82", "ttft_s": 0.8803304789980757, "first_content_s": 2.651041837001685, "elapsed_s": 2.811309409997193, "decode_tps_approx": 36.74879746387076, "metrics_delta": {"vllm:prefix_cache_queries_total": 8173.0, "vllm:prefix_cache_hits_total": 6400.0}}
|
||||
{"label": "eager", "time": 1789664035.7953794, "event": "REQUEST_START", "name": "prefill-32768-2", "profile": false, "prompt_tokens_preflight": 32745}
|
||||
{"label": "eager", "time": 1789664053.153125, "event": "RESULT", "name": "prefill-32768-2", "profile": false, "correct": true, "content": "\n\n90812", "finish_reason": "stop", "usage": {"prompt_tokens": 32745, "total_tokens": 32814, "completion_tokens": 69, "completion_tokens_details": {"reasoning_tokens": 61}}, "max_tokens": 2048, "prompt_sha256": "d8796cea8608150048ce4f253e451ba96da5416ccfea4a6560dc2f6bb2412004", "output_sha256": "8b503330d6b70b8845bab80540601fc48c8f840d54d07f38d80250c8167de89f", "ttft_s": 15.41581593700539, "first_content_s": 17.18937433300016, "elapsed_s": 17.35202403000585, "decode_tps_approx": 36.713832650404036, "metrics_delta": {"vllm:prefix_cache_queries_total": 32745.0, "vllm:prefix_cache_hits_total": 0.0}}
|
||||
{"label": "eager", "time": 1789664053.1986973, "event": "REQUEST_START", "name": "reuse-32768-2", "profile": false, "prompt_tokens_preflight": 32745}
|
||||
{"label": "eager", "time": 1789664056.5784104, "event": "RESULT", "name": "reuse-32768-2", "profile": false, "correct": true, "content": "\n\n90812", "finish_reason": "stop", "usage": {"prompt_tokens": 32745, "total_tokens": 32814, "completion_tokens": 69, "completion_tokens_details": {"reasoning_tokens": 61}}, "max_tokens": 2048, "prompt_sha256": "d8796cea8608150048ce4f253e451ba96da5416ccfea4a6560dc2f6bb2412004", "output_sha256": "8b503330d6b70b8845bab80540601fc48c8f840d54d07f38d80250c8167de89f", "ttft_s": 1.357998088002205, "first_content_s": 3.2952792900032364, "elapsed_s": 3.3747847780032316, "decode_tps_approx": 33.71865583479496, "metrics_delta": {"vllm:prefix_cache_queries_total": 32745.0, "vllm:prefix_cache_hits_total": 30400.0}}
|
||||
{"label": "eager", "time": 1789664056.5835583, "event": "REQUEST_START", "name": "trace-short", "profile": true, "prompt_tokens_preflight": 65}
|
||||
{"label": "eager", "time": 1789664067.1052518, "event": "RESULT", "name": "trace-short", "profile": true, "correct": true, "content": "\n\n数据库索引类似于书籍的目录,它通过B+树等数据结构将数据按特定字段有序组织,使查询时能快速定位目标记录,避免全表扫描,从而大幅提升读取效率。然而,索引对写入操作有负面影响:每次插入、更新或删除数据时,数据库都需要同步维护索引结构,这会增加额外的I/O和计算开销,导致写入性能下降。因此,索引是\"以空间换时间\"的典型策略,需根据实际读写比例合理设计。", "finish_reason": "stop", "usage": {"prompt_tokens": 65, "total_tokens": 279, "completion_tokens": 214, "completion_tokens_details": {"reasoning_tokens": 111}}, "max_tokens": 512, "prompt_sha256": "40e7d1df3e96ad760d661cfaff312f5d732cb2edd5eab555b4b5c43cea4512ca", "output_sha256": "126232af3bc1889efc9c0dad67f2d47fc60d3c4ff56e699a83f1576c3763c274", "ttft_s": 0.28372384200338274, "first_content_s": 3.9164422460016795, "elapsed_s": 10.166879840006004, "decode_tps_approx": 29.24182266761644, "metrics_delta": {"vllm:prefix_cache_queries_total": 65.0, "vllm:prefix_cache_hits_total": 0.0}}
|
||||
{"label": "eager", "time": 1789664067.137696, "event": "REQUEST_START", "name": "trace-prefill-32768", "profile": true, "prompt_tokens_preflight": 32742}
|
||||
{"label": "eager", "time": 1789664086.9158142, "event": "RESULT", "name": "trace-prefill-32768", "profile": true, "correct": true, "content": "\n\n90812", "finish_reason": "stop", "usage": {"prompt_tokens": 32742, "total_tokens": 32809, "completion_tokens": 67, "completion_tokens_details": {"reasoning_tokens": 59}}, "max_tokens": 2048, "prompt_sha256": "464e9da0976ca56ddece006a237605bca8b60d3dd0aab538abbc0aaa1696815c", "output_sha256": "ca2e8f89b165bee277a12cac7bb1cbc8ce34eef2ee29ecd1a8a46aa71c502cfa", "ttft_s": 15.503820653997536, "first_content_s": 17.303895275996183, "elapsed_s": 19.650969089998398, "decode_tps_approx": 35.058139901273435, "metrics_delta": {"vllm:prefix_cache_queries_total": 32742.0, "vllm:prefix_cache_hits_total": 0.0}}
|
||||
{"label": "eager", "time": 1789664086.9506161, "event": "REQUEST_START", "name": "trace-reuse-32768", "profile": true, "prompt_tokens_preflight": 32742}
|
||||
{"label": "eager", "time": 1789664093.2510297, "event": "RESULT", "name": "trace-reuse-32768", "profile": true, "correct": true, "content": "\n\n90812", "finish_reason": "stop", "usage": {"prompt_tokens": 32742, "total_tokens": 32810, "completion_tokens": 68, "completion_tokens_details": {"reasoning_tokens": 60}}, "max_tokens": 2048, "prompt_sha256": "464e9da0976ca56ddece006a237605bca8b60d3dd0aab538abbc0aaa1696815c", "output_sha256": "ce351e2ffc5a1b2641f170a3fa56c81034b4d81226c524d10936eb65c7eb3926", "ttft_s": 1.385693889002141, "first_content_s": 3.091133230998821, "elapsed_s": 6.173271200997988, "decode_tps_approx": 37.499826794181516, "metrics_delta": {"vllm:prefix_cache_queries_total": 32742.0, "vllm:prefix_cache_hits_total": 30400.0}}
|
||||
{"label": "eager", "time": 1789664093.251096, "event": "PROFILE_SUITE_PASS"}
|
||||
File diff suppressed because it is too large
Load Diff
@@ -0,0 +1,823 @@
|
||||
{
|
||||
"file": "trace.2.sqlite",
|
||||
"gpu_activity_span_s": 17.45799776,
|
||||
"gpu_busy_union_s": 16.999974145,
|
||||
"gpu_idle_within_span_s": 0.4580236150000019,
|
||||
"kernel_count": 107899,
|
||||
"graph_node_kernel_count": 0,
|
||||
"api": {
|
||||
"cudaStreamSynchronize_v3020": [
|
||||
92,
|
||||
6.388727359999995
|
||||
],
|
||||
"cuLaunchKernelEx": [
|
||||
30313,
|
||||
2.942097295999861
|
||||
],
|
||||
"cudaLaunchKernelExC_v11060": [
|
||||
27262,
|
||||
2.8157289279998983
|
||||
],
|
||||
"cudaLaunchKernel_v7000": [
|
||||
34403,
|
||||
2.1073674239999285
|
||||
],
|
||||
"cuLaunchKernel": [
|
||||
15921,
|
||||
0.18051391999999974
|
||||
],
|
||||
"cudaMemcpyAsync_v3020": [
|
||||
2359,
|
||||
0.02657009600000002
|
||||
],
|
||||
"cudaEventSynchronize_v3020": [
|
||||
46,
|
||||
0.01590248
|
||||
],
|
||||
"cudaEventRecordWithFlags_v11010": [
|
||||
3413,
|
||||
0.006630880000000002
|
||||
],
|
||||
"cudaMemsetAsync_v3020": [
|
||||
736,
|
||||
0.0055193599999999975
|
||||
],
|
||||
"cuTensorMapEncodeTiled": [
|
||||
27416,
|
||||
0.004886224000000271
|
||||
],
|
||||
"cuKernelGetFunction": [
|
||||
15921,
|
||||
0.0038959359999997194
|
||||
],
|
||||
"cudaGetDriverEntryPointByVersion_v12050": [
|
||||
26496,
|
||||
0.0037329280000005572
|
||||
],
|
||||
"cuKernelGetName": [
|
||||
35159,
|
||||
0.0033410560000001516
|
||||
],
|
||||
"cudaStreamWaitEvent_v3020": [
|
||||
2588,
|
||||
0.0021758880000000004
|
||||
],
|
||||
"cudaEventQuery_v3020": [
|
||||
1494,
|
||||
0.0019392639999999957
|
||||
],
|
||||
"cudaStreamIsCapturing_v10000": [
|
||||
6960,
|
||||
0.0014392959999999955
|
||||
],
|
||||
"cudaEventCreateWithFlags_v3020": [
|
||||
92,
|
||||
0.000234256
|
||||
],
|
||||
"cudaEventDestroy_v3020": [
|
||||
92,
|
||||
0.00018534400000000007
|
||||
]
|
||||
},
|
||||
"nvtx": {
|
||||
"PLE_lookup_tokens=1600": {
|
||||
"count": 20,
|
||||
"sum_s": 5.58014968,
|
||||
"union_s": 5.58014968,
|
||||
"gpu_overlap_s": 5.318215141
|
||||
},
|
||||
"PLE_wait_for_GPU": {
|
||||
"count": 46,
|
||||
"sum_s": 6.389229296,
|
||||
"union_s": 6.389229296,
|
||||
"gpu_overlap_s": 6.341070402
|
||||
},
|
||||
"PLE_CPU_dedup": {
|
||||
"count": 46,
|
||||
"sum_s": 0.01914024,
|
||||
"union_s": 0.01914024,
|
||||
"gpu_overlap_s": 0.0
|
||||
},
|
||||
"PLE_CPU_gather": {
|
||||
"count": 46,
|
||||
"sum_s": 0.275209776,
|
||||
"union_s": 0.275209776,
|
||||
"gpu_overlap_s": 0.0
|
||||
},
|
||||
"PLE_lookup_tokens=742": {
|
||||
"count": 1,
|
||||
"sum_s": 0.29896912,
|
||||
"union_s": 0.29896912,
|
||||
"gpu_overlap_s": 0.278459202
|
||||
},
|
||||
"PLE_lookup_tokens=3": {
|
||||
"count": 25,
|
||||
"sum_s": 0.820145456,
|
||||
"union_s": 0.820145456,
|
||||
"gpu_overlap_s": 0.746778411
|
||||
}
|
||||
},
|
||||
"dispatch": {
|
||||
"CG_DISPATCH:actual=1600:mode=NONE:padded=1600": 40,
|
||||
"CG_DISPATCH:actual=1:mode=NONE:padded=1": 46,
|
||||
"CG_DISPATCH:actual=742:mode=NONE:padded=742": 2,
|
||||
"CG_DISPATCH:actual=3:mode=NONE:padded=3": 50
|
||||
},
|
||||
"dispatch_timeline": [
|
||||
{
|
||||
"start_ns": 45828672,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 415343408,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 441920176,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 467464272,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 1117406208,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 1146479600,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 1170640896,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 1839789152,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 1870939456,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 1896334032,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 2568046544,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 2598479856,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 2624039184,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 3305750304,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 3335311984,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 3361130160,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 4046256096,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 4076056912,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 4101620656,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 4786128624,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 4816276720,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 4842056960,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 5522623936,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 5552300720,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 5576921440,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 6259346144,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 6290320240,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 6315131568,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 7002476032,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 7032926192,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 7058327584,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 7753655264,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 7785387856,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 7811813408,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 8504434640,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 8535191696,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 8561506848,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 9253677696,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 9285233312,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 9311217936,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 10009614064,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 10040506016,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 10066581360,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 10757969120,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 10789570400,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 10815237984,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 11519444752,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 11550823200,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 11575788896,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 12278226144,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 12309497040,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 12335515504,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 13034051120,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 13066055552,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 13091963328,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 13790163952,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 13821175296,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 13846090352,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 14557613040,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 14589908560,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 14615422640,
|
||||
"name": "CG_DISPATCH:actual=742:mode=NONE:padded=742"
|
||||
},
|
||||
{
|
||||
"start_ns": 15203740016,
|
||||
"name": "CG_DISPATCH:actual=742:mode=NONE:padded=742"
|
||||
},
|
||||
{
|
||||
"start_ns": 15220760336,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 15234841296,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 15491074784,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 15493643856,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 15497135184,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 15564983824,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 15566982016,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 15570203408,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 15644290848,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 15646296832,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 15649329520,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 15729121168,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 15731019920,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 15734338320,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 15813356656,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 15815902416,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 15819264448,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 15894042800,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 15895960816,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 15898832688,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 15982235440,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 15984929392,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 15988404656,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 16064077552,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 16066769024,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 16070591120,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 16140542304,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 16142492784,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 16145392256,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 16230150416,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 16232897584,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 16236422000,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 16305893904,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 16307780160,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 16310441504,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 16394138896,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 16396838944,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 16400267168,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 16476712080,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 16479446496,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 16482904144,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 16559135840,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 16561788800,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 16565117088,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 16632650032,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 16634497504,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 16637262864,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 16722767040,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 16725243152,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 16728652272,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 16807455696,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 16810202752,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 16813561920,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 16894260000,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 16898134688,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 16903894256,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 16982500864,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 16985299840,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 16988503344,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 17054985488,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 17056901872,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 17059848448,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 17143682800,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 17146548688,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 17150240800,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 17226658560,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 17229352608,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 17233024096,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 17306935664,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 17308968096,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 17311992480,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 17388634448,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 17391419008,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 17394878992,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 17463051952,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 17464913120,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
}
|
||||
],
|
||||
"top_kernels": [
|
||||
[
|
||||
"device_kernel",
|
||||
[
|
||||
4416,
|
||||
5.785133960999994
|
||||
]
|
||||
],
|
||||
[
|
||||
"Kernel2",
|
||||
[
|
||||
15921,
|
||||
1.6569721249999942
|
||||
]
|
||||
],
|
||||
[
|
||||
"_qsa_sparse_paged_gqa_splitk_kernel",
|
||||
[
|
||||
644,
|
||||
1.481612926999999
|
||||
]
|
||||
],
|
||||
[
|
||||
"nvjet_sm121_tst_mma_128x208x64_2_32x104x64_tmaAB_bz_TNNN",
|
||||
[
|
||||
1000,
|
||||
1.3622323199999982
|
||||
]
|
||||
],
|
||||
[
|
||||
"_hc_combine_norm_kernel",
|
||||
[
|
||||
4646,
|
||||
0.8981812170000064
|
||||
]
|
||||
],
|
||||
[
|
||||
"kernel",
|
||||
[
|
||||
1510,
|
||||
0.8489803810000001
|
||||
]
|
||||
],
|
||||
[
|
||||
"_hc_gate_mix_kernel",
|
||||
[
|
||||
4738,
|
||||
0.6790835519999859
|
||||
]
|
||||
],
|
||||
[
|
||||
"nvjet_sm121_tst_mma_128x240x64_2_32x120x64_tmaAB_bz_TNNN",
|
||||
[
|
||||
980,
|
||||
0.5484992310000006
|
||||
]
|
||||
],
|
||||
[
|
||||
"nvjet_sm121_tst_mma_128x128x64_3_32x64x64_tmaAB_bz_TNNN",
|
||||
[
|
||||
2000,
|
||||
0.484882112
|
||||
]
|
||||
],
|
||||
[
|
||||
"finalizeMoeRoutingKernel",
|
||||
[
|
||||
2208,
|
||||
0.4100567680000026
|
||||
]
|
||||
],
|
||||
[
|
||||
"_causal_conv1d_fwd_kernel",
|
||||
[
|
||||
756,
|
||||
0.2423829440000001
|
||||
]
|
||||
],
|
||||
[
|
||||
"_fused_post_conv_kernel",
|
||||
[
|
||||
756,
|
||||
0.23299334399999985
|
||||
]
|
||||
],
|
||||
[
|
||||
"doActivationKernel",
|
||||
[
|
||||
2208,
|
||||
0.2290338470000007
|
||||
]
|
||||
],
|
||||
[
|
||||
"elementwise_kernel",
|
||||
[
|
||||
6060,
|
||||
0.22898364799999965
|
||||
]
|
||||
],
|
||||
[
|
||||
"kernel_cutlass_kernel_flashinfergdn_kernelsdelta_rule_dsldelta_rule_sm120_FullyFusedDeltaRuleSm120_object_at__tensorptrf32gmemalign16o1_tensorptrf32gmemalign16o1_CopyAtom_ThrID10_TVLayout_0",
|
||||
[
|
||||
756,
|
||||
0.19385532799999977
|
||||
]
|
||||
],
|
||||
[
|
||||
"vectorized_elementwise_kernel",
|
||||
[
|
||||
10976,
|
||||
0.18249009600000335
|
||||
]
|
||||
],
|
||||
[
|
||||
"expandInputRowsKernel",
|
||||
[
|
||||
2208,
|
||||
0.17563680399999979
|
||||
]
|
||||
],
|
||||
[
|
||||
"layer_norm_fwd_kernel",
|
||||
[
|
||||
756,
|
||||
0.17546492700000002
|
||||
]
|
||||
],
|
||||
[
|
||||
"sm120_fp8_fp4_gemm_1d1d_impl",
|
||||
[
|
||||
184,
|
||||
0.16484399999999996
|
||||
]
|
||||
],
|
||||
[
|
||||
"nvjet_sm121_tst_mma_128x192x64_2_32x96x64_tmaAB_bz_TNNN",
|
||||
[
|
||||
980,
|
||||
0.13451315200000008
|
||||
]
|
||||
]
|
||||
],
|
||||
"notes": "GPU busy is interval union, not utilization counter. CPU/API/NVTX times overlap GPU work and each other; do not sum them as exclusive costs."
|
||||
}
|
||||
@@ -0,0 +1,583 @@
|
||||
{
|
||||
"file": "trace.3.sqlite",
|
||||
"gpu_activity_span_s": 3.245924224,
|
||||
"gpu_busy_union_s": 3.090907949,
|
||||
"gpu_idle_within_span_s": 0.15501627499999993,
|
||||
"kernel_count": 58496,
|
||||
"graph_node_kernel_count": 0,
|
||||
"api": {
|
||||
"cudaStreamSynchronize_v3020": [
|
||||
52,
|
||||
1.0780331839999997
|
||||
],
|
||||
"cudaLaunchKernelExC_v11060": [
|
||||
16696,
|
||||
0.3360265120000019
|
||||
],
|
||||
"cuLaunchKernelEx": [
|
||||
13138,
|
||||
0.22730212800000002
|
||||
],
|
||||
"cudaLaunchKernel_v7000": [
|
||||
16229,
|
||||
0.21477230400000039
|
||||
],
|
||||
"cuLaunchKernel": [
|
||||
12433,
|
||||
0.05239439999999998
|
||||
],
|
||||
"cudaMemcpyAsync_v3020": [
|
||||
939,
|
||||
0.00983091199999999
|
||||
],
|
||||
"cudaEventSynchronize_v3020": [
|
||||
26,
|
||||
0.0089596
|
||||
],
|
||||
"cudaEventRecordWithFlags_v11010": [
|
||||
2877,
|
||||
0.004890783999999996
|
||||
],
|
||||
"cuKernelGetFunction": [
|
||||
12433,
|
||||
0.0027621279999998907
|
||||
],
|
||||
"cuTensorMapEncodeTiled": [
|
||||
15496,
|
||||
0.002756256000000075
|
||||
],
|
||||
"cudaMemsetAsync_v3020": [
|
||||
416,
|
||||
0.0023939039999999988
|
||||
],
|
||||
"cudaGetDriverEntryPointByVersion_v12050": [
|
||||
14976,
|
||||
0.001993504000000201
|
||||
],
|
||||
"cudaStreamWaitEvent_v3020": [
|
||||
2430,
|
||||
0.0018610560000000033
|
||||
],
|
||||
"cuKernelGetName": [
|
||||
16301,
|
||||
0.0014150560000001082
|
||||
],
|
||||
"cudaStreamIsCapturing_v10000": [
|
||||
5810,
|
||||
0.0011180959999999769
|
||||
],
|
||||
"cudaEventQuery_v3020": [
|
||||
798,
|
||||
0.001018544
|
||||
],
|
||||
"cudaEventCreateWithFlags_v3020": [
|
||||
52,
|
||||
0.000134192
|
||||
],
|
||||
"cudaEventDestroy_v3020": [
|
||||
52,
|
||||
0.00010068799999999997
|
||||
]
|
||||
},
|
||||
"nvtx": {
|
||||
"PLE_lookup_tokens=1600": {
|
||||
"count": 1,
|
||||
"sum_s": 0.018957376,
|
||||
"union_s": 0.018957376,
|
||||
"gpu_overlap_s": 0.011426528
|
||||
},
|
||||
"PLE_wait_for_GPU": {
|
||||
"count": 26,
|
||||
"sum_s": 1.078286144,
|
||||
"union_s": 1.078286144,
|
||||
"gpu_overlap_s": 1.058767637
|
||||
},
|
||||
"PLE_CPU_dedup": {
|
||||
"count": 26,
|
||||
"sum_s": 0.002243088,
|
||||
"union_s": 0.002243088,
|
||||
"gpu_overlap_s": 0.0
|
||||
},
|
||||
"PLE_CPU_gather": {
|
||||
"count": 26,
|
||||
"sum_s": 0.062812192,
|
||||
"union_s": 0.062812192,
|
||||
"gpu_overlap_s": 0.0
|
||||
},
|
||||
"PLE_lookup_tokens=742": {
|
||||
"count": 1,
|
||||
"sum_s": 0.292337536,
|
||||
"union_s": 0.292337536,
|
||||
"gpu_overlap_s": 0.279791319
|
||||
},
|
||||
"PLE_lookup_tokens=3": {
|
||||
"count": 24,
|
||||
"sum_s": 0.842259408,
|
||||
"union_s": 0.842259408,
|
||||
"gpu_overlap_s": 0.771610318
|
||||
}
|
||||
},
|
||||
"dispatch": {
|
||||
"CG_DISPATCH:actual=1600:mode=NONE:padded=1600": 2,
|
||||
"CG_DISPATCH:actual=1:mode=NONE:padded=1": 26,
|
||||
"CG_DISPATCH:actual=742:mode=NONE:padded=742": 2,
|
||||
"CG_DISPATCH:actual=3:mode=NONE:padded=3": 48
|
||||
},
|
||||
"dispatch_timeline": [
|
||||
{
|
||||
"start_ns": 46439136,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 440098864,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 472554240,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 498029616,
|
||||
"name": "CG_DISPATCH:actual=742:mode=NONE:padded=742"
|
||||
},
|
||||
{
|
||||
"start_ns": 1080640368,
|
||||
"name": "CG_DISPATCH:actual=742:mode=NONE:padded=742"
|
||||
},
|
||||
{
|
||||
"start_ns": 1097925088,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 1109689376,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 1374741120,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 1377837856,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 1381724192,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 1449909984,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 1452632352,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 1456060352,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 1537442544,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 1540106512,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 1543362000,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 1622817104,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 1625489424,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 1628952064,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 1697539568,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 1699380000,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 1702138000,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 1777756128,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 1779549280,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 1782099440,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 1865886816,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 1868546304,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 1872017632,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 1939645984,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 1941472672,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 1944250112,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 2021839712,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 2023852576,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 2026597104,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 2103480432,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 2105282848,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 2107900064,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 2193835440,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 2196580896,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 2199840944,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 2282745504,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 2285692352,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 2289777504,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 2362513824,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 2365324800,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 2368818832,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 2436200000,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 2438187600,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 2441020928,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 2516737184,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 2518684848,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 2521416768,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 2596916928,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 2598750864,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 2601295488,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 2678072112,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 2679938368,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 2682459200,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 2756745440,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 2758542912,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 2761074960,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 2838330208,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 2840177632,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 2842655760,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 2918308832,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 2920073040,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 2922520640,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 3006922960,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 3008938128,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 3011760224,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 3094064288,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 3096710352,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 3100043552,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 3174526160,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 3177165152,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 3180415120,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 3257416576,
|
||||
"name": "CG_DISPATCH:actual=3:mode=NONE:padded=3"
|
||||
},
|
||||
{
|
||||
"start_ns": 3260041984,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
}
|
||||
],
|
||||
"top_kernels": [
|
||||
[
|
||||
"Kernel2",
|
||||
[
|
||||
12433,
|
||||
1.083779682999996
|
||||
]
|
||||
],
|
||||
[
|
||||
"device_kernel",
|
||||
[
|
||||
2496,
|
||||
0.8852412400000004
|
||||
]
|
||||
],
|
||||
[
|
||||
"kernel",
|
||||
[
|
||||
360,
|
||||
0.38886832
|
||||
]
|
||||
],
|
||||
[
|
||||
"_qsa_sparse_paged_gqa_splitk_kernel",
|
||||
[
|
||||
364,
|
||||
0.12547817500000008
|
||||
]
|
||||
],
|
||||
[
|
||||
"_hc_combine_norm_kernel",
|
||||
[
|
||||
2626,
|
||||
0.06954325899999982
|
||||
]
|
||||
],
|
||||
[
|
||||
"nvjet_sm121_tst_mma_128x208x64_2_32x104x64_tmaAB_bz_TNNN",
|
||||
[
|
||||
50,
|
||||
0.068864416
|
||||
]
|
||||
],
|
||||
[
|
||||
"_hc_gate_mix_kernel",
|
||||
[
|
||||
2678,
|
||||
0.04872185600000148
|
||||
]
|
||||
],
|
||||
[
|
||||
"finalizeMoeRoutingKernel",
|
||||
[
|
||||
1248,
|
||||
0.03866553600000007
|
||||
]
|
||||
],
|
||||
[
|
||||
"sm120_fp8_fp4_gemm_1d1d_impl",
|
||||
[
|
||||
104,
|
||||
0.038346688000000004
|
||||
]
|
||||
],
|
||||
[
|
||||
"nvjet_sm121_tst_mma_192x160x64_2_48x80x64_tmaAB_bz_TNNN",
|
||||
[
|
||||
51,
|
||||
0.034701248
|
||||
]
|
||||
],
|
||||
[
|
||||
"nvjet_sm121_tst_mma_128x240x64_2_32x120x64_tmaAB_bz_TNNN",
|
||||
[
|
||||
49,
|
||||
0.02691910399999999
|
||||
]
|
||||
],
|
||||
[
|
||||
"nvjet_sm121_tst_mma_128x128x64_3_32x64x64_tmaAB_bz_TNNN",
|
||||
[
|
||||
100,
|
||||
0.024421632000000002
|
||||
]
|
||||
],
|
||||
[
|
||||
"gdn_decode_post_conv_mtp_kernel",
|
||||
[
|
||||
864,
|
||||
0.023706272000000014
|
||||
]
|
||||
],
|
||||
[
|
||||
"expandInputRowsKernel",
|
||||
[
|
||||
1248,
|
||||
0.022278243999999982
|
||||
]
|
||||
],
|
||||
[
|
||||
"doActivationKernel",
|
||||
[
|
||||
1248,
|
||||
0.02122110900000002
|
||||
]
|
||||
],
|
||||
[
|
||||
"elementwise_kernel",
|
||||
[
|
||||
2004,
|
||||
0.020444831999999843
|
||||
]
|
||||
],
|
||||
[
|
||||
"vectorized_elementwise_kernel",
|
||||
[
|
||||
5307,
|
||||
0.018513366999999715
|
||||
]
|
||||
],
|
||||
[
|
||||
"_causal_conv1d_fwd_kernel",
|
||||
[
|
||||
72,
|
||||
0.017434495999999997
|
||||
]
|
||||
],
|
||||
[
|
||||
"nvjet_sm121_tst_mma_128x224x64_2_64x56x64_tmaAB_bz_TNNN",
|
||||
[
|
||||
50,
|
||||
0.016634912
|
||||
]
|
||||
],
|
||||
[
|
||||
"_fused_post_conv_kernel",
|
||||
[
|
||||
72,
|
||||
0.015161888000000002
|
||||
]
|
||||
]
|
||||
],
|
||||
"notes": "GPU busy is interval union, not utilization counter. CPU/API/NVTX times overlap GPU work and each other; do not sum them as exclusive costs."
|
||||
}
|
||||
@@ -0,0 +1,60 @@
|
||||
{
|
||||
"context_limit": 131072,
|
||||
"batch_token_limit": 2048,
|
||||
"observed_attention_block_size": 1600,
|
||||
"graph_capture_sizes": [
|
||||
1,
|
||||
2,
|
||||
4,
|
||||
8,
|
||||
12
|
||||
],
|
||||
"nsys_version": "NVIDIA Nsight Systems version 2025.3.2.474-253236389321v0",
|
||||
"trace_cases": {
|
||||
"1": "short",
|
||||
"2": "fresh-32768",
|
||||
"3": "reused-32768"
|
||||
},
|
||||
"files": {
|
||||
"graph/trace.2.nsys-rep": {
|
||||
"bytes": 7688024,
|
||||
"sha256": "443a6e052586d3f4bfe7c6f1492500e495374e8e52e6de5766a437e53d73d146"
|
||||
},
|
||||
"graph/trace.3.nsys-rep": {
|
||||
"bytes": 4115977,
|
||||
"sha256": "15e7602b8dd8d92e9f295fcc1a24df4a6245a36339555c83f04dda32c7d5dc86"
|
||||
},
|
||||
"graph/trace.1.nsys-rep": {
|
||||
"bytes": 9244219,
|
||||
"sha256": "4eceb2f5241678ac387064369ccfb485d741dec140d0a1ace46866da0d70e687"
|
||||
},
|
||||
"eager/trace.2.nsys-rep": {
|
||||
"bytes": 8540396,
|
||||
"sha256": "543a9c1a5698de7e2ce7fe9e25dab546d64236e0f6f2014648223f664ff967a1"
|
||||
},
|
||||
"eager/trace.3.nsys-rep": {
|
||||
"bytes": 4635497,
|
||||
"sha256": "af251e99de398302611f691f39073e1a48ae766bd9473931b2437c46e2e42227"
|
||||
},
|
||||
"eager/trace.1.nsys-rep": {
|
||||
"bytes": 15842503,
|
||||
"sha256": "0390fdffffaeee0b78809d64cbf8696e8301d5730de36c2d67da4e1aa18bc5d0"
|
||||
},
|
||||
"vllm_ple_mmap.original.py": {
|
||||
"bytes": 40227,
|
||||
"sha256": "02adb76b789f9fccf472e922e3be7b4b2a95962ae7d7c3eb99bfeac6052dccad"
|
||||
},
|
||||
"cudagraph_utils.original.py": {
|
||||
"bytes": 46652,
|
||||
"sha256": "e123e808a8b6a616f69cc97f8b0d25d77f0eb35d8ebc91c4cdbc505f649927c0"
|
||||
},
|
||||
"vllm_ple_mmap.profile.py": {
|
||||
"bytes": 40773,
|
||||
"sha256": "2dff4fe71a5f2f5a8579af09eebf798d49da0d1554a56819ab475b86aebe17e4"
|
||||
},
|
||||
"cudagraph_utils.profile.py": {
|
||||
"bytes": 47119,
|
||||
"sha256": "ef6dccbd27c6b109691a2ab2b9a74cc08c1d3055498ff91638f78af1fe5faf66"
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,39 @@
|
||||
{"label": "graph", "time": 1789664773.0911076, "event": "REQUEST_START", "name": "warmup", "profile": false, "prompt_tokens_preflight": 51}
|
||||
{"label": "graph", "time": 1789664774.6022928, "event": "RESULT", "name": "warmup", "profile": false, "correct": true, "content": "\n\n323", "finish_reason": "stop", "usage": {"prompt_tokens": 51, "total_tokens": 83, "completion_tokens": 32, "completion_tokens_details": {"reasoning_tokens": 26}}, "max_tokens": 2048, "prompt_sha256": "4181e1a44a0bed10d00c74e822377bd75a6c6c9a1ea458f5170b8193e9fce959", "output_sha256": "4477e82dcd9f3f8da1dd56be073d72e4ba20121d6ce6fbc97ee54258fb4b2959", "ttft_s": 0.6364480190022732, "first_content_s": 1.4175457309975172, "elapsed_s": 1.5047177029991872, "decode_tps_approx": 39.6877362766987, "metrics_delta": {"vllm:prefix_cache_queries_total": 51.0, "vllm:prefix_cache_hits_total": 0.0}}
|
||||
{"label": "graph", "time": 1789664774.614013, "event": "REQUEST_START", "name": "short-0", "profile": false, "prompt_tokens_preflight": 65}
|
||||
{"label": "graph", "time": 1789664781.980255, "event": "RESULT", "name": "short-0", "profile": false, "correct": true, "content": "\n\n数据库索引类似于书籍的目录,它通过B+树等数据结构将数据按特定字段有序组织,使查询时能快速定位目标记录,避免全表扫描,从而大幅提升读取效率。然而,索引对写入操作有负面影响:每次插入、更新或删除数据时,数据库都需要同步维护索引结构,这会增加额外的I/O和计算开销,导致写入性能下降。因此,索引是\"以空间换时间\"的典型策略,需根据实际读写比例合理设计。", "finish_reason": "stop", "usage": {"prompt_tokens": 65, "total_tokens": 279, "completion_tokens": 214, "completion_tokens_details": {"reasoning_tokens": 111}}, "max_tokens": 512, "prompt_sha256": "40e7d1df3e96ad760d661cfaff312f5d732cb2edd5eab555b4b5c43cea4512ca", "output_sha256": "126232af3bc1889efc9c0dad67f2d47fc60d3c4ff56e699a83f1576c3763c274", "ttft_s": 0.32573997599683935, "first_content_s": 3.8834300599992275, "elapsed_s": 7.360314821999054, "decode_tps_approx": 30.279974628892248, "metrics_delta": {"vllm:prefix_cache_queries_total": 65.0, "vllm:prefix_cache_hits_total": 0.0}}
|
||||
{"label": "graph", "time": 1789664781.9975877, "event": "REQUEST_START", "name": "prefill-8192-0", "profile": false, "prompt_tokens_preflight": 8173}
|
||||
{"label": "graph", "time": 1789664788.1946163, "event": "RESULT", "name": "prefill-8192-0", "profile": false, "correct": true, "content": "\n\n21573", "finish_reason": "stop", "usage": {"prompt_tokens": 8173, "total_tokens": 8247, "completion_tokens": 74, "completion_tokens_details": {"reasoning_tokens": 66}}, "max_tokens": 2048, "prompt_sha256": "f59f163e1c7bb8a79d08df830d94d6d2e98fec457e9de75e3c09f6fca7ab0813", "output_sha256": "c0d6109d7f5d6c6f5486e269f63ea0338562a90ab0ceeede5f228b020e701eea", "ttft_s": 3.921214998998039, "first_content_s": 6.03858627199952, "elapsed_s": 6.191493427999376, "decode_tps_approx": 33.338899955373805, "metrics_delta": {"vllm:prefix_cache_queries_total": 8173.0, "vllm:prefix_cache_hits_total": 0.0}}
|
||||
{"label": "graph", "time": 1789664788.2102215, "event": "REQUEST_START", "name": "reuse-8192-0", "profile": false, "prompt_tokens_preflight": 8173}
|
||||
{"label": "graph", "time": 1789664791.047786, "event": "RESULT", "name": "reuse-8192-0", "profile": false, "correct": true, "content": "\n\n21573", "finish_reason": "stop", "usage": {"prompt_tokens": 8173, "total_tokens": 8238, "completion_tokens": 65, "completion_tokens_details": {"reasoning_tokens": 57}}, "max_tokens": 2048, "prompt_sha256": "f59f163e1c7bb8a79d08df830d94d6d2e98fec457e9de75e3c09f6fca7ab0813", "output_sha256": "50530d055152bcaa11d1779a66b456a8ad5175a5343fb48f268c163f397d0089", "ttft_s": 0.8737782960015466, "first_content_s": 2.671151395996276, "elapsed_s": 2.8319971849996364, "decode_tps_approx": 34.055335354489465, "metrics_delta": {"vllm:prefix_cache_queries_total": 8173.0, "vllm:prefix_cache_hits_total": 6400.0}}
|
||||
{"label": "graph", "time": 1789664791.096218, "event": "REQUEST_START", "name": "prefill-32768-0", "profile": false, "prompt_tokens_preflight": 32745}
|
||||
{"label": "graph", "time": 1789664808.7936926, "event": "RESULT", "name": "prefill-32768-0", "profile": false, "correct": true, "content": "\n\n90812", "finish_reason": "stop", "usage": {"prompt_tokens": 32745, "total_tokens": 32815, "completion_tokens": 70, "completion_tokens_details": {"reasoning_tokens": 62}}, "max_tokens": 2048, "prompt_sha256": "605f7b492252f50eaf60e8d128dde3f2d2e1b50dd26f0b45db88dac00e36a5ca", "output_sha256": "7eec57291da7ee95ae6d9214f8ed940a405e64beb9f49a4693387362aac3e14c", "ttft_s": 15.511759164000978, "first_content_s": 17.52370634300314, "elapsed_s": 17.691939716001798, "decode_tps_approx": 32.87688689931954, "metrics_delta": {"vllm:prefix_cache_queries_total": 32745.0, "vllm:prefix_cache_hits_total": 0.0}}
|
||||
{"label": "graph", "time": 1789664808.8368387, "event": "REQUEST_START", "name": "reuse-32768-0", "profile": false, "prompt_tokens_preflight": 32745}
|
||||
{"label": "graph", "time": 1789664812.1356192, "event": "RESULT", "name": "reuse-32768-0", "profile": false, "correct": true, "content": "\n\n90812", "finish_reason": "stop", "usage": {"prompt_tokens": 32745, "total_tokens": 32814, "completion_tokens": 69, "completion_tokens_details": {"reasoning_tokens": 61}}, "max_tokens": 2048, "prompt_sha256": "605f7b492252f50eaf60e8d128dde3f2d2e1b50dd26f0b45db88dac00e36a5ca", "output_sha256": "8b503330d6b70b8845bab80540601fc48c8f840d54d07f38d80250c8167de89f", "ttft_s": 1.325252010996337, "first_content_s": 3.1243166819986072, "elapsed_s": 3.293290553992847, "decode_tps_approx": 36.2184980346298, "metrics_delta": {"vllm:prefix_cache_queries_total": 32745.0, "vllm:prefix_cache_hits_total": 30400.0}}
|
||||
{"label": "graph", "time": 1789664812.1429493, "event": "REQUEST_START", "name": "short-1", "profile": false, "prompt_tokens_preflight": 65}
|
||||
{"label": "graph", "time": 1789664819.2091334, "event": "RESULT", "name": "short-1", "profile": false, "correct": true, "content": "\n\n数据库索引类似于书籍的目录,它通过B+树等数据结构将数据按特定字段有序组织,使查询时能快速定位目标记录,避免全表扫描,从而大幅提升读取效率。然而,索引对写入操作有负面影响:每次插入、更新或删除数据时,数据库都需要同步维护索引结构,这会增加额外的I/O和计算开销,导致写入性能下降。因此,索引是\"以空间换时间\"的典型策略,需根据实际读写比例合理设计。", "finish_reason": "stop", "usage": {"prompt_tokens": 65, "total_tokens": 279, "completion_tokens": 214, "completion_tokens_details": {"reasoning_tokens": 111}}, "max_tokens": 512, "prompt_sha256": "40e7d1df3e96ad760d661cfaff312f5d732cb2edd5eab555b4b5c43cea4512ca", "output_sha256": "126232af3bc1889efc9c0dad67f2d47fc60d3c4ff56e699a83f1576c3763c274", "ttft_s": 0.2600282569983392, "first_content_s": 3.7012813049950637, "elapsed_s": 7.060739258995454, "decode_tps_approx": 31.32107217692348, "metrics_delta": {"vllm:prefix_cache_queries_total": 65.0, "vllm:prefix_cache_hits_total": 0.0}}
|
||||
{"label": "graph", "time": 1789664819.2247505, "event": "REQUEST_START", "name": "prefill-8192-1", "profile": false, "prompt_tokens_preflight": 8173}
|
||||
{"label": "graph", "time": 1789664824.82708, "event": "RESULT", "name": "prefill-8192-1", "profile": false, "correct": true, "content": "\n\n21573", "finish_reason": "stop", "usage": {"prompt_tokens": 8173, "total_tokens": 8236, "completion_tokens": 63, "completion_tokens_details": {"reasoning_tokens": 55}}, "max_tokens": 2048, "prompt_sha256": "6c13deb2d4ddee54db95d7ebc8a1b87d24ed8a4112c33644c18b6b5d61ad7a16", "output_sha256": "09d17d110b06ec13d532382c41ea29955b7edfa8dd4bd34163a08fa746e56610", "ttft_s": 3.7701075279983343, "first_content_s": 5.436177673000202, "elapsed_s": 5.596617848997994, "decode_tps_approx": 35.516432222041566, "metrics_delta": {"vllm:prefix_cache_queries_total": 8173.0, "vllm:prefix_cache_hits_total": 0.0}}
|
||||
{"label": "graph", "time": 1789664824.8427408, "event": "REQUEST_START", "name": "reuse-8192-1", "profile": false, "prompt_tokens_preflight": 8173}
|
||||
{"label": "graph", "time": 1789664827.8556995, "event": "RESULT", "name": "reuse-8192-1", "profile": false, "correct": true, "content": "\n\n21573", "finish_reason": "stop", "usage": {"prompt_tokens": 8173, "total_tokens": 8247, "completion_tokens": 74, "completion_tokens_details": {"reasoning_tokens": 66}}, "max_tokens": 2048, "prompt_sha256": "6c13deb2d4ddee54db95d7ebc8a1b87d24ed8a4112c33644c18b6b5d61ad7a16", "output_sha256": "c0d6109d7f5d6c6f5486e269f63ea0338562a90ab0ceeede5f228b020e701eea", "ttft_s": 0.8690901810041396, "first_content_s": 2.8428866180038312, "elapsed_s": 3.0074559670028975, "decode_tps_approx": 35.510793077788634, "metrics_delta": {"vllm:prefix_cache_queries_total": 8173.0, "vllm:prefix_cache_hits_total": 6400.0}}
|
||||
{"label": "graph", "time": 1789664827.8988104, "event": "REQUEST_START", "name": "prefill-32768-1", "profile": false, "prompt_tokens_preflight": 32745}
|
||||
{"label": "graph", "time": 1789664845.1446679, "event": "RESULT", "name": "prefill-32768-1", "profile": false, "correct": true, "content": "\n\n90812", "finish_reason": "stop", "usage": {"prompt_tokens": 32745, "total_tokens": 32814, "completion_tokens": 69, "completion_tokens_details": {"reasoning_tokens": 61}}, "max_tokens": 2048, "prompt_sha256": "0a5e4ba33ff264d180b8a56d89184c0fefad59a96fdd348d655bd54fe5b5ed45", "output_sha256": "8b503330d6b70b8845bab80540601fc48c8f840d54d07f38d80250c8167de89f", "ttft_s": 15.357147006005107, "first_content_s": 17.08089536800253, "elapsed_s": 17.239325528003974, "decode_tps_approx": 37.728914829613814, "metrics_delta": {"vllm:prefix_cache_queries_total": 32745.0, "vllm:prefix_cache_hits_total": 0.0}}
|
||||
{"label": "graph", "time": 1789664845.1907535, "event": "REQUEST_START", "name": "reuse-32768-1", "profile": false, "prompt_tokens_preflight": 32745}
|
||||
{"label": "graph", "time": 1789664848.6512141, "event": "RESULT", "name": "reuse-32768-1", "profile": false, "correct": true, "content": "\n\n90812", "finish_reason": "stop", "usage": {"prompt_tokens": 32745, "total_tokens": 32815, "completion_tokens": 70, "completion_tokens_details": {"reasoning_tokens": 62}}, "max_tokens": 2048, "prompt_sha256": "0a5e4ba33ff264d180b8a56d89184c0fefad59a96fdd348d655bd54fe5b5ed45", "output_sha256": "7eec57291da7ee95ae6d9214f8ed940a405e64beb9f49a4693387362aac3e14c", "ttft_s": 1.3272808809997514, "first_content_s": 3.296925946000556, "elapsed_s": 3.4540864679947845, "decode_tps_approx": 33.71980234098872, "metrics_delta": {"vllm:prefix_cache_queries_total": 32745.0, "vllm:prefix_cache_hits_total": 30400.0}}
|
||||
{"label": "graph", "time": 1789664848.660781, "event": "REQUEST_START", "name": "short-2", "profile": false, "prompt_tokens_preflight": 65}
|
||||
{"label": "graph", "time": 1789664855.73135, "event": "RESULT", "name": "short-2", "profile": false, "correct": true, "content": "\n\n数据库索引类似于书籍的目录,它通过B+树等数据结构将数据按特定字段有序组织,使查询时能快速定位目标记录,避免全表扫描,从而大幅提升读取效率。然而,索引对写入操作有负面影响:每次插入、更新或删除数据时,数据库都需要同步维护索引结构,这会增加额外的I/O和计算开销,导致写入性能下降。因此,索引是\"以空间换时间\"的典型策略,需根据实际读写比例合理设计。", "finish_reason": "stop", "usage": {"prompt_tokens": 65, "total_tokens": 279, "completion_tokens": 214, "completion_tokens_details": {"reasoning_tokens": 111}}, "max_tokens": 512, "prompt_sha256": "40e7d1df3e96ad760d661cfaff312f5d732cb2edd5eab555b4b5c43cea4512ca", "output_sha256": "126232af3bc1889efc9c0dad67f2d47fc60d3c4ff56e699a83f1576c3763c274", "ttft_s": 0.2623967160034226, "first_content_s": 3.7147180649990332, "elapsed_s": 7.0642776130043785, "decode_tps_approx": 31.315723764990278, "metrics_delta": {"vllm:prefix_cache_queries_total": 65.0, "vllm:prefix_cache_hits_total": 0.0}}
|
||||
{"label": "graph", "time": 1789664855.7496016, "event": "REQUEST_START", "name": "prefill-8192-2", "profile": false, "prompt_tokens_preflight": 8173}
|
||||
{"label": "graph", "time": 1789664861.4250884, "event": "RESULT", "name": "prefill-8192-2", "profile": false, "correct": true, "content": "\n\n21573", "finish_reason": "stop", "usage": {"prompt_tokens": 8173, "total_tokens": 8242, "completion_tokens": 69, "completion_tokens_details": {"reasoning_tokens": 61}}, "max_tokens": 2048, "prompt_sha256": "b49f6dbbf80f49ad860a60b0a7c58dfc2122e38bb8493ae155790f22305d8f7a", "output_sha256": "f72eba21a095aa0a43356b25cb120346b064d7f0a253a63cee7579d2c3422b82", "ttft_s": 3.768740819999948, "first_content_s": 5.512005377000605, "elapsed_s": 5.670499198000471, "decode_tps_approx": 37.32072215199369, "metrics_delta": {"vllm:prefix_cache_queries_total": 8173.0, "vllm:prefix_cache_hits_total": 0.0}}
|
||||
{"label": "graph", "time": 1789664861.4389272, "event": "REQUEST_START", "name": "reuse-8192-2", "profile": false, "prompt_tokens_preflight": 8173}
|
||||
{"label": "graph", "time": 1789664864.1561685, "event": "RESULT", "name": "reuse-8192-2", "profile": false, "correct": true, "content": "\n\n21573", "finish_reason": "stop", "usage": {"prompt_tokens": 8173, "total_tokens": 8236, "completion_tokens": 63, "completion_tokens_details": {"reasoning_tokens": 55}}, "max_tokens": 2048, "prompt_sha256": "b49f6dbbf80f49ad860a60b0a7c58dfc2122e38bb8493ae155790f22305d8f7a", "output_sha256": "09d17d110b06ec13d532382c41ea29955b7edfa8dd4bd34163a08fa746e56610", "ttft_s": 0.8816422219970264, "first_content_s": 2.543782171997009, "elapsed_s": 2.711807350999152, "decode_tps_approx": 35.522341967132114, "metrics_delta": {"vllm:prefix_cache_queries_total": 8173.0, "vllm:prefix_cache_hits_total": 6400.0}}
|
||||
{"label": "graph", "time": 1789664864.196416, "event": "REQUEST_START", "name": "prefill-32768-2", "profile": false, "prompt_tokens_preflight": 32745}
|
||||
{"label": "graph", "time": 1789664881.472875, "event": "RESULT", "name": "prefill-32768-2", "profile": false, "correct": true, "content": "\n\n90812", "finish_reason": "stop", "usage": {"prompt_tokens": 32745, "total_tokens": 32814, "completion_tokens": 69, "completion_tokens_details": {"reasoning_tokens": 61}}, "max_tokens": 2048, "prompt_sha256": "d8796cea8608150048ce4f253e451ba96da5416ccfea4a6560dc2f6bb2412004", "output_sha256": "8b503330d6b70b8845bab80540601fc48c8f840d54d07f38d80250c8167de89f", "ttft_s": 15.375977513001999, "first_content_s": 17.108021717001975, "elapsed_s": 17.268514374001825, "decode_tps_approx": 37.54862656190945, "metrics_delta": {"vllm:prefix_cache_queries_total": 32745.0, "vllm:prefix_cache_hits_total": 0.0}}
|
||||
{"label": "graph", "time": 1789664881.5214229, "event": "REQUEST_START", "name": "reuse-32768-2", "profile": false, "prompt_tokens_preflight": 32745}
|
||||
{"label": "graph", "time": 1789664884.7351754, "event": "RESULT", "name": "reuse-32768-2", "profile": false, "correct": true, "content": "\n\n90812", "finish_reason": "stop", "usage": {"prompt_tokens": 32745, "total_tokens": 32814, "completion_tokens": 69, "completion_tokens_details": {"reasoning_tokens": 61}}, "max_tokens": 2048, "prompt_sha256": "d8796cea8608150048ce4f253e451ba96da5416ccfea4a6560dc2f6bb2412004", "output_sha256": "8b503330d6b70b8845bab80540601fc48c8f840d54d07f38d80250c8167de89f", "ttft_s": 1.3270910260034725, "first_content_s": 3.045836498000426, "elapsed_s": 3.2060768180017476, "decode_tps_approx": 37.82214005937634, "metrics_delta": {"vllm:prefix_cache_queries_total": 32745.0, "vllm:prefix_cache_hits_total": 30400.0}}
|
||||
{"label": "graph", "time": 1789664884.7446704, "event": "REQUEST_START", "name": "trace-short", "profile": true, "prompt_tokens_preflight": 65}
|
||||
{"label": "graph", "time": 1789664894.3422408, "event": "RESULT", "name": "trace-short", "profile": true, "correct": true, "content": "\n\n数据库索引类似于书籍的目录,它通过B+树等数据结构将数据按特定字段有序组织,使查询时能快速定位目标记录,避免全表扫描,从而大幅提升读取效率。然而,索引对写入操作有负面影响:每次插入、更新或删除数据时,数据库都需要同步维护索引结构,这会增加额外的I/O和计算开销,导致写入性能下降。因此,索引是\"以空间换时间\"的典型策略,需根据实际读写比例合理设计。", "finish_reason": "stop", "usage": {"prompt_tokens": 65, "total_tokens": 279, "completion_tokens": 214, "completion_tokens_details": {"reasoning_tokens": 111}}, "max_tokens": 512, "prompt_sha256": "40e7d1df3e96ad760d661cfaff312f5d732cb2edd5eab555b4b5c43cea4512ca", "output_sha256": "126232af3bc1889efc9c0dad67f2d47fc60d3c4ff56e699a83f1576c3763c274", "ttft_s": 0.23192091599776177, "first_content_s": 3.732794447001652, "elapsed_s": 9.19204898099997, "decode_tps_approx": 30.944518883030874, "metrics_delta": {"vllm:prefix_cache_queries_total": 65.0, "vllm:prefix_cache_hits_total": 0.0}}
|
||||
{"label": "graph", "time": 1789664894.3764791, "event": "REQUEST_START", "name": "trace-prefill-32768", "profile": true, "prompt_tokens_preflight": 32742}
|
||||
{"label": "graph", "time": 1789664914.0692568, "event": "RESULT", "name": "trace-prefill-32768", "profile": true, "correct": true, "content": "\n\n90812", "finish_reason": "stop", "usage": {"prompt_tokens": 32742, "total_tokens": 32812, "completion_tokens": 70, "completion_tokens_details": {"reasoning_tokens": 62}}, "max_tokens": 2048, "prompt_sha256": "464e9da0976ca56ddece006a237605bca8b60d3dd0aab538abbc0aaa1696815c", "output_sha256": "7eec57291da7ee95ae6d9214f8ed940a405e64beb9f49a4693387362aac3e14c", "ttft_s": 15.522943222997128, "first_content_s": 17.421024231000047, "elapsed_s": 19.56250263000402, "decode_tps_approx": 34.818013845987785, "metrics_delta": {"vllm:prefix_cache_queries_total": 32742.0, "vllm:prefix_cache_hits_total": 0.0}}
|
||||
{"label": "graph", "time": 1789664914.1044817, "event": "REQUEST_START", "name": "trace-reuse-32768", "profile": true, "prompt_tokens_preflight": 32742}
|
||||
{"label": "graph", "time": 1789664920.2165241, "event": "RESULT", "name": "trace-reuse-32768", "profile": true, "correct": true, "content": "\n\n90812", "finish_reason": "stop", "usage": {"prompt_tokens": 32742, "total_tokens": 32816, "completion_tokens": 74, "completion_tokens_details": {"reasoning_tokens": 66}}, "max_tokens": 2048, "prompt_sha256": "464e9da0976ca56ddece006a237605bca8b60d3dd0aab538abbc0aaa1696815c", "output_sha256": "997501e1996b1c37af707622b1c4caf8dba93113e868f28210bd47ba700b5804", "ttft_s": 1.3475809320007102, "first_content_s": 3.358952419999696, "elapsed_s": 5.983986291001202, "decode_tps_approx": 34.79999075012843, "metrics_delta": {"vllm:prefix_cache_queries_total": 32742.0, "vllm:prefix_cache_hits_total": 30400.0}}
|
||||
{"label": "graph", "time": 1789664920.2165873, "event": "PROFILE_SUITE_PASS"}
|
||||
File diff suppressed because it is too large
Load Diff
@@ -0,0 +1,852 @@
|
||||
{
|
||||
"file": "trace.2.sqlite",
|
||||
"gpu_activity_span_s": 17.613957408,
|
||||
"gpu_busy_union_s": 17.173015684,
|
||||
"gpu_idle_within_span_s": 0.44094172400000176,
|
||||
"kernel_count": 111055,
|
||||
"graph_node_kernel_count": 49140,
|
||||
"api": {
|
||||
"cudaStreamSynchronize_v3020": [
|
||||
96,
|
||||
7.291123311999998
|
||||
],
|
||||
"cuLaunchKernelEx": [
|
||||
22540,
|
||||
2.9715556959999576
|
||||
],
|
||||
"cudaLaunchKernelExC_v11060": [
|
||||
11197,
|
||||
2.7721689919999997
|
||||
],
|
||||
"cudaLaunchKernel_v7000": [
|
||||
24666,
|
||||
2.074785632000026
|
||||
],
|
||||
"cuLaunchKernel": [
|
||||
3512,
|
||||
0.1402584800000001
|
||||
],
|
||||
"cudaGraphLaunch_v10000": [
|
||||
1431,
|
||||
0.13174172799999972
|
||||
],
|
||||
"cudaEventSynchronize_v3020": [
|
||||
48,
|
||||
0.042627568
|
||||
],
|
||||
"cudaMemcpyAsync_v3020": [
|
||||
2398,
|
||||
0.02786942400000001
|
||||
],
|
||||
"cudaMemsetAsync_v3020": [
|
||||
768,
|
||||
0.005161456000000002
|
||||
],
|
||||
"cuKernelGetName": [
|
||||
25422,
|
||||
0.002872591999999989
|
||||
],
|
||||
"cudaEventQuery_v3020": [
|
||||
1557,
|
||||
0.0023473760000000004
|
||||
],
|
||||
"cuTensorMapEncodeTiled": [
|
||||
12786,
|
||||
0.002259360000000043
|
||||
],
|
||||
"cudaEventRecordWithFlags_v11010": [
|
||||
1003,
|
||||
0.001927552000000001
|
||||
],
|
||||
"cudaGetDriverEntryPointByVersion_v12050": [
|
||||
12096,
|
||||
0.0017841920000000863
|
||||
],
|
||||
"cuKernelGetFunction": [
|
||||
3512,
|
||||
0.0013231039999999851
|
||||
],
|
||||
"cudaStreamIsCapturing_v10000": [
|
||||
2144,
|
||||
0.0006217759999999976
|
||||
],
|
||||
"cudaStreamWaitEvent_v3020": [
|
||||
144,
|
||||
0.00041838400000000016
|
||||
],
|
||||
"cudaEventCreateWithFlags_v3020": [
|
||||
96,
|
||||
0.00024513599999999997
|
||||
],
|
||||
"cudaEventDestroy_v3020": [
|
||||
96,
|
||||
0.00022665600000000007
|
||||
]
|
||||
},
|
||||
"nvtx": {
|
||||
"PLE_lookup_tokens=1600": {
|
||||
"count": 20,
|
||||
"sum_s": 5.607371088,
|
||||
"union_s": 5.607371088,
|
||||
"gpu_overlap_s": 5.336825187
|
||||
},
|
||||
"PLE_wait_for_GPU": {
|
||||
"count": 48,
|
||||
"sum_s": 7.291963728,
|
||||
"union_s": 7.291963728,
|
||||
"gpu_overlap_s": 7.233423306
|
||||
},
|
||||
"PLE_CPU_dedup": {
|
||||
"count": 48,
|
||||
"sum_s": 0.020901856,
|
||||
"union_s": 0.020901856,
|
||||
"gpu_overlap_s": 0.0
|
||||
},
|
||||
"PLE_CPU_gather": {
|
||||
"count": 48,
|
||||
"sum_s": 0.292111936,
|
||||
"union_s": 0.292111936,
|
||||
"gpu_overlap_s": 0.0
|
||||
},
|
||||
"PLE_lookup_tokens=742": {
|
||||
"count": 1,
|
||||
"sum_s": 0.29688864,
|
||||
"union_s": 0.29688864,
|
||||
"gpu_overlap_s": 0.282970556
|
||||
},
|
||||
"PLE_lookup_tokens=4": {
|
||||
"count": 27,
|
||||
"sum_s": 1.71834584,
|
||||
"union_s": 1.71834584,
|
||||
"gpu_overlap_s": 1.61634065
|
||||
}
|
||||
},
|
||||
"dispatch": {
|
||||
"CG_DISPATCH:actual=1600:mode=NONE:padded=1600": 40,
|
||||
"CG_DISPATCH:actual=1:mode=NONE:padded=1": 48,
|
||||
"CG_DISPATCH:actual=742:mode=NONE:padded=742": 2,
|
||||
"CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4": 27,
|
||||
"CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4": 27
|
||||
},
|
||||
"dispatch_timeline": [
|
||||
{
|
||||
"start_ns": 46743952,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 416451360,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 443082544,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 469310784,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 1124006496,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 1154683696,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 1179774000,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 1845801456,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 1877195728,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 1902001888,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 2578899920,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 2609916288,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 2635831856,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 3318901952,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 3349740016,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 3376088736,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 4054732640,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 4085130432,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 4110746784,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 4797770512,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 4829017936,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 4854209488,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 5551719760,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 5581307632,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 5607134000,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 6295334464,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 6325098224,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 6349994272,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 7042285248,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 7072015024,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 7097530960,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 7785043856,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 7816741056,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 7842952592,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 8537987536,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 8569536560,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 8595452192,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 9285771248,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 9317233072,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 9342219952,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 10060629616,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 10091559456,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 10116812048,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 10807242976,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 10837982160,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 10864521920,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 11562733424,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 11593737408,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 11619654832,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 12323127600,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 12354216560,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 12380982848,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 13081498400,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 13113506512,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 13138598000,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 13833930640,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 13865480864,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 13890978272,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 14599979088,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 14632423408,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 14657285504,
|
||||
"name": "CG_DISPATCH:actual=742:mode=NONE:padded=742"
|
||||
},
|
||||
{
|
||||
"start_ns": 15244687408,
|
||||
"name": "CG_DISPATCH:actual=742:mode=NONE:padded=742"
|
||||
},
|
||||
{
|
||||
"start_ns": 15261168608,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 15272860704,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 15510425328,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 15511343584,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 15516520016,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 15591171712,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 15592218224,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 15597515056,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 15670668032,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 15671484256,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 15675816752,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 15748911968,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 15749703280,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 15753801520,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 15829628848,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 15830463168,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 15834544656,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 15911548208,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 15912496096,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 15917334144,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 15989715552,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 15990467984,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 15994717936,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 16067760256,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 16068616608,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 16072918640,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 16151936192,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 16153013376,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 16158295088,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 16231124064,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 16232113808,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 16237373136,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 16308737952,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 16309646016,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 16314212128,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 16388651696,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 16389407904,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 16393690448,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 16470779408,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 16471838880,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 16476836288,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 16549986896,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 16550865424,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 16555256672,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 16632067984,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 16632894512,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 16637253408,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 16709708944,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 16710541136,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 16714847184,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 16789609392,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 16790634928,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 16795643776,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 16869435120,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 16870399200,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 16875445984,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 16952330656,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 16953311568,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 16958431376,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 17031362304,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 17032114688,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 17036354400,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 17111019568,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 17112056224,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 17117105808,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 17188540768,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 17189324992,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 17193788208,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 17269626256,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 17270693728,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 17276082128,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 17347956512,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 17348732480,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 17352868096,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 17427032704,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 17427739648,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 17431571536,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 17508960384,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 17509972320,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 17514870544,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 17592660304,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 17593603200,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
}
|
||||
],
|
||||
"top_kernels": [
|
||||
[
|
||||
"device_kernel",
|
||||
[
|
||||
4608,
|
||||
5.806803458000012
|
||||
]
|
||||
],
|
||||
[
|
||||
"Kernel2",
|
||||
[
|
||||
16931,
|
||||
1.7378030570000023
|
||||
]
|
||||
],
|
||||
[
|
||||
"_qsa_sparse_paged_gqa_splitk_kernel",
|
||||
[
|
||||
672,
|
||||
1.4925525440000014
|
||||
]
|
||||
],
|
||||
[
|
||||
"nvjet_sm121_tst_mma_128x208x64_2_32x104x64_tmaAB_bz_TNNN",
|
||||
[
|
||||
1000,
|
||||
1.3766302719999985
|
||||
]
|
||||
],
|
||||
[
|
||||
"_hc_combine_norm_kernel",
|
||||
[
|
||||
4848,
|
||||
0.9208744449999874
|
||||
]
|
||||
],
|
||||
[
|
||||
"kernel",
|
||||
[
|
||||
2853,
|
||||
0.8909366829999992
|
||||
]
|
||||
],
|
||||
[
|
||||
"_hc_gate_mix_kernel",
|
||||
[
|
||||
4944,
|
||||
0.6871558350000507
|
||||
]
|
||||
],
|
||||
[
|
||||
"nvjet_sm121_tst_mma_128x240x64_2_32x120x64_tmaAB_bz_TNNN",
|
||||
[
|
||||
980,
|
||||
0.5473217280000005
|
||||
]
|
||||
],
|
||||
[
|
||||
"nvjet_sm121_tst_mma_128x128x64_3_32x64x64_tmaAB_bz_TNNN",
|
||||
[
|
||||
2000,
|
||||
0.4623586880000001
|
||||
]
|
||||
],
|
||||
[
|
||||
"finalizeMoeRoutingKernel",
|
||||
[
|
||||
2304,
|
||||
0.4111630399999976
|
||||
]
|
||||
],
|
||||
[
|
||||
"_causal_conv1d_fwd_kernel",
|
||||
[
|
||||
756,
|
||||
0.24460185600000015
|
||||
]
|
||||
],
|
||||
[
|
||||
"doActivationKernel",
|
||||
[
|
||||
2304,
|
||||
0.23128571500000175
|
||||
]
|
||||
],
|
||||
[
|
||||
"_fused_post_conv_kernel",
|
||||
[
|
||||
756,
|
||||
0.23053734400000014
|
||||
]
|
||||
],
|
||||
[
|
||||
"elementwise_kernel",
|
||||
[
|
||||
6192,
|
||||
0.22792899099999814
|
||||
]
|
||||
],
|
||||
[
|
||||
"kernel_cutlass_kernel_flashinfergdn_kernelsdelta_rule_dsldelta_rule_sm120_FullyFusedDeltaRuleSm120_object_at__tensorptrf32gmemalign16o1_tensorptrf32gmemalign16o1_CopyAtom_ThrID10_TVLayout_0",
|
||||
[
|
||||
756,
|
||||
0.19290607999999995
|
||||
]
|
||||
],
|
||||
[
|
||||
"vectorized_elementwise_kernel",
|
||||
[
|
||||
11370,
|
||||
0.18636042799999758
|
||||
]
|
||||
],
|
||||
[
|
||||
"layer_norm_fwd_kernel",
|
||||
[
|
||||
756,
|
||||
0.18163276799999997
|
||||
]
|
||||
],
|
||||
[
|
||||
"sm120_fp8_fp4_gemm_1d1d_impl",
|
||||
[
|
||||
192,
|
||||
0.17354873599999998
|
||||
]
|
||||
],
|
||||
[
|
||||
"expandInputRowsKernel",
|
||||
[
|
||||
2304,
|
||||
0.1723688820000004
|
||||
]
|
||||
],
|
||||
[
|
||||
"nvjet_sm121_tst_mma_128x192x64_2_32x96x64_tmaAB_bz_TNNN",
|
||||
[
|
||||
980,
|
||||
0.13571574399999997
|
||||
]
|
||||
]
|
||||
],
|
||||
"notes": "GPU busy is interval union, not utilization counter. CPU/API/NVTX times overlap GPU work and each other; do not sum them as exclusive costs."
|
||||
}
|
||||
@@ -0,0 +1,636 @@
|
||||
{
|
||||
"file": "trace.3.sqlite",
|
||||
"gpu_activity_span_s": 3.549492992,
|
||||
"gpu_busy_union_s": 3.382135698,
|
||||
"gpu_idle_within_span_s": 0.16735729399999988,
|
||||
"kernel_count": 66056,
|
||||
"graph_node_kernel_count": 50960,
|
||||
"api": {
|
||||
"cudaStreamSynchronize_v3020": [
|
||||
60,
|
||||
2.063916384
|
||||
],
|
||||
"cudaLaunchKernelExC_v11060": [
|
||||
1284,
|
||||
0.28009695999999934
|
||||
],
|
||||
"cuLaunchKernelEx": [
|
||||
5990,
|
||||
0.2070799999999992
|
||||
],
|
||||
"cudaLaunchKernel_v7000": [
|
||||
7285,
|
||||
0.1831771039999994
|
||||
],
|
||||
"cudaGraphLaunch_v10000": [
|
||||
1484,
|
||||
0.119649248
|
||||
],
|
||||
"cudaEventSynchronize_v3020": [
|
||||
30,
|
||||
0.040196352
|
||||
],
|
||||
"cuLaunchKernel": [
|
||||
537,
|
||||
0.01219283200000001
|
||||
],
|
||||
"cudaMemcpyAsync_v3020": [
|
||||
1043,
|
||||
0.010008064000000006
|
||||
],
|
||||
"cudaMemsetAsync_v3020": [
|
||||
480,
|
||||
0.0025478400000000017
|
||||
],
|
||||
"cudaEventQuery_v3020": [
|
||||
921,
|
||||
0.0011545439999999998
|
||||
],
|
||||
"cuKernelGetName": [
|
||||
7357,
|
||||
0.00088769599999998
|
||||
],
|
||||
"cudaEventRecordWithFlags_v11010": [
|
||||
605,
|
||||
0.0008087840000000007
|
||||
],
|
||||
"cudaStreamIsCapturing_v10000": [
|
||||
1274,
|
||||
0.00032420800000000107
|
||||
],
|
||||
"cuTensorMapEncodeTiled": [
|
||||
1472,
|
||||
0.0003082879999999999
|
||||
],
|
||||
"cuKernelGetFunction": [
|
||||
537,
|
||||
0.0002906560000000001
|
||||
],
|
||||
"cudaGetDriverEntryPointByVersion_v12050": [
|
||||
1152,
|
||||
0.00016971199999999957
|
||||
],
|
||||
"cudaStreamWaitEvent_v3020": [
|
||||
90,
|
||||
0.000155392
|
||||
],
|
||||
"cudaEventCreateWithFlags_v3020": [
|
||||
60,
|
||||
0.00014443199999999998
|
||||
],
|
||||
"cudaEventDestroy_v3020": [
|
||||
60,
|
||||
8.001600000000003e-05
|
||||
]
|
||||
},
|
||||
"nvtx": {
|
||||
"PLE_lookup_tokens=1600": {
|
||||
"count": 1,
|
||||
"sum_s": 0.019309872,
|
||||
"union_s": 0.019309872,
|
||||
"gpu_overlap_s": 0.010934608
|
||||
},
|
||||
"PLE_wait_for_GPU": {
|
||||
"count": 30,
|
||||
"sum_s": 2.064303616,
|
||||
"union_s": 2.064303616,
|
||||
"gpu_overlap_s": 2.035614378
|
||||
},
|
||||
"PLE_CPU_dedup": {
|
||||
"count": 30,
|
||||
"sum_s": 0.00226768,
|
||||
"union_s": 0.00226768,
|
||||
"gpu_overlap_s": 0.0
|
||||
},
|
||||
"PLE_CPU_gather": {
|
||||
"count": 30,
|
||||
"sum_s": 0.109426208,
|
||||
"union_s": 0.109426208,
|
||||
"gpu_overlap_s": 0.0
|
||||
},
|
||||
"PLE_lookup_tokens=742": {
|
||||
"count": 1,
|
||||
"sum_s": 0.293855216,
|
||||
"union_s": 0.293855216,
|
||||
"gpu_overlap_s": 0.27970459
|
||||
},
|
||||
"PLE_lookup_tokens=4": {
|
||||
"count": 28,
|
||||
"sum_s": 1.874469744,
|
||||
"union_s": 1.874469744,
|
||||
"gpu_overlap_s": 1.748828475
|
||||
}
|
||||
},
|
||||
"dispatch": {
|
||||
"CG_DISPATCH:actual=1600:mode=NONE:padded=1600": 2,
|
||||
"CG_DISPATCH:actual=1:mode=NONE:padded=1": 30,
|
||||
"CG_DISPATCH:actual=742:mode=NONE:padded=742": 2,
|
||||
"CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4": 28,
|
||||
"CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4": 28
|
||||
},
|
||||
"dispatch_timeline": [
|
||||
{
|
||||
"start_ns": 47667312,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 442471520,
|
||||
"name": "CG_DISPATCH:actual=1600:mode=NONE:padded=1600"
|
||||
},
|
||||
{
|
||||
"start_ns": 473807600,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 498595616,
|
||||
"name": "CG_DISPATCH:actual=742:mode=NONE:padded=742"
|
||||
},
|
||||
{
|
||||
"start_ns": 1085215856,
|
||||
"name": "CG_DISPATCH:actual=742:mode=NONE:padded=742"
|
||||
},
|
||||
{
|
||||
"start_ns": 1101603472,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 1113435424,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 1335575776,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 1336548688,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 1340993280,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 1412886784,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 1413839472,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 1418416416,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 1496932144,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 1497849376,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 1502437168,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 1574235248,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 1574893632,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 1578152608,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 1659975184,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 1660978304,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 1666040464,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 1739695136,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 1740582992,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 1745368080,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 1819362544,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 1820259120,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 1825002016,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 1898444080,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 1899320480,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 1903978080,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 1973992688,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 1974642976,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 1977460352,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 2058854000,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 2059781888,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 2064550128,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 2140294400,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 2141166416,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 2145782752,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 2215803760,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 2216437328,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 2219274592,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 2295308880,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 2295905872,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 2298529488,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 2373456000,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 2374001712,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 2376400016,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 2451495056,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 2452018624,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 2454446720,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 2534005728,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 2534498160,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 2536805632,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 2617391424,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 2618055024,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 2620944928,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 2698279104,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 2699135648,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 2704056256,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 2781484592,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 2782408128,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 2787209600,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 2872938352,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 2873856272,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 2878412464,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 2955704400,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 2956657904,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 2961494272,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 3046465232,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 3047422016,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 3052108288,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 3123689648,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 3124649664,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 3129308048,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 3199193712,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 3199819264,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 3202578560,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 3285661472,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 3286559472,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 3291397856,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 3363509456,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 3364085808,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 3366681840,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 3447715488,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 3448689632,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
},
|
||||
{
|
||||
"start_ns": 3453541344,
|
||||
"name": "CG_DISPATCH:actual=3:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 3529683168,
|
||||
"name": "CG_DISPATCH:actual=4:mode=PIECEWISE:padded=4"
|
||||
},
|
||||
{
|
||||
"start_ns": 3530662560,
|
||||
"name": "CG_DISPATCH:actual=1:mode=NONE:padded=1"
|
||||
}
|
||||
],
|
||||
"top_kernels": [
|
||||
[
|
||||
"Kernel2",
|
||||
[
|
||||
14453,
|
||||
1.2634071650000012
|
||||
]
|
||||
],
|
||||
[
|
||||
"device_kernel",
|
||||
[
|
||||
2880,
|
||||
0.9274175719999982
|
||||
]
|
||||
],
|
||||
[
|
||||
"kernel",
|
||||
[
|
||||
1772,
|
||||
0.4575362240000001
|
||||
]
|
||||
],
|
||||
[
|
||||
"_qsa_sparse_paged_gqa_splitk_kernel",
|
||||
[
|
||||
420,
|
||||
0.13202313600000007
|
||||
]
|
||||
],
|
||||
[
|
||||
"_hc_combine_norm_kernel",
|
||||
[
|
||||
3030,
|
||||
0.07136511900000105
|
||||
]
|
||||
],
|
||||
[
|
||||
"nvjet_sm121_tst_mma_128x208x64_2_32x104x64_tmaAB_bz_TNNN",
|
||||
[
|
||||
50,
|
||||
0.06789817599999999
|
||||
]
|
||||
],
|
||||
[
|
||||
"sm120_fp8_fp4_gemm_1d1d_impl",
|
||||
[
|
||||
120,
|
||||
0.05055315199999998
|
||||
]
|
||||
],
|
||||
[
|
||||
"_hc_gate_mix_kernel",
|
||||
[
|
||||
3090,
|
||||
0.049511773999999786
|
||||
]
|
||||
],
|
||||
[
|
||||
"finalizeMoeRoutingKernel",
|
||||
[
|
||||
1440,
|
||||
0.04024281600000026
|
||||
]
|
||||
],
|
||||
[
|
||||
"nvjet_sm121_tst_mma_192x160x64_2_48x80x64_tmaAB_bz_TNNN",
|
||||
[
|
||||
51,
|
||||
0.035206783999999984
|
||||
]
|
||||
],
|
||||
[
|
||||
"gdn_decode_post_conv_mtp_kernel",
|
||||
[
|
||||
1008,
|
||||
0.027670688000000006
|
||||
]
|
||||
],
|
||||
[
|
||||
"nvjet_sm121_tst_mma_128x240x64_2_32x120x64_tmaAB_bz_TNNN",
|
||||
[
|
||||
49,
|
||||
0.027358240000000006
|
||||
]
|
||||
],
|
||||
[
|
||||
"elementwise_kernel",
|
||||
[
|
||||
2268,
|
||||
0.02362060799999982
|
||||
]
|
||||
],
|
||||
[
|
||||
"nvjet_sm121_tst_mma_128x128x64_3_32x64x64_tmaAB_bz_TNNN",
|
||||
[
|
||||
100,
|
||||
0.023191840000000002
|
||||
]
|
||||
],
|
||||
[
|
||||
"vectorized_elementwise_kernel",
|
||||
[
|
||||
6095,
|
||||
0.02246828099999996
|
||||
]
|
||||
],
|
||||
[
|
||||
"doActivationKernel",
|
||||
[
|
||||
1440,
|
||||
0.02160813899999983
|
||||
]
|
||||
],
|
||||
[
|
||||
"expandInputRowsKernel",
|
||||
[
|
||||
1440,
|
||||
0.020183961000000014
|
||||
]
|
||||
],
|
||||
[
|
||||
"_causal_conv1d_fwd_kernel",
|
||||
[
|
||||
72,
|
||||
0.017587615999999997
|
||||
]
|
||||
],
|
||||
[
|
||||
"nvjet_sm121_tst_mma_128x224x64_2_64x56x64_tmaAB_bz_TNNN",
|
||||
[
|
||||
50,
|
||||
0.016199904
|
||||
]
|
||||
],
|
||||
[
|
||||
"_fused_post_conv_kernel",
|
||||
[
|
||||
72,
|
||||
0.014694784
|
||||
]
|
||||
]
|
||||
],
|
||||
"notes": "GPU busy is interval union, not utilization counter. CPU/API/NVTX times overlap GPU work and each other; do not sum them as exclusive costs."
|
||||
}
|
||||
@@ -0,0 +1,6 @@
|
||||
{"models": {"object": "list", "data": [{"id": "qwen3.8-flash-next", "object": "model", "created": 1789665662, "owned_by": "vllm", "root": "/root/.cache/huggingface/hub/models--nvidia--Qwen3.8-Flash-Next-NVFP4/snapshots/fc694b54fb0174e0913e6adf86691ef85a4ead47", "parent": null, "max_model_len": 262144, "permission": [{"id": "modelperm-b7d422ef1689806e", "object": "model_permission", "created": 1789665662, "allow_create_engine": false, "allow_sampling": true, "allow_logprobs": true, "allow_search_indices": false, "allow_view": true, "allow_fine_tuning": false, "organization": "*", "group": null, "is_blocking": false}]}]}}
|
||||
{"test": "arithmetic", "seconds": 2.43, "content": "\n\n323", "finish_reason": "stop", "usage": {"prompt_tokens": 53, "total_tokens": 109, "completion_tokens": 56, "prompt_tokens_details": null, "completion_tokens_details": {"reasoning_tokens": 50}}}
|
||||
{"test": "chinese", "seconds": 5.48, "content": "\n\n太阳光包含各种颜色的光(红、橙、黄、绿、蓝、紫等)。当阳光进入大气层时,会与空气中的氮、氧等小分子发生散射。\n\n短波长的光(蓝、紫)比长波长的光(红、橙)更容易被散射,这种现象称为**瑞利散射**。蓝光被散射到四面八方,充满整个天空,所以我们抬头看到的天空就是蓝色的。\n\n(紫光其实散射更强,但人眼对紫光不敏感,且部分紫光被高层大气吸收,所以天空呈现蓝色而非紫色。)", "finish_reason": "stop", "usage": {"prompt_tokens": 49, "total_tokens": 199, "completion_tokens": 150, "prompt_tokens_details": null, "completion_tokens_details": {"reasoning_tokens": 31}}}
|
||||
{"test": "tool_call", "seconds": 2.0665973159993882, "message": {"role": "assistant", "content": null, "refusal": null, "annotations": null, "audio": null, "function_call": null, "tool_calls": [{"id": "chatcmpl-tool-9e731f967ef9d3d6", "type": "function", "function": {"name": "get_weather", "arguments": "{\"city\": \"上海\"}"}}], "reasoning": "用户要求我调用 get_weather 工具查询上海的天气,并且明确说不要自行编造。我需要直接调用这个工具。\n"}}
|
||||
TOOL_TEST_PASS
|
||||
SMOKE_TEST_PASS
|
||||
@@ -0,0 +1,14 @@
|
||||
{
|
||||
"time": "2026-09-18T01:21:40.503543+08:00",
|
||||
"image": "local/qwen38-flash-spark:prefix-eager-0bfc7a15",
|
||||
"health": "healthy",
|
||||
"restart_policy": "unless-stopped",
|
||||
"compose_unchanged": true,
|
||||
"profile_mounts": 0,
|
||||
"profiler_arg_present": false,
|
||||
"host_oom_kills": 25,
|
||||
"available_GiB": 18.851730346679688,
|
||||
"system_failed_units": 0,
|
||||
"user_failed_units": 0,
|
||||
"smoke_pass": true
|
||||
}
|
||||
File diff suppressed because it is too large
Load Diff
@@ -0,0 +1,38 @@
|
||||
# Nsight Systems 诊断脚本
|
||||
|
||||
固定 nightly `0bfc7a15`、当前 PLE 适配,单 GPU、串行请求。
|
||||
方案见 [plan.md](plan.md)。这些脚本是实验工具,不是默认启动入口。
|
||||
|
||||
- `prepare.py prepare --data ...`:使用正在运行的同一模型生成 8K/32K 固定检索数据并预检 token 数。
|
||||
- `instrument.py`:读取同目录 `vllm_ple_mmap.original.py` 和 `cudagraph_utils.original.py`,
|
||||
精确匹配源码后生成临时 NVTX 版本;保持算子签名与计算不变。原始文件来自容器,不提交重复副本。
|
||||
- `client.py --label eager|graph`:在测试容器内运行,数据挂到 `/profiles/data`,密钥读挂载文件。
|
||||
三轮正常计时和三次独立 trace;所有请求相同,temperature=0、seed=42、reasoning_effort=low。
|
||||
- `monitor.py`:在 Spark 宿主机运行,写资源采样到 stdout;以脚本目录中的文件发出停止信号。
|
||||
必须在启动实验前启动监控,并配套退出后恢复原配置的 runner。保护是尽力而为,不能保证拦截瞬时 OOM。
|
||||
- `summarize_trace.py <trace.sqlite> ...`:导出 CUDA/NVTX 汇总,不读取进程命令行和环境。
|
||||
|
||||
Nsight 使用主机安装目录只读挂载到 `/opt/nsight`。临时服务添加:
|
||||
|
||||
```text
|
||||
--max-model-len 131072 --profiler-config '{"profiler":"cuda"}'
|
||||
```
|
||||
|
||||
在原有 `vllm serve ...` 前包裹:
|
||||
|
||||
```text
|
||||
/opt/nsight/bin/nsys profile --sample=none --trace=cuda,nvtx,osrt \
|
||||
--capture-range=cudaProfilerApi --capture-range-end=repeat:3 --kill=none \
|
||||
--trace-fork-before-exec=true --cuda-graph-trace=node --force-overwrite=true \
|
||||
-o /profiles/<label>/trace vllm serve ...
|
||||
```
|
||||
|
||||
临时将服务绑定到容器回环地址,客户端使用 `docker exec`;保留已有缓存、代理和 secret 挂载。
|
||||
PLE 和 CudaGraphManager 的标记文件分别只读覆盖原模块,不修改镜像或权重。
|
||||
测试时自动重启关闭;原 Compose 保留备份,退出时恢复并做健康及 smoke 检查。
|
||||
|
||||
计时阶段没有启动 capture,但 Nsight launcher 和 NVTX 包装仍存在;本轮不是完全卸载 profiler 的独立测速。
|
||||
node 级采样可能增加开销,因此 trace 时长只用作诊断,不能代替无采样阶段排名。
|
||||
CPU API、NVTX 和 GPU 时间存在重叠;GPU busy 是 kernel/copy/memset 区间并集,
|
||||
分母为首次至末次 GPU 活动的时间跨度,不是 SM 利用率、带宽利用率或完整 HTTP 延迟。
|
||||
原始 `.nsys-rep` / SQLite 可能包含进程参数,只保留在 Spark 的受限目录,不提交 Git。
|
||||
@@ -0,0 +1,67 @@
|
||||
import argparse,json,hashlib,time,urllib.request,signal
|
||||
from pathlib import Path
|
||||
parser=argparse.ArgumentParser();parser.add_argument('--label',required=True);args=parser.parse_args()
|
||||
base='http://127.0.0.1:8000'
|
||||
key=Path('/run/secrets/qwen_api_key').read_text().strip()
|
||||
headers={'Authorization':'Bearer '+key,'Content-Type':'application/json'}
|
||||
def emit(r):print(json.dumps({'label':args.label,'time':time.time(),**r},ensure_ascii=False),flush=True)
|
||||
def request(path,payload=None):
|
||||
req=urllib.request.Request(base+path,headers=headers,data=json.dumps(payload).encode() if payload is not None else None)
|
||||
with urllib.request.urlopen(req,timeout=60) as res:return res.read()
|
||||
def metrics():
|
||||
result={}
|
||||
for l in request('/metrics').decode().splitlines():
|
||||
if l.startswith('vllm:ple_mmap_') or l.startswith('vllm:prefix_cache_'):
|
||||
name=l.split('{')[0].split()[0]
|
||||
if name.endswith('_total'):result[name]=result.get(name,0)+float(l.split()[-1])
|
||||
return result
|
||||
def deadline(*_):raise TimeoutError('Request exceeded 180 seconds')
|
||||
signal.signal(signal.SIGALRM,deadline)
|
||||
def run(name,prompt,expected=None,profile=False):
|
||||
m=[{'role':'user','content':prompt}]
|
||||
count=json.loads(request('/tokenize',{'model':'qwen3.8-flash-next','messages':m,'chat_template_kwargs':{'reasoning_effort':'low'}}))['count']
|
||||
assert count+2048<=131072
|
||||
before=metrics();payload={'model':'qwen3.8-flash-next','messages':m,'temperature':0,'seed':42,'reasoning_effort':'low','max_tokens':2048 if expected else 512,'stream':True,'stream_options':{'include_usage':True}}
|
||||
req=urllib.request.Request(base+'/v1/chat/completions',headers=headers,data=json.dumps(payload,ensure_ascii=False).encode())
|
||||
emit({'event':'REQUEST_START','name':name,'profile':profile,'prompt_tokens_preflight':count})
|
||||
if profile:request('/start_profile',{})
|
||||
t=time.monotonic();first=last=visible=None;content=reasoning='';usage={};finish=None
|
||||
signal.alarm(180)
|
||||
try:
|
||||
with urllib.request.urlopen(req,timeout=180) as res:
|
||||
for raw in res:
|
||||
line=raw.decode().strip()
|
||||
if not line.startswith('data: ') or line=='data: [DONE]':continue
|
||||
x=json.loads(line[6:])
|
||||
if x.get('error'):raise RuntimeError('Streaming error')
|
||||
if x.get('usage'):usage=x['usage']
|
||||
for c in x.get('choices',[]):
|
||||
d=c.get('delta',{});a=d.get('content') or '';r=d.get('reasoning') or d.get('reasoning_content') or ''
|
||||
if a or r:
|
||||
now=time.monotonic();first=first if first is not None else now;last=now
|
||||
if a:visible=visible if visible is not None else now
|
||||
content+=a;reasoning+=r;finish=c.get('finish_reason') or finish
|
||||
finally:
|
||||
signal.alarm(0)
|
||||
if profile:request('/stop_profile',{})
|
||||
elapsed=time.monotonic()-t
|
||||
after=metrics()
|
||||
correct=bool(content.strip()) and (expected is None or content.strip()==expected) and finish=='stop'
|
||||
emit({'event':'RESULT','name':name,'profile':profile,'correct':correct,'content':content,'finish_reason':finish,'usage':usage,'max_tokens':payload['max_tokens'],'prompt_sha256':hashlib.sha256(prompt.encode()).hexdigest(),'output_sha256':hashlib.sha256((reasoning+'\0'+content).encode()).hexdigest(),'ttft_s':first-t if first else None,'first_content_s':visible-t if visible else None,'elapsed_s':elapsed,'decode_tps_approx':(usage.get('completion_tokens',0)-1)/(last-first) if last and first and last>first else None,'metrics_delta':{k:after[k]-before.get(k,0) for k in after}})
|
||||
assert correct and usage['prompt_tokens']==count,name
|
||||
return prompt
|
||||
short='用中文写一段约150字的说明,解释数据库索引为什么能加快查询,以及它对写入有什么影响。'
|
||||
run('warmup','计算17乘19,只输出结果。','323')
|
||||
# Same prefix fixtures across configurations; phase/round identity is before the document.
|
||||
fixtures={n:json.loads(Path(f'/profiles/data/ctx-{n}-1.json').read_text())['cases'][1] for n in [8192,32768]}
|
||||
for i in range(3):
|
||||
run(f'short-{i}',short)
|
||||
for n,c in fixtures.items():
|
||||
prompt=f'测试文档批次 BENCH-{i}-{n}。\n'+c['messages'][0]['content']
|
||||
run(f'prefill-{n}-{i}',prompt,c['expected'])
|
||||
run(f'reuse-{n}-{i}',prompt,c['expected'])
|
||||
run('trace-short',short,profile=True)
|
||||
c=fixtures[32768];prompt='测试文档批次 TRACE-32768。\n'+c['messages'][0]['content']
|
||||
run('trace-prefill-32768',prompt,c['expected'],profile=True)
|
||||
run('trace-reuse-32768',prompt,c['expected'],profile=True)
|
||||
emit({'event':'PROFILE_SUITE_PASS'})
|
||||
@@ -0,0 +1,29 @@
|
||||
from pathlib import Path
|
||||
import ast
|
||||
p=Path(__file__).resolve().parent
|
||||
s=(p/'vllm_ple_mmap.original.py').read_text()
|
||||
replacements=[
|
||||
(' torch.cuda.current_stream(ids.device).synchronize()', ' with torch.cuda.nvtx.range("PLE_wait_for_GPU"):\n torch.cuda.current_stream(ids.device).synchronize()'),
|
||||
(' rows = table.gather(uniq) # uint8 [U, row_bytes], fresh & writable',' with torch.cuda.nvtx.range("PLE_CPU_gather"):\n rows = table.gather(uniq) # uint8 [U, row_bytes], fresh & writable'),
|
||||
(' uniq, inverse = np.unique(ids_np, return_inverse=True)', ' with torch.cuda.nvtx.range("PLE_CPU_dedup"):\n uniq, inverse = np.unique(ids_np, return_inverse=True)')]
|
||||
for before,after in replacements:
|
||||
assert s.count(before)==1,(before,s.count(before))
|
||||
s=s.replace(before,after)
|
||||
s+='''\n# Temporary diagnostic annotations; no tensor or arithmetic changes.
|
||||
_profile_original_lookup_ids = _lookup_ids_impl
|
||||
def _lookup_ids_impl(ngram_ids: torch.Tensor, output: torch.Tensor, layer_name: str) -> None:
|
||||
with torch.cuda.nvtx.range("PLE_lookup_tokens=" + str(ngram_ids.shape[0])):
|
||||
return _profile_original_lookup_ids(ngram_ids, output, layer_name)
|
||||
'''
|
||||
ast.parse(s);(p/'vllm_ple_mmap.profile.py').write_text(s)
|
||||
s=(p/'cudagraph_utils.original.py').read_text()
|
||||
s+='''\n# Diagnostic dispatch marker; does not change graph selection.
|
||||
_profile_original_dispatch = CudaGraphManager.dispatch
|
||||
def _profile_dispatch(self, *args, **kwargs):
|
||||
desc = _profile_original_dispatch(self, *args, **kwargs)
|
||||
actual = args[1] if len(args) > 1 else kwargs.get("num_tokens", -1)
|
||||
torch.cuda.nvtx.mark(f"CG_DISPATCH:actual={actual}:mode={desc.cg_mode.name}:padded={desc.num_tokens}")
|
||||
return desc
|
||||
CudaGraphManager.dispatch = _profile_dispatch
|
||||
'''
|
||||
ast.parse(s);(p/'cudagraph_utils.profile.py').write_text(s)
|
||||
@@ -0,0 +1,41 @@
|
||||
import json,time,subprocess,os
|
||||
from pathlib import Path
|
||||
out=Path(__file__).resolve().parent
|
||||
low=[]
|
||||
critical=[]
|
||||
initial_oom=None
|
||||
while not (out/'monitor.stop').exists():
|
||||
mem={l.split(':')[0]:int(l.split()[1]) for l in Path('/proc/meminfo').read_text().splitlines()}
|
||||
vm=dict(l.split() for l in Path('/proc/vmstat').read_text().splitlines())
|
||||
state=subprocess.run(['docker','inspect','--format','{{json .State}}','qwen38-flash-vllm'],capture_output=True,text=True)
|
||||
status=json.loads(state.stdout) if state.returncode==0 else {}
|
||||
if initial_oom is None: initial_oom=int(vm.get('oom_kill',0))
|
||||
swapped=(int(vm['pswpin'])+int(vm['pswpout']))*os.sysconf('SC_PAGE_SIZE')
|
||||
rec={'time':time.time(),'available_kib':mem['MemAvailable'],'swap_used_kib':mem['SwapTotal']-mem['SwapFree'],
|
||||
'pswpin':int(vm['pswpin']),'pswpout':int(vm['pswpout']),'status':status.get('Status'),
|
||||
'oom_killed':status.get('OOMKilled'),'host_oom_kills':int(vm.get('oom_kill',0)),'memory_pressure':Path('/proc/pressure/memory').read_text().strip()}
|
||||
print(json.dumps(rec),flush=True)
|
||||
if mem['MemAvailable']<1024*1024:
|
||||
(out/'GUARD_STOP').write_text('MemAvailable below 1 GiB')
|
||||
subprocess.run(['docker','stop','-t','2','qwen38-flash-vllm'])
|
||||
break
|
||||
if mem['MemAvailable']<2*1024*1024:
|
||||
critical.append(time.monotonic())
|
||||
else:
|
||||
critical=[]
|
||||
if int(vm.get('oom_kill',0))>initial_oom or (len(critical)>=2 and critical[-1]-critical[0]>=5):
|
||||
(out/'GUARD_STOP').write_text('Host OOM counter increased or MemAvailable below 2 GiB for 5s')
|
||||
subprocess.run(['docker','stop','-t','5','qwen38-flash-vllm'])
|
||||
break
|
||||
if mem['MemAvailable']<4*1024*1024:
|
||||
low.append((time.monotonic(),swapped))
|
||||
else:
|
||||
low=[]
|
||||
if len(low)>=7 and low[-1][0]-low[0][0]>=30 and low[-1][1]-low[0][1]>512*1024**2:
|
||||
(out/'GUARD_STOP').write_text('MemAvailable below 4 GiB for 30s with >512 MiB swapping')
|
||||
subprocess.run(['docker','stop','-t','5','qwen38-flash-vllm'])
|
||||
break
|
||||
if status.get('OOMKilled'):
|
||||
(out/'GUARD_STOP').write_text('Docker reported OOMKilled')
|
||||
break
|
||||
time.sleep(1)
|
||||
@@ -0,0 +1,17 @@
|
||||
# CUDA Graph 性能剖析方案(2026-09-18)
|
||||
|
||||
目的:定位当前 CUDA Graph 约 4.4% 收益的限制,确认捕获范围与 PLE 图外路径的影响。
|
||||
先采证据,再决定是否扩大图覆盖;不将推断写成瓶颈结论。
|
||||
|
||||
1. 小型 CUDA 程序验证主机 Nsight Systems 在当前镜像中可采集 CUDA 时间线。
|
||||
2. 临时服务上下文上限 131072,串行输入不超过 128K;原镜像、模型、MTP=2、BF16 KV、内存比例 0.80 不变。
|
||||
3. 对照 eager、小尺寸图 [1,2,4,8,12]。若证据支持,再试含 2048-token 预填充尺寸的图配置。
|
||||
4. 固定合成短问答、8K/32K 首次/重复输入;记录延迟、输出/思考 tokens、PLE 分项计数、图重放和内存。
|
||||
5. Nsight 使用 API 控制的短时间窗口,只采选定请求;计时基准在采样关闭时运行,trace 不作性能定量排名。
|
||||
6. 用 CPU CUDA API、GPU 活动、同步、PLE 指标判断瓶颈;这些时间可能重叠,不简单相加当百分比。
|
||||
7. 主机 OOM 计数增加、MemAvailable 单次低于 1 GiB 或持续低于 2 GiB 即停止。
|
||||
保留并检查内核 OOM 日志;测试后恢复原服务和 smoke test。128K 仅在候选通过后做资源/正确性验证。
|
||||
8. 原始大型 trace 留在 Spark,仓库保存方案、配置、脚本、摘要及证据索引,不提交凭据。
|
||||
|
||||
基础事实:当前图模式为 breakable PIECEWISE,PLE 哈希+CPU mmap 查表在 eager break 中运行。
|
||||
大于已捕获 token 数的 batch 无匹配图时回退无图执行。扩大图尺寸仍需实测兼容性与额外内存。
|
||||
@@ -0,0 +1,198 @@
|
||||
"""Generate reproducible long-context retrieval fixtures and test a running vLLM.
|
||||
|
||||
Run inside the model container. Never prints the mounted API key.
|
||||
"""
|
||||
import argparse
|
||||
import hashlib
|
||||
import json
|
||||
import random
|
||||
import signal
|
||||
import time
|
||||
import urllib.request
|
||||
from pathlib import Path
|
||||
|
||||
SNAPSHOT = '/root/.cache/huggingface/hub/models--nvidia--Qwen3.8-Flash-Next-NVFP4/snapshots/fc694b54fb0174e0913e6adf86691ef85a4ead47'
|
||||
MODEL = 'qwen3.8-flash-next'
|
||||
MAX_OUTPUT = 2048
|
||||
MAX_CONTEXT = 131072
|
||||
BASE = 'http://127.0.0.1:8000'
|
||||
|
||||
|
||||
def emit(record):
|
||||
print(json.dumps(record, ensure_ascii=False), flush=True)
|
||||
|
||||
|
||||
def digest(value):
|
||||
return hashlib.sha256(value.encode()).hexdigest()
|
||||
|
||||
|
||||
def headers():
|
||||
key = Path('/run/secrets/qwen_api_key').read_text().strip()
|
||||
return {'Authorization': 'Bearer ' + key, 'Content-Type': 'application/json'}
|
||||
|
||||
|
||||
def request(path, payload=None, timeout=120):
|
||||
req = urllib.request.Request(BASE + path, headers=headers(),
|
||||
data=json.dumps(payload, ensure_ascii=False).encode() if payload is not None else None)
|
||||
with urllib.request.urlopen(req, timeout=timeout) as response:
|
||||
return json.load(response)
|
||||
|
||||
|
||||
def metrics():
|
||||
req = urllib.request.Request(BASE + '/metrics', headers=headers())
|
||||
with urllib.request.urlopen(req, timeout=15) as response:
|
||||
lines = response.read().decode().splitlines()
|
||||
names = ['vllm:prefix_cache_hits_total', 'vllm:prefix_cache_queries_total',
|
||||
'vllm:num_preemptions_total']
|
||||
return {name: sum(float(l.split()[-1]) for l in lines if l.split('{')[0].split(' ')[0] == name)
|
||||
for name in names}
|
||||
|
||||
|
||||
def messages(document, project):
|
||||
prompt = ('以下是一份虚构项目预算档案,请按项目编号检索其中的数据。\n' + document +
|
||||
'\n档案结束。\n问题:项目编号 ' + project + ' 的核定预算是多少元?只输出数字。')
|
||||
return [{'role': 'user', 'content': prompt}]
|
||||
|
||||
|
||||
def prepare(folder):
|
||||
from transformers import AutoTokenizer
|
||||
tok = AutoTokenizer.from_pretrained(SNAPSHOT, local_files_only=True)
|
||||
folder.mkdir(parents=True, exist_ok=True)
|
||||
for target in [8192, 32768]:
|
||||
for sample in [1, 2]:
|
||||
name = f'ctx-{target}-{sample}'
|
||||
rng = random.Random(target * 10 + sample)
|
||||
amounts = [str(rng.randrange(10000, 99999)) for _ in range(3)]
|
||||
projects = [f'TARGET-{position}-{sample}' for position in ['FRONT', 'MIDDLE', 'END']]
|
||||
filler = [f'项目编号 BG-{i:06d};核定预算 {rng.randrange(10000,99999)} 元;类别 C{i%19:02d}。' for i in range(20000)]
|
||||
|
||||
def build(n):
|
||||
rows = filler[:n].copy()
|
||||
positions = [int(n * f) for f in [0.05, 0.5, 0.95]]
|
||||
for pos, project, amount in zip(positions, projects, amounts):
|
||||
rows[pos] = f'项目编号 {project};核定预算 {amount} 元;类别 C20。'
|
||||
return '文档编号:PROFILE128K-' + name + '\n' + '\n'.join(rows), positions, rows
|
||||
|
||||
low, high = 20, len(filler)
|
||||
while low < high:
|
||||
n = (low + high + 1) // 2
|
||||
doc, _, _ = build(n)
|
||||
count = len(tok.apply_chat_template(messages(doc, projects[1]), tokenize=True, add_generation_prompt=True, return_dict=False))
|
||||
if count <= target:
|
||||
low = n
|
||||
else:
|
||||
high = n - 1
|
||||
doc, positions, rows = build(low)
|
||||
cases = []
|
||||
for label, index in [('front', 0), ('middle', 1), ('end', 2), ('front_repeat', 0)]:
|
||||
msgs = messages(doc, projects[index])
|
||||
count = request('/tokenize', {'model': MODEL, 'messages': msgs, 'add_generation_prompt': True, 'chat_template_kwargs': {'reasoning_effort': 'low'}})['count']
|
||||
assert target - 128 <= count <= target + 32, (name, count, target)
|
||||
assert count + MAX_OUTPUT <= MAX_CONTEXT, (name, count)
|
||||
cases.append({'case': label, 'messages': msgs, 'expected': amounts[index],
|
||||
'prompt_tokens_preflight': count, 'prompt_sha256': digest(msgs[0]['content']),
|
||||
'max_tokens': MAX_OUTPUT})
|
||||
fractions = [round(len(tok.encode('文档编号:PROFILE128K-'+name+'\n'+'\n'.join(rows[:pos]), add_special_tokens=False)) / cases[0]['prompt_tokens_preflight'], 4)
|
||||
for pos in positions]
|
||||
fixture = {'document': name, 'target_input_tokens': target, 'sample': sample,
|
||||
'document_sha256': digest(doc), 'needle_token_fractions': fractions, 'cases': cases}
|
||||
(folder / (name + '.json')).write_text(json.dumps(fixture, ensure_ascii=False))
|
||||
emit({'event': 'FIXTURE_READY', **{k:v for k,v in fixture.items() if k != 'cases'},
|
||||
'prompt_tokens': [c['prompt_tokens_preflight'] for c in cases]})
|
||||
|
||||
|
||||
def deadline(signum, frame):
|
||||
raise TimeoutError('Request exceeded the 1200 second wall-clock limit')
|
||||
|
||||
|
||||
def run_case(label, fixture, case, attempt):
|
||||
before = metrics()
|
||||
payload = {'model': MODEL, 'messages': case['messages'], 'temperature': 0, 'seed': 42,
|
||||
'reasoning_effort': 'low', 'max_tokens': MAX_OUTPUT,
|
||||
'stream': True, 'stream_options': {'include_usage': True}}
|
||||
req = urllib.request.Request(BASE + '/v1/chat/completions', headers=headers(),
|
||||
data=json.dumps(payload, ensure_ascii=False).encode())
|
||||
start = time.monotonic()
|
||||
first = first_content = last = None
|
||||
content = reasoning = ''
|
||||
usage = {}
|
||||
finish = None
|
||||
signal.alarm(1200)
|
||||
try:
|
||||
with urllib.request.urlopen(req, timeout=1200) as response:
|
||||
for raw in response:
|
||||
line = raw.decode().strip()
|
||||
if not line.startswith('data: ') or line == 'data: [DONE]':
|
||||
continue
|
||||
item = json.loads(line[6:])
|
||||
if item.get('error'):
|
||||
raise RuntimeError('Server returned a streaming error')
|
||||
if item.get('usage'):
|
||||
usage = item['usage']
|
||||
for choice in item.get('choices', []):
|
||||
delta = choice.get('delta', {})
|
||||
answer = delta.get('content') or ''
|
||||
thought = delta.get('reasoning') or delta.get('reasoning_content') or ''
|
||||
if answer or thought:
|
||||
now = time.monotonic()
|
||||
first = first if first is not None else now
|
||||
last = now
|
||||
if answer:
|
||||
first_content = first_content if first_content is not None else now
|
||||
content += answer
|
||||
reasoning += thought
|
||||
finish = choice.get('finish_reason') or finish
|
||||
finally:
|
||||
signal.alarm(0)
|
||||
if finish is None or not usage:
|
||||
raise RuntimeError('Incomplete response or missing usage')
|
||||
after = metrics()
|
||||
elapsed = time.monotonic() - start
|
||||
record = {'event': 'RESULT', 'label': label, 'document': fixture['document'],
|
||||
'target_input_tokens': fixture['target_input_tokens'], 'sample': fixture['sample'],
|
||||
'case': case['case'], 'attempt': attempt, 'expected': case['expected'],
|
||||
'content': content, 'correct': content.strip() == case['expected'] and finish == 'stop',
|
||||
'finish_reason': finish, 'usage': usage, 'max_tokens': MAX_OUTPUT,
|
||||
'prompt_tokens_preflight': case['prompt_tokens_preflight'], 'prompt_sha256': case['prompt_sha256'],
|
||||
'output_sha256': digest(reasoning + '\0' + content),
|
||||
'ttft_s': round(first-start, 4) if first else None,
|
||||
'first_content_s': round(first_content-start, 4) if first_content else None,
|
||||
'elapsed_s': round(elapsed, 4),
|
||||
'decode_tps_approx': round((usage['completion_tokens']-1)/(last-first), 3) if last and first and last>first else None,
|
||||
'metrics_delta': {k: after[k]-before[k] for k in after}}
|
||||
emit(record)
|
||||
assert usage['prompt_tokens'] + MAX_OUTPUT <= MAX_CONTEXT
|
||||
assert usage['prompt_tokens'] == case['prompt_tokens_preflight'], 'Server tokenization differs from preflight'
|
||||
return record
|
||||
|
||||
|
||||
def run(folder, label):
|
||||
signal.signal(signal.SIGALRM, deadline)
|
||||
failures = []
|
||||
for target in [8192, 32768]:
|
||||
emit({'event': 'STAGE_START', 'label': label, 'target_input_tokens': target})
|
||||
for sample in [1, 2]:
|
||||
fixture = json.loads((folder / f'ctx-{target}-{sample}.json').read_text())
|
||||
emit({'event': 'DOCUMENT_START', 'label': label, 'document': fixture['document'],
|
||||
'document_sha256': fixture['document_sha256'], 'needle_token_fractions': fixture['needle_token_fractions']})
|
||||
for case in fixture['cases']:
|
||||
emit({'event': 'REQUEST_START', 'label': label, 'document': fixture['document'], 'case': case['case']})
|
||||
result = run_case(label, fixture, case, 1)
|
||||
if not result['correct']:
|
||||
failures.append([fixture['document'], case['case']])
|
||||
run_case(label, fixture, case, 2)
|
||||
emit({'event': 'STAGE_COMPLETED', 'label': label, 'target_input_tokens': target})
|
||||
emit({'event': 'CONTEXT_SUITE_COMPLETED', 'label': label, 'first_attempt_failures': failures})
|
||||
# Quality failures are reported above, not hidden; infrastructure failures raise.
|
||||
|
||||
|
||||
if __name__ == '__main__':
|
||||
parser = argparse.ArgumentParser()
|
||||
parser.add_argument('mode', choices=['prepare', 'run'])
|
||||
parser.add_argument('--data', type=Path, default=Path('/tmp/context262k-data'))
|
||||
parser.add_argument('--label', default='current')
|
||||
args = parser.parse_args()
|
||||
if args.mode == 'prepare':
|
||||
prepare(args.data)
|
||||
else:
|
||||
run(args.data, args.label)
|
||||
@@ -0,0 +1,20 @@
|
||||
import json,statistics,sys
|
||||
from pathlib import Path
|
||||
result={}
|
||||
for fn in sys.argv[1:]:
|
||||
rows=[]
|
||||
for line in Path(fn).read_text().splitlines():
|
||||
try:r=json.loads(line)
|
||||
except ValueError:continue
|
||||
if r.get('event')=='RESULT':rows.append(r)
|
||||
label=rows[0]['label'];groups={}
|
||||
for name in ['short','prefill-8192','reuse-8192','prefill-32768','reuse-32768']:
|
||||
rr=[r for r in rows if not r['profile'] and r['name'].rsplit('-',1)[0]==name]
|
||||
if not rr:continue
|
||||
groups[name]={'n':len(rr),'correct':sum(r['correct'] for r in rr)}
|
||||
for k in ['ttft_s','elapsed_s','decode_tps_approx']:
|
||||
values=[r[k] for r in rr];groups[name][k]={'median':statistics.median(values),'values':values}
|
||||
groups[name]['output_tokens']=[r['usage']['completion_tokens'] for r in rr]
|
||||
groups[name]['reasoning_tokens']=[r['usage'].get('completion_tokens_details',{}).get('reasoning_tokens') for r in rr]
|
||||
result[label]={'total_requests':len(rows),'correct':sum(r['correct'] for r in rows),'groups':groups}
|
||||
print(json.dumps(result,indent=2))
|
||||
@@ -0,0 +1,52 @@
|
||||
"""Export only CUDA/NVTX aggregates; omit process arguments and environment."""
|
||||
import sqlite3,json,sys,collections
|
||||
from pathlib import Path
|
||||
|
||||
def merge(intervals):
|
||||
out=[]
|
||||
for a,b in sorted(intervals):
|
||||
if b<=a:continue
|
||||
if out and a<=out[-1][1]:out[-1]=(out[-1][0],max(out[-1][1],b))
|
||||
else:out.append((a,b))
|
||||
return out
|
||||
|
||||
def duration(xs):return sum(b-a for a,b in xs)/1e9
|
||||
|
||||
def overlap(a,b):
|
||||
i=j=0;total=0
|
||||
while i<len(a) and j<len(b):
|
||||
total+=max(0,min(a[i][1],b[j][1])-max(a[i][0],b[j][0]))
|
||||
if a[i][1]<b[j][1]:i+=1
|
||||
else:j+=1
|
||||
return total/1e9
|
||||
|
||||
for fn in sys.argv[1:]:
|
||||
db=sqlite3.connect(fn);db.row_factory=sqlite3.Row
|
||||
tables={r[0] for r in db.execute("select name from sqlite_master where type='table'")}
|
||||
strings=dict(db.execute('select id,value from StringIds'))
|
||||
gpu=[];kerns=[];api=collections.defaultdict(lambda:[0,0]);nv=collections.defaultdict(list);dispatch=collections.Counter();dispatch_timeline=[]
|
||||
for table in ['CUPTI_ACTIVITY_KIND_KERNEL','CUPTI_ACTIVITY_KIND_MEMCPY','CUPTI_ACTIVITY_KIND_MEMSET']:
|
||||
if table in tables:
|
||||
rows=list(db.execute('select * from '+table));gpu.extend((r['start'],r['end']) for r in rows)
|
||||
if table.endswith('KERNEL'):kerns=rows
|
||||
for table in ['CUPTI_ACTIVITY_KIND_RUNTIME','CUPTI_ACTIVITY_KIND_DRIVER']:
|
||||
if table in tables:
|
||||
for r in db.execute('select start,end,nameId from '+table):
|
||||
name=strings.get(r['nameId'],str(r['nameId']));api[name][0]+=1;api[name][1]+=(r['end']-r['start'])/1e9
|
||||
if 'NVTX_EVENTS' in tables:
|
||||
for r in db.execute('select * from NVTX_EVENTS'):
|
||||
keys=r.keys();name=r['text'] if 'text' in keys else None
|
||||
if not name and 'textId' in keys:name=strings.get(r['textId'],'')
|
||||
name=name or ''
|
||||
if name.startswith('CG_DISPATCH:'):
|
||||
dispatch[name]+=1;dispatch_timeline.append({'start_ns':r['start'],'name':name})
|
||||
if name.startswith('PLE_') and r['end'] is not None:nv[name].append((r['start'],r['end']))
|
||||
merged=merge(gpu);span=(merged[-1][1]-merged[0][0])/1e9 if merged else 0
|
||||
nvout={}
|
||||
for name,ranges in nv.items():
|
||||
u=merge(ranges);nvout[name]={'count':len(ranges),'sum_s':duration(ranges),'union_s':duration(u),'gpu_overlap_s':overlap(u,merged)}
|
||||
kernel_names=collections.defaultdict(lambda:[0,0])
|
||||
for r in kerns:
|
||||
name=strings.get(r['shortName'],'?');kernel_names[name][0]+=1;kernel_names[name][1]+=(r['end']-r['start'])/1e9
|
||||
result={'file':Path(fn).name,'gpu_activity_span_s':span,'gpu_busy_union_s':duration(merged),'gpu_idle_within_span_s':span-duration(merged),'kernel_count':len(kerns),'graph_node_kernel_count':sum(bool(r['graphNodeId']) for r in kerns),'api':dict(sorted(api.items(),key=lambda x:-x[1][1])),'nvtx':nvout,'dispatch':dict(dispatch),'dispatch_timeline':dispatch_timeline,'top_kernels':sorted(kernel_names.items(),key=lambda x:-x[1][1])[:20],'notes':'GPU busy is interval union, not utilization counter. CPU/API/NVTX times overlap GPU work and each other; do not sum them as exclusive costs.'}
|
||||
dest=Path(fn).with_suffix('.summary.json');dest.write_text(json.dumps(result,indent=2));print(dest)
|
||||
Reference in New Issue
Block a user