Compare commits
3 Commits
8189942353
...
9e56401384
| Author | SHA1 | Date | |
|---|---|---|---|
|
9e56401384 | ||
|
|
5ed30006a5 | ||
|
|
1e8c36b7d1 |
@ -0,0 +1,40 @@
|
|||||||
|
# GLM-5.3-NVFP4 TP4PP2+MTP 深入攻坚（2026-09-07 第三轮）
|
||||||
|
|
||||||
|
机器：174.1.60.5（8×RTX 6000D）。基础补丁 = 09-06 的 r33（sglang_patch2 七文件挂载 + nightly-dev-cu13-20260901-07c8f729），本轮在其上做竞态二分、修复尝试与 decode 轮量化。主报告（结论与工程清单）：`results/analysis/GLM53_TP4PP2_Profile与优化方向_2026-09-06.md` §9。
|
||||||
|
|
||||||
|
## scripts/
|
||||||
|
|
||||||
|
| 文件 | 作用 |
|
||||||
|
|---|---|
|
||||||
|
| eagle_worker_v2_mask.py | r33 的位掩码版：7 处阶段边界同步改为 `/tmp/pp_sync_mask` 热开关（bit0-6 = site1-7：draftstep/prefill-ext/ext-pre/extlogits/draft-tree/accept/replay-in），缺省 127 ≡ r33 |
|
||||||
|
| scheduler_pp_mixin_r34.py | 修复候选：`_pp_commit_comm_work` 延迟一轮释放 P2PWork（实验结果：反而加速竞态触发，弃用） |
|
||||||
|
| deploy_ppmtp_mask.sh | 部署掩码构建（挂 mask 版 eagle_worker_v2） |
|
||||||
|
| deploy_ppmtp_r34.sh | 部署 r34（挂 r34 mixin + 健康后写掩码文件 SYNCMASK） |
|
||||||
|
| deploy_ppmtp_r34_nocache.sh | r34 + PYTORCH_NO_CUDA_MEMORY_CACHING=1（死于启动期 graph capture，实验不可行） |
|
||||||
|
| race_run.sh | 竞态单点运行器：热写掩码 → conc_test big → 监控容器/健康死亡并收割日志 |
|
||||||
|
| prof_mtp_decode{,2,3}.sh | torch profiler 采集编排（12 步 spec round，无 stack） |
|
||||||
|
| prof_mtp_analyze.py | trace 分析：轮时拆解、cudaStreamSynchronize/cudaLaunchKernel 宿主成本、NCCL 分桶、gap 归因、长宿主 op |
|
||||||
|
| prof_ann_census.py | user_annotation 相位普查（TARGET_VERIFY/draft/DECODE/边界收发时长 + gloo 操作画像） |
|
||||||
|
|
||||||
|
## results/ops/（实验原始记录）
|
||||||
|
|
||||||
|
- race_r0.log：mask=127 对照，conc_test 8×16384 PASS 64.4s
|
||||||
|
- race_r1.log / race_r1b_ladder.log / race_r1c_big.log：mask=0 conc_test 3/3 PASS（30.7/27.5/26.6s）——**假信心**，bench 才触发竞态
|
||||||
|
- crash_mask0_full.log：r33+mask0 完整 bench cc16 崩溃（prefill 爬坡段，PG2 watchdog IMA，表面 dsa init/allocator insert_events 同 09-06）
|
||||||
|
- crash_r34_mask0.log：r34+mask0 conc_test 23.6s 崩（commit wait 也是隐性遮罩的直接证据）
|
||||||
|
- deploy_r34.log / deploy_r34nc.log：部署记录（后者 = NO_CACHING 死于 cudaErrorStreamCaptureInvalidated）
|
||||||
|
- prof_mtp_decode.log / prof_mtp_decode3.log：采集编排日志
|
||||||
|
- prof_mtp_analysis.txt：8 rank trace 全量分析输出
|
||||||
|
|
||||||
|
## 核心事实（详见主报告 §9）
|
||||||
|
|
||||||
|
1. 三源证伪：7 同步 4.9ms/轮、eager launch 3.3ms/轮、边界 SendRecv 0.17ms/轮——合计 <10ms，解释不了 240ms 回归。
|
||||||
|
2. 真凶：每轮 3 次串行 GLOO 元数据 rendezvous（parallel_state.send_object = 2 gloo isend/字典），相位漂移下发送端 CPU 阻塞 300-540ms → TP 到达偏斜 → AR 自旋 132ms/轮 + GPU 空洞 ~130ms/轮（293.5ms/轮的 72% 是 TARGET_VERIFY 的等待态）。
|
||||||
|
3. 竞态为真且与停顿去除耦合：mask=0 在 bench 剖面崩；r34 去 commit wait 更快崩。根因定位只能靠 compute-sanitizer（NO_CACHING/LAUNCH_BLOCKING 均引入同步污染）。
|
||||||
|
4. 16k 输入下 KV 池把 decode 并发钉在 ≤8：cc16 压测的 decode 实际是 bs 7-8 混合态，"稳态 bs16"不存在。
|
||||||
|
|
||||||
|
## 远端留存（60.5）
|
||||||
|
|
||||||
|
- /root/sglang_patch2/（r33 七文件）、/root/eagle_worker_v2_mask.py、/root/scheduler_pp_mixin_r34.py、deploy_ppmtp_{mask,r34,r34_nocache}.sh、race_run.sh、prof_mtp_*.sh/py
|
||||||
|
- /data/hf_models/ppmtp_profiles/M_decode/（8 rank × 12 步 trace，~48MB）
|
||||||
|
- 崩溃与竞态日志：/root/crash_*.log、/root/race_*.log
|
||||||
@ -0,0 +1,267 @@
|
|||||||
|
# GLM-5.3-NVFP4 TP4PP2 Profile 与优化方向
|
||||||
|
|
||||||
|
日期：2026-09-06（profile）/ 2026-09-07（第二轮优化实验） | 机器：174.1.60.5（8×RTX 6000D 85.6GB，SM120，PCIe Gen5，无 NVLink） | 模型：GLM-5.3-NVFP4（78 层 + 1 nextn，hidden 6144，256 专家，first_k_dense 3，kv_lora_rank 512，DSA index_topk 2048）
|
||||||
|
|
||||||
|
## 0. 一句话结论
|
||||||
|
|
||||||
|
TP4PP2 在 prefill 主导的负载下**没有教科书式 PP 气泡**（chunked prefill 天然形成微批流水，两 stage 同忙 96%）；真正的墙是 **TP4 AllReduce 占 prefill 关键路径 30%**。09-07 第二轮实验把配置层面的快赢**全部排除**：NCCL 协议切换（env 到不了 rank + AR 已近 PCIe 带宽极限）、chunk 16384（MoE 工作区 OOM + 带宽受限下无次数红利）、现成通信开关（量化通信 NPU 专用、AR 融合 SM90/100 门控、symm-mem/NVLS 是 NVLink 向）——**现有配置（chunk 8192 / mem-frac 0.88 / 默认 NCCL）已是本栈的配置最优**；128k 场景的意外头号热点 DSA indexer（28%，1.6× 稀疏注意力本体）经核算属**算法固有成本**（无缓存复用机会、无可切换内核、模型已做跨层 topk 共享）。进一步提升全部是开发投入：优先级 = **PP+MTP 上游化（decode，TPOT -60%+ 潜力）> AR 分块重叠/量化 AR 内核开发（prefill，-15~25% 潜力）**。
|
||||||
|
|
||||||
|
## 1. 环境与部署忠实性
|
||||||
|
|
||||||
|
部署命令（与 09-05 对比报告获胜配置逐字一致）：
|
||||||
|
|
||||||
|
```
|
||||||
|
bash /root/deploy_par.sh "--tp 4 --pp-size 2 --disable-overlap-schedule --max-prefill-tokens 16384" nomtp
|
||||||
|
```
|
||||||
|
|
||||||
|
- 镜像 `lmsysorg/sglang:nightly-dev-20260828-daf63171`，容器 glm53-nvfp4:30000，权重 /data/hf_models/GLM-5.3-NVFP4
|
||||||
|
- KV fp8_e4m3 自动；SM120 三件套 `--disable-shared-experts-fusion --moe-runner-backend flashinfer_cutlass --disable-flashinfer-autotune`；hicache-ratio 3；mem-frac 0.88；chunked-prefill-size 8192
|
||||||
|
- PP 模式 nightly 硬约束：nomtp + `--disable-overlap-schedule`
|
||||||
|
- 系统库 NCCL 2.28.3（CUDA 13.0 build），容器未设任何 NCCL 环境变量（默认调参）
|
||||||
|
|
||||||
|
**忠实性验证**：干净无 profiler 复测 cc16 16k/512 —— 输出 104.7 tok/s、TPOT 50.9ms，与 09-05 基准（103.3 tok/s / 55.4ms）一致 → 本机部署复现了对比报告的性能，profile 结论可回推到基准数据。
|
||||||
|
|
||||||
|
| 指标 | 09-05 基准 | 本次干净复测 |
|
||||||
|
|---|---|---|
|
||||||
|
| 输入吞吐 | 3306 tok/s | 3350 tok/s |
|
||||||
|
| 输出吞吐 | 103.3 tok/s | 104.7 tok/s |
|
||||||
|
| TTFT | 29.2 s | ~31 s |
|
||||||
|
| TPOT | 55.4 ms | 50.9 ms |
|
||||||
|
|
||||||
|
## 2. 采集协议
|
||||||
|
|
||||||
|
torch profiler（sglang `POST /start_profile`，CPU+GPU activities，无 stack），8 rank 全采，精读 PP0-TP0 / PP1-TP0（TP1 做一致性对照）。跨 rank 墙钟对齐用 trace 内 baseTimeNanoseconds。采集前 flush cache + 微量预热 + 再 flush；bench 一律换新 seed（run-id 隔离）。
|
||||||
|
|
||||||
|
| 场景 | 负载 | 采集窗口 | 步数 |
|
||||||
|
|---|---|---|---|
|
||||||
|
| A prefill | 16×(16384→1) cc16 | 全部 prefill forward | ~34 前向 |
|
||||||
|
| B decode | 16×(16384→512) cc16 | TTFT 后新 "Decode batch" 起 | 60 步 |
|
||||||
|
| C mapping | with_stack 小采集 | — | **两次均失败**（见 §4.5） |
|
||||||
|
| D 128k/90% | 8×131072（90% 共享前缀）cc8 | 后缀 prefill 窗口 | 16 前向 |
|
||||||
|
|
||||||
|
trace 留存：`60.5:/data/hf_models/pp2tp4_profiles/{A_prefill,B_decode,D_128k}/`；分析输出：`.../analysis/`；脚本：`60.5:/root/prof_tools/`（本地副本 `D:\sskj\prof_tools\`）。
|
||||||
|
|
||||||
|
## 3. 核心发现
|
||||||
|
|
||||||
|
### 3.1 Prefill（A：16×16384→1，cc16，墙钟 35.2s）
|
||||||
|
|
||||||
|
**PP 气泡假说被证伪。** 跨 stage 1ms 分箱对齐：两 stage 同时忙碌 **96.0%**，仅 PP0 忙 1.1%，仅 PP1 忙 2.9%，双闲 0.0%。chunked prefill 把 16 个请求切成 ~34 个 8192-token 前向，天然形成微批流水，PP1 的计算完全藏在 PP0 之后。09-05 报告里 TP4PP2 相对 TP8 的收益**不是**以气泡为代价换来的。
|
||||||
|
|
||||||
|
**墙钟 = 最慢 stage 的关键链路：**
|
||||||
|
|
||||||
|
| PP0 链路（决定墙钟） | 耗时 | 占比 |
|
||||||
|
|---|---|---|
|
||||||
|
| GPU 计算 | 22.8 s | 65% |
|
||||||
|
| TP4 AllReduce（2686 次 × p50 3.69ms） | 10.4 s | 30% |
|
||||||
|
| 其余（launch 间隙等） | ~2 s | 5% |
|
||||||
|
|
||||||
|
PP1 链路（计算 21.2s + AR 10.0s + AllGather 0.24s）被流水完全隐藏。级间 SendRecv 真传输 p50 3.33ms×5 次/前向；p90 678ms 的空转档是稳态背压信号，被流水线吸收、不进墙钟。
|
||||||
|
|
||||||
|
**计算热点（两 stage 一致）：** sparse_mla_prefill 8.5s（~25%）；cutlass NVFP4 MoE GEMM ~5.5s（~17%）；deep_gemm fp8 密集 GEMM ~4s（PP0 ~12%）；MoE 路由胶水（expand/finalize/activation/topk）~2s（~6%）；DSA indexer（mqa_logits+topk）~3%；Memcpy DtoH 568ms（n=166）。
|
||||||
|
|
||||||
|
**结构性观察：每前向仅 8192 token。** max-prefill-tokens 虽为 16384，但 chunked-prefill-size 8192 封顶了每个前向的有效 token 数——AR 次数（79/前向）、流水线填充、launch 等固定开销只按 8192 token 摊薄。
|
||||||
|
|
||||||
|
### 3.2 Decode（B：16k KV，bs 7-16，步长 21-26ms，cuda graph 重放）
|
||||||
|
|
||||||
|
| PP0 每步构成 | 耗时 | 占比 |
|
||||||
|
|---|---|---|
|
||||||
|
| MoE NVFP4 GEMM（72 次小 batch，均值 106µs） | 7.65 ms | 35% |
|
||||||
|
| 级间 SendRecv（~5 次 × 0.8ms） | 4 ms | 18% |
|
||||||
|
| dense bf16 GEMM | ~2.3 ms | 10% |
|
||||||
|
| TP4 AllReduce（79 次 × 18µs） | 1.4 ms | 6% |
|
||||||
|
| sparse_mla_decode | 0.8 ms | 4% |
|
||||||
|
|
||||||
|
PP1 每步：MoE GEMM ~9.2ms（42%）、AR 1.9ms、**AllGather 1.4ms（273 次 launch 走 cuLaunchKernelExC，不在 cuda graph 内）**、SendRecv ~1ms。步间 CPU 间隙 p50 2.3-2.6ms（overlap-schedule 被禁的代价，~10%）。
|
||||||
|
|
||||||
|
- 无 MTP：每步 1 token/请求。干净 TPOT 50.9ms 与步长 21-26ms 的差值 = 交错进入 + 低并发期 + 尾部（最后一个请求独占解码，~28 tok/s）+ CPU 间隙。
|
||||||
|
- 采集期内仅 7-8 个请求同时处于 decode（TTFT 错峰进入），说明 cc16 下 decode 相位从未满batch。
|
||||||
|
- 带 profiler 时 TPOT 虚高到 97.8ms（2×）——**profiler 采集期的 bench 指标一律作废**，干净指标必须单独跑。
|
||||||
|
|
||||||
|
### 3.3 128k / 90% 前缀命中（D：cc8，后缀 prefill 窗口 19.9s）
|
||||||
|
|
||||||
|
日志核实命中率 0.8999。后缀有效 prefill 速率 **~5.2k tok/s**（同机 TP8 仅 ~3.2k——TP4PP2 的 KV 容量 2.37× 使全部 KV GPU 常驻，227k < 589k 容量，无需 TP8 必需的 hicache 往返）。
|
||||||
|
|
||||||
|
| 热点 | PP0 | PP1 |
|
||||||
|
|---|---|---|
|
||||||
|
| **DSA indexer `deep_gemm::sm120_fp8_mqa_logits`** | **28%** | **23%** |
|
||||||
|
| sparse_mla（稀疏注意力本体） | 17% | 17% |
|
||||||
|
| TP4 AllReduce | 20% | 20% |
|
||||||
|
| hicache/radix 恢复间隙（p50 77.5ms/前向） | ~5% | ~5% |
|
||||||
|
|
||||||
|
**意外发现：indexer 的 MQA logits GEMM 是它所服务的稀疏注意力的 1.6 倍。** DSA 选索引时对 ~118k 已缓存前缀键逐层算 q·k logits——前缀命中场景下这部分键对每个请求都相同，但当前实现对整个前缀重算。这是 128k 场景 TTFT 的第一大头，也是最值得攻的算法级优化点。
|
||||||
|
|
||||||
|
### 3.4 NCCL 取证（A）
|
||||||
|
|
||||||
|
- AllReduce 内核名 `ncclDevKernel_AllReduce_Sum_bf16_RING_LL`：~96MB 消息（8192×6144×2B）每次 3.69ms（极均匀，p10-p99 都在 3.68-4.18ms）。
|
||||||
|
- 每前向 79 次 AR ≈ 292ms ≈ **29% 的 1.03s 前向**；整个 A 窗口 2686 次 ≈ 10.4s。
|
||||||
|
- **带宽核算：96MB / 3.7ms ⇒ algbw 26GB/s，环形 busbw ≈ 39GB/s**，约为 PCIe Gen5 x16 实际大包峰值的 ~75%。真跑半带宽效率的 LL 线格式需要 ~78GB/s 原始速率——超过 PCIe Gen5 理论值，物理不可能。结论：内核名里的 "_LL" 是 NCCL 2.28 混合 plan 的命名，不代表线上是 LL 格式；**AR 实际已以接近最优的有效带宽运行，协议层没有免费午餐**（§5 的强制 Simple 实验证实了这一点）。
|
||||||
|
- SendRecv 双峰：p50 3.33ms 真传输 + p90 678ms 空转（流水背压，非墙钟成本）。
|
||||||
|
- 环境核查：容器与 launcher 进程均无 NCCL 调参变量（NCCL 2.28.3 全默认调参）。
|
||||||
|
|
||||||
|
### 3.5 工程发现（profiler 操作性结论，本 nightly）
|
||||||
|
|
||||||
|
1. **profiler flush 阻塞整个流水线**：60 步无栈 decode trace 的 flush 让服务停顿 11s（日志 14:57:28 stop → 14:57:39 done）。采集步数控制在 ≤60 步、bench 指标单独干净跑。
|
||||||
|
2. **with_stack 病态（两次复现）**：16k prefill 上开 with_stack → 单 rank RSS 145-173GB、容器 841GB/1008GB、health=000、容器僵尸（docker restart 失败），必须宿主机 kill -9 全部 sglang 进程 + docker rm -f 重部署；2k 小 probe 上开 with_stack → stop 已确认但 flush 5 分钟内不出文件、服务再次拖死。**本版本 with_stack 一律不用**，kernel→源码映射改用经验+符号名。
|
||||||
|
3. 采集自动停条件是 forward_ct ≥ start+num_steps+1（+1 守卫步），步数给多了 profiler 会一直武装着等——需要手动 stop_profile（该 POST 会阻塞到 flush 完，不能放在 rsh 链内联）。
|
||||||
|
|
||||||
|
## 4. 优化方向排序（09-07 第二轮实验后更新）
|
||||||
|
|
||||||
|
| # | 方向 | 状态（09-07 更新） | 预期收益 | 落地难度 |
|
||||||
|
|---|---|---|---|---|
|
||||||
|
| 1 | **PP+MTP（decode 最大杠杆）** | 唯一剩余大杠杆。步长 21-26ms 无投机；MTP accept ~2.5；楔死根因已定位（09-05 补丁），需上游化 | TPOT 等效 ~9-10ms/token（-60%+） | 高（开发项目） |
|
||||||
|
| 2 | **prefill AR 通信-计算重叠 / 量化 AR** | **无现成开关**（§6.3 扫描：quant-comm NPU 专用、AR 融合 SM90/100 门控、symm-mem/NVLS 是 NVLink 向），需上游/内核开发 | AR 逼近零暴露：prefill 墙钟 -15~25% | 中高（分块重叠的 TP linear 改造，或将 AR 融合内核扩展到 SM120+大 token 数） |
|
||||||
|
| 3 | MoE 胶水融合 + PP1 AllGather 入图 | 上游开发 | decode -2~4% | 中 |
|
||||||
|
| 4 | decode 低并发/尾部 | overlap-sched 被 PP 禁是 nightly 硬约束 | 长尾 P99 改善 | 中 |
|
||||||
|
|
||||||
|
**已证伪/关闭（记录在案，勿再投入）**：
|
||||||
|
- **NCCL 协议/算法调参**（§5）：容器 env 到不了 rank 进程，且 AR 已在带宽极限附近（busbw≈39GB/s ≈ PCIe Gen5 实际峰值 ~75%）；即便进程内强制生效上限也只有 prefill ~6%。
|
||||||
|
- **chunk 16384**（§6.2）：MoE 工作区随 chunk 增长 OOM（3.08GB > 2.11GB 空闲，需 mem-frac ≤0.85）；且 AR 带宽受限使"次数减半"无红利、indexer 功不变——剩余收益 ≤3%（噪声内）却要付 KV -5% + Max ITL 翻倍。TP8 口径"chunk 8192 勿调大"在 TP4PP2 同样成立。
|
||||||
|
- **128k DSA indexer**（§6.4）：28% 属算法固有成本（每请求后缀 q 独有、无跨请求复用；SM120 仅 deepgemm 一个内核后端；模型已做 index_topk_freq=4 跨层共享；kernel ~219 TFLOPS/rank 属 paged-gather 合理量级）。逃逸路径（分层索引）是影响精度的上游研究项。
|
||||||
|
|
||||||
|
优先级依据：#1 是唯一数量级杠杆且根因/补丁已就绪；#2 是 prefill 唯一剩余的结构性方向（协议红利已排除）；#3-#4 是小额累积项。
|
||||||
|
|
||||||
|
## 5. 验证实验：NCCL_PROTO=Simple（假设证伪）
|
||||||
|
|
||||||
|
Profile 提出的第一假设是"prefill AR 内核名带 `_RING_LL` → NCCL 在 96MB 大消息上误选 LL 协议（半带宽）→ 强制 Simple 可省 ~11% prefill 墙钟"。为此做了同配置 +1 环境变量的对照实验：
|
||||||
|
|
||||||
|
**设置**：`deploy_par.sh` 逐字相同，仅 docker run 追加 `-e NCCL_PROTO=Simple`（容器 config env 与 launcher 进程 environ 均核实含有该变量）。
|
||||||
|
|
||||||
|
**结果**：
|
||||||
|
|
||||||
|
| 指标 | 原样（默认调参） | 强制 Simple | 判读 |
|
||||||
|
|---|---|---|---|
|
||||||
|
| 输入吞吐（cc16 16k/512） | 3350 tok/s | 3248 tok/s | 无改善，单轮噪声内 |
|
||||||
|
| 输出吞吐 | 104.7 tok/s | 101.5 tok/s | 同上 |
|
||||||
|
| Mean TTFT | 31.36 s | 30.71 s | 持平 |
|
||||||
|
| Mean TPOT | 50.87 ms | 57.54 ms | 单轮波动（chunked prefill 干扰下 TPOT 本身有较大 run 间方差） |
|
||||||
|
| AR 内核（2 步 prefill trace 取证） | `..._RING_LL` p50 3.69ms ×79/前向 | `..._RING_LL` p50 3.70ms ×79/前向 | **协议未变** |
|
||||||
|
|
||||||
|
**根因：环境变量根本没到 rank 进程。** 逐一核查 `/proc/<pid>/environ`：launcher（`python3 -m sglang.launch_server`）与 multiprocessing resource_tracker 均含 `NCCL_PROTO=Simple`，但**全部 8 个 scheduler rank（PP0/PP1 × TP0-3，真正跑 NCCL 的进程）都没有**——sglang 的 worker spawn 重建了环境，丢弃继承来的 NCCL 调参变量（只保留自己注入的 `NCCL_CUMEM_ENABLE=0`、`NCCL_NVLS_ENABLE=0`）。因此内核选择与时长不变是必然，e2e 差值纯属噪声。
|
||||||
|
|
||||||
|
**即便变量送达也救不了这个方向（带宽核算）**：96MB/3.7ms ⇒ busbw≈39GB/s ≈ PCIe Gen5 x16 实际峰值的 ~75%；若真是半带宽 LL 线格式，需要 ~78GB/s 原始速率——超 PCIe Gen5 理论上限，物理不可能。故 "_LL" 命名不代表线上跑 LL 格式（NCCL 2.28 混合 plan 命名），AR 已在近最优有效带宽上运行。协议切换的理论上限 = 把 75% 提到 ~100% ⇒ AR 3.7→2.9ms ⇒ prefill 墙钟最多 -6%。
|
||||||
|
|
||||||
|
**结论**：协议/算法调参作为"一行环境变量的快赢"**证伪**。留在桌面上的通信优化只有两条真路：①通信-计算重叠（把 30% 的 AR 藏进计算）；②量化/压缩 AR（把载荷减半）。若仍想穷尽协议实验，正确做法是在 rank 进程内、NCCL 初始化前设置（例如通过 deploy 已挂载的 sglang_patch 在 layer_setup.py 中 `os.environ["NCCL_PROTO"]="Simple"`），预期收益 ≤6% prefill。
|
||||||
|
|
||||||
|
trace 留存：`60.5:/data/hf_models/pp2tp4_profiles/E_nccl_simple/`（8 rank × 2 前向）；bench：`/data/hf_models/bs_results/nccl_simple_cc16.json`；实验部署脚本（已用后删除，命令记录于本节）：`sed 's|--ipc=host|--ipc=host -e NCCL_PROTO=Simple|' /root/deploy_par.sh`。
|
||||||
|
|
||||||
|
## 6. 第二轮优化实验（2026-09-07）：配置级方向全部证伪
|
||||||
|
|
||||||
|
按 §4 排序推进"能落地的先落地"，结果是三个方向全部被证据关闭——这本身是本轮最有价值的产出：**TP4PP2 现有配置已是本栈配置最优**，后续提升必须走开发投入。详细证据归档：sskj-review `results/analysis/optimization_round2_findings.md`。
|
||||||
|
|
||||||
|
### 6.1 干净 128k/90% 基线（补缺口）
|
||||||
|
|
||||||
|
此前 D 场景的 bench 数字带 profiler 作废，TP4PP2 一直没有干净 128k 指标。原样配置补测（cc8，8×131072/512，hit 0.8999）：**TTFT p50 14.70s**（min 4.64 / max 20.02），TPOT 均值 51.7ms，输出 104.0 tok/s，e2e 均值 39.24s，retract 0。
|
||||||
|
|
||||||
|
### 6.2 chunk 16384 证伪（方向④关闭）
|
||||||
|
|
||||||
|
变体（`deploy_par.sh` 第 3 位置参数）启动成功、冒烟通过，但**首个 16k 请求让全部 8 个调度器 CUDA OOM 崩溃**：flashinfer cutlass MoE 工作区随 chunk 规模增长，16384-token 前向需 3.08GB，mem-frac 0.88 下每卡仅 2.11GB 空闲。跑通需 mem-frac ≤0.85（KV 池 -5%），但物理核算表明收益主体不存在——AR 纯带宽受限（每次调用固定开销≈0，"次数减半"省不出带宽部分）、indexer 的 q×k 总功与 chunk 切分无关；剩余收益（调度/launch 摊薄）≤3%，在单轮 bench 噪声（±3%）内，却要再付 Max ITL 翻倍（混合相位单次 prefill 阻塞 1.03s→2.06s）。附带发现：prefill 分块图（backend=breakable）只捕获 ≤2048 token 档位，8192/16384 的 extend 均走 eager_runner 逐层重放图块。
|
||||||
|
|
||||||
|
### 6.3 AR 现成开关扫描（方向①降级为纯上游开发）
|
||||||
|
|
||||||
|
| 开关 | 判定 |
|
||||||
|
|---|---|
|
||||||
|
| `--enable-quant-communications` | **NPU 专用**（server_args 校验非 NPU 直接 ValueError） |
|
||||||
|
| flashinfer allreduce fusion | 自动启用门控 `is_sm90 or is_sm100`——**SM120 被排除**；trtllm/mnnvl 后端为 NVLink/小 batch 设计 |
|
||||||
|
| `--enable-symm-mem` / `--enable-torch-symm-mem` | one-shot 小消息/decode 导向，96MB prefill AR 在 PCIe 不适用 |
|
||||||
|
| `--enable-nccl-nvls` | NVLink 专用 |
|
||||||
|
| `--pp-async-batch-depth` | PP 气泡仅 ~4%，非瓶颈 |
|
||||||
|
|
||||||
|
**本栈不存在 prefill AR 重叠/量化的现成路径。** 开发落点：分块 AR 与后续 GEMM 重叠（async-TP 式 TP linear 改造）；或将 flashinfer AR 融合（fp8 量化 AR+residual+norm）扩展到 SM120 与大 token 数。
|
||||||
|
|
||||||
|
### 6.4 indexer 关闭（方向③：算法固有成本，修正原设想）
|
||||||
|
|
||||||
|
- **修正 §3.3 的原始设想**："前缀 logits 增量缓存"不成立——每请求后缀 q 独有，其对共享前缀键的 logits 是该请求固有功；前缀 index-K 已由 fused_store_index_k_cache 缓存，无跨请求复用。
|
||||||
|
- 既有 skip 仅适用 kv_len ≤ index_topk=2048（全选即正确），128k 不适用；CUDA decode 侧 skip 还被 HIP-only 门控。
|
||||||
|
- SM120 上内核后端无选择（deepgemm 唯一，cutedsl=SM100 门控、aiter=ROCm）；模型已做 `index_topk_freq=4` 跨层 topk 共享。
|
||||||
|
- roofline：9.37ms/次 ≈ 8192 新 token × ~122k kv × 2048 FLOP/pair（每 rank，32 头×128 维×2 ÷ TP4）≈ 2.05 TFLOP/rank → **~219 TFLOPS/rank fp8**，paged-gather 内核的合理量级，无 5-10× 调优空间。
|
||||||
|
|
||||||
|
结论：28% 是 DSA 在 128k 语境扫全 kv 选 top-2048 的固有成本；逃逸路径（分层/粗到细索引）是影响精度的上游研究项。
|
||||||
|
|
||||||
|
## 7. 复现与留存
|
||||||
|
|
||||||
|
```
|
||||||
|
# 部署（获胜配置）
|
||||||
|
bash /root/deploy_par.sh "--tp 4 --pp-size 2 --disable-overlap-schedule --max-prefill-tokens 16384" nomtp
|
||||||
|
|
||||||
|
# 采集（60.5:/root/prof_tools/）
|
||||||
|
bash profile_cap.sh # A/B/C 编排：warm_and_flush → arm → bench → wait_traces
|
||||||
|
python3 bench_128k_hit90_prof.py --concurrency 8 --num-requests 8 --run-id <新> \
|
||||||
|
--arm-profile $BASE/D_128k --profile-steps 16
|
||||||
|
|
||||||
|
# 分析
|
||||||
|
python3 pp_align.py <trace_dir> # 跨 rank 对齐 + 每 rank 报告 + 1ms 分箱
|
||||||
|
python3 nccl_forensics.py <trace> # AR/SendRecv 分位 + 时间聚类
|
||||||
|
python3 census.py <trace> # 事件类别普查
|
||||||
|
# triage 三表：skill 脚本 docker cp 进容器运行（analyze_llm_torch_profile.py triage）
|
||||||
|
```
|
||||||
|
|
||||||
|
- trace：`60.5:/data/hf_models/pp2tp4_profiles/{A_prefill,B_decode,D_128k}/`（8 rank × 3 场景；A 的 PP0 ~70-77MB gz）
|
||||||
|
- 分析输出：`.../pp2tp4_profiles/analysis/`
|
||||||
|
- 部署/采集日志：`/root/deploy_pp2tp4_0906*.log`、`/root/prof_cap_*.log`、`/root/phase2.log`、`/root/bs_results/cleanB.log`
|
||||||
|
- 脚本本地副本：`D:\sskj\prof_tools\`；归档：sskj-review `experiments/glm53-tp4pp2-profile-20260906/`
|
||||||
|
|
||||||
|
## 8. 60.5 事后状态
|
||||||
|
|
||||||
|
- 当前运行：**TP4PP2 原样配置**（容器 glm53-nvfp4:30000，chunk 8192 / mem-frac 0.88；09-07 二轮实验后已恢复并验收 health=200）。第二轮实验证明该配置即本栈配置最优，可直接作为后续开发（PP+MTP / AR 内核）的基线。
|
||||||
|
- 若需恢复 GLM-5.3-Flash kt 服务：`bash /root/deploy_glm53_flash_kt_gpu_graphs.sh`（恢复前 kill -9 本容器 sglang 进程并等显存归零）。
|
||||||
|
|
||||||
|
## 9. PP+MTP 深入（2026-09-07 第三轮，进行中）
|
||||||
|
|
||||||
|
### 9.1 复现现场（Step 1a/1b）
|
||||||
|
|
||||||
|
补丁原样重部署（`deploy_ppmtp_bench.sh`，7 文件挂载 + 0901 镜像，健康 ~420s）。smoke 输出与 wiki 报告字节一致（Paris. Distance … Euro.），16k 单请求 31.0s 在报告区间。cc16 压测（种子 6117）：**输入 1598 tok/s（=基线 3305.5 的 0.48×）、TPOT 133.7ms（2.4×基线）**，与 09-06 报告（0.49×/145.7ms）一致——回归稳定复现。
|
||||||
|
|
||||||
|
### 9.2 三源假设实测证伪 + 真凶定位（Step 1c，profiler 无 stack 12 步）
|
||||||
|
|
||||||
|
解码轮结构（TP-0 双 stage，bs 7-8 过渡窗，12 轮 span 3.522s，**293.5ms/轮**，相位标注账目闭合）：
|
||||||
|
|
||||||
|
| 相位 | 每轮 | 占比 |
|
||||||
|
|---|---|---|
|
||||||
|
| TARGET_VERIFY（verify 前向 39 层） | ~210ms | 72% |
|
||||||
|
| draft（3 步 eager） | ~30ms | 10% |
|
||||||
|
| DECODE 步 | ~13ms | 4% |
|
||||||
|
| draft_extend | ~5ms | 2% |
|
||||||
|
| 边界收发 + 调度 | ~35ms | 12% |
|
||||||
|
|
||||||
|
**原三源实测全部出局**：7 处主机同步合计仅 **4.9ms/轮**（120 次 cudaStreamSynchronize 均值 0.49ms）；eager launch 256 次/轮 × 13μs ≈ **3.3ms/轮**；边界 SendRecv GPU 时长仅 **0.17ms/轮**（整卡发送在此尺度无感）。三者合计 <10ms/轮，解释不了 240ms 的回归量。
|
||||||
|
|
||||||
|
**真凶（两级传播链）**：
|
||||||
|
|
||||||
|
1. **边界 GLOO 元数据通道的同步 CPU rendezvous**：`parallel_state.send_object` 每条边界字典 = 2 次 gloo isend（8B size + pickle 体），收端 `recv_object` 两次 irecv 均 `work.wait()` 真阻塞。对端迟到时发送端 gloo:send **自旋阻塞 300-540ms**（trace user_annotation 实测，~每 2 轮一次）；trace 中最大 GPU 空洞 240ms 与之对应。
|
||||||
|
2. **TP 到达偏斜 → AllReduce 自旋**：被 gloo 卡住的 rank 在后续每个 AR 迟到，其余 3 rank 的 AR 内核空转。PP0 AR 合计 **132ms/轮（占轮时 45%、GPU busy 的 82%）**；TP2-PP0 busy 仅 30%（掉队者空闲最多），其余 rank busy 55-62%（AR 等待撑大）。
|
||||||
|
|
||||||
|
结论：**回归主体是协议层（边界通道停顿 + 其 TP 传播），不是执行层**。Step 3 优先级据此改写：边界通道重构（去热路径 GLOO / 预挂 recv）>> 图捕获（仅省 ~20-30ms 计算与 launch）>> 拆同步（≈0）>> 分片发送（≈0）。
|
||||||
|
|
||||||
|
### 9.3 竞态二分结果（Step 2，掩码热切换实测）
|
||||||
|
|
||||||
|
工具：`eagle_worker_v2_mask.py`（7 处同步改为 `/tmp/pp_sync_mask` 位掩码热开关，默认 127 ≡ r33）+ `race_run.sh`。实验矩阵：
|
||||||
|
|
||||||
|
| 构建 | 掩码 | 负载 | 结果 |
|
||||||
|
|---|---|---|---|
|
||||||
|
| r33 | 127 | conc_test 8×16384 | PASS 64.4s（对照，≡09-06 基线） |
|
||||||
|
| r33 | 0 | conc_test 8×16384 ×3 + 完整阶梯 | **3/3 PASS**（30.7/27.5/26.6s）+ smoke 字节一致 |
|
||||||
|
| r33 | 0 | **bench cc16 完整压测** | **CRASH**：prefill 爬坡段（16 并发+队列 churn），PG2(tp:device) NCCL watchdog 报 IMA；表面 dsa init_forward_metadata / allocator insert_events——与 09-06 记录同款 |
|
||||||
|
| r34（延迟 commit） | 0 | conc_test 8×16384 | **CRASH 23.6s**（比 r33+bench 更快） |
|
||||||
|
|
||||||
|
三条结论：
|
||||||
|
|
||||||
|
1. **conc_test 是不合格触发器**：8 并发杀手负载 3/3 通过给了假信心；bench（16 并发、512 输出长 decode、队列 admission churn）才触发。09-06 以全套压测为验证标准是对的。
|
||||||
|
2. **7 处同步在 bench 剖面下确属必要遮罩**（09-06 结论维持），但其成本仅 4.9ms/轮——拆它没有收益，竞态才是封锁项。
|
||||||
|
3. **commit 时的 work.wait() 本身也是隐性时序遮罩**：r34 把它延迟一轮后竞态反而更快触发（23.6s）——发送端节奏是竞态防护的一部分。任何"去停顿"改动（延迟 commit、预挂 recv、去 GLOO）都会扰动同一时序织物，**在竞态根因定位前一律不安全**。
|
||||||
|
|
||||||
|
r34 崩溃样本的表面仍为异步粘性错误（dsa metadata init 的 H2D 拷贝 + allocator insert_events），内核元凶未定位。
|
||||||
|
|
||||||
|
**判别实验不可行性（工具矩阵终点）**：`PYTORCH_NO_CUDA_MEMORY_CACHING=1` 部署死于启动期 CUDA graph capture（cudaErrorStreamCaptureInvalidated——capture 依赖缓存分配器的内存池）；即使能启动，cudaFree 自带的设备同步 ≈ 部分版 CUDA_LAUNCH_BLOCKING（09-06 已证其为遮罩），判别结果必然失真。结论：本竞态的根因定位只能靠 **compute-sanitizer memcheck 全 rank 包裹启动**（唯一不引入同步污染的定位器，~20min 启动成本），归入上游工程。
|
||||||
|
|
||||||
|
### 9.4 终局判断与上游工程清单
|
||||||
|
|
||||||
|
**本轮净结论**：PP+MTP 的 0.41-0.49× 回归不是"eager 无图 + 7 同步 + 4× 线路字节"的执行层问题（三者实测合计 <10ms/轮），而是**边界控制通道的协议层问题**：每轮 3 次串行 GLOO CPU rendezvous 在相位漂移下把发送端 CPU 卡死 300-540ms，传播为 TP 到达偏斜 → AR 自旋（132ms/轮）+ GPU 空洞（~130ms/轮）。同时该竞态（去遮罩即 async IMA）与停顿去除**共享同一时序织物**——7 同步与 commit wait 都是隐性屏障，任何性能手术在根因定位前都不安全。
|
||||||
|
|
||||||
|
上游工程清单（按依赖序）：
|
||||||
|
|
||||||
|
1. **竞态根因定位与修复（解锁一切的前提）**：compute-sanitizer memcheck 包裹启动复现；嫌疑面已收窄到边界 P2P 生命周期/跨流序（record_stream 的流指向、`work.wait()` 事件落在 wait 时刻的当前流而消费在 forward/copy 流）。修复形态候选：P2P 专用 NCCL communicator + 显式跨流 event 协议（发送：生产流→NCCL 流；接收：NCCL 流→消费流）。
|
||||||
|
2. **边界控制通道去 GLOO**（性能主杠杆，依赖 1）：实测停顿链见 9.2；方向 = 固定尺寸元数据走 NCCL P2P / 接收端预挂 recv（后台线程）/ 合并每轮 3 个字典为 1。
|
||||||
|
3. **PP+spec 图捕获**（依赖 1，图回放会跳过 Python 同步）：上游 `_hosts_draft` 半成品落地后的下一步；本栈实测 prefill 分段图仅覆盖 ≤2048 token，decode 图捕获是收益 ~20-30ms/轮的第二杠杆。
|
||||||
|
4. 已证伪/无需：拆同步（4.9ms/轮，且是必要遮罩）、分片发送（SendRecv 0.17ms/轮）、NCCL_PROTO（环回测试已关）、chunk 16384（OOM + 无收益）。
|
||||||
|
|
||||||
|
**60.5 终态**：TP4PP2 无 MTP 基线恢复部署（生产口径不变）。MTP 实验产物（r33/r34/掩码构建、崩溃日志、trace、脚本）留档见 §7 与 sskj-review。
|
||||||
@ -0,0 +1,62 @@
|
|||||||
|
<h1>1. 背景与结论</h1>
|
||||||
|
<p>上游 sglang nightly 镜像（lmsysorg/sglang:nightly-dev-cu13-20260901-07c8f729）对 GLM-5.3-NVFP4 的 TP4PP2 + EAGLE(MTP) 组合只有部分实现：非最后 PP stage 的 scheduler 缺少 draft/verify 结果的回环链路，spec 解码在 PP 下无法工作。本次在 174.1.60.5（8×RTX 6000D，容器 glm53-nvfp4）上自行补全该链路，使 TP4PP2+MTP 服务完整跑通并完成与基线同参的压测。实验全程仅在 60.5 上进行。</p>
|
||||||
|
<callout emoji="🏁" background-color="light-blue" border-color="blue">
|
||||||
|
<p><b>结论：</b>功能层面已达成——输出正确、两 stage 副本锁步一致，楔死与崩溃分别根治/稳定化，连续负载与全套压测零故障。但当前形态下 MTP 是净性能回归：输入吞吐为 TP4PP2 无 MTP 基线的 0.41–0.49×，TPOT 为基线的 2.3–3.4×。回归不来自推测质量（accept len ≈ 2.4–2.5，接受率健康），而来自执行开销：eager draft/verify 无 CUDA 图、每轮 7 次主机同步、边界全张量发送约 4× 线路字节。三项均有明确的后续优化路径。</p>
|
||||||
|
</callout>
|
||||||
|
<h1>2. 实现内容</h1>
|
||||||
|
<p>补丁以挂载文件形式注入容器（/root/sglang_patch2/，只读挂载），核心是两个文件：scheduler_pp_mixin.py（PP 边界通信协议）与 eagle_worker_v2.py（draft/verify 执行与重放）。启动方式：</p>
|
||||||
|
<pre lang="bash" caption="60.5 上的启动命令（deploy_ppmtp.sh，健康约 420s）"><code>bash /root/deploy_ppmtp.sh '--tp 4 --pp-size 2 --disable-overlap-schedule --max-prefill-tokens 16384' mtp
|
||||||
|
# MTP 参数: --speculative-algorithm EAGLE --speculative-num-steps 3
|
||||||
|
# --speculative-eagle-topk 1 --speculative-num-draft-tokens 4</code></pre>
|
||||||
|
<p>协议要点：PP0→PP1 方向由调度器把 draft/verify 所需代理张量（proxy tensors）随 batch 启动发出，经 launch_event 以 wait_event 门控，保证后级在 forward 前拿到；PP1（最后 stage）→PP0 方向把 verify 结果（hidden states、accept_lens、next_token_ids、verify_slots 等）回环给非最后 stage，由 q_event 门控；PP0 再将整理后的结果回传 PP1 供下一轮 draft。元数据先走 GLOO，数据张量再逐个走 NCCL P2P；接收端按消息类型做 typed demux（proxy/output），错序到达暂存 inbox。verify 阶段分配的 draft-KV slots 在 _forward_isolation 回滚后无法再从 batch 上读到，改为在 verify 出口显式捕获并随结果中继，非最后 stage 重放 draft-extend 时显式回填 batch.out_cache_loc。</p>
|
||||||
|
<p>调试与强制项：SGLANG_PP_SPEC_DEBUG 门控全部分布式探针日志；SGLANG_PP_SPEC_FORCE_EAGER_DRAFT=1 与 SGLANG_PP_FORCE_EAGER_VERIFY=1 强制 draft/verify 走 eager——上游没有实现 PP+spec 组合下的图捕获路径（decode 日志可见 cuda graph: False），这是性能回归的第一大头。</p>
|
||||||
|
<h1>3. 两个故障的根因链</h1>
|
||||||
|
<h2>3.1 分块 prefill 确定性楔死（已根治）</h2>
|
||||||
|
<p>现象：10684 token 的 16k 复现请求（分块 [8192, 2492]，两块）在第 2 块边界 7/7 确定性死锁，小 prompt 从不楔死。证据链：py-spy 显示 8 个 scheduler rank 全部停在自己 post-recv 的同步点、GPU 100% 空转；NCCL COLL 日志显示每个 rank 的全部 Send/Recv（各 350+ 个）实际运行在同一条 stream 上——parallel_state.send_tensor_dict/recv_tensor_dict 走的是 torch.distributed.isend/irecv（ProcessGroupNCCL），代码里 with torch.cuda.stream(...) 的包裹并不改变实际执行流。于是两端同时 send-first 时，每个 rank 的 irecv 在 FIFO 上排在自己的 isend 之后，两侧各挂一个等待对端 recv 的大发送，构成环。</p>
|
||||||
|
<p>为什么只有 spec 楔死：spec 回环携带 chunk 尺度的 hidden states（8192 token × 6144 hidden × bf16 ≈ 100MB），NCCL 对该尺寸走 rendezvous，发送必须等对端 recv 就位才完成；非 spec 与 smoke 场景的回环是小张量，走 buffered/LL 路径，发送不等对端也能完成，环不成立。</p>
|
||||||
|
<p>修复：按 pp_rank 奇偶交替排序边界 P2P——偶数 rank 先 send 后 recv，奇数 rank 先 recv 后 send，任何环形等待都被该排序解开。这条规则代码里原本就有，但只对 blocking-send 后端生效（XPU 路径注释写明），本次把它应用到 CUDA。修复后 16k 复现 3/3 通过（39.7s / 22.0s / 31.7s），同步探针收发计数 1072/1072 平衡。</p>
|
||||||
|
<h2>3.2 并发负载竞态（已稳定化，底层竞态对未定位）</h2>
|
||||||
|
<p>现象：debug 探针关闭后，8×16384 并发负载必然触发 async CUDA illegal memory access，崩溃表面点随机（allocator insert_events、dsa_backend init_forward_metadata 等），PP0 为首发方。定性：CUDA_LAUNCH_BLOCKING=1 下同一负载通过（78.1s），是纯竞态；消融实验证明 recv 后无条件 drain schedule stream 不能替代遮罩（36.4s 即崩），真正的遮罩是 eagle_worker_v2 中 7 个 debug 探针块的 .item()/.tolist() 主机同步。</p>
|
||||||
|
<p>修复：把这 7 处阶段边界的主机同步改为无条件执行（日志仍由 debug 门控），8×16384 复现 3/3 通过（60.7s / 30.1s / 28.1s），全套压测零崩溃。另有两项加固：NCCL Work.wait() 默认模式下是 host no-op，"提交"后即释放的张量可能在 isend/irecv 内核仍在读时被分配器复用，故对边界全部 CUDA 张量在发送侧（schedule stream + 当前流）与接收侧（forward + copy stream）补了 record_stream()；最后 stage 的发送数据门控同时 wait_event 到实际执行流（schedule_stream）。</p>
|
||||||
|
<p><b>诚实的限制：</b>7 次无条件同步是遮罩不是根治——底层竞态对未定位（已排查并排除 launch 流门控顺序、分配器复用等嫌疑），定位之前不建议拆除同步。</p>
|
||||||
|
<h1>4. 正确性与稳定性验证</h1>
|
||||||
|
<p>正确性：标准 smoke prompt 每次运行输出字节一致（"Paris. Distance from London to Paris is 343 km, while distance from Manchester to Paris is 605 km. The currency in France is Euro."），PP0/PP1 两 stage 的 decode 日志逐行一致，副本锁步。温度 0 下随机 garbage prompt 的输出存在轻微漂移，为 batch shape 变化引入的数值非确定性（近似值 argmax 翻转），不是正确性缺陷。</p>
|
||||||
|
<p>稳定性阶梯：16k 单请求复现 3/3；conc_test 阶梯（8×64 → 8×2048 → 8×16384）全过，其中杀手负载 8×16384 为 3/3；bench_serving 全套 cc8/16/32（含 warmup 与 flush-cache）零崩溃，completed 数与请求数一致。</p>
|
||||||
|
<h1>5. 压测结果</h1>
|
||||||
|
<p>方法与基线完全同参：容器内 sglang.bench_serving，random-ids 数据集，输入 16384 / 输出 512、range-ratio 1.0，温度 0，--flush-cache，warmup 1，cc=8/16/32；MTP 侧种子 5301/5302/5303，基线种子 7408/7416/7432。基线取 TP4PP2（无 MTP）与 TP8/EP8 两组，详见 <cite type="doc" doc-id="GCHPdWg3WokhCbxjRdbcSj99ncc"></cite>。</p>
|
||||||
|
<table>
|
||||||
|
<thead><tr><th background-color="light-gray">并发</th><th background-color="light-gray">配置</th><th background-color="light-gray">输入吞吐 (tok/s)</th><th background-color="light-gray">输出吞吐 (tok/s)</th><th background-color="light-gray">平均 TTFT (ms)</th><th background-color="light-gray">平均 TPOT (ms)</th></tr></thead>
|
||||||
|
<tbody>
|
||||||
|
<tr><td>8</td><td>TP4PP2+MTP</td><td>1167.63</td><td>36.49</td><td>18323.2</td><td>127.1</td></tr>
|
||||||
|
<tr><td>8</td><td>TP4PP2（基线）</td><td>2829.07</td><td>88.41</td><td>17646.9</td><td>56.0</td></tr>
|
||||||
|
<tr><td>8</td><td>TP8/EP8</td><td>2775.3</td><td>86.73</td><td>29053.7</td><td>28.8</td></tr>
|
||||||
|
<tr><td>16</td><td>TP4PP2+MTP</td><td>1625.79</td><td>50.81</td><td>34744.6</td><td>145.7</td></tr>
|
||||||
|
<tr><td>16</td><td>TP4PP2（基线）</td><td>3305.5</td><td>103.3</td><td>29161.2</td><td>55.4</td></tr>
|
||||||
|
<tr><td>16</td><td>TP8/EP8</td><td>3253.14</td><td>101.66</td><td>46970.8</td><td>54.3</td></tr>
|
||||||
|
<tr><td>32</td><td>TP4PP2+MTP</td><td>1692.36</td><td>52.89</td><td>82031.7</td><td>168.3</td></tr>
|
||||||
|
<tr><td>32</td><td>TP4PP2（基线）</td><td>3846.42</td><td>120.2</td><td>60479.0</td><td>49.6</td></tr>
|
||||||
|
<tr><td>32</td><td>TP8/EP8</td><td>3432.52</td><td>107.27</td><td>80889.2</td><td>73.5</td></tr>
|
||||||
|
</tbody>
|
||||||
|
</table>
|
||||||
|
<p>解读：MTP 侧输入吞吐为 TP4PP2 基线的 0.41–0.49×，输出吞吐同比例；TPOT 为 TP4PP2 基线的 2.3–3.4×（cc8 对 EP8 达 4.4×）；TTFT 相对受损最小，对 EP8 甚至持平或略优（cc8：18.3s vs 29.1s），说明回归集中在 decode 循环，prefill 侧的 spec 开销不大。decode 日志给出 accept len 2.48 / accept rate 0.49——每轮 verify 平均落袋 2.48 个 token，推测质量健康，回归是纯执行开销。归因排序：一是 draft/verify 全 eager 无 CUDA 图，逐层 Python 调度与 kernel launch 开销；二是每轮 7 次无条件主机同步；三是边界张量整卡发送（未按 TP 分片聚合），线路字节约为分片方案的 4 倍。</p>
|
||||||
|
<h1>6. 限制与后续</h1>
|
||||||
|
<p>当前实现的已知限制：</p>
|
||||||
|
<ul>
|
||||||
|
<li>仅覆盖文本路径，mm_input_embeds 未中继，多模态输入不支持；</li>
|
||||||
|
<li>draft/verify 强制 eager（上游无 PP+spec 图捕获路径），decode 无法用图；</li>
|
||||||
|
<li>每轮 7 次无条件主机同步是竞态遮罩，底层竞态对未定位，拆除有崩溃风险；</li>
|
||||||
|
<li>边界张量整卡发送，线路字节约 4×；</li>
|
||||||
|
<li>SGLANG_SPEC_SKIP_ZERO_STEP_DRAFT_EXTEND 开关不支持；</li>
|
||||||
|
<li>温度 0 下输出存在 batch-shape 级数值非确定性（见第 4 节）。</li>
|
||||||
|
</ul>
|
||||||
|
<p>后续优化按收益排序：一是为 draft/verify 实现图捕获（边界 P2P 留在图外、以图间事件衔接），这是回归最大头；二是定位 3.2 节的竞态对后拆除 7 次同步；三是边界按 TP 分片发送或低精度传输压线路字节。三项落地后该组合才有与基线对齐甚至反超的空间——accept len 2.48 说明上限存在。</p>
|
||||||
|
<h1>7. 产物位置</h1>
|
||||||
|
<table>
|
||||||
|
<thead><tr><th background-color="light-gray">产物</th><th background-color="light-gray">服务器（174.1.60.5）</th><th background-color="light-gray">本机（D:\sskj）</th></tr></thead>
|
||||||
|
<tbody>
|
||||||
|
<tr><td>scheduler_pp_mixin.py（边界协议 + 奇偶修复 + 生命周期加固）</td><td>/root/sglang_patch2/</td><td>scheduler_pp_mixin_r27.py</td></tr>
|
||||||
|
<tr><td>eagle_worker_v2.py（7 处阶段边界同步 + 探针）</td><td>/root/sglang_patch2/</td><td>eagle_worker_v2_r33.py</td></tr>
|
||||||
|
<tr><td>部署/压测/复现脚本</td><td>/root/deploy_ppmtp*.sh、/root/bench_ppmtp.sh、/root/conc_test.py、/root/smoke_repro_r25.py</td><td>bench_ppmtp.sh、conc_test.py、wedge_capture.sh</td></tr>
|
||||||
|
<tr><td>压测原始 JSON（cc8/16/32）</td><td>/data/hf_models/bs_results/m5_ppmtp_cc*.json（基线在 /root/bs_results/）</td><td>—</td></tr>
|
||||||
|
</tbody>
|
||||||
|
</table>
|
||||||
@ -0,0 +1,216 @@
|
|||||||
|
== pm-1788716023.1756692-TP-0-PP-0.trace.json.gz ==
|
||||||
|
span 3.522s n_steps~12 round~293.5ms
|
||||||
|
GPU-busy 1939.1ms (55.1%) idle 1582.4ms (44.9%) kernels n=20106
|
||||||
|
host cudaStreamSynchronize: n=120 total=58.6ms (per-call avg 0.49ms)
|
||||||
|
host cudaEventSynchronize: n=12 total=0.1ms cudaEventRecord n=4164
|
||||||
|
host cudaLaunchKernel: n=3072 cpu=39.9ms (avg 13us) memcpy n=314 gpu=0.3ms
|
||||||
|
gaps>0.3ms: n=839 total=932.7ms top5=['239.5', '49.7', '46.1', '42.1', '39.1']ms
|
||||||
|
gap-us attributed to cudaStreamSynchronize overlap: 39.6ms
|
||||||
|
long host ops >1ms: n=1039 total=35837.5ms; top10:
|
||||||
|
+ -21.0ms 3546.8ms [Trace] PyTorch Profiler (0)
|
||||||
|
+ -8.6ms 537.4ms [user_annotation] gloo:send
|
||||||
|
+ 293.5ms 517.0ms [user_annotation] gloo:send
|
||||||
|
+ 275.2ms 513.9ms [user_annotation] gloo:send
|
||||||
|
+ 275.2ms 513.8ms [user_annotation] gloo:send
|
||||||
|
+ 2074.4ms 324.6ms [user_annotation] gloo:send
|
||||||
|
+ 2074.4ms 324.6ms [user_annotation] gloo:send
|
||||||
|
+ 2632.4ms 319.9ms [user_annotation] gloo:send
|
||||||
|
+ 2632.4ms 319.9ms [user_annotation] gloo:send
|
||||||
|
+ 3184.2ms 313.3ms [user_annotation] gloo:send
|
||||||
|
nccl_AllReduce 1588.1ms n= 1056 (45.1% span, 81.9% busy)
|
||||||
|
cutlass_moe 289.5ms n= 9612 (8.2% span, 14.9% busy)
|
||||||
|
sparse_mla 26.9ms n= 1008 (0.8% span, 1.4% busy)
|
||||||
|
other 17.5ms n= 7986 (0.5% span, 0.9% busy)
|
||||||
|
nccl_other 14.1ms n= 36 (0.4% span, 0.7% busy)
|
||||||
|
mqa_logits 3.7ms n= 204 (0.1% span, 0.2% busy)
|
||||||
|
nccl_SendRecv 2.0ms n= 204 (0.1% span, 0.1% busy)
|
||||||
|
|
||||||
|
== pm-1788716023.1756692-TP-0-PP-1.trace.json.gz ==
|
||||||
|
span 11.041s n_steps~12 round~920.1ms
|
||||||
|
GPU-busy 947.6ms (8.6%) idle 10093.5ms (91.4%) kernels n=20457
|
||||||
|
host cudaStreamSynchronize: n=120 total=36.2ms (per-call avg 0.30ms)
|
||||||
|
host cudaEventSynchronize: n=12 total=0.1ms cudaEventRecord n=4374
|
||||||
|
host cudaLaunchKernel: n=3240 cpu=38.9ms (avg 12us) memcpy n=314 gpu=0.3ms
|
||||||
|
gaps>0.3ms: n=1432 total=8788.7ms top5=['7772.2', '23.8', '15.3', '15.2', '14.4']ms
|
||||||
|
gap-us attributed to cudaStreamSynchronize overlap: 7870.4ms
|
||||||
|
long host ops >1ms: n=526 total=51095.8ms; top10:
|
||||||
|
+ -8.3ms 11056.9ms [Trace] PyTorch Profiler (0)
|
||||||
|
+ 3250.1ms 7773.9ms [user_annotation] recv_res_dict_from_prev_stage
|
||||||
|
+ 3250.3ms 7770.3ms [user_annotation] gloo:recv
|
||||||
|
+ 2966.7ms 283.1ms [user_annotation] gloo:send
|
||||||
|
+ 2966.8ms 283.1ms [user_annotation] gloo:send
|
||||||
|
+ 1305.0ms 277.6ms [user_annotation] gloo:send
|
||||||
|
+ 1305.1ms 277.5ms [user_annotation] gloo:send
|
||||||
|
+ 754.3ms 276.8ms [user_annotation] gloo:send
|
||||||
|
+ 754.4ms 276.7ms [user_annotation] gloo:send
|
||||||
|
+ 1869.2ms 275.5ms [user_annotation] gloo:send
|
||||||
|
nccl_AllReduce 535.6ms n= 1044 (4.9% span, 56.5% busy)
|
||||||
|
cutlass_moe 350.3ms n= 9876 (3.2% span, 37.0% busy)
|
||||||
|
sparse_mla 27.4ms n= 1008 (0.2% span, 2.9% busy)
|
||||||
|
other 17.5ms n= 8094 (0.2% span, 1.8% busy)
|
||||||
|
nccl_other 12.1ms n= 48 (0.1% span, 1.3% busy)
|
||||||
|
mqa_logits 3.0ms n= 180 (0.0% span, 0.3% busy)
|
||||||
|
nccl_SendRecv 2.7ms n= 207 (0.0% span, 0.3% busy)
|
||||||
|
|
||||||
|
== pm-1788716023.1756692-TP-1-PP-0.trace.json.gz ==
|
||||||
|
span 3.522s n_steps~12 round~293.5ms
|
||||||
|
GPU-busy 2179.6ms (61.9%) idle 1342.0ms (38.1%) kernels n=20106
|
||||||
|
host cudaStreamSynchronize: n=120 total=94.4ms (per-call avg 0.79ms)
|
||||||
|
host cudaEventSynchronize: n=12 total=0.1ms cudaEventRecord n=4164
|
||||||
|
host cudaLaunchKernel: n=3072 cpu=37.6ms (avg 12us) memcpy n=314 gpu=0.3ms
|
||||||
|
gaps>0.3ms: n=659 total=837.0ms top5=['240.0', '49.5', '45.9', '42.3', '39.6']ms
|
||||||
|
gap-us attributed to cudaStreamSynchronize overlap: 34.6ms
|
||||||
|
long host ops >1ms: n=1138 total=32034.5ms; top10:
|
||||||
|
+ -17.7ms 3543.5ms [Trace] PyTorch Profiler (0)
|
||||||
|
+ -8.5ms 531.8ms [user_annotation] gloo:send
|
||||||
|
+ 279.8ms 507.9ms [user_annotation] gloo:send
|
||||||
|
+ 279.8ms 507.9ms [user_annotation] gloo:send
|
||||||
|
+ 2067.8ms 320.2ms [user_annotation] gloo:send
|
||||||
|
+ 2067.8ms 320.2ms [user_annotation] gloo:send
|
||||||
|
+ 2628.5ms 316.7ms [user_annotation] gloo:send
|
||||||
|
+ 2628.5ms 316.7ms [user_annotation] gloo:send
|
||||||
|
+ 3177.8ms 314.3ms [user_annotation] gloo:send
|
||||||
|
+ 3177.8ms 314.3ms [user_annotation] gloo:send
|
||||||
|
nccl_AllReduce 1828.6ms n= 1056 (51.9% span, 83.9% busy)
|
||||||
|
cutlass_moe 288.3ms n= 9612 (8.2% span, 13.2% busy)
|
||||||
|
sparse_mla 26.7ms n= 1008 (0.8% span, 1.2% busy)
|
||||||
|
other 17.2ms n= 7986 (0.5% span, 0.8% busy)
|
||||||
|
nccl_other 16.2ms n= 36 (0.5% span, 0.7% busy)
|
||||||
|
mqa_logits 3.6ms n= 204 (0.1% span, 0.2% busy)
|
||||||
|
nccl_SendRecv 2.0ms n= 204 (0.1% span, 0.1% busy)
|
||||||
|
|
||||||
|
== pm-1788716023.1756692-TP-1-PP-1.trace.json.gz ==
|
||||||
|
span 11.046s n_steps~12 round~920.5ms
|
||||||
|
GPU-busy 2059.8ms (18.6%) idle 8986.2ms (81.4%) kernels n=20457
|
||||||
|
host cudaStreamSynchronize: n=120 total=86.0ms (per-call avg 0.72ms)
|
||||||
|
host cudaEventSynchronize: n=12 total=0.1ms cudaEventRecord n=4374
|
||||||
|
host cudaLaunchKernel: n=3240 cpu=37.2ms (avg 11us) memcpy n=314 gpu=0.3ms
|
||||||
|
gaps>0.3ms: n=827 total=8437.5ms top5=['7771.2', '22.5', '20.4', '15.5', '15.4']ms
|
||||||
|
gap-us attributed to cudaStreamSynchronize overlap: 7836.1ms
|
||||||
|
long host ops >1ms: n=1091 total=52267.7ms; top10:
|
||||||
|
+ -8.5ms 11062.3ms [Trace] PyTorch Profiler (0)
|
||||||
|
+ 3255.5ms 7772.7ms [user_annotation] recv_res_dict_from_prev_stage
|
||||||
|
+ 3255.6ms 7769.1ms [user_annotation] gloo:recv
|
||||||
|
+ 2965.9ms 289.2ms [user_annotation] gloo:send
|
||||||
|
+ 2966.0ms 289.2ms [user_annotation] gloo:send
|
||||||
|
+ 1862.4ms 287.6ms [user_annotation] gloo:send
|
||||||
|
+ 1862.4ms 287.5ms [user_annotation] gloo:send
|
||||||
|
+ 1302.8ms 285.2ms [user_annotation] gloo:send
|
||||||
|
+ 1302.8ms 285.2ms [user_annotation] gloo:send
|
||||||
|
+ 759.5ms 276.8ms [user_annotation] gloo:send
|
||||||
|
nccl_AllReduce 1638.2ms n= 1044 (14.8% span, 79.5% busy)
|
||||||
|
cutlass_moe 349.2ms n= 9876 (3.2% span, 17.0% busy)
|
||||||
|
sparse_mla 27.4ms n= 1008 (0.2% span, 1.3% busy)
|
||||||
|
nccl_other 23.7ms n= 48 (0.2% span, 1.2% busy)
|
||||||
|
other 17.7ms n= 8094 (0.2% span, 0.9% busy)
|
||||||
|
mqa_logits 3.0ms n= 180 (0.0% span, 0.1% busy)
|
||||||
|
nccl_SendRecv 2.6ms n= 207 (0.0% span, 0.1% busy)
|
||||||
|
|
||||||
|
== pm-1788716023.1756692-TP-2-PP-0.trace.json.gz ==
|
||||||
|
span 3.522s n_steps~12 round~293.5ms
|
||||||
|
GPU-busy 1056.4ms (30.0%) idle 2465.4ms (70.0%) kernels n=20106
|
||||||
|
host cudaStreamSynchronize: n=120 total=58.8ms (per-call avg 0.49ms)
|
||||||
|
host cudaEventSynchronize: n=12 total=0.1ms cudaEventRecord n=4164
|
||||||
|
host cudaLaunchKernel: n=3072 cpu=41.0ms (avg 13us) memcpy n=314 gpu=0.3ms
|
||||||
|
gaps>0.3ms: n=1278 total=1208.5ms top5=['240.0', '47.7', '44.1', '40.6', '37.9']ms
|
||||||
|
gap-us attributed to cudaStreamSynchronize overlap: 36.0ms
|
||||||
|
long host ops >1ms: n=621 total=31027.9ms; top10:
|
||||||
|
+ -11.2ms 3537.6ms [Trace] PyTorch Profiler (0)
|
||||||
|
+ 258.3ms 538.5ms [user_annotation] gloo:send
|
||||||
|
+ 258.4ms 538.5ms [user_annotation] gloo:send
|
||||||
|
+ -8.4ms 537.8ms [user_annotation] gloo:send
|
||||||
|
+ 2077.2ms 322.5ms [user_annotation] gloo:send
|
||||||
|
+ 2077.2ms 322.5ms [user_annotation] gloo:send
|
||||||
|
+ 3185.7ms 319.6ms [user_annotation] gloo:send
|
||||||
|
+ 3185.8ms 319.6ms [user_annotation] gloo:send
|
||||||
|
+ 2635.6ms 317.1ms [user_annotation] gloo:send
|
||||||
|
+ 2635.6ms 317.1ms [user_annotation] gloo:send
|
||||||
|
nccl_AllReduce 707.7ms n= 1056 (20.1% span, 67.0% busy)
|
||||||
|
cutlass_moe 289.3ms n= 9612 (8.2% span, 27.4% busy)
|
||||||
|
sparse_mla 27.0ms n= 1008 (0.8% span, 2.6% busy)
|
||||||
|
other 17.0ms n= 7986 (0.5% span, 1.6% busy)
|
||||||
|
nccl_other 10.1ms n= 36 (0.3% span, 1.0% busy)
|
||||||
|
mqa_logits 3.7ms n= 204 (0.1% span, 0.3% busy)
|
||||||
|
nccl_SendRecv 2.8ms n= 204 (0.1% span, 0.3% busy)
|
||||||
|
|
||||||
|
== pm-1788716023.1756692-TP-2-PP-1.trace.json.gz ==
|
||||||
|
span 11.040s n_steps~12 round~920.0ms
|
||||||
|
GPU-busy 2123.5ms (19.2%) idle 8916.7ms (80.8%) kernels n=20457
|
||||||
|
host cudaStreamSynchronize: n=120 total=134.5ms (per-call avg 1.12ms)
|
||||||
|
host cudaEventSynchronize: n=12 total=0.1ms cudaEventRecord n=4374
|
||||||
|
host cudaLaunchKernel: n=3240 cpu=37.8ms (avg 12us) memcpy n=314 gpu=0.3ms
|
||||||
|
gaps>0.3ms: n=744 total=8389.1ms top5=['7776.1', '24.2', '23.4', '20.7', '14.6']ms
|
||||||
|
gap-us attributed to cudaStreamSynchronize overlap: 7893.3ms
|
||||||
|
long host ops >1ms: n=1082 total=52177.2ms; top10:
|
||||||
|
+ -8.1ms 11055.7ms [Trace] PyTorch Profiler (0)
|
||||||
|
+ 3249.7ms 7777.6ms [user_annotation] recv_res_dict_from_prev_stage
|
||||||
|
+ 3249.8ms 7774.0ms [user_annotation] gloo:recv
|
||||||
|
+ 1869.1ms 275.0ms [user_annotation] gloo:send
|
||||||
|
+ 1869.2ms 274.9ms [user_annotation] gloo:send
|
||||||
|
+ 2974.5ms 274.9ms [user_annotation] gloo:send
|
||||||
|
+ 2974.6ms 274.8ms [user_annotation] gloo:send
|
||||||
|
+ 758.2ms 272.3ms [user_annotation] gloo:send
|
||||||
|
+ 758.2ms 272.3ms [user_annotation] gloo:send
|
||||||
|
+ 1312.0ms 270.1ms [user_annotation] gloo:send
|
||||||
|
nccl_AllReduce 1703.9ms n= 1044 (15.4% span, 80.2% busy)
|
||||||
|
cutlass_moe 348.0ms n= 9876 (3.2% span, 16.4% busy)
|
||||||
|
sparse_mla 27.1ms n= 1008 (0.2% span, 1.3% busy)
|
||||||
|
nccl_other 23.5ms n= 48 (0.2% span, 1.1% busy)
|
||||||
|
other 17.4ms n= 8094 (0.2% span, 0.8% busy)
|
||||||
|
mqa_logits 3.0ms n= 180 (0.0% span, 0.1% busy)
|
||||||
|
nccl_SendRecv 2.5ms n= 207 (0.0% span, 0.1% busy)
|
||||||
|
|
||||||
|
== pm-1788716023.1756692-TP-3-PP-0.trace.json.gz ==
|
||||||
|
span 3.522s n_steps~12 round~293.5ms
|
||||||
|
GPU-busy 2075.2ms (58.9%) idle 1446.3ms (41.1%) kernels n=20106
|
||||||
|
host cudaStreamSynchronize: n=120 total=86.1ms (per-call avg 0.72ms)
|
||||||
|
host cudaEventSynchronize: n=12 total=0.1ms cudaEventRecord n=4164
|
||||||
|
host cudaLaunchKernel: n=3072 cpu=35.9ms (avg 12us) memcpy n=314 gpu=0.3ms
|
||||||
|
gaps>0.3ms: n=806 total=908.2ms top5=['239.9', '49.7', '45.9', '42.3', '39.7']ms
|
||||||
|
gap-us attributed to cudaStreamSynchronize overlap: 34.7ms
|
||||||
|
long host ops >1ms: n=1110 total=31972.9ms; top10:
|
||||||
|
+ -30.9ms 3556.9ms [Trace] PyTorch Profiler (0)
|
||||||
|
+ -8.6ms 538.8ms [user_annotation] gloo:send
|
||||||
|
+ 273.2ms 516.4ms [user_annotation] gloo:send
|
||||||
|
+ 273.3ms 516.4ms [user_annotation] gloo:send
|
||||||
|
+ 2067.2ms 325.5ms [user_annotation] gloo:send
|
||||||
|
+ 2067.2ms 325.5ms [user_annotation] gloo:send
|
||||||
|
+ 2625.7ms 323.3ms [user_annotation] gloo:send
|
||||||
|
+ 2625.7ms 323.2ms [user_annotation] gloo:send
|
||||||
|
+ 3177.4ms 316.5ms [user_annotation] gloo:send
|
||||||
|
+ 3177.4ms 316.5ms [user_annotation] gloo:send
|
||||||
|
nccl_AllReduce 1726.3ms n= 1056 (49.0% span, 83.2% busy)
|
||||||
|
cutlass_moe 289.0ms n= 9612 (8.2% span, 13.9% busy)
|
||||||
|
sparse_mla 26.8ms n= 1008 (0.8% span, 1.3% busy)
|
||||||
|
other 17.4ms n= 7986 (0.5% span, 0.8% busy)
|
||||||
|
nccl_other 13.1ms n= 36 (0.4% span, 0.6% busy)
|
||||||
|
mqa_logits 3.7ms n= 204 (0.1% span, 0.2% busy)
|
||||||
|
nccl_SendRecv 2.0ms n= 204 (0.1% span, 0.1% busy)
|
||||||
|
|
||||||
|
== pm-1788716023.1756692-TP-3-PP-1.trace.json.gz ==
|
||||||
|
span 11.040s n_steps~12 round~920.0ms
|
||||||
|
GPU-busy 2171.9ms (19.7%) idle 8868.2ms (80.3%) kernels n=20457
|
||||||
|
host cudaStreamSynchronize: n=120 total=111.6ms (per-call avg 0.93ms)
|
||||||
|
host cudaEventSynchronize: n=12 total=0.1ms cudaEventRecord n=4374
|
||||||
|
host cudaLaunchKernel: n=3240 cpu=38.8ms (avg 12us) memcpy n=314 gpu=0.3ms
|
||||||
|
gaps>0.3ms: n=692 total=8361.0ms top5=['7770.1', '21.1', '16.9', '15.9', '14.1']ms
|
||||||
|
gap-us attributed to cudaStreamSynchronize overlap: 7846.8ms
|
||||||
|
long host ops >1ms: n=1083 total=52334.1ms; top10:
|
||||||
|
+ -7.5ms 11055.2ms [Trace] PyTorch Profiler (0)
|
||||||
|
+ 3249.5ms 7771.5ms [user_annotation] recv_res_dict_from_prev_stage
|
||||||
|
+ 3249.7ms 7768.0ms [user_annotation] gloo:recv
|
||||||
|
+ 2961.9ms 287.2ms [user_annotation] gloo:send
|
||||||
|
+ 2962.0ms 287.2ms [user_annotation] gloo:send
|
||||||
|
+ 1298.2ms 283.7ms [user_annotation] gloo:send
|
||||||
|
+ 1298.2ms 283.7ms [user_annotation] gloo:send
|
||||||
|
+ 1860.6ms 283.3ms [user_annotation] gloo:send
|
||||||
|
+ 1860.7ms 283.3ms [user_annotation] gloo:send
|
||||||
|
+ 753.5ms 276.9ms [user_annotation] gloo:send
|
||||||
|
nccl_AllReduce 1753.4ms n= 1044 (15.9% span, 80.7% busy)
|
||||||
|
cutlass_moe 348.1ms n= 9876 (3.2% span, 16.0% busy)
|
||||||
|
sparse_mla 27.1ms n= 1008 (0.2% span, 1.2% busy)
|
||||||
|
nccl_other 22.3ms n= 48 (0.2% span, 1.0% busy)
|
||||||
|
other 17.5ms n= 8094 (0.2% span, 0.8% busy)
|
||||||
|
mqa_logits 3.0ms n= 180 (0.0% span, 0.1% busy)
|
||||||
|
nccl_SendRecv 2.4ms n= 207 (0.0% span, 0.1% busy)
|
||||||
|
|
||||||
@ -0,0 +1,70 @@
|
|||||||
|
#!/bin/bash
|
||||||
|
# PP+MTP 竞态二分实验部署（60.5, nightly-dev-cu13-20260901-07c8f729）
|
||||||
|
# 与 deploy_ppmtp_bench.sh 唯一差异：eagle_worker_v2 挂载掩码版（eagle_worker_v2_mask.py）。
|
||||||
|
# 同步掩码运行时读 /tmp/pp_sync_mask（16进制或10进制，缺省=127 全同步），
|
||||||
|
# 可在容器外热改: docker exec glm53-nvfp4 sh -c 'echo 0x40 > /tmp/pp_sync_mask'
|
||||||
|
# 用法: bash deploy_ppmtp_mask.sh "<并行参数>" [mtp|nomtp] [chunk] [memfrac]
|
||||||
|
PAR=${1:?usage: deploy_ppmtp_mask.sh "<flags>" [mtp|nomtp] [chunk] [memfrac]}
|
||||||
|
MTPMODE=${2:-nomtp}
|
||||||
|
CHUNK=${3:-8192}
|
||||||
|
MEMFRAC=${4:-0.88}
|
||||||
|
|
||||||
|
MTPARGS=""
|
||||||
|
if [ "$MTPMODE" = "mtp" ]; then
|
||||||
|
MTPARGS="--speculative-algorithm EAGLE --speculative-num-steps 3 --speculative-eagle-topk 1 --speculative-num-draft-tokens 4"
|
||||||
|
fi
|
||||||
|
|
||||||
|
docker update --restart=no glm53-nvfp4 >/dev/null 2>&1
|
||||||
|
for i in 1 2 3 4 5; do
|
||||||
|
docker rm -f glm53-nvfp4 >/dev/null 2>&1
|
||||||
|
sleep 2
|
||||||
|
docker ps -a --format '{{.Names}}' 2>/dev/null | grep -q '^glm53-nvfp4$' || break
|
||||||
|
done
|
||||||
|
if docker ps -a --format '{{.Names}}' 2>/dev/null | grep -q '^glm53-nvfp4$'; then
|
||||||
|
echo "ERROR: old container cannot be removed"; exit 1
|
||||||
|
fi
|
||||||
|
for i in $(seq 1 15); do ss -ltn 2>/dev/null | grep -q ":30000 " || break; sleep 2; done
|
||||||
|
|
||||||
|
docker run -d --name glm53-nvfp4 --gpus all --shm-size 64g --ipc=host --cap-add SYS_PTRACE \
|
||||||
|
-v /root/sglang_patch2/layer_setup.py:/sgl-workspace/sglang/python/sglang/srt/model_executor/model_runner_components/layer_setup.py:ro \
|
||||||
|
-v /root/sglang_patch2/validation_hook.py:/sgl-workspace/sglang/python/sglang/srt/arg_groups/validation_hook.py:ro \
|
||||||
|
-v /root/eagle_worker_v2_mask.py:/sgl-workspace/sglang/python/sglang/srt/speculative/eagle_worker_v2.py:ro \
|
||||||
|
-v /root/sglang_patch2/eagle_worker_common.py:/sgl-workspace/sglang/python/sglang/srt/speculative/eagle_worker_common.py:ro \
|
||||||
|
-v /root/sglang_patch2/deepseek_nextn.py:/sgl-workspace/sglang/python/sglang/srt/models/deepseek_nextn.py:ro \
|
||||||
|
-v /root/sglang_patch2/scheduler_pp_mixin.py:/sgl-workspace/sglang/python/sglang/srt/managers/scheduler_pp_mixin.py:ro \
|
||||||
|
-v /root/sglang_patch2/deepseek_v2.py:/sgl-workspace/sglang/python/sglang/srt/models/deepseek_v2.py:ro \
|
||||||
|
-e SGLANG_PP_SPEC_DEBUG=0 \
|
||||||
|
-e SGLANG_PP_SPEC_FORCE_EAGER_DRAFT=1 \
|
||||||
|
-e SGLANG_PP_FORCE_EAGER_VERIFY=1 \
|
||||||
|
--restart no -p 30000:30000 \
|
||||||
|
-v /data/hf_models:/data/hf_models \
|
||||||
|
lmsysorg/sglang:nightly-dev-cu13-20260901-07c8f729 \
|
||||||
|
python3 -m sglang.launch_server \
|
||||||
|
--model-path /data/hf_models/GLM-5.3-NVFP4 \
|
||||||
|
--tp 8 \
|
||||||
|
--mem-fraction-static ${MEMFRAC} \
|
||||||
|
--max-running-requests 16 \
|
||||||
|
--chunked-prefill-size ${CHUNK} \
|
||||||
|
--disable-shared-experts-fusion \
|
||||||
|
--moe-runner-backend flashinfer_cutlass \
|
||||||
|
--disable-flashinfer-autotune \
|
||||||
|
--reasoning-parser glm45 --tool-call-parser glm47 \
|
||||||
|
--enable-hierarchical-cache --hicache-ratio 3 \
|
||||||
|
${MTPARGS} \
|
||||||
|
${PAR} \
|
||||||
|
--host 0.0.0.0 --port 30000
|
||||||
|
|
||||||
|
echo "deployed: par=[${PAR}] mtp=${MTPMODE} mask-file=/tmp/pp_sync_mask (default 127); waiting for health..."
|
||||||
|
for i in $(seq 10 10 1800); do
|
||||||
|
code=$(curl -s -o /dev/null -m3 -w '%{http_code}' http://127.0.0.1:30000/health 2>/dev/null)
|
||||||
|
if [ "$code" = "200" ]; then
|
||||||
|
echo "healthy after ${i}s"
|
||||||
|
docker logs glm53-nvfp4 2>&1 | grep -oE "max_total_num_tokens = [0-9]+" | head -1
|
||||||
|
exit 0
|
||||||
|
fi
|
||||||
|
if ! docker ps --format '{{.Names}}' | grep -q '^glm53-nvfp4$'; then
|
||||||
|
echo "CONTAINER DIED after ${i}s"; docker logs --tail 60 glm53-nvfp4 2>&1 | grep -iE "error|assert|not support|incompatible" | tail -8; exit 1
|
||||||
|
fi
|
||||||
|
sleep 10
|
||||||
|
done
|
||||||
|
echo "TIMEOUT waiting for health"; exit 1
|
||||||
@ -0,0 +1,70 @@
|
|||||||
|
#!/bin/bash
|
||||||
|
# PP+MTP 竞态二分实验部署（60.5, nightly-dev-cu13-20260901-07c8f729）
|
||||||
|
# 与 deploy_ppmtp_bench.sh 唯一差异：eagle_worker_v2 挂载掩码版（eagle_worker_v2_mask.py）。
|
||||||
|
# 同步掩码运行时读 /tmp/pp_sync_mask（16进制或10进制，缺省=127 全同步），
|
||||||
|
# 可在容器外热改: docker exec glm53-nvfp4 sh -c 'echo 0x40 > /tmp/pp_sync_mask'
|
||||||
|
# 用法: bash deploy_ppmtp_mask.sh "<并行参数>" [mtp|nomtp] [chunk] [memfrac]
|
||||||
|
PAR=${1:?usage: deploy_ppmtp_mask.sh "<flags>" [mtp|nomtp] [chunk] [memfrac]}
|
||||||
|
MTPMODE=${2:-nomtp}
|
||||||
|
CHUNK=${3:-8192}
|
||||||
|
MEMFRAC=${4:-0.88}
|
||||||
|
|
||||||
|
MTPARGS=""
|
||||||
|
if [ "$MTPMODE" = "mtp" ]; then
|
||||||
|
MTPARGS="--speculative-algorithm EAGLE --speculative-num-steps 3 --speculative-eagle-topk 1 --speculative-num-draft-tokens 4"
|
||||||
|
fi
|
||||||
|
|
||||||
|
docker update --restart=no glm53-nvfp4 >/dev/null 2>&1
|
||||||
|
for i in 1 2 3 4 5; do
|
||||||
|
docker rm -f glm53-nvfp4 >/dev/null 2>&1
|
||||||
|
sleep 2
|
||||||
|
docker ps -a --format '{{.Names}}' 2>/dev/null | grep -q '^glm53-nvfp4$' || break
|
||||||
|
done
|
||||||
|
if docker ps -a --format '{{.Names}}' 2>/dev/null | grep -q '^glm53-nvfp4$'; then
|
||||||
|
echo "ERROR: old container cannot be removed"; exit 1
|
||||||
|
fi
|
||||||
|
for i in $(seq 1 15); do ss -ltn 2>/dev/null | grep -q ":30000 " || break; sleep 2; done
|
||||||
|
|
||||||
|
docker run -d --name glm53-nvfp4 --gpus all --shm-size 64g --ipc=host --cap-add SYS_PTRACE \
|
||||||
|
-v /root/sglang_patch2/layer_setup.py:/sgl-workspace/sglang/python/sglang/srt/model_executor/model_runner_components/layer_setup.py:ro \
|
||||||
|
-v /root/sglang_patch2/validation_hook.py:/sgl-workspace/sglang/python/sglang/srt/arg_groups/validation_hook.py:ro \
|
||||||
|
-v /root/eagle_worker_v2_mask.py:/sgl-workspace/sglang/python/sglang/srt/speculative/eagle_worker_v2.py:ro \
|
||||||
|
-v /root/sglang_patch2/eagle_worker_common.py:/sgl-workspace/sglang/python/sglang/srt/speculative/eagle_worker_common.py:ro \
|
||||||
|
-v /root/sglang_patch2/deepseek_nextn.py:/sgl-workspace/sglang/python/sglang/srt/models/deepseek_nextn.py:ro \
|
||||||
|
-v /root/scheduler_pp_mixin_r34.py:/sgl-workspace/sglang/python/sglang/srt/managers/scheduler_pp_mixin.py:ro \
|
||||||
|
-v /root/sglang_patch2/deepseek_v2.py:/sgl-workspace/sglang/python/sglang/srt/models/deepseek_v2.py:ro \
|
||||||
|
-e SGLANG_PP_SPEC_DEBUG=0 \
|
||||||
|
-e SGLANG_PP_SPEC_FORCE_EAGER_DRAFT=1 \
|
||||||
|
-e SGLANG_PP_FORCE_EAGER_VERIFY=1 \
|
||||||
|
--restart no -p 30000:30000 \
|
||||||
|
-v /data/hf_models:/data/hf_models \
|
||||||
|
lmsysorg/sglang:nightly-dev-cu13-20260901-07c8f729 \
|
||||||
|
python3 -m sglang.launch_server \
|
||||||
|
--model-path /data/hf_models/GLM-5.3-NVFP4 \
|
||||||
|
--tp 8 \
|
||||||
|
--mem-fraction-static ${MEMFRAC} \
|
||||||
|
--max-running-requests 16 \
|
||||||
|
--chunked-prefill-size ${CHUNK} \
|
||||||
|
--disable-shared-experts-fusion \
|
||||||
|
--moe-runner-backend flashinfer_cutlass \
|
||||||
|
--disable-flashinfer-autotune \
|
||||||
|
--reasoning-parser glm45 --tool-call-parser glm47 \
|
||||||
|
--enable-hierarchical-cache --hicache-ratio 3 \
|
||||||
|
${MTPARGS} \
|
||||||
|
${PAR} \
|
||||||
|
--host 0.0.0.0 --port 30000
|
||||||
|
|
||||||
|
echo "deployed: par=[${PAR}] mtp=${MTPMODE} mask-file=/tmp/pp_sync_mask (default 127); waiting for health..."
|
||||||
|
for i in $(seq 10 10 1800); do
|
||||||
|
code=$(curl -s -o /dev/null -m3 -w '%{http_code}' http://127.0.0.1:30000/health 2>/dev/null)
|
||||||
|
if [ "$code" = "200" ]; then
|
||||||
|
echo "healthy after ${i}s"; docker exec glm53-nvfp4 sh -c "echo ${SYNCMASK:-127} > /tmp/pp_sync_mask"
|
||||||
|
docker logs glm53-nvfp4 2>&1 | grep -oE "max_total_num_tokens = [0-9]+" | head -1
|
||||||
|
exit 0
|
||||||
|
fi
|
||||||
|
if ! docker ps --format '{{.Names}}' | grep -q '^glm53-nvfp4$'; then
|
||||||
|
echo "CONTAINER DIED after ${i}s"; docker logs --tail 60 glm53-nvfp4 2>&1 | grep -iE "error|assert|not support|incompatible" | tail -8; exit 1
|
||||||
|
fi
|
||||||
|
sleep 10
|
||||||
|
done
|
||||||
|
echo "TIMEOUT waiting for health"; exit 1
|
||||||
@ -0,0 +1,70 @@
|
|||||||
|
#!/bin/bash
|
||||||
|
# PP+MTP 竞态二分实验部署（60.5, nightly-dev-cu13-20260901-07c8f729）
|
||||||
|
# 与 deploy_ppmtp_bench.sh 唯一差异：eagle_worker_v2 挂载掩码版（eagle_worker_v2_mask.py）。
|
||||||
|
# 同步掩码运行时读 /tmp/pp_sync_mask（16进制或10进制，缺省=127 全同步），
|
||||||
|
# 可在容器外热改: docker exec glm53-nvfp4 sh -c 'echo 0x40 > /tmp/pp_sync_mask'
|
||||||
|
# 用法: bash deploy_ppmtp_mask.sh "<并行参数>" [mtp|nomtp] [chunk] [memfrac]
|
||||||
|
PAR=${1:?usage: deploy_ppmtp_mask.sh "<flags>" [mtp|nomtp] [chunk] [memfrac]}
|
||||||
|
MTPMODE=${2:-nomtp}
|
||||||
|
CHUNK=${3:-8192}
|
||||||
|
MEMFRAC=${4:-0.88}
|
||||||
|
|
||||||
|
MTPARGS=""
|
||||||
|
if [ "$MTPMODE" = "mtp" ]; then
|
||||||
|
MTPARGS="--speculative-algorithm EAGLE --speculative-num-steps 3 --speculative-eagle-topk 1 --speculative-num-draft-tokens 4"
|
||||||
|
fi
|
||||||
|
|
||||||
|
docker update --restart=no glm53-nvfp4 >/dev/null 2>&1
|
||||||
|
for i in 1 2 3 4 5; do
|
||||||
|
docker rm -f glm53-nvfp4 >/dev/null 2>&1
|
||||||
|
sleep 2
|
||||||
|
docker ps -a --format '{{.Names}}' 2>/dev/null | grep -q '^glm53-nvfp4$' || break
|
||||||
|
done
|
||||||
|
if docker ps -a --format '{{.Names}}' 2>/dev/null | grep -q '^glm53-nvfp4$'; then
|
||||||
|
echo "ERROR: old container cannot be removed"; exit 1
|
||||||
|
fi
|
||||||
|
for i in $(seq 1 15); do ss -ltn 2>/dev/null | grep -q ":30000 " || break; sleep 2; done
|
||||||
|
|
||||||
|
docker run -d --name glm53-nvfp4 --gpus all --shm-size 64g --ipc=host --cap-add SYS_PTRACE \
|
||||||
|
-v /root/sglang_patch2/layer_setup.py:/sgl-workspace/sglang/python/sglang/srt/model_executor/model_runner_components/layer_setup.py:ro \
|
||||||
|
-v /root/sglang_patch2/validation_hook.py:/sgl-workspace/sglang/python/sglang/srt/arg_groups/validation_hook.py:ro \
|
||||||
|
-v /root/eagle_worker_v2_mask.py:/sgl-workspace/sglang/python/sglang/srt/speculative/eagle_worker_v2.py:ro \
|
||||||
|
-v /root/sglang_patch2/eagle_worker_common.py:/sgl-workspace/sglang/python/sglang/srt/speculative/eagle_worker_common.py:ro \
|
||||||
|
-v /root/sglang_patch2/deepseek_nextn.py:/sgl-workspace/sglang/python/sglang/srt/models/deepseek_nextn.py:ro \
|
||||||
|
-v /root/scheduler_pp_mixin_r34.py:/sgl-workspace/sglang/python/sglang/srt/managers/scheduler_pp_mixin.py:ro \
|
||||||
|
-v /root/sglang_patch2/deepseek_v2.py:/sgl-workspace/sglang/python/sglang/srt/models/deepseek_v2.py:ro \
|
||||||
|
-e SGLANG_PP_SPEC_DEBUG=0 \
|
||||||
|
-e SGLANG_PP_SPEC_FORCE_EAGER_DRAFT=1 \
|
||||||
|
-e SGLANG_PP_FORCE_EAGER_VERIFY=1 -e PYTORCH_NO_CUDA_MEMORY_CACHING=1 \
|
||||||
|
--restart no -p 30000:30000 \
|
||||||
|
-v /data/hf_models:/data/hf_models \
|
||||||
|
lmsysorg/sglang:nightly-dev-cu13-20260901-07c8f729 \
|
||||||
|
python3 -m sglang.launch_server \
|
||||||
|
--model-path /data/hf_models/GLM-5.3-NVFP4 \
|
||||||
|
--tp 8 \
|
||||||
|
--mem-fraction-static ${MEMFRAC} \
|
||||||
|
--max-running-requests 16 \
|
||||||
|
--chunked-prefill-size ${CHUNK} \
|
||||||
|
--disable-shared-experts-fusion \
|
||||||
|
--moe-runner-backend flashinfer_cutlass \
|
||||||
|
--disable-flashinfer-autotune \
|
||||||
|
--reasoning-parser glm45 --tool-call-parser glm47 \
|
||||||
|
--enable-hierarchical-cache --hicache-ratio 3 \
|
||||||
|
${MTPARGS} \
|
||||||
|
${PAR} \
|
||||||
|
--host 0.0.0.0 --port 30000
|
||||||
|
|
||||||
|
echo "deployed: par=[${PAR}] mtp=${MTPMODE} mask-file=/tmp/pp_sync_mask (default 127); waiting for health..."
|
||||||
|
for i in $(seq 10 10 1800); do
|
||||||
|
code=$(curl -s -o /dev/null -m3 -w '%{http_code}' http://127.0.0.1:30000/health 2>/dev/null)
|
||||||
|
if [ "$code" = "200" ]; then
|
||||||
|
echo "healthy after ${i}s"; docker exec glm53-nvfp4 sh -c "echo ${SYNCMASK:-127} > /tmp/pp_sync_mask"
|
||||||
|
docker logs glm53-nvfp4 2>&1 | grep -oE "max_total_num_tokens = [0-9]+" | head -1
|
||||||
|
exit 0
|
||||||
|
fi
|
||||||
|
if ! docker ps --format '{{.Names}}' | grep -q '^glm53-nvfp4$'; then
|
||||||
|
echo "CONTAINER DIED after ${i}s"; docker logs --tail 60 glm53-nvfp4 2>&1 | grep -iE "error|assert|not support|incompatible" | tail -8; exit 1
|
||||||
|
fi
|
||||||
|
sleep 10
|
||||||
|
done
|
||||||
|
echo "TIMEOUT waiting for health"; exit 1
|
||||||
File diff suppressed because it is too large
Load Diff
@ -0,0 +1,31 @@
|
|||||||
|
import gzip, json, sys
|
||||||
|
from collections import defaultdict
|
||||||
|
|
||||||
|
def census(path):
|
||||||
|
with gzip.open(path, 'rt') as f:
|
||||||
|
data = json.load(f)
|
||||||
|
ann = defaultdict(lambda: [0, 0.0])
|
||||||
|
gloo = defaultdict(lambda: [0, 0.0, 0.0]) # name -> [n, total, max]
|
||||||
|
for e in data['traceEvents']:
|
||||||
|
if e.get('ph') != 'X': continue
|
||||||
|
nm = e.get('name', ''); cat = e.get('cat', ''); d = e.get('dur', 0)
|
||||||
|
if cat == 'user_annotation':
|
||||||
|
ann[nm][0] += 1; ann[nm][1] += d
|
||||||
|
elif nm.startswith('gloo:') and cat not in ('kernel',):
|
||||||
|
gloo[nm][2] = max(gloo[nm][2], d)
|
||||||
|
gloo[nm][0] += 1; gloo[nm][1] += d
|
||||||
|
print(f"== {path.split('/')[-1]} ==")
|
||||||
|
print(" user_annotations:")
|
||||||
|
for nm in sorted(ann, key=lambda x: -ann[x][1]):
|
||||||
|
n, t = ann[nm]
|
||||||
|
print(f" {t/1e3:9.1f}ms n={n:5d} {nm[:100]}")
|
||||||
|
print(" gloo ops:")
|
||||||
|
for nm in sorted(gloo, key=lambda x: -gloo[x][1]):
|
||||||
|
n, t, mx = gloo[nm]
|
||||||
|
print(f" {t/1e3:9.1f}ms n={n:5d} max={mx/1e3:8.1f}ms {nm}")
|
||||||
|
print()
|
||||||
|
|
||||||
|
if __name__ == '__main__':
|
||||||
|
import glob
|
||||||
|
for p in sorted(glob.glob(sys.argv[1])):
|
||||||
|
census(p)
|
||||||
@ -0,0 +1,104 @@
|
|||||||
|
import gzip, json, sys
|
||||||
|
from collections import defaultdict
|
||||||
|
|
||||||
|
def merged(intervals):
|
||||||
|
ivs = sorted((s, e) for s, e in intervals if e > s)
|
||||||
|
out = []
|
||||||
|
for s, e in ivs:
|
||||||
|
if out and s <= out[-1][1]:
|
||||||
|
if e > out[-1][1]: out[-1][1] = e
|
||||||
|
else:
|
||||||
|
out.append([s, e])
|
||||||
|
return out
|
||||||
|
|
||||||
|
def analyze(path, n_steps=12):
|
||||||
|
with gzip.open(path, 'rt') as f:
|
||||||
|
data = json.load(f)
|
||||||
|
ev = [e for e in data['traceEvents'] if e.get('ph') == 'X']
|
||||||
|
kernels = [] # (ts, end, name) us
|
||||||
|
memcpy = []
|
||||||
|
cpu_rt = defaultdict(lambda: [0, 0.0]) # cuda_runtime name -> [n, total_us]
|
||||||
|
cpu_sync_windows = [] # cudaStreamSynchronize (ts, end)
|
||||||
|
launches = []
|
||||||
|
for e in ev:
|
||||||
|
cat = e.get('cat', '')
|
||||||
|
ts = e['ts']; dur = e.get('dur', 0)
|
||||||
|
if cat == 'kernel':
|
||||||
|
kernels.append((ts, ts + dur, e.get('name', '')))
|
||||||
|
elif cat == 'gpu_memcpy':
|
||||||
|
memcpy.append((ts, ts + dur, e.get('name', '')))
|
||||||
|
elif cat == 'cuda_runtime':
|
||||||
|
nm = e.get('name', '')
|
||||||
|
cpu_rt[nm][0] += 1; cpu_rt[nm][1] += dur
|
||||||
|
if nm == 'cudaStreamSynchronize':
|
||||||
|
cpu_sync_windows.append((ts, ts + dur))
|
||||||
|
elif nm == 'cudaLaunchKernel':
|
||||||
|
launches.append((ts, dur))
|
||||||
|
if not kernels:
|
||||||
|
print(f"{path}: no kernels"); return
|
||||||
|
t0 = min(k[0] for k in kernels); t1 = max(k[1] for k in kernels)
|
||||||
|
# long host-side ops (any ph=X outside cuda_runtime): gloo recv, work.wait, python fns
|
||||||
|
long_cpu = []
|
||||||
|
for e in ev:
|
||||||
|
cat = e.get('cat', '')
|
||||||
|
if cat in ('cuda_runtime', 'kernel', 'gpu_memcpy'): continue
|
||||||
|
d = e.get('dur', 0)
|
||||||
|
if d > 1000:
|
||||||
|
long_cpu.append((e['ts'] - t0, d, e.get('name', '')[:110], cat))
|
||||||
|
long_cpu.sort(key=lambda x: -x[1])
|
||||||
|
span = t1 - t0
|
||||||
|
busy_iv = merged([(k[0], k[1]) for k in kernels])
|
||||||
|
busy = sum(e - s for s, e in busy_iv)
|
||||||
|
idle = span - busy
|
||||||
|
# gaps
|
||||||
|
gaps = []
|
||||||
|
for (s1, e1), (s2, e2) in zip(busy_iv, busy_iv[1:]):
|
||||||
|
if s2 - e1 > 0: gaps.append((e1, s2 - e1))
|
||||||
|
# attribute gaps >0.3ms: overlapping a cudaStreamSynchronize window?
|
||||||
|
sync_set = merged(cpu_sync_windows)
|
||||||
|
def overlaps_sync(gts):
|
||||||
|
for s, e in sync_set:
|
||||||
|
if s < gts and e > gts - 0: # any overlap
|
||||||
|
return True
|
||||||
|
return False
|
||||||
|
big = [(gs, gd) for gs, gd in gaps if gd > 300]
|
||||||
|
sync_attr = sum(gd for gs, gd in big if any(s < gs + gd and e > gs for s, e in sync_set))
|
||||||
|
# kernel buckets
|
||||||
|
buckets = defaultdict(lambda: [0, 0.0])
|
||||||
|
for ts, te, nm in kernels:
|
||||||
|
b = ('nccl_SendRecv' if 'SendRecv' in nm else
|
||||||
|
'nccl_AllReduce' if 'AllReduce' in nm else
|
||||||
|
'nccl_other' if 'ncclDevKernel' in nm else
|
||||||
|
'mqa_logits' if 'mqa_logits' in nm else
|
||||||
|
'sparse_mla' if 'sparse_mla' in nm else
|
||||||
|
'cutlass_moe' if 'cutlass' in nm.lower() else
|
||||||
|
'other')
|
||||||
|
buckets[b][0] += 1; buckets[b][1] += te - ts
|
||||||
|
mc = sum(1 for _ in memcpy); mt = sum(e - s for s, e, _ in memcpy)
|
||||||
|
rt = lambda n: cpu_rt.get(n, [0, 0.0])
|
||||||
|
n_sync, sync_us = rt('cudaStreamSynchronize')
|
||||||
|
n_lk, lk_us = rt('cudaLaunchKernel')
|
||||||
|
n_ev, ev_us = rt('cudaEventSynchronize')
|
||||||
|
n_ecr, ecr_us = rt('cudaEventRecord')
|
||||||
|
print(f"== {path.split('/')[-1]} ==")
|
||||||
|
print(f" span {span/1e6:.3f}s n_steps~{n_steps} round~{span/1e3/n_steps:.1f}ms")
|
||||||
|
print(f" GPU-busy {busy/1e3:.1f}ms ({100*busy/span:.1f}%) idle {idle/1e3:.1f}ms ({100*idle/span:.1f}%) kernels n={len(kernels)}")
|
||||||
|
print(f" host cudaStreamSynchronize: n={n_sync} total={sync_us/1e3:.1f}ms (per-call avg {sync_us/max(n_sync,1)/1e3:.2f}ms)")
|
||||||
|
print(f" host cudaEventSynchronize: n={n_ev} total={ev_us/1e3:.1f}ms cudaEventRecord n={n_ecr}")
|
||||||
|
print(f" host cudaLaunchKernel: n={n_lk} cpu={lk_us/1e3:.1f}ms (avg {lk_us/max(n_lk,1):.0f}us) memcpy n={mc} gpu={mt/1e3:.1f}ms")
|
||||||
|
print(f" gaps>0.3ms: n={len(big)} total={sum(gd for _,gd in big)/1e3:.1f}ms top5={[f'{gd/1e3:.1f}' for _,gd in sorted(big,key=lambda x:-x[1])[:5]]}ms")
|
||||||
|
print(f" gap-us attributed to cudaStreamSynchronize overlap: {sync_attr/1e3:.1f}ms")
|
||||||
|
print(f" long host ops >1ms: n={len(long_cpu)} total={sum(d for _,d,_,_ in long_cpu)/1e3:.1f}ms; top10:")
|
||||||
|
for ts, d, nm, cat in long_cpu[:10]:
|
||||||
|
print(f" +{ts/1e3:9.1f}ms {d/1e3:7.1f}ms [{cat}] {nm}")
|
||||||
|
for b in sorted(buckets, key=lambda x: -buckets[x][1]):
|
||||||
|
n, d = buckets[b]
|
||||||
|
if d > 1000:
|
||||||
|
print(f" {b:14s} {d/1e3:9.1f}ms n={n:6d} ({100*d/span:.1f}% span, {100*d/busy:.1f}% busy)")
|
||||||
|
print()
|
||||||
|
|
||||||
|
if __name__ == '__main__':
|
||||||
|
base = sys.argv[1] if len(sys.argv) > 1 else '/data/hf_models/ppmtp_profiles/M_decode'
|
||||||
|
import glob
|
||||||
|
for p in sorted(glob.glob(base + '/pm-*.trace.json.gz')):
|
||||||
|
analyze(p)
|
||||||
@ -0,0 +1,47 @@
|
|||||||
|
#!/bin/bash
|
||||||
|
# PP+MTP spec-decode round 采集（60.5, glm53-nvfp4:30000, TP4PP2+MTP 全 eager）
|
||||||
|
# 用法: bash prof_mtp_decode.sh
|
||||||
|
# 产出: /data/hf_models/ppmtp_profiles/M_decode/ 8 ranks x 12 steps trace
|
||||||
|
set -u
|
||||||
|
BASE=/data/hf_models/ppmtp_profiles
|
||||||
|
URL=http://127.0.0.1:30000
|
||||||
|
mkdir -p "$BASE/M_decode" /root/bs_results /data/hf_models/bs_results
|
||||||
|
|
||||||
|
health() { [ "$(curl -s -o /dev/null -m3 -w '%{http_code}' $URL/health)" = "200" ]; }
|
||||||
|
health || { echo "server not healthy"; exit 1; }
|
||||||
|
|
||||||
|
SEED=6221
|
||||||
|
curl -s -X POST $URL/flush_cache >/dev/null; sleep 3
|
||||||
|
|
||||||
|
D0=$(docker logs glm53-nvfp4 2>&1 | grep -c "Decode batch")
|
||||||
|
docker exec glm53-nvfp4 python3 -m sglang.bench_serving \
|
||||||
|
--backend sglang --host 127.0.0.1 --port 30000 \
|
||||||
|
--dataset-name random-ids --tokenizer /data/hf_models/GLM-5.3-NVFP4 \
|
||||||
|
--num-prompts 16 --random-input-len 16384 --random-output-len 512 \
|
||||||
|
--random-range-ratio 1.0 --max-concurrency 16 --temperature 0.0 \
|
||||||
|
--warmup-requests 0 --seed $SEED \
|
||||||
|
--output-file /data/hf_models/bs_results/prof_M_mtp_${SEED}.json \
|
||||||
|
> /root/bs_results/prof_M_bench.log 2>&1 &
|
||||||
|
BPID=$!
|
||||||
|
|
||||||
|
for i in $(seq 1 120); do
|
||||||
|
D=$(docker logs glm53-nvfp4 2>&1 | grep -c "Decode batch")
|
||||||
|
[ "$D" -gt "$D0" ] && break
|
||||||
|
sleep 2
|
||||||
|
done
|
||||||
|
echo "[M] decode phase after ~$((i*2))s (decode-lines D0=$D0 -> D=$D)"
|
||||||
|
|
||||||
|
curl -s -X POST $URL/start_profile -H 'Content-Type: application/json' \
|
||||||
|
-d "{\"output_dir\":\"$BASE/M_decode\",\"num_steps\":12,\"activities\":[\"CPU\",\"GPU\"],\"with_stack\":false,\"record_shapes\":false,\"profile_prefix\":\"pm\"}"
|
||||||
|
echo; echo "[arm] 12 steps at $(date +%T)"
|
||||||
|
|
||||||
|
wait $BPID; echo "[bench] exit=$? (数字仅参考, 采集期内)"
|
||||||
|
grep -E "Output token throughput|Mean TPOT" /root/bs_results/prof_M_bench.log | tail -2
|
||||||
|
|
||||||
|
n=0
|
||||||
|
for i in $(seq 1 90); do
|
||||||
|
n=$(ls "$BASE/M_decode"/*.trace.json.gz 2>/dev/null | wc -l)
|
||||||
|
[ "$n" -ge 8 ] && { echo "[traces] $n files after ~$((i*5))s"; ls -la "$BASE/M_decode" | tail -10; exit 0; }
|
||||||
|
sleep 5
|
||||||
|
done
|
||||||
|
echo "[traces] TIMEOUT only $n files"; ls -la "$BASE/M_decode" 2>/dev/null
|
||||||
@ -0,0 +1,43 @@
|
|||||||
|
#!/bin/bash
|
||||||
|
# PP+MTP spec-decode 稳态采集 v2：等全部 16 请求进入 decode（#running-req: 16 连续 5 行）后 arm。
|
||||||
|
set -u
|
||||||
|
BASE=/data/hf_models/ppmtp_profiles
|
||||||
|
URL=http://127.0.0.1:30000
|
||||||
|
mkdir -p "$BASE/M2_decode" /root/bs_results /data/hf_models/bs_results
|
||||||
|
health() { [ "$(curl -s -o /dev/null -m3 -w '%{http_code}' $URL/health)" = "200" ]; }
|
||||||
|
health || { echo "server not healthy"; exit 1; }
|
||||||
|
|
||||||
|
SEED=6231
|
||||||
|
curl -s -X POST $URL/flush_cache >/dev/null; sleep 3
|
||||||
|
|
||||||
|
D0=$(docker logs glm53-nvfp4 2>&1 | grep -c "#running-req: 16")
|
||||||
|
docker exec glm53-nvfp4 python3 -m sglang.bench_serving \
|
||||||
|
--backend sglang --host 127.0.0.1 --port 30000 \
|
||||||
|
--dataset-name random-ids --tokenizer /data/hf_models/GLM-5.3-NVFP4 \
|
||||||
|
--num-prompts 16 --random-input-len 16384 --random-output-len 512 \
|
||||||
|
--random-range-ratio 1.0 --max-concurrency 16 --temperature 0.0 \
|
||||||
|
--warmup-requests 0 --seed $SEED \
|
||||||
|
--output-file /data/hf_models/bs_results/prof_M2_mtp_${SEED}.json \
|
||||||
|
> /root/bs_results/prof_M2_bench.log 2>&1 &
|
||||||
|
BPID=$!
|
||||||
|
|
||||||
|
echo "[M2] waiting for steady decode (#running-req: 16 x5)..."
|
||||||
|
for i in $(seq 1 240); do
|
||||||
|
D=$(docker logs glm53-nvfp4 2>&1 | grep -c "#running-req: 16")
|
||||||
|
[ "$D" -ge $((D0 + 5)) ] && break
|
||||||
|
sleep 2
|
||||||
|
done
|
||||||
|
echo "[M2] steady decode after ~$((i*2))s (D0=$D0 -> D=$D)"
|
||||||
|
|
||||||
|
curl -s -X POST $URL/start_profile -H 'Content-Type: application/json' \
|
||||||
|
-d "{\"output_dir\":\"$BASE/M2_decode\",\"num_steps\":12,\"activities\":[\"CPU\",\"GPU\"],\"with_stack\":false,\"record_shapes\":false,\"profile_prefix\":\"pm2\"}"
|
||||||
|
echo; echo "[arm] 12 steps at $(date +%T)"
|
||||||
|
|
||||||
|
wait $BPID; echo "[bench] exit=$?"
|
||||||
|
n=0
|
||||||
|
for i in $(seq 1 90); do
|
||||||
|
n=$(ls "$BASE/M2_decode"/*.trace.json.gz 2>/dev/null | wc -l)
|
||||||
|
[ "$n" -ge 8 ] && { echo "[traces] $n files after ~$((i*5))s"; ls -la "$BASE/M2_decode" | tail -10; exit 0; }
|
||||||
|
sleep 5
|
||||||
|
done
|
||||||
|
echo "[traces] TIMEOUT only $n files"; ls -la "$BASE/M2_decode" 2>/dev/null
|
||||||
@ -0,0 +1,48 @@
|
|||||||
|
#!/bin/bash
|
||||||
|
# PP+MTP 稳态采集 v3: 16k in / 2048 out 拉长 decode 段, 等 running-req>=12 连续5行后 arm 12 步。
|
||||||
|
set -u
|
||||||
|
BASE=/data/hf_models/ppmtp_profiles
|
||||||
|
URL=http://127.0.0.1:30000
|
||||||
|
mkdir -p "$BASE/M3_decode" /root/bs_results /data/hf_models/bs_results
|
||||||
|
rmdir "$BASE/M2_decode" 2>/dev/null
|
||||||
|
health() { [ "$(curl -s -o /dev/null -m3 -w '%{http_code}' $URL/health)" = "200" ]; }
|
||||||
|
health || { echo "server not healthy"; exit 1; }
|
||||||
|
|
||||||
|
SEED=6241
|
||||||
|
curl -s -X POST $URL/flush_cache >/dev/null; sleep 3
|
||||||
|
|
||||||
|
D0=$(docker logs glm53-nvfp4 2>&1 | grep -cE '#running-req: 1[2-9],|#running-req: 16,')
|
||||||
|
docker exec glm53-nvfp4 python3 -m sglang.bench_serving \
|
||||||
|
--backend sglang --host 127.0.0.1 --port 30000 \
|
||||||
|
--dataset-name random-ids --tokenizer /data/hf_models/GLM-5.3-NVFP4 \
|
||||||
|
--num-prompts 16 --random-input-len 16384 --random-output-len 2048 \
|
||||||
|
--random-range-ratio 1.0 --max-concurrency 16 --temperature 0.0 \
|
||||||
|
--warmup-requests 0 --seed $SEED \
|
||||||
|
--output-file /data/hf_models/bs_results/prof_M3_mtp_${SEED}.json \
|
||||||
|
> /root/bs_results/prof_M3_bench.log 2>&1 &
|
||||||
|
BPID=$!
|
||||||
|
|
||||||
|
echo "[M3] waiting for high-bs decode (running-req>=12 x5)..."
|
||||||
|
for i in $(seq 1 300); do
|
||||||
|
D=$(docker logs glm53-nvfp4 2>&1 | grep -cE '#running-req: 1[2-9],|#running-req: 16,')
|
||||||
|
[ "$D" -ge $((D0 + 5)) ] && break
|
||||||
|
sleep 2
|
||||||
|
done
|
||||||
|
echo "[M3] high-bs decode after ~$((i*2))s (D0=$D0 -> D=$D)"
|
||||||
|
|
||||||
|
curl -s -X POST $URL/start_profile -H 'Content-Type: application/json' \
|
||||||
|
-d "{\"output_dir\":\"$BASE/M3_decode\",\"num_steps\":12,\"activities\":[\"CPU\",\"GPU\"],\"with_stack\":false,\"record_shapes\":false,\"profile_prefix\":\"pm3\"}"
|
||||||
|
echo; echo "[arm] 12 steps at $(date +%T)"
|
||||||
|
|
||||||
|
# 不等 bench 完成: arm 后 12 步 ~4s + flush ~11s, 之后再等 bench
|
||||||
|
sleep 40
|
||||||
|
n=0
|
||||||
|
for i in $(seq 1 60); do
|
||||||
|
n=$(ls "$BASE/M3_decode"/*.trace.json.gz 2>/dev/null | wc -l)
|
||||||
|
[ "$n" -ge 8 ] && { echo "[traces] $n files after ~$((i*5))s"; break; }
|
||||||
|
sleep 5
|
||||||
|
done
|
||||||
|
[ "$n" -ge 8 ] || { echo "[traces] only $n files"; ls -la "$BASE/M3_decode" 2>/dev/null; }
|
||||||
|
kill $BPID 2>/dev/null
|
||||||
|
wait $BPID 2>/dev/null
|
||||||
|
echo DONE
|
||||||
@ -0,0 +1,46 @@
|
|||||||
|
#!/bin/bash
|
||||||
|
# 竞态二分单点运行器: race_run.sh <mask> <label> [timeout_s]
|
||||||
|
# 1) 热写掩码 2) 跑 conc_test big (8x16384 杀手负载)
|
||||||
|
# 3) 监控容器死亡/健康失效, 及时收割 docker 日志
|
||||||
|
# 退出码: 0=PASS 2=容器死 3=健康失效 4=客户端超时/失败
|
||||||
|
MASK=$1; LABEL=$2; TMO=${3:-420}
|
||||||
|
URL=http://127.0.0.1:30000
|
||||||
|
|
||||||
|
docker exec glm53-nvfp4 sh -c "echo $MASK > /tmp/pp_sync_mask"
|
||||||
|
sleep 1
|
||||||
|
echo "MASK=$MASK LABEL=$LABEL start $(date +%T)"
|
||||||
|
|
||||||
|
timeout $TMO python3 /root/conc_test.py big > /root/race_${LABEL}.log 2>&1 &
|
||||||
|
CPID=$!
|
||||||
|
|
||||||
|
BADHEALTH=0
|
||||||
|
for i in $(seq 1 $((TMO/10))); do
|
||||||
|
sleep 10
|
||||||
|
if ! docker ps --format '{{.Names}}' 2>/dev/null | grep -q '^glm53-nvfp4$'; then
|
||||||
|
echo "VERDICT=CONTAINER-DIED at +$((i*10))s"
|
||||||
|
kill $CPID 2>/dev/null
|
||||||
|
docker logs --tail 120 glm53-nvfp4 > /root/race_${LABEL}_docker.log 2>&1
|
||||||
|
grep -nE "illegal|CUDA error|RuntimeError|Traceback|assert" /root/race_${LABEL}_docker.log | sed 's/closed by/closed-by/g' | tail -25
|
||||||
|
exit 2
|
||||||
|
fi
|
||||||
|
CODE=$(curl -s -o /dev/null -m3 -w '%{http_code}' $URL/health 2>/dev/null)
|
||||||
|
if [ "$CODE" != "200" ]; then
|
||||||
|
BADHEALTH=$((BADHEALTH+1))
|
||||||
|
if [ $BADHEALTH -ge 3 ]; then
|
||||||
|
echo "VERDICT=HEALTH-DEAD at +$((i*10))s (container up, service dead)"
|
||||||
|
kill $CPID 2>/dev/null
|
||||||
|
docker logs --tail 120 glm53-nvfp4 > /root/race_${LABEL}_docker.log 2>&1
|
||||||
|
grep -nE "illegal|CUDA error|RuntimeError|Traceback|assert" /root/race_${LABEL}_docker.log | sed 's/closed by/closed-by/g' | tail -25
|
||||||
|
exit 3
|
||||||
|
fi
|
||||||
|
else
|
||||||
|
BADHEALTH=0
|
||||||
|
fi
|
||||||
|
kill -0 $CPID 2>/dev/null || break
|
||||||
|
done
|
||||||
|
|
||||||
|
wait $CPID; RC=$?
|
||||||
|
echo "MASK=$MASK LABEL=$LABEL client_rc=$RC $(date +%T)"
|
||||||
|
tail -4 /root/race_${LABEL}.log
|
||||||
|
if [ $RC -eq 0 ]; then echo VERDICT=PASS; exit 0; fi
|
||||||
|
echo VERDICT=CLIENT-FAIL; exit 4
|
||||||
File diff suppressed because it is too large
Load Diff
@ -0,0 +1,82 @@
|
|||||||
|
# GLM-5.3-NVFP4 (Pro6000D×8) SGLang TP4PP2 Torch-Profiler Profile 与优化方向
|
||||||
|
|
||||||
|
- 日期：2026-09-06 机器：174.1.60.5（8×RTX 6000D 85.6GB, SM120, PCIe Gen5, 无 NVLink）
|
||||||
|
- 完整报告：`D:\sskj\reports\GLM53_TP4PP2_Profile与优化方向_2026-09-06.md`（结论、排名表、NCCL 验证实验全在该文件）
|
||||||
|
- 上游依据：09-05 并行方案对比（TP8 / TP4PP2 / EP8 / DP）确认 TP4PP2 为最优配置
|
||||||
|
|
||||||
|
## 目的
|
||||||
|
|
||||||
|
对获胜配置 TP4PP2 做 prefill / decode / 128k 前缀命中三场景 torch profiler 采集，
|
||||||
|
量化 kernel 热点、PP 流水线气泡、NCCL 通信占比、CPU 调度空隙，产出优化方向排序；
|
||||||
|
并用对照实验检验 profile 提出的第一假设（NCCL LL 协议误选）。
|
||||||
|
|
||||||
|
## 部署（与 09-05 获胜配置逐字一致）
|
||||||
|
|
||||||
|
```
|
||||||
|
bash /root/deploy_par.sh "--tp 4 --pp-size 2 --disable-overlap-schedule --max-prefill-tokens 16384" nomtp
|
||||||
|
```
|
||||||
|
|
||||||
|
镜像 `lmsysorg/sglang:nightly-dev-20260828-daf63171`；容器 glm53-nvfp4:30000；
|
||||||
|
KV fp8_e4m3 自动；SM120 三件套（`--disable-shared-experts-fusion --moe-runner-backend flashinfer_cutlass --disable-flashinfer-autotune`）；
|
||||||
|
hicache-ratio 3；mem-frac 0.88；chunked-prefill 8192。忠实性：干净复测 cc16 16k/512
|
||||||
|
输出 104.7 tok/s / TPOT 50.87ms，与 09-05 基准（103.3 / 55.4）一致。
|
||||||
|
|
||||||
|
## 采集协议
|
||||||
|
|
||||||
|
| 场景 | 负载 | 说明 |
|
||||||
|
|---|---|---|
|
||||||
|
| A prefill | 16×(16384→1) cc16 | ~34 个 8192-token chunk 前向 |
|
||||||
|
| B decode | 16×(16384→512) cc16 | TTFT 过后 arm 60 步 |
|
||||||
|
| D 128k/90% | 8×131072 cc8 | 共享前缀 117968；后缀 prefill 窗口 arm 16 前向 |
|
||||||
|
| E 协议对照 | 16k probe ×2 前向 | NCCL_PROTO=Simple 变体下的 AR kernel 取证 |
|
||||||
|
|
||||||
|
- `POST /start_profile`，CPU+GPU activities，无 stack；采集前 flush + 微量预热 + 再 flush；bench 每轮换 seed
|
||||||
|
- **with_stack 在本 nightly 病态（两次复现：16k 直接 OOM 僵尸；2k probe flush 挂死）——禁用**
|
||||||
|
|
||||||
|
## 关键结论（详见报告）
|
||||||
|
|
||||||
|
1. **PP 气泡证伪**：chunked prefill 形成天然微批流水，两 stage 同忙 96.0%，经典气泡仅 ~4%
|
||||||
|
2. **prefill 墙 = PP0 链路**：计算 22.8s(65%) + TP4 AR 10.4s(30%)，35.2s 墙钟
|
||||||
|
3. **AR 已近带宽极限**：96MB/3.7ms ⇒ busbw≈39GB/s ≈ PCIe Gen5 实际峰值 ~75%；"_RING_LL" 内核名是 NCCL 2.28 混合 plan 命名，非半带宽 LL 线格式
|
||||||
|
4. **NCCL_PROTO 强制实验证伪快赢**：docker `-e` 到不了 rank 进程（sglang worker spawn 丢弃 NCCL env），AR p50 3.69→3.70ms 不变，e2e 噪声内
|
||||||
|
5. **decode**：步长 21-26ms（MoE GEMM 35% + SendRecv 18%），无 MTP；PP1 AllGather 273 次不在 cuda graph 内
|
||||||
|
6. **128k/90%**：DSA indexer `sm120_fp8_mqa_logits` 占 28%（PP0）/23%（PP1），是其服务的稀疏注意力本体（17%）的 1.6×；后缀有效速率 5.2k tok/s
|
||||||
|
7. **profiler 工程发现**：flush 阻塞全流水线 11s；带 profiler 的 bench 指标作废
|
||||||
|
|
||||||
|
## 第二轮优化实验（2026-09-07）：配置级方向全部证伪
|
||||||
|
|
||||||
|
详见 `results/analysis/optimization_round2_findings.md`。干净 128k/90% 基线补测（TTFT p50 14.70s / TPOT 51.7ms）；
|
||||||
|
chunk 16384 因 MoE 工作区 OOM（3.08GB > 2.11GB 空闲）证伪且收益主体不存在（AR 带宽受限、indexer 功不变）；
|
||||||
|
AR 现成开关扫描全部不适用（quant-comm=NPU、AR 融合=SM90/100 门控、symm-mem/NVLS=NVLink 向）；
|
||||||
|
indexer 28% 属 DSA 算法固有成本（无复用机会、SM120 仅 deepgemm、模型已做 index_topk_freq=4 跨层共享）。
|
||||||
|
**结论：现有配置即本栈配置最优；剩余提升均为开发投入（PP+MTP > AR 重叠/量化内核）。**
|
||||||
|
|
||||||
|
## 优化方向（排序见报告 §4，09-07 更新）
|
||||||
|
|
||||||
|
① prefill AR 通信-计算重叠 / 量化 AR（协议红利已排除）；② PP+MTP（decode 最大杠杆）；
|
||||||
|
③ 128k 场景 DSA indexer 前缀复用/调优；④ chunk 8192→16384 重测；⑤ MoE 胶水融合 + PP1 AG 入图；⑥ decode 尾部。
|
||||||
|
|
||||||
|
## 目录
|
||||||
|
|
||||||
|
```
|
||||||
|
scripts/ deploy_par_605.sh（部署）、profile_cap.sh（A/B/C 编排）、run_phase2.sh（冒烟+干净bench+D+C）、
|
||||||
|
bench_128k_hit90_prof.py（D 驱动+arm）、probe_c*.py（小探针）、
|
||||||
|
pp_align.py（跨rank墙钟对齐+分箱）、nccl_forensics.py / nccl_names.py（NCCL 取证）、census.py（事件普查）
|
||||||
|
results/ analysis/（pp_align/triage/forensics/census 输出）、bench/（干净 bench 与采集期 bench 日志）、ops/（部署与采集日志）
|
||||||
|
```
|
||||||
|
|
||||||
|
trace 原始文件（体积大，不入库）：`60.5:/data/hf_models/pp2tp4_profiles/{A_prefill,B_decode,D_128k,E_nccl_simple}/`
|
||||||
|
（8 rank × 4 场景；A 的 PP0 ~70-77MB gz，其余 ~2-9.5MB）。
|
||||||
|
|
||||||
|
## 复现
|
||||||
|
|
||||||
|
1. `bash scripts/deploy_par_605.sh "<同上参数>" nomtp`（60.5 上原路径 /root/deploy_par.sh）
|
||||||
|
2. 等 /health=200，GSM8K 冒烟
|
||||||
|
3. `bash scripts/profile_cap.sh`（A/B）；`python3 scripts/bench_128k_hit90_prof.py --concurrency 8 --num-requests 8 --run-id <新> --arm-profile <dir> --profile-steps 16`（D）
|
||||||
|
4. `python3 scripts/pp_align.py <trace_dir>`；`python3 scripts/nccl_forensics.py <trace>`；`python3 scripts/nccl_names.py <trace>`
|
||||||
|
5. triage 三表用 zcode skill `llm-torch-profiler-analysis` 的 analyze 脚本（docker cp 进容器跑）
|
||||||
|
|
||||||
|
## 事后状态
|
||||||
|
|
||||||
|
60.5 当前运行原样 TP4PP2（glm53-nvfp4:30000）。恢复 GLM-5.3-Flash kt：
|
||||||
|
`bash /root/deploy_glm53_flash_kt_gpu_graphs.sh`（先 kill -9 本容器 sglang 进程并等显存归零）。
|
||||||
@ -0,0 +1,19 @@
|
|||||||
|
# GLM-5.3-NVFP4 TP4PP2 profile 实验配置（174.1.60.5, 2026-09-06）
|
||||||
|
HOST=174.1.60.5
|
||||||
|
MODEL=/data/hf_models/GLM-5.3-NVFP4
|
||||||
|
IMAGE=lmsysorg/sglang:nightly-dev-20260828-daf63171
|
||||||
|
CONTAINER=glm53-nvfp4
|
||||||
|
PORT=30000
|
||||||
|
|
||||||
|
# 部署（与 09-05 获胜配置逐字一致）
|
||||||
|
DEPLOY="bash /root/deploy_par.sh --tp 4 --pp-size 2 --disable-overlap-schedule --max-prefill-tokens 16384 nomtp"
|
||||||
|
EXTRA_SERVER_DEFAULTS="kv fp8_e4m3 auto; --disable-shared-experts-fusion --moe-runner-backend flashinfer_cutlass --disable-flashinfer-autotune; hicache-ratio 3; mem-frac 0.88; chunked-prefill 8192"
|
||||||
|
|
||||||
|
# 采集
|
||||||
|
PROFILER=sglang /start_profile (CPU+GPU, no stack; with_stack 在本 nightly 病态禁用)
|
||||||
|
TRACE_BASE=/data/hf_models/pp2tp4_profiles
|
||||||
|
SCENARIOS=A_prefill(16x16384->1 cc16), B_decode(16x16384->512 cc16, 60 steps), D_128k(8x131072 cc8, 90% shared, 16 steps), E_nccl_simple(16k probe x2 fwd)
|
||||||
|
|
||||||
|
# 基准（干净无 profiler，seed 每轮更换）
|
||||||
|
BENCH_CC16="python3 -m sglang.bench_serving --backend sglang --dataset-name random-ids --num-prompts 16 --random-input-len 16384 --random-output-len 512 --max-concurrency 16 --temperature 0.0 --flush-cache --warmup-requests 1"
|
||||||
|
BASELINE_CLEAN="input 3350 tok/s; output 104.7 tok/s; TTFT 31.36s; TPOT 50.87ms"
|
||||||
@ -0,0 +1,24 @@
|
|||||||
|
|
||||||
|
== pa-1788706112.6645906-TP-0-PP-0.trace.json.gz ==
|
||||||
|
user_annotation n=851444 total= 864258.2ms
|
||||||
|
Trace n=1 total= 238118.8ms
|
||||||
|
gpu_user_annotation n=3109 total= 99906.2ms
|
||||||
|
kernel n=46252 total= 55472.0ms
|
||||||
|
cpu_op n=2774288 total= 42373.2ms
|
||||||
|
cuda_runtime n=170665 total= 25753.1ms
|
||||||
|
overhead n=2644 total= 4244.5ms
|
||||||
|
cuda_driver n=17680 total= 1394.8ms
|
||||||
|
gpu_memcpy n=1356 total= 597.9ms
|
||||||
|
gpu_memset n=444 total= 0.5ms
|
||||||
|
|
||||||
|
== pa-1788706112.6645906-TP-0-PP-1.trace.json.gz ==
|
||||||
|
user_annotation n=681892 total= 947007.0ms
|
||||||
|
Trace n=1 total= 345621.9ms
|
||||||
|
gpu_user_annotation n=3177 total= 101222.1ms
|
||||||
|
kernel n=47531 total= 34473.8ms
|
||||||
|
cuda_runtime n=190764 total= 26457.3ms
|
||||||
|
cpu_op n=1927437 total= 21763.8ms
|
||||||
|
overhead n=2897 total= 4953.8ms
|
||||||
|
cuda_driver n=17577 total= 1152.1ms
|
||||||
|
gpu_memcpy n=1308 total= 599.9ms
|
||||||
|
gpu_memset n=379 total= 0.4ms
|
||||||
@ -0,0 +1,10 @@
|
|||||||
|
== pa-1788706112.6645906-TP-0-PP-0.trace.json.gz ==
|
||||||
|
AllReduce: n=2686 total=10419.5ms min=3.03 p10=3.66 p50=3.69 p90=3.84 p99=9.39 max=21.46 (ms)
|
||||||
|
span 34.2s; top 2s-buckets (offset_s -> ms): 12s->634ms, 20s->633ms, 14s->632ms, 28s->632ms, 10s->631ms, 26s->629ms
|
||||||
|
SendRecv : n=170 total=22601.8ms min=0.01 p10=0.01 p50=3.33 p90=678.24 p99=714.23 max=718.95 (ms)
|
||||||
|
span 34.2s; top 2s-buckets (offset_s -> ms): 16s->1836ms, 6s->1430ms, 8s->1426ms, 4s->1418ms, 10s->1416ms, 22s->1409ms
|
||||||
|
== pa-1788706112.6645906-TP-0-PP-1.trace.json.gz ==
|
||||||
|
AllReduce: n=2652 total=9968.0ms min=3.04 p10=3.68 p50=3.70 p90=3.84 p99=4.16 max=98.21 (ms)
|
||||||
|
span 34.2s; top 2s-buckets (offset_s -> ms): 12s->606ms, 14s->603ms, 20s->602ms, 22s->601ms, 10s->600ms, 18s->597ms
|
||||||
|
SendRecv : n=170 total=3590.7ms min=0.00 p10=0.00 p50=1.33 p90=83.62 p99=555.66 max=656.03 (ms)
|
||||||
|
span 34.9s; top 2s-buckets (offset_s -> ms): 0s->784ms, 2s->559ms, 16s->342ms, 18s->156ms, 26s->153ms, 6s->151ms
|
||||||
@ -0,0 +1,67 @@
|
|||||||
|
== PP0 rank ==
|
||||||
|
file: /data/hf_models/pp2tp4_profiles/A_prefill/pa-1788706112.6645906-TP-0-PP-0.trace.json.gz
|
||||||
|
baseTimeNanoseconds present: True
|
||||||
|
span 35.176s GPU-busy 34.065s (96.8%) kernels 48052
|
||||||
|
cudaMemcpy/Set GPU time: 598.4ms cudaLaunchKernel calls: 4606 (421.1ms CPU)
|
||||||
|
NCCL total: 33021.3ms (96.9% of busy, 93.9% of span)
|
||||||
|
SendRecv 22601.8ms n=170
|
||||||
|
AllReduce 10419.5ms n=2686
|
||||||
|
top kernels (of GPU-busy):
|
||||||
|
66.3% 22601.8ms n=170 ncclDevKernel_SendRecv(ncclDevKernelArgsStorage
|
||||||
|
30.6% 10419.5ms n=2686 ncclDevKernel_AllReduce_Sum_bf16_RING_LL(ncclDevKernelArgsStorage
|
||||||
|
24.9% 8483.5ms n=1326 void sparse_mla_prefill_kernel
|
||||||
|
16.6% 5661.4ms n=11322 void cutlass::Kernel2
|
||||||
|
11.9% 4044.2ms n=2448 _ZN7cutlass13device_kernelINS_4gemm6kernel13GemmUniversalINS1_17GroupProblemShap
|
||||||
|
2.6% 877.3ms n=1224 void tensorrt_llm::kernels::cutlass_kernels::finalizeMoeRoutingKernel
|
||||||
|
2.4% 819.2ms n=408 void deep_gemm::sm120_fp8_mqa_logits
|
||||||
|
2.0% 677.7ms n=2618 kernel_cutlass_kernel_flashinfernormkernelsfused_add_rmsnormFusedAddRMSNormKerne
|
||||||
|
1.8% 612.9ms n=1224 void tensorrt_llm::kernels::cutlass_kernels::expandInputRowsKernel
|
||||||
|
1.8% 598.4ms n=1800 ::
|
||||||
|
0.8% 271.0ms n=1326 void sglang::concat_mla_absorb_q_kernel
|
||||||
|
0.8% 261.2ms n=2047 void at::native::vectorized_elementwise_kernel
|
||||||
|
0.7% 236.9ms n=1224 void tensorrt_llm::kernels::cutlass_kernels::doActivationKernel
|
||||||
|
0.6% 221.1ms n=408 (anonymous namespace)::topk_transform_prefill_kernel((anonymous namespace)::Fast
|
||||||
|
0.2% 54.1ms n=2652 kernel_cutlass_kernel_flashinfernormkernelsrmsnormRMSNormKernel_object_at__tenso
|
||||||
|
forwards: 5 (gap>4.0ms split)
|
||||||
|
wall ms: mean 6831.0 p50 1977.3 min 21.5 max 17604.3
|
||||||
|
small forwards (<30% of p50): 1 (warmup/mixed steps)
|
||||||
|
GPU-covered fraction inside main forwards: mean 99.4%
|
||||||
|
inter-forward gap ms: mean 255.2 p50 243.9 p90 355.2 max 355.2
|
||||||
|
|
||||||
|
== PP1 rank ==
|
||||||
|
file: /data/hf_models/pp2tp4_profiles/A_prefill/pa-1788706112.6645906-TP-0-PP-1.trace.json.gz
|
||||||
|
baseTimeNanoseconds present: True
|
||||||
|
span 35.181s GPU-busy 34.609s (98.4%) kernels 49218
|
||||||
|
cudaMemcpy/Set GPU time: 600.3ms cudaLaunchKernel calls: 5240 (1209.2ms CPU)
|
||||||
|
NCCL total: 13796.9ms (39.9% of busy, 39.2% of span)
|
||||||
|
AllReduce 9968.0ms n=2652
|
||||||
|
SendRecv 3590.7ms n=170
|
||||||
|
AG_RS 238.2ms n=136
|
||||||
|
top kernels (of GPU-busy):
|
||||||
|
28.8% 9968.0ms n=2652 ncclDevKernel_AllReduce_Sum_bf16_RING_LL(ncclDevKernelArgsStorage
|
||||||
|
24.6% 8505.8ms n=1326 void sparse_mla_prefill_kernel
|
||||||
|
15.4% 5323.7ms n=11235 void cutlass::Kernel2
|
||||||
|
10.4% 3590.7ms n=170 ncclDevKernel_SendRecv(ncclDevKernelArgsStorage
|
||||||
|
8.5% 2930.6ms n=2652 _ZN7cutlass13device_kernelINS_4gemm6kernel13GemmUniversalINS1_17GroupProblemShap
|
||||||
|
2.7% 949.5ms n=1326 void tensorrt_llm::kernels::cutlass_kernels::finalizeMoeRoutingKernel
|
||||||
|
2.0% 697.6ms n=2686 kernel_cutlass_kernel_flashinfernormkernelsfused_add_rmsnormFusedAddRMSNormKerne
|
||||||
|
1.9% 668.9ms n=1326 void tensorrt_llm::kernels::cutlass_kernels::expandInputRowsKernel
|
||||||
|
1.7% 600.3ms n=1687 ::
|
||||||
|
1.0% 355.1ms n=324 void deep_gemm::sm120_fp8_mqa_logits
|
||||||
|
0.8% 278.8ms n=2169 void at::native::vectorized_elementwise_kernel
|
||||||
|
0.8% 270.7ms n=1326 void sglang::concat_mla_absorb_q_kernel
|
||||||
|
0.7% 245.0ms n=1326 void tensorrt_llm::kernels::cutlass_kernels::doActivationKernel
|
||||||
|
0.7% 238.2ms n=136 ncclDevKernel_AllGather_RING_LL(ncclDevKernelArgsStorage
|
||||||
|
0.5% 171.8ms n=324 (anonymous namespace)::topk_transform_prefill_kernel((anonymous namespace)::Fast
|
||||||
|
forwards: 10 (gap>4.0ms split)
|
||||||
|
wall ms: mean 3482.4 p50 3786.0 min 4.3 max 9922.5
|
||||||
|
small forwards (<30% of p50): 4 (warmup/mixed steps)
|
||||||
|
GPU-covered fraction inside main forwards: mean 99.5%
|
||||||
|
inter-forward gap ms: mean 39.7 p50 7.5 p90 277.8 max 277.8
|
||||||
|
|
||||||
|
== cross-stage bins (1.0ms, total 35.18s) ==
|
||||||
|
both busy : 33.77s 96.0%
|
||||||
|
only PP0 : 0.38s 1.1% <- PP1 bubble
|
||||||
|
only PP1 : 1.03s 2.9% <- PP0 bubble
|
||||||
|
neither : 0.01s 0.0% <- sched/CPU idle
|
||||||
|
pipeline efficiency: both/(any busy) = 96.0%
|
||||||
@ -0,0 +1,32 @@
|
|||||||
|
Triage View
|
||||||
|
Mode: single-trace
|
||||||
|
Framework: SGLang
|
||||||
|
Input traces: /data/hf_models/pp2tp4_profiles/A_prefill/pa-1788706112.6645906-TP-0-PP-0.trace.json.gz
|
||||||
|
|
||||||
|
Kernel Table
|
||||||
|
##### extend/prefill
|
||||||
|
| Kernel | Category | GPU time | Share | Launches | Python location (site share) | CPU op |
|
||||||
|
| --- | --- | ---: | ---: | ---: | --- | --- |
|
||||||
|
| ncclDevKernel_SendRecv(ncclDevKernelArgsStorage<4096ul>) | communication | 22601.84 ms | 40.3% | 170 | unresolved | cuLaunchKernelEx |
|
||||||
|
| ncclDevKernel_AllReduce_Sum_bf16_RING_LL(ncclDevKernelArgsStorage<4096ul>) | communication | 10419.48 ms | 18.6% | 2686 | unresolved | cuLaunchKernelEx |
|
||||||
|
| void sparse_mla_prefill_kernel<(ModelType)2, (ComputeMode)0, 16, 2048, 64> | memory | 8483.48 ms | 15.1% | 1326 | unresolved | cudaLaunchKernelExC |
|
||||||
|
| _ZN7cutlass13device_kernelINS_4gemm6kernel13GemmUniversalINS1_17GroupProblemShapeIN4cute5tupleIJlllEEEEENS1_10collective13CollectiveMmaINS1_47MainloopSm120ArrayTmaWarpSpecializedBlockScaledILi2ELi3ENS6_IJNS5_1CILi1EEESD_SD_EEENS1_59KernelPtrArrayTmaWarpSpecializedCooperativeBlockScaledSm120ILi3EEEEENS6_IJNSC_ILi128EEESI_NSC_ILi256EEEEEENS6_IJNS_12float_e2m1_tENS_13float_ue4m3_tEEEENS6_IJPNS6_IJlSD_NSC_ILi0EEEEEEPNS5_6LayoutINS6_IJNS6_IJNS6_IJNSC_ILi32EEENSC_ILi4EEEEEEiEEENS6_IJNS6_IJNSC_ILi16EEEST_EEEiEEENS6_IJSD_iEEEEEENS6_IJSY_NS6_IJNS6_IJSO_SD_EEENSC_ILi512EEEEEENS6_IJSO_iEEEEEEEEEEESN_S18_NS5_8TiledMMAINS5_8MMA_AtomIJNS5_5SM12011BLOCKSCALED19SM120_16x8x64_TN_VSISL_SL_fSM_Li16EEEEEENSR_INS6_IJST_NSC_ILi2EEESD_EEENS6_IJSD_ST_SO_EEEEENS6_IJSI_NSR_INS6_IJNSC_ILi8EEES1G_S1G_EEENS6_IJSD_SW_S1K_EEEEENSC_ILi64EEEEEEEENS6_IJNS5_13SM90_TMA_LOADES1R_EEENS6_IJNS5_14ComposedLayoutINS5_7SwizzleILi3ELi4ELi3EEENS5_18smem_ptr_flag_bitsILi4EEENSR_INS6_IJS1K_SJ_EEENS6_IJSJ_SD_EEEEEEENSR_INS6_IJNS6_IJSU_SD_EEENS6_IJSX_SD_ST_EEEEEENS6_IJNS6_IJSX_S12_EEENS6_IJS11_ST_S12_EEEEEEEEEEENS6_IJNS5_9Copy_AtomIJNS5_17SM75_U32x4_LDSM_NENS_15integer_subbyteILi4ELb0EEEEEENS2A_IJNS5_13UniversalCopyISM_SM_EESM_EEEEEENS5_8identityES1S_S29_S2I_S2J_EENS_8epilogue10collective18CollectiveEpilogueINS2L_30Sm90PtrArrayTmaWarpSpecializedILi3ELi2ELi4ELb1ELb0ELi2EEEJSK_NS6_IJS1O_SS_EEENS_10bfloat16_tESQ_S2R_SQ_NS2L_6fusion15FusionCallbacksINS2L_31Sm120PtrArrayTmaWarpSpecializedILi3ELi2ELi4ELb1ELb0ELi2EEENS2S_17LinearCombinationIS2R_fS2R_fLNS_15FloatRoundStyleE2EEESK_S2Q_JEEES1R_NS1T_INS1U_ILi2ELi4ELi3EEENS1W_ILi16EEENSR_INS6_IJS1K_SS_EEENS6_IJSS_SD_EEEEEEENS5_17SM75_U32x2_LDSM_NENS5_14SM90_TMA_STOREES35_NS5_17SM90_U32x2_STSM_NENS2A_IJS38_NS_6half_tEEEEvEEEvvEEEEvNT_6ParamsE | moe | 4044.20 ms | 7.2% | 2448 | unresolved | cudaLaunchKernelExC |
|
||||||
|
| void cutlass::Kernel2<cutlass_80_tensorop_bf16_s16816gemm_relu_bf16_64x256_32x4_tn_align8> | gemm | 3380.00 ms | 6.0% | 3180 | unresolved | cuLaunchKernel |
|
||||||
|
| void cutlass::Kernel2<cutlass_80_tensorop_bf16_s16816gemm_relu_bf16_256x128_32x3_tn_align8> | gemm | 1422.33 ms | 2.5% | 5559 | unresolved | cuLaunchKernel |
|
||||||
|
| void tensorrt_llm::kernels::cutlass_kernels::finalizeMoeRoutingKernel<__nv_bfloat16, __nv_bfloat16, __nv_bfloat16, (tensorrt_llm::kernels::cutlass_kernels::ScaleMode)1> | moe | 877.25 ms | 1.6% | 1224 | unresolved | cudaLaunchKernelExC |
|
||||||
|
| void deep_gemm::sm120_fp8_mqa_logits<32u, 128u, false, 4u, 128u, 2u, 3u, 156u, 128u, 256u, float> | gemm | 819.24 ms | 1.5% | 408 | unresolved | cuLaunchKernelEx |
|
||||||
|
| kernel_cutlass_kernel_flashinfernormkernelsfused_add_rmsnormFusedAddRMSNormKernel_object_at__tensorptrbf16gmemalign128oi64614461441_tensorptrbf16gmemalign128oi64614461441_tensorptrbf16gme_0 | gemm | 677.69 ms | 1.2% | 2618 | unresolved | cudaLaunchKernelExC |
|
||||||
|
| void tensorrt_llm::kernels::cutlass_kernels::expandInputRowsKernel<__nv_bfloat16, __nv_fp4_e2m1, (tensorrt_llm::kernels::cutlass_kernels::TmaWarpSpecializedGroupedGemmInput::FpXBlockScalingType)1, false, false, std::integral_constant<bool, false> > | gemm | 612.91 ms | 1.1% | 1224 | unresolved | cudaLaunchKernelExC |
|
||||||
|
| Memcpy DtoH (Device -> Pinned) | memory | 568.62 ms | 1.0% | 166 | unresolved | - |
|
||||||
|
|
||||||
|
Overlap Opportunity Table
|
||||||
|
| Priority | Verdict | Kernel | Python scope | Formal signal | Dep risk | Recommendation |
|
||||||
|
| --- | --- | --- | --- | --- | --- | --- |
|
||||||
|
| - | - | No rows cleared the 1.0% reporting bar. Use mapping/formal mode for overlap attribution. | - | - | - | - |
|
||||||
|
|
||||||
|
Fuse Opportunity Table
|
||||||
|
##### extend/prefill
|
||||||
|
| Pattern | Confidence | Related GPU time | Share | Evidence kernels | Current kernel Python location | Candidate fused Python path | Rationale |
|
||||||
|
| --- | --- | ---: | ---: | --- | --- | --- | --- |
|
||||||
|
| NSA fused quantize + indexed K-cache store | Confirmed | 4119.01 ms | 7.3% | _ZN7cutlass13device_kernelINS_4gemm6kernel13GemmUniversalINS1_17GroupProblemShapeIN4cute5tupleIJlllEEEEENS1_10collective13CollectiveMmaINS1_47MainloopSm120ArrayTmaWarpSpecializedBlockScaledILi2ELi3ENS6_IJNS5_1CILi1EEESD_SD_EEENS1_59KernelPtrArrayTmaWarpSpecializedCooperativeBlockScaledSm120ILi3EEEEENS6_IJNSC_ILi128EEESI_NSC_ILi256EEEEEENS6_IJNS_12float_e2m1_tENS_13float_ue4m3_tEEEENS6_IJPNS6_IJlSD_NSC_ILi0EEEEEEPNS5_6LayoutINS6_IJNS6_IJNS6_IJNSC_ILi32EEENSC_ILi4EEEEEEiEEENS6_IJNS6_IJNSC_ILi16EEEST_EEEiEEENS6_IJSD_iEEEEEENS6_IJSY_NS6_IJNS6_IJSO_SD_EEENSC_ILi512EEEEEENS6_IJSO_iEEEEEEEEEEESN_S18_NS5_8TiledMMAINS5_8MMA_AtomIJNS5_5SM12011BLOCKSCALED19SM120_16x8x64_TN_VSISL_SL_fSM_Li16EEEEEENSR_INS6_IJST_NSC_ILi2EEESD_EEENS6_IJSD_ST_SO_EEEEENS6_IJSI_NSR_INS6_IJNSC_ILi8EEES1G_S1G_EEENS6_IJSD_SW_S1K_EEEEENSC_ILi64EEEEEEEENS6_IJNS5_13SM90_TMA_LOADES1R_EEENS6_IJNS5_14ComposedLayoutINS5_7SwizzleILi3ELi4ELi3EEENS5_18smem_ptr_flag_bitsILi4EEENSR_INS6_IJS1K_SJ_EEENS6_IJSJ_SD_EEEEEEENSR_INS6_IJNS6_IJSU_SD_EEENS6_IJSX_SD_ST_EEEEEENS6_IJNS6_IJSX_S12_EEENS6_IJS11_ST_S12_EEEEEEEEEEENS6_IJNS5_9Copy_AtomIJNS5_17SM75_U32x4_LDSM_NENS_15integer_subbyteILi4ELb0EEEEEENS2A_IJNS5_13UniversalCopyISM_SM_EESM_EEEEEENS5_8identityES1S_S29_S2I_S2J_EENS_8epilogue10collective18CollectiveEpilogueINS2L_30Sm90PtrArrayTmaWarpSpecializedILi3ELi2ELi4ELb1ELb0ELi2EEEJSK_NS6_IJS1O_SS_EEENS_10bfloat16_tESQ_S2R_SQ_NS2L_6fusion15FusionCallbacksINS2L_31Sm120PtrArrayTmaWarpSpecializedILi3ELi2ELi4ELb1ELb0ELi2EEENS2S_17LinearCombinationIS2R_fS2R_fLNS_15FloatRoundStyleE2EEESK_S2Q_JEEES1R_NS1T_INS1U_ILi2ELi4ELi3EEENS1W_ILi16EEENSR_INS6_IJS1K_SS_EEENS6_IJSS_SD_EEEEEEENS5_17SM75_U32x2_LDSM_NENS5_14SM90_TMA_STOREES35_NS5_17SM90_U32x2_STSM_NENS2A_IJS38_NS_6half_tEEEEvEEEvvEEEEvNT_6ParamsE (7.2%) | unresolved | python/sglang/kernels/ops/attention/fused_store_index_cache.py<br>python/sglang/srt/layers/attention/nsa/nsa_indexer.py | Split kernels in this family take 7.3% of GPU time. This tree already has a matching path. NSA already has a fused quantize-and-indexed-store kernel family. |
|
||||||
|
| Fused residual add + RMSNorm | Confirmed | 677.69 ms | 1.2% | kernel_cutlass_kernel_flashinfernormkernelsfused_add_rmsnormFusedAddRMSNormKernel_object_at__tensorptrbf16gmemalign128oi64614461441_tensorptrbf16gmemalign128oi64614461441_tensorptrbf16gme_0 (1.2%) | unresolved | python/sglang/srt/layers/layernorm.py<br>python/sglang/srt/layers/quantization/modelslim/modelslim.py | `Fused residual add + RMSNorm` is present in this trace (1.2% related GPU time). Residual add plus RMSNorm already has fused implementations across several backends. |
|
||||||
@ -0,0 +1,33 @@
|
|||||||
|
Triage View
|
||||||
|
Mode: single-trace
|
||||||
|
Framework: SGLang
|
||||||
|
Input traces: /data/hf_models/pp2tp4_profiles/A_prefill/pa-1788706112.6645906-TP-0-PP-1.trace.json.gz
|
||||||
|
|
||||||
|
Kernel Table
|
||||||
|
##### extend/prefill
|
||||||
|
| Kernel | Category | GPU time | Share | Launches | Python location (site share) | CPU op |
|
||||||
|
| --- | --- | ---: | ---: | ---: | --- | --- |
|
||||||
|
| ncclDevKernel_AllReduce_Sum_bf16_RING_LL(ncclDevKernelArgsStorage<4096ul>) | communication | 9968.03 ms | 28.4% | 2652 | unresolved | cuLaunchKernelEx |
|
||||||
|
| void sparse_mla_prefill_kernel<(ModelType)2, (ComputeMode)0, 16, 2048, 64> | memory | 8505.81 ms | 24.3% | 1326 | unresolved | cudaLaunchKernelExC |
|
||||||
|
| ncclDevKernel_SendRecv(ncclDevKernelArgsStorage<4096ul>) | communication | 3590.66 ms | 10.2% | 170 | unresolved | cuLaunchKernelEx |
|
||||||
|
| void cutlass::Kernel2<cutlass_80_tensorop_bf16_s16816gemm_relu_bf16_64x256_32x4_tn_align8> | gemm | 3013.03 ms | 8.6% | 2880 | unresolved | cuLaunchKernel |
|
||||||
|
| _ZN7cutlass13device_kernelINS_4gemm6kernel13GemmUniversalINS1_17GroupProblemShapeIN4cute5tupleIJlllEEEEENS1_10collective13CollectiveMmaINS1_47MainloopSm120ArrayTmaWarpSpecializedBlockScaledILi2ELi3ENS6_IJNS5_1CILi1EEESD_SD_EEENS1_59KernelPtrArrayTmaWarpSpecializedCooperativeBlockScaledSm120ILi3EEEEENS6_IJNSC_ILi128EEESI_NSC_ILi256EEEEEENS6_IJNS_12float_e2m1_tENS_13float_ue4m3_tEEEENS6_IJPNS6_IJlSD_NSC_ILi0EEEEEEPNS5_6LayoutINS6_IJNS6_IJNS6_IJNSC_ILi32EEENSC_ILi4EEEEEEiEEENS6_IJNS6_IJNSC_ILi16EEEST_EEEiEEENS6_IJSD_iEEEEEENS6_IJSY_NS6_IJNS6_IJSO_SD_EEENSC_ILi512EEEEEENS6_IJSO_iEEEEEEEEEEESN_S18_NS5_8TiledMMAINS5_8MMA_AtomIJNS5_5SM12011BLOCKSCALED19SM120_16x8x64_TN_VSISL_SL_fSM_Li16EEEEEENSR_INS6_IJST_NSC_ILi2EEESD_EEENS6_IJSD_ST_SO_EEEEENS6_IJSI_NSR_INS6_IJNSC_ILi8EEES1G_S1G_EEENS6_IJSD_SW_S1K_EEEEENSC_ILi64EEEEEEEENS6_IJNS5_13SM90_TMA_LOADES1R_EEENS6_IJNS5_14ComposedLayoutINS5_7SwizzleILi3ELi4ELi3EEENS5_18smem_ptr_flag_bitsILi4EEENSR_INS6_IJS1K_SJ_EEENS6_IJSJ_SD_EEEEEEENSR_INS6_IJNS6_IJSU_SD_EEENS6_IJSX_SD_ST_EEEEEENS6_IJNS6_IJSX_S12_EEENS6_IJS11_ST_S12_EEEEEEEEEEENS6_IJNS5_9Copy_AtomIJNS5_17SM75_U32x4_LDSM_NENS_15integer_subbyteILi4ELb0EEEEEENS2A_IJNS5_13UniversalCopyISM_SM_EESM_EEEEEENS5_8identityES1S_S29_S2I_S2J_EENS_8epilogue10collective18CollectiveEpilogueINS2L_30Sm90PtrArrayTmaWarpSpecializedILi3ELi2ELi4ELb1ELb0ELi2EEEJSK_NS6_IJS1O_SS_EEENS_10bfloat16_tESQ_S2R_SQ_NS2L_6fusion15FusionCallbacksINS2L_31Sm120PtrArrayTmaWarpSpecializedILi3ELi2ELi4ELb1ELb0ELi2EEENS2S_17LinearCombinationIS2R_fS2R_fLNS_15FloatRoundStyleE2EEESK_S2Q_JEEES1R_NS1T_INS1U_ILi2ELi4ELi3EEENS1W_ILi16EEENSR_INS6_IJS1K_SS_EEENS6_IJSS_SD_EEEEEEENS5_17SM75_U32x2_LDSM_NENS5_14SM90_TMA_STOREES35_NS5_17SM90_U32x2_STSM_NENS2A_IJS38_NS_6half_tEEEEvEEEvvEEEEvNT_6ParamsE | moe | 2930.57 ms | 8.4% | 2652 | unresolved | cudaLaunchKernelExC |
|
||||||
|
| void cutlass::Kernel2<cutlass_80_tensorop_bf16_s16816gemm_relu_bf16_256x128_32x3_tn_align8> | gemm | 1394.28 ms | 4.0% | 5562 | unresolved | cuLaunchKernel |
|
||||||
|
| void tensorrt_llm::kernels::cutlass_kernels::finalizeMoeRoutingKernel<__nv_bfloat16, __nv_bfloat16, __nv_bfloat16, (tensorrt_llm::kernels::cutlass_kernels::ScaleMode)1> | moe | 949.47 ms | 2.7% | 1326 | unresolved | cudaLaunchKernelExC |
|
||||||
|
| kernel_cutlass_kernel_flashinfernormkernelsfused_add_rmsnormFusedAddRMSNormKernel_object_at__tensorptrbf16gmemalign128oi64614461441_tensorptrbf16gmemalign128oi64614461441_tensorptrbf16gme_0 | gemm | 697.65 ms | 2.0% | 2686 | unresolved | cudaLaunchKernelExC |
|
||||||
|
| void tensorrt_llm::kernels::cutlass_kernels::expandInputRowsKernel<__nv_bfloat16, __nv_fp4_e2m1, (tensorrt_llm::kernels::cutlass_kernels::TmaWarpSpecializedGroupedGemmInput::FpXBlockScalingType)1, false, false, std::integral_constant<bool, false> > | gemm | 668.88 ms | 1.9% | 1326 | unresolved | cudaLaunchKernelExC |
|
||||||
|
| void cutlass::Kernel2<cutlass_80_tensorop_bf16_s16816gemm_relu_bf16_64x128_64x3_tn_align8> | gemm | 597.53 ms | 1.7% | 1326 | unresolved | cuLaunchKernel |
|
||||||
|
| Memcpy DtoH (Device -> Pinned) | memory | 577.18 ms | 1.6% | 166 | unresolved | - |
|
||||||
|
| void deep_gemm::sm120_fp8_mqa_logits<32u, 128u, false, 4u, 128u, 2u, 3u, 156u, 128u, 256u, float> | gemm | 355.07 ms | 1.0% | 324 | unresolved | cuLaunchKernelEx |
|
||||||
|
|
||||||
|
Overlap Opportunity Table
|
||||||
|
| Priority | Verdict | Kernel | Python scope | Formal signal | Dep risk | Recommendation |
|
||||||
|
| --- | --- | --- | --- | --- | --- | --- |
|
||||||
|
| - | - | No rows cleared the 1.0% reporting bar. Use mapping/formal mode for overlap attribution. | - | - | - | - |
|
||||||
|
|
||||||
|
Fuse Opportunity Table
|
||||||
|
##### extend/prefill
|
||||||
|
| Pattern | Confidence | Related GPU time | Share | Evidence kernels | Current kernel Python location | Candidate fused Python path | Rationale |
|
||||||
|
| --- | --- | ---: | ---: | --- | --- | --- | --- |
|
||||||
|
| NSA fused quantize + indexed K-cache store | Confirmed | 3000.27 ms | 8.6% | _ZN7cutlass13device_kernelINS_4gemm6kernel13GemmUniversalINS1_17GroupProblemShapeIN4cute5tupleIJlllEEEEENS1_10collective13CollectiveMmaINS1_47MainloopSm120ArrayTmaWarpSpecializedBlockScaledILi2ELi3ENS6_IJNS5_1CILi1EEESD_SD_EEENS1_59KernelPtrArrayTmaWarpSpecializedCooperativeBlockScaledSm120ILi3EEEEENS6_IJNSC_ILi128EEESI_NSC_ILi256EEEEEENS6_IJNS_12float_e2m1_tENS_13float_ue4m3_tEEEENS6_IJPNS6_IJlSD_NSC_ILi0EEEEEEPNS5_6LayoutINS6_IJNS6_IJNS6_IJNSC_ILi32EEENSC_ILi4EEEEEEiEEENS6_IJNS6_IJNSC_ILi16EEEST_EEEiEEENS6_IJSD_iEEEEEENS6_IJSY_NS6_IJNS6_IJSO_SD_EEENSC_ILi512EEEEEENS6_IJSO_iEEEEEEEEEEESN_S18_NS5_8TiledMMAINS5_8MMA_AtomIJNS5_5SM12011BLOCKSCALED19SM120_16x8x64_TN_VSISL_SL_fSM_Li16EEEEEENSR_INS6_IJST_NSC_ILi2EEESD_EEENS6_IJSD_ST_SO_EEEEENS6_IJSI_NSR_INS6_IJNSC_ILi8EEES1G_S1G_EEENS6_IJSD_SW_S1K_EEEEENSC_ILi64EEEEEEEENS6_IJNS5_13SM90_TMA_LOADES1R_EEENS6_IJNS5_14ComposedLayoutINS5_7SwizzleILi3ELi4ELi3EEENS5_18smem_ptr_flag_bitsILi4EEENSR_INS6_IJS1K_SJ_EEENS6_IJSJ_SD_EEEEEEENSR_INS6_IJNS6_IJSU_SD_EEENS6_IJSX_SD_ST_EEEEEENS6_IJNS6_IJSX_S12_EEENS6_IJS11_ST_S12_EEEEEEEEEEENS6_IJNS5_9Copy_AtomIJNS5_17SM75_U32x4_LDSM_NENS_15integer_subbyteILi4ELb0EEEEEENS2A_IJNS5_13UniversalCopyISM_SM_EESM_EEEEEENS5_8identityES1S_S29_S2I_S2J_EENS_8epilogue10collective18CollectiveEpilogueINS2L_30Sm90PtrArrayTmaWarpSpecializedILi3ELi2ELi4ELb1ELb0ELi2EEEJSK_NS6_IJS1O_SS_EEENS_10bfloat16_tESQ_S2R_SQ_NS2L_6fusion15FusionCallbacksINS2L_31Sm120PtrArrayTmaWarpSpecializedILi3ELi2ELi4ELb1ELb0ELi2EEENS2S_17LinearCombinationIS2R_fS2R_fLNS_15FloatRoundStyleE2EEESK_S2Q_JEEES1R_NS1T_INS1U_ILi2ELi4ELi3EEENS1W_ILi16EEENSR_INS6_IJS1K_SS_EEENS6_IJSS_SD_EEEEEEENS5_17SM75_U32x2_LDSM_NENS5_14SM90_TMA_STOREES35_NS5_17SM90_U32x2_STSM_NENS2A_IJS38_NS_6half_tEEEEvEEEvvEEEEvNT_6ParamsE (8.4%) | unresolved | python/sglang/kernels/ops/attention/fused_store_index_cache.py<br>python/sglang/srt/layers/attention/nsa/nsa_indexer.py | Split kernels in this family take 8.6% of GPU time. This tree already has a matching path. NSA already has a fused quantize-and-indexed-store kernel family. |
|
||||||
|
| Fused residual add + RMSNorm | Confirmed | 697.65 ms | 2.0% | kernel_cutlass_kernel_flashinfernormkernelsfused_add_rmsnormFusedAddRMSNormKernel_object_at__tensorptrbf16gmemalign128oi64614461441_tensorptrbf16gmemalign128oi64614461441_tensorptrbf16gme_0 (2.0%) | unresolved | python/sglang/srt/layers/layernorm.py<br>python/sglang/srt/layers/quantization/modelslim/modelslim.py | `Fused residual add + RMSNorm` is present in this trace (2.0% related GPU time). Residual add plus RMSNorm already has fused implementations across several backends. |
|
||||||
@ -0,0 +1,68 @@
|
|||||||
|
== PP0 rank ==
|
||||||
|
file: /data/hf_models/pp2tp4_profiles/B_decode/pb-1788706647.2466695-TP-0-PP-0.trace.json.gz
|
||||||
|
baseTimeNanoseconds present: True
|
||||||
|
span 1.579s GPU-busy 1.167s (73.9%) kernels 80049
|
||||||
|
cudaMemcpy/Set GPU time: 0.3ms cudaLaunchKernel calls: 757 (10.0ms CPU)
|
||||||
|
NCCL total: 324.5ms (27.8% of busy, 20.5% of span)
|
||||||
|
SendRecv 239.3ms n=300
|
||||||
|
AllReduce 85.0ms n=4740
|
||||||
|
AG_RS 0.3ms n=30
|
||||||
|
top kernels (of GPU-busy):
|
||||||
|
39.3% 459.2ms n=4320 _ZN7cutlass13device_kernelINS_4gemm6kernel13GemmUniversalINS1_17GroupProblemShap
|
||||||
|
20.5% 239.3ms n=300 ncclDevKernel_SendRecv(ncclDevKernelArgsStorage
|
||||||
|
18.1% 211.7ms n=13140 void cutlass::Kernel2
|
||||||
|
7.3% 85.0ms n=4740 ncclDevKernel_AllReduce_Sum_bf16_RING_LL(ncclDevKernelArgsStorage
|
||||||
|
5.9% 68.4ms n=2340 kernel_cutlass__dsv3_fused_a_gemm_kernel_tensorptri32gmemo2624307230721_tensorpt
|
||||||
|
4.1% 47.8ms n=2340 void flashinfer::sparse_mla_sm120::sparse_mla_decode_dsv3_2_kernel
|
||||||
|
3.0% 35.1ms n=2340 kernel_cutlass__dsv3_fused_a_gemm_kernel_tensorptri32gmemo4096102410241_tensorpt
|
||||||
|
2.2% 25.8ms n=2160 void tensorrt_llm::kernels::cutlass_kernels::finalizeMoeRoutingKernel
|
||||||
|
1.4% 16.1ms n=4620 kernel_cutlass_kernel_flashinfernormkernelsfused_add_rmsnormFusedAddRMSNormKerne
|
||||||
|
1.3% 14.7ms n=2160 void sglang::router_gemm_kernel
|
||||||
|
1.0% 11.7ms n=720 void deep_gemm::sm120_fp8_paged_mqa_logits
|
||||||
|
0.9% 10.5ms n=2160 void tensorrt_llm::kernels::cutlass_kernels::expandInputRowsKernel
|
||||||
|
0.8% 9.4ms n=4680 kernel_cutlass_kernel_flashinfernormkernelsrmsnormRMSNormKernel_object_at__tenso
|
||||||
|
0.8% 9.0ms n=2160 void tensorrt_llm::kernels::cutlass_kernels::doActivationKernel
|
||||||
|
0.7% 8.1ms n=2160 void tensorrt_llm::kernels::cutlass_kernels::fusedBuildExpertMapsSortFirstTokenK
|
||||||
|
forwards: 12 (gap>2.0ms split)
|
||||||
|
wall ms: mean 124.1 p50 54.0 min 10.7 max 386.7
|
||||||
|
small forwards (<30% of p50): 1 (warmup/mixed steps)
|
||||||
|
GPU-covered fraction inside main forwards: mean 76.5%
|
||||||
|
inter-forward gap ms: mean 8.2 p50 2.3 p90 4.4 max 64.0
|
||||||
|
|
||||||
|
== PP1 rank ==
|
||||||
|
file: /data/hf_models/pp2tp4_profiles/B_decode/pb-1788706647.2466695-TP-0-PP-1.trace.json.gz
|
||||||
|
baseTimeNanoseconds present: True
|
||||||
|
span 12.096s GPU-busy 1.300s (10.7%) kernels 81315
|
||||||
|
cudaMemcpy/Set GPU time: 0.2ms cudaLaunchKernel calls: 997 (11.2ms CPU)
|
||||||
|
NCCL total: 255.8ms (19.7% of busy, 2.1% of span)
|
||||||
|
AllReduce 112.4ms n=4680
|
||||||
|
AG_RS 85.7ms n=273
|
||||||
|
SendRecv 57.7ms n=303
|
||||||
|
top kernels (of GPU-busy):
|
||||||
|
42.4% 551.1ms n=4680 _ZN7cutlass13device_kernelINS_4gemm6kernel13GemmUniversalINS1_17GroupProblemShap
|
||||||
|
16.7% 216.9ms n=12840 void cutlass::Kernel2
|
||||||
|
8.6% 112.4ms n=4680 ncclDevKernel_AllReduce_Sum_bf16_RING_LL(ncclDevKernelArgsStorage
|
||||||
|
6.6% 85.7ms n=273 ncclDevKernel_AllGather_RING_LL(ncclDevKernelArgsStorage
|
||||||
|
5.3% 68.5ms n=2340 kernel_cutlass__dsv3_fused_a_gemm_kernel_tensorptri32gmemo2624307230721_tensorpt
|
||||||
|
4.4% 57.7ms n=303 ncclDevKernel_SendRecv(ncclDevKernelArgsStorage
|
||||||
|
3.7% 47.7ms n=2340 void flashinfer::sparse_mla_sm120::sparse_mla_decode_dsv3_2_kernel
|
||||||
|
2.7% 35.3ms n=2340 kernel_cutlass__dsv3_fused_a_gemm_kernel_tensorptri32gmemo4096102410241_tensorpt
|
||||||
|
2.2% 28.0ms n=2340 void tensorrt_llm::kernels::cutlass_kernels::finalizeMoeRoutingKernel
|
||||||
|
1.3% 16.5ms n=4740 kernel_cutlass_kernel_flashinfernormkernelsfused_add_rmsnormFusedAddRMSNormKerne
|
||||||
|
1.2% 16.0ms n=2340 void sglang::router_gemm_kernel
|
||||||
|
0.9% 11.4ms n=2340 void tensorrt_llm::kernels::cutlass_kernels::expandInputRowsKernel
|
||||||
|
0.8% 10.0ms n=2340 void tensorrt_llm::kernels::cutlass_kernels::doActivationKernel
|
||||||
|
0.7% 9.3ms n=4680 kernel_cutlass_kernel_flashinfernormkernelsrmsnormRMSNormKernel_object_at__tenso
|
||||||
|
0.7% 8.8ms n=2340 void tensorrt_llm::kernels::cutlass_kernels::fusedBuildExpertMapsSortFirstTokenK
|
||||||
|
forwards: 11 (gap>2.0ms split)
|
||||||
|
wall ms: mean 142.6 p50 56.0 min 5.6 max 346.6
|
||||||
|
small forwards (<30% of p50): 2 (warmup/mixed steps)
|
||||||
|
GPU-covered fraction inside main forwards: mean 84.5%
|
||||||
|
inter-forward gap ms: mean 1052.8 p50 2.6 p90 10503.3 max 10503.3
|
||||||
|
|
||||||
|
== cross-stage bins (1.0ms, total 12.17s) ==
|
||||||
|
both busy : 1.44s 11.8%
|
||||||
|
only PP0 : 0.03s 0.3% <- PP1 bubble
|
||||||
|
only PP1 : 0.13s 1.1% <- PP0 bubble
|
||||||
|
neither : 10.56s 86.8% <- sched/CPU idle
|
||||||
|
pipeline efficiency: both/(any busy) = 89.6%
|
||||||
@ -0,0 +1,34 @@
|
|||||||
|
Triage View
|
||||||
|
Mode: single-trace
|
||||||
|
Framework: SGLang
|
||||||
|
Input traces: /data/hf_models/pp2tp4_profiles/B_decode/pb-1788706647.2466695-TP-0-PP-0.trace.json.gz
|
||||||
|
|
||||||
|
Kernel Table
|
||||||
|
##### decode
|
||||||
|
| Kernel | Category | GPU time | Share | Launches | Python location (site share) | CPU op |
|
||||||
|
| --- | --- | ---: | ---: | ---: | --- | --- |
|
||||||
|
| _ZN7cutlass13device_kernelINS_4gemm6kernel13GemmUniversalINS1_17GroupProblemShapeIN4cute5tupleIJlllEEEEENS1_10collective13CollectiveMmaINS1_47MainloopSm120ArrayTmaWarpSpecializedBlockScaledILi2ELi3ENS6_IJNS5_1CILi1EEESD_SD_EEENS1_59KernelPtrArrayTmaWarpSpecializedCooperativeBlockScaledSm120ILi3EEEEENS6_IJNSC_ILi128EEESI_NSC_ILi256EEEEEENS6_IJNS_12float_e2m1_tENS_13float_ue4m3_tEEEENS6_IJPNS6_IJlSD_NSC_ILi0EEEEEEPNS5_6LayoutINS6_IJNS6_IJNS6_IJNSC_ILi32EEENSC_ILi4EEEEEEiEEENS6_IJNS6_IJNSC_ILi16EEEST_EEEiEEENS6_IJSD_iEEEEEENS6_IJSY_NS6_IJNS6_IJSO_SD_EEENSC_ILi512EEEEEENS6_IJSO_iEEEEEEEEEEESN_S18_NS5_8TiledMMAINS5_8MMA_AtomIJNS5_5SM12011BLOCKSCALED19SM120_16x8x64_TN_VSISL_SL_fSM_Li16EEEEEENSR_INS6_IJST_NSC_ILi2EEESD_EEENS6_IJSD_ST_SO_EEEEENS6_IJSI_NSR_INS6_IJNSC_ILi8EEES1G_S1G_EEENS6_IJSD_SW_S1K_EEEEENSC_ILi64EEEEEEEENS6_IJNS5_13SM90_TMA_LOADES1R_EEENS6_IJNS5_14ComposedLayoutINS5_7SwizzleILi3ELi4ELi3EEENS5_18smem_ptr_flag_bitsILi4EEENSR_INS6_IJS1K_SJ_EEENS6_IJSJ_SD_EEEEEEENSR_INS6_IJNS6_IJSU_SD_EEENS6_IJSX_SD_ST_EEEEEENS6_IJNS6_IJSX_S12_EEENS6_IJS11_ST_S12_EEEEEEEEEEENS6_IJNS5_9Copy_AtomIJNS5_17SM75_U32x4_LDSM_NENS_15integer_subbyteILi4ELb0EEEEEENS2A_IJNS5_13UniversalCopyISM_SM_EESM_EEEEEENS5_8identityES1S_S29_S2I_S2J_EENS_8epilogue10collective18CollectiveEpilogueINS2L_30Sm90PtrArrayTmaWarpSpecializedILi3ELi2ELi4ELb1ELb0ELi2EEEJSK_NS6_IJS1O_SS_EEENS_10bfloat16_tESQ_S2R_SQ_NS2L_6fusion15FusionCallbacksINS2L_31Sm120PtrArrayTmaWarpSpecializedILi3ELi2ELi4ELb1ELb0ELi2EEENS2S_17LinearCombinationIS2R_fS2R_fLNS_15FloatRoundStyleE2EEESK_S2Q_JEEES1R_NS1T_INS1U_ILi2ELi4ELi3EEENS1W_ILi16EEENSR_INS6_IJS1K_SS_EEENS6_IJSS_SD_EEEEEEENS5_17SM75_U32x2_LDSM_NENS5_14SM90_TMA_STOREES35_NS5_17SM90_U32x2_STSM_NENS2A_IJS38_NS_6half_tEEEEvEEEvvEEEEvNT_6ParamsE | moe | 459.23 ms | 34.8% | 4320 | unresolved | cudaGraphLaunch |
|
||||||
|
| ncclDevKernel_SendRecv(ncclDevKernelArgsStorage<4096ul>) | communication | 239.28 ms | 18.1% | 300 | unresolved | cuLaunchKernelEx |
|
||||||
|
| void cutlass::Kernel2<cutlass_80_wmma_tensorop_bf16_s161616gemm_bf16_16x16_128x2_tn_align8> | gemm | 142.38 ms | 10.8% | 5400 | unresolved | cudaGraphLaunch |
|
||||||
|
| ncclDevKernel_AllReduce_Sum_bf16_RING_LL(ncclDevKernelArgsStorage<4096ul>) | communication | 84.97 ms | 6.4% | 4740 | unresolved | cudaGraphLaunch |
|
||||||
|
| kernel_cutlass__dsv3_fused_a_gemm_kernel_tensorptri32gmemo2624307230721_tensorptri32gmemo16307230721_tensorptrbf16gmemo16262426241__24_8_8_0 | gemm | 68.37 ms | 5.2% | 2340 | unresolved | cudaGraphLaunch |
|
||||||
|
| void flashinfer::sparse_mla_sm120::sparse_mla_decode_dsv3_2_kernel<(ModelType)2, 16, 2048, 64> | attention | 47.80 ms | 3.6% | 2340 | unresolved | cudaGraphLaunch |
|
||||||
|
| kernel_cutlass__dsv3_fused_a_gemm_kernel_tensorptri32gmemo4096102410241_tensorptri32gmemo16102410241_tensorptrbf16gmemo16409640961__8_8_8_0 | gemm | 35.11 ms | 2.7% | 2340 | unresolved | cudaGraphLaunch |
|
||||||
|
| void cutlass::Kernel2<cutlass_80_wmma_tensorop_bf16_s161616gemm_bf16_32x32_64x2_tn_align8> | gemm | 28.22 ms | 2.1% | 2160 | unresolved | cudaGraphLaunch |
|
||||||
|
| void tensorrt_llm::kernels::cutlass_kernels::finalizeMoeRoutingKernel<__nv_bfloat16, __nv_bfloat16, __nv_bfloat16, (tensorrt_llm::kernels::cutlass_kernels::ScaleMode)1> | moe | 25.78 ms | 2.0% | 2160 | unresolved | cudaGraphLaunch |
|
||||||
|
| void cutlass::Kernel2<cutlass_80_wmma_tensorop_bf16_s161616gemm_bf16_32x32_128x2_tn_align8> | gemm | 18.78 ms | 1.4% | 3060 | unresolved | cudaGraphLaunch |
|
||||||
|
| kernel_cutlass_kernel_flashinfernormkernelsfused_add_rmsnormFusedAddRMSNormKernel_object_at__tensorptrbf16gmemalign128oi64614461441_tensorptrbf16gmemalign128oi64614461441_tensorptrbf16gme_0 | gemm | 16.06 ms | 1.2% | 4620 | unresolved | cudaGraphLaunch |
|
||||||
|
| void sglang::router_gemm_kernel<__nv_bfloat16, float, 128, 8, 8, 256, 6144, true> | gemm | 14.71 ms | 1.1% | 2160 | unresolved | cudaGraphLaunch |
|
||||||
|
|
||||||
|
Overlap Opportunity Table
|
||||||
|
| Priority | Verdict | Kernel | Python scope | Formal signal | Dep risk | Recommendation |
|
||||||
|
| --- | --- | --- | --- | --- | --- | --- |
|
||||||
|
| - | - | No rows cleared the 1.0% reporting bar. Use mapping/formal mode for overlap attribution. | - | - | - | - |
|
||||||
|
|
||||||
|
Fuse Opportunity Table
|
||||||
|
##### decode
|
||||||
|
| Pattern | Confidence | Related GPU time | Share | Evidence kernels | Current kernel Python location | Candidate fused Python path | Rationale |
|
||||||
|
| --- | --- | ---: | ---: | --- | --- | --- | --- |
|
||||||
|
| Fused MoE grouped-topk / gate kernels | Confirmed | 29.29 ms | 2.2% | void sglang::router_gemm_kernel<__nv_bfloat16, float, 128, 8, 8, 256, 6144, true> (1.1%) | unresolved | python/sglang/srt/layers/moe/topk.py | Split kernels in this family take 2.2% of GPU time. This tree already has a matching path. Grouped-topk, bias handling, and routed scaling already have fused gate kernels. |
|
||||||
|
| NSA fused quantize + indexed K-cache store | Confirmed | 465.20 ms | 35.2% | _ZN7cutlass13device_kernelINS_4gemm6kernel13GemmUniversalINS1_17GroupProblemShapeIN4cute5tupleIJlllEEEEENS1_10collective13CollectiveMmaINS1_47MainloopSm120ArrayTmaWarpSpecializedBlockScaledILi2ELi3ENS6_IJNS5_1CILi1EEESD_SD_EEENS1_59KernelPtrArrayTmaWarpSpecializedCooperativeBlockScaledSm120ILi3EEEEENS6_IJNSC_ILi128EEESI_NSC_ILi256EEEEEENS6_IJNS_12float_e2m1_tENS_13float_ue4m3_tEEEENS6_IJPNS6_IJlSD_NSC_ILi0EEEEEEPNS5_6LayoutINS6_IJNS6_IJNS6_IJNSC_ILi32EEENSC_ILi4EEEEEEiEEENS6_IJNS6_IJNSC_ILi16EEEST_EEEiEEENS6_IJSD_iEEEEEENS6_IJSY_NS6_IJNS6_IJSO_SD_EEENSC_ILi512EEEEEENS6_IJSO_iEEEEEEEEEEESN_S18_NS5_8TiledMMAINS5_8MMA_AtomIJNS5_5SM12011BLOCKSCALED19SM120_16x8x64_TN_VSISL_SL_fSM_Li16EEEEEENSR_INS6_IJST_NSC_ILi2EEESD_EEENS6_IJSD_ST_SO_EEEEENS6_IJSI_NSR_INS6_IJNSC_ILi8EEES1G_S1G_EEENS6_IJSD_SW_S1K_EEEEENSC_ILi64EEEEEEEENS6_IJNS5_13SM90_TMA_LOADES1R_EEENS6_IJNS5_14ComposedLayoutINS5_7SwizzleILi3ELi4ELi3EEENS5_18smem_ptr_flag_bitsILi4EEENSR_INS6_IJS1K_SJ_EEENS6_IJSJ_SD_EEEEEEENSR_INS6_IJNS6_IJSU_SD_EEENS6_IJSX_SD_ST_EEEEEENS6_IJNS6_IJSX_S12_EEENS6_IJS11_ST_S12_EEEEEEEEEEENS6_IJNS5_9Copy_AtomIJNS5_17SM75_U32x4_LDSM_NENS_15integer_subbyteILi4ELb0EEEEEENS2A_IJNS5_13UniversalCopyISM_SM_EESM_EEEEEENS5_8identityES1S_S29_S2I_S2J_EENS_8epilogue10collective18CollectiveEpilogueINS2L_30Sm90PtrArrayTmaWarpSpecializedILi3ELi2ELi4ELb1ELb0ELi2EEEJSK_NS6_IJS1O_SS_EEENS_10bfloat16_tESQ_S2R_SQ_NS2L_6fusion15FusionCallbacksINS2L_31Sm120PtrArrayTmaWarpSpecializedILi3ELi2ELi4ELb1ELb0ELi2EEENS2S_17LinearCombinationIS2R_fS2R_fLNS_15FloatRoundStyleE2EEESK_S2Q_JEEES1R_NS1T_INS1U_ILi2ELi4ELi3EEENS1W_ILi16EEENSR_INS6_IJS1K_SS_EEENS6_IJSS_SD_EEEEEEENS5_17SM75_U32x2_LDSM_NENS5_14SM90_TMA_STOREES35_NS5_17SM90_U32x2_STSM_NENS2A_IJS38_NS_6half_tEEEEvEEEvvEEEEvNT_6ParamsE (34.8%) | unresolved | python/sglang/kernels/ops/attention/fused_store_index_cache.py<br>python/sglang/srt/layers/attention/nsa/nsa_indexer.py | Split kernels in this family take 35.2% of GPU time. This tree already has a matching path. NSA already has a fused quantize-and-indexed-store kernel family. |
|
||||||
|
| Fused residual add + RMSNorm | Confirmed | 16.06 ms | 1.2% | kernel_cutlass_kernel_flashinfernormkernelsfused_add_rmsnormFusedAddRMSNormKernel_object_at__tensorptrbf16gmemalign128oi64614461441_tensorptrbf16gmemalign128oi64614461441_tensorptrbf16gme_0 (1.2%) | unresolved | python/sglang/srt/layers/layernorm.py<br>python/sglang/srt/layers/quantization/modelslim/modelslim.py | `Fused residual add + RMSNorm` is present in this trace (1.2% related GPU time). Residual add plus RMSNorm already has fused implementations across several backends. |
|
||||||
@ -0,0 +1,36 @@
|
|||||||
|
Triage View
|
||||||
|
Mode: single-trace
|
||||||
|
Framework: SGLang
|
||||||
|
Input traces: /data/hf_models/pp2tp4_profiles/B_decode/pb-1788706647.2466695-TP-0-PP-1.trace.json.gz
|
||||||
|
|
||||||
|
Kernel Table
|
||||||
|
##### decode
|
||||||
|
| Kernel | Category | GPU time | Share | Launches | Python location (site share) | CPU op |
|
||||||
|
| --- | --- | ---: | ---: | ---: | --- | --- |
|
||||||
|
| _ZN7cutlass13device_kernelINS_4gemm6kernel13GemmUniversalINS1_17GroupProblemShapeIN4cute5tupleIJlllEEEEENS1_10collective13CollectiveMmaINS1_47MainloopSm120ArrayTmaWarpSpecializedBlockScaledILi2ELi3ENS6_IJNS5_1CILi1EEESD_SD_EEENS1_59KernelPtrArrayTmaWarpSpecializedCooperativeBlockScaledSm120ILi3EEEEENS6_IJNSC_ILi128EEESI_NSC_ILi256EEEEEENS6_IJNS_12float_e2m1_tENS_13float_ue4m3_tEEEENS6_IJPNS6_IJlSD_NSC_ILi0EEEEEEPNS5_6LayoutINS6_IJNS6_IJNS6_IJNSC_ILi32EEENSC_ILi4EEEEEEiEEENS6_IJNS6_IJNSC_ILi16EEEST_EEEiEEENS6_IJSD_iEEEEEENS6_IJSY_NS6_IJNS6_IJSO_SD_EEENSC_ILi512EEEEEENS6_IJSO_iEEEEEEEEEEESN_S18_NS5_8TiledMMAINS5_8MMA_AtomIJNS5_5SM12011BLOCKSCALED19SM120_16x8x64_TN_VSISL_SL_fSM_Li16EEEEEENSR_INS6_IJST_NSC_ILi2EEESD_EEENS6_IJSD_ST_SO_EEEEENS6_IJSI_NSR_INS6_IJNSC_ILi8EEES1G_S1G_EEENS6_IJSD_SW_S1K_EEEEENSC_ILi64EEEEEEEENS6_IJNS5_13SM90_TMA_LOADES1R_EEENS6_IJNS5_14ComposedLayoutINS5_7SwizzleILi3ELi4ELi3EEENS5_18smem_ptr_flag_bitsILi4EEENSR_INS6_IJS1K_SJ_EEENS6_IJSJ_SD_EEEEEEENSR_INS6_IJNS6_IJSU_SD_EEENS6_IJSX_SD_ST_EEEEEENS6_IJNS6_IJSX_S12_EEENS6_IJS11_ST_S12_EEEEEEEEEEENS6_IJNS5_9Copy_AtomIJNS5_17SM75_U32x4_LDSM_NENS_15integer_subbyteILi4ELb0EEEEEENS2A_IJNS5_13UniversalCopyISM_SM_EESM_EEEEEENS5_8identityES1S_S29_S2I_S2J_EENS_8epilogue10collective18CollectiveEpilogueINS2L_30Sm90PtrArrayTmaWarpSpecializedILi3ELi2ELi4ELb1ELb0ELi2EEEJSK_NS6_IJS1O_SS_EEENS_10bfloat16_tESQ_S2R_SQ_NS2L_6fusion15FusionCallbacksINS2L_31Sm120PtrArrayTmaWarpSpecializedILi3ELi2ELi4ELb1ELb0ELi2EEENS2S_17LinearCombinationIS2R_fS2R_fLNS_15FloatRoundStyleE2EEESK_S2Q_JEEES1R_NS1T_INS1U_ILi2ELi4ELi3EEENS1W_ILi16EEENSR_INS6_IJS1K_SS_EEENS6_IJSS_SD_EEEEEEENS5_17SM75_U32x2_LDSM_NENS5_14SM90_TMA_STOREES35_NS5_17SM90_U32x2_STSM_NENS2A_IJS38_NS_6half_tEEEEvEEEvvEEEEvNT_6ParamsE | moe | 551.07 ms | 40.7% | 4680 | unresolved | cudaGraphLaunch |
|
||||||
|
| void cutlass::Kernel2<cutlass_80_wmma_tensorop_bf16_s161616gemm_bf16_16x16_128x2_tn_align8> | gemm | 136.35 ms | 10.1% | 5220 | unresolved | cudaGraphLaunch |
|
||||||
|
| ncclDevKernel_AllReduce_Sum_bf16_RING_LL(ncclDevKernelArgsStorage<4096ul>) | communication | 112.42 ms | 8.3% | 4680 | unresolved | cudaGraphLaunch |
|
||||||
|
| ncclDevKernel_AllGather_RING_LL(ncclDevKernelArgsStorage<4096ul>) | communication | 85.71 ms | 6.3% | 273 | unresolved | cuLaunchKernelEx |
|
||||||
|
| kernel_cutlass__dsv3_fused_a_gemm_kernel_tensorptri32gmemo2624307230721_tensorptri32gmemo16307230721_tensorptrbf16gmemo16262426241__24_8_8_0 | gemm | 68.51 ms | 5.1% | 2340 | unresolved | cudaGraphLaunch |
|
||||||
|
| ncclDevKernel_SendRecv(ncclDevKernelArgsStorage<4096ul>) | communication | 57.72 ms | 4.3% | 303 | unresolved | cuLaunchKernelEx |
|
||||||
|
| void flashinfer::sparse_mla_sm120::sparse_mla_decode_dsv3_2_kernel<(ModelType)2, 16, 2048, 64> | attention | 47.68 ms | 3.5% | 2340 | unresolved | cudaGraphLaunch |
|
||||||
|
| kernel_cutlass__dsv3_fused_a_gemm_kernel_tensorptri32gmemo4096102410241_tensorptri32gmemo16102410241_tensorptrbf16gmemo16409640961__8_8_8_0 | gemm | 35.26 ms | 2.6% | 2340 | unresolved | cudaGraphLaunch |
|
||||||
|
| void cutlass::Kernel2<cutlass_80_wmma_tensorop_bf16_s161616gemm_bf16_32x32_64x2_tn_align8> | gemm | 30.64 ms | 2.3% | 2340 | unresolved | cudaGraphLaunch |
|
||||||
|
| void tensorrt_llm::kernels::cutlass_kernels::finalizeMoeRoutingKernel<__nv_bfloat16, __nv_bfloat16, __nv_bfloat16, (tensorrt_llm::kernels::cutlass_kernels::ScaleMode)1> | moe | 27.98 ms | 2.1% | 2340 | unresolved | cudaGraphLaunch |
|
||||||
|
| void cutlass::Kernel2<cutlass_80_tensorop_bf16_s16816gemm_relu_bf16_64x64_64x6_tn_align8> | gemm | 21.88 ms | 1.6% | 60 | unresolved | cudaGraphLaunch |
|
||||||
|
| void cutlass::Kernel2<cutlass_80_wmma_tensorop_bf16_s161616gemm_bf16_32x32_128x2_tn_align8> | gemm | 17.44 ms | 1.3% | 2880 | unresolved | cudaGraphLaunch |
|
||||||
|
| kernel_cutlass_kernel_flashinfernormkernelsfused_add_rmsnormFusedAddRMSNormKernel_object_at__tensorptrbf16gmemalign128oi64614461441_tensorptrbf16gmemalign128oi64614461441_tensorptrbf16gme_0 | gemm | 16.46 ms | 1.2% | 4740 | unresolved | cudaGraphLaunch |
|
||||||
|
| void sglang::router_gemm_kernel<__nv_bfloat16, float, 128, 8, 8, 256, 6144, true> | gemm | 16.02 ms | 1.2% | 2340 | unresolved | cudaGraphLaunch |
|
||||||
|
|
||||||
|
Overlap Opportunity Table
|
||||||
|
| Priority | Verdict | Kernel | Python scope | Formal signal | Dep risk | Recommendation |
|
||||||
|
| --- | --- | --- | --- | --- | --- | --- |
|
||||||
|
| - | - | No rows cleared the 1.0% reporting bar. Use mapping/formal mode for overlap attribution. | - | - | - | - |
|
||||||
|
|
||||||
|
Fuse Opportunity Table
|
||||||
|
##### decode
|
||||||
|
| Pattern | Confidence | Related GPU time | Share | Evidence kernels | Current kernel Python location | Candidate fused Python path | Rationale |
|
||||||
|
| --- | --- | ---: | ---: | --- | --- | --- | --- |
|
||||||
|
| Fused MoE grouped-topk / gate kernels | Confirmed | 29.89 ms | 2.2% | void sglang::router_gemm_kernel<__nv_bfloat16, float, 128, 8, 8, 256, 6144, true> (1.2%) | unresolved | python/sglang/srt/layers/moe/topk.py | Split kernels in this family take 2.2% of GPU time. This tree already has a matching path. Grouped-topk, bias handling, and routed scaling already have fused gate kernels. |
|
||||||
|
| NSA fused quantize + indexed K-cache store | Confirmed | 556.77 ms | 41.2% | _ZN7cutlass13device_kernelINS_4gemm6kernel13GemmUniversalINS1_17GroupProblemShapeIN4cute5tupleIJlllEEEEENS1_10collective13CollectiveMmaINS1_47MainloopSm120ArrayTmaWarpSpecializedBlockScaledILi2ELi3ENS6_IJNS5_1CILi1EEESD_SD_EEENS1_59KernelPtrArrayTmaWarpSpecializedCooperativeBlockScaledSm120ILi3EEEEENS6_IJNSC_ILi128EEESI_NSC_ILi256EEEEEENS6_IJNS_12float_e2m1_tENS_13float_ue4m3_tEEEENS6_IJPNS6_IJlSD_NSC_ILi0EEEEEEPNS5_6LayoutINS6_IJNS6_IJNS6_IJNSC_ILi32EEENSC_ILi4EEEEEEiEEENS6_IJNS6_IJNSC_ILi16EEEST_EEEiEEENS6_IJSD_iEEEEEENS6_IJSY_NS6_IJNS6_IJSO_SD_EEENSC_ILi512EEEEEENS6_IJSO_iEEEEEEEEEEESN_S18_NS5_8TiledMMAINS5_8MMA_AtomIJNS5_5SM12011BLOCKSCALED19SM120_16x8x64_TN_VSISL_SL_fSM_Li16EEEEEENSR_INS6_IJST_NSC_ILi2EEESD_EEENS6_IJSD_ST_SO_EEEEENS6_IJSI_NSR_INS6_IJNSC_ILi8EEES1G_S1G_EEENS6_IJSD_SW_S1K_EEEEENSC_ILi64EEEEEEEENS6_IJNS5_13SM90_TMA_LOADES1R_EEENS6_IJNS5_14ComposedLayoutINS5_7SwizzleILi3ELi4ELi3EEENS5_18smem_ptr_flag_bitsILi4EEENSR_INS6_IJS1K_SJ_EEENS6_IJSJ_SD_EEEEEEENSR_INS6_IJNS6_IJSU_SD_EEENS6_IJSX_SD_ST_EEEEEENS6_IJNS6_IJSX_S12_EEENS6_IJS11_ST_S12_EEEEEEEEEEENS6_IJNS5_9Copy_AtomIJNS5_17SM75_U32x4_LDSM_NENS_15integer_subbyteILi4ELb0EEEEEENS2A_IJNS5_13UniversalCopyISM_SM_EESM_EEEEEENS5_8identityES1S_S29_S2I_S2J_EENS_8epilogue10collective18CollectiveEpilogueINS2L_30Sm90PtrArrayTmaWarpSpecializedILi3ELi2ELi4ELb1ELb0ELi2EEEJSK_NS6_IJS1O_SS_EEENS_10bfloat16_tESQ_S2R_SQ_NS2L_6fusion15FusionCallbacksINS2L_31Sm120PtrArrayTmaWarpSpecializedILi3ELi2ELi4ELb1ELb0ELi2EEENS2S_17LinearCombinationIS2R_fS2R_fLNS_15FloatRoundStyleE2EEESK_S2Q_JEEES1R_NS1T_INS1U_ILi2ELi4ELi3EEENS1W_ILi16EEENSR_INS6_IJS1K_SS_EEENS6_IJSS_SD_EEEEEEENS5_17SM75_U32x2_LDSM_NENS5_14SM90_TMA_STOREES35_NS5_17SM90_U32x2_STSM_NENS2A_IJS38_NS_6half_tEEEEvEEEvvEEEEvNT_6ParamsE (40.7%) | unresolved | python/sglang/kernels/ops/attention/fused_store_index_cache.py<br>python/sglang/srt/layers/attention/nsa/nsa_indexer.py | Split kernels in this family take 41.2% of GPU time. This tree already has a matching path. NSA already has a fused quantize-and-indexed-store kernel family. |
|
||||||
|
| Fused residual add + RMSNorm | Confirmed | 16.46 ms | 1.2% | kernel_cutlass_kernel_flashinfernormkernelsfused_add_rmsnormFusedAddRMSNormKernel_object_at__tensorptrbf16gmemalign128oi64614461441_tensorptrbf16gmemalign128oi64614461441_tensorptrbf16gme_0 (1.2%) | unresolved | python/sglang/srt/layers/layernorm.py<br>python/sglang/srt/layers/quantization/modelslim/modelslim.py | `Fused residual add + RMSNorm` is present in this trace (1.2% related GPU time). Residual add plus RMSNorm already has fused implementations across several backends. |
|
||||||
@ -0,0 +1,68 @@
|
|||||||
|
== PP0 rank ==
|
||||||
|
file: /data/hf_models/pp2tp4_profiles/D_128k/pd-1788708582.9504747-TP-0-PP-0.trace.json.gz
|
||||||
|
baseTimeNanoseconds present: True
|
||||||
|
span 19.900s GPU-busy 19.311s (97.0%) kernels 30279
|
||||||
|
cudaMemcpy/Set GPU time: 212.3ms cudaLaunchKernel calls: 8256 (1914.1ms CPU)
|
||||||
|
NCCL total: 10898.6ms (56.4% of busy, 54.8% of span)
|
||||||
|
SendRecv 6970.9ms n=77
|
||||||
|
AllReduce 3927.6ms n=1264
|
||||||
|
AG_RS 0.0ms n=2
|
||||||
|
top kernels (of GPU-busy):
|
||||||
|
36.1% 6970.9ms n=77 ncclDevKernel_SendRecv(ncclDevKernelArgsStorage
|
||||||
|
28.0% 5397.8ms n=576 void deep_gemm::sm120_fp8_mqa_logits
|
||||||
|
20.3% 3927.6ms n=1264 ncclDevKernel_AllReduce_Sum_bf16_RING_LL(ncclDevKernelArgsStorage
|
||||||
|
17.3% 3347.8ms n=507 void sparse_mla_prefill_kernel
|
||||||
|
10.9% 2106.7ms n=4986 void cutlass::Kernel2
|
||||||
|
6.7% 1300.1ms n=1152 _ZN7cutlass13device_kernelINS_4gemm6kernel13GemmUniversalINS1_17GroupProblemShap
|
||||||
|
3.1% 608.0ms n=576 (anonymous namespace)::topk_transform_prefill_kernel((anonymous namespace)::Fast
|
||||||
|
2.5% 486.4ms n=583 void at::native::vectorized_gather_kernel
|
||||||
|
1.7% 334.3ms n=576 void tensorrt_llm::kernels::cutlass_kernels::finalizeMoeRoutingKernel
|
||||||
|
1.3% 257.0ms n=1232 kernel_cutlass_kernel_flashinfernormkernelsfused_add_rmsnormFusedAddRMSNormKerne
|
||||||
|
1.2% 232.6ms n=576 void tensorrt_llm::kernels::cutlass_kernels::expandInputRowsKernel
|
||||||
|
1.1% 212.3ms n=1735 ::
|
||||||
|
0.5% 102.2ms n=624 void sglang::concat_mla_absorb_q_kernel
|
||||||
|
0.5% 99.5ms n=1635 void at::native::vectorized_elementwise_kernel
|
||||||
|
0.5% 87.3ms n=576 void tensorrt_llm::kernels::cutlass_kernels::doActivationKernel
|
||||||
|
forwards: 7 (gap>4.0ms split)
|
||||||
|
wall ms: mean 2769.3 p50 583.1 min 3.1 max 11507.6
|
||||||
|
small forwards (<30% of p50): 3 (warmup/mixed steps)
|
||||||
|
GPU-covered fraction inside main forwards: mean 99.1%
|
||||||
|
inter-forward gap ms: mean 85.8 p50 77.5 p90 231.2 max 231.2
|
||||||
|
|
||||||
|
== PP1 rank ==
|
||||||
|
file: /data/hf_models/pp2tp4_profiles/D_128k/pd-1788708582.9504747-TP-0-PP-1.trace.json.gz
|
||||||
|
baseTimeNanoseconds present: True
|
||||||
|
span 30.873s GPU-busy 19.218s (62.2%) kernels 36752
|
||||||
|
cudaMemcpy/Set GPU time: 222.0ms cudaLaunchKernel calls: 13111 (5216.2ms CPU)
|
||||||
|
NCCL total: 6021.4ms (31.3% of busy, 19.5% of span)
|
||||||
|
AllReduce 3757.5ms n=1248
|
||||||
|
SendRecv 2191.8ms n=82
|
||||||
|
AG_RS 72.1ms n=69
|
||||||
|
top kernels (of GPU-busy):
|
||||||
|
23.1% 4429.8ms n=1008 void deep_gemm::sm120_fp8_mqa_logits
|
||||||
|
19.6% 3757.5ms n=1248 ncclDevKernel_AllReduce_Sum_bf16_RING_LL(ncclDevKernelArgsStorage
|
||||||
|
17.4% 3340.8ms n=507 void sparse_mla_prefill_kernel
|
||||||
|
11.4% 2191.8ms n=82 ncclDevKernel_SendRecv(ncclDevKernelArgsStorage
|
||||||
|
10.3% 1973.9ms n=4939 void cutlass::Kernel2
|
||||||
|
5.8% 1113.6ms n=1248 _ZN7cutlass13device_kernelINS_4gemm6kernel13GemmUniversalINS1_17GroupProblemShap
|
||||||
|
2.9% 547.9ms n=1008 (anonymous namespace)::topk_transform_prefill_kernel((anonymous namespace)::Fast
|
||||||
|
1.9% 361.8ms n=624 void tensorrt_llm::kernels::cutlass_kernels::finalizeMoeRoutingKernel
|
||||||
|
1.9% 355.7ms n=1022 void at::native::vectorized_gather_kernel
|
||||||
|
1.4% 264.8ms n=1264 kernel_cutlass_kernel_flashinfernormkernelsfused_add_rmsnormFusedAddRMSNormKerne
|
||||||
|
1.3% 252.9ms n=624 void tensorrt_llm::kernels::cutlass_kernels::expandInputRowsKernel
|
||||||
|
1.2% 222.0ms n=2576 ::
|
||||||
|
0.6% 106.7ms n=2063 void at::native::vectorized_elementwise_kernel
|
||||||
|
0.5% 102.4ms n=624 void sglang::concat_mla_absorb_q_kernel
|
||||||
|
0.5% 91.4ms n=624 void tensorrt_llm::kernels::cutlass_kernels::doActivationKernel
|
||||||
|
forwards: 4 (gap>4.0ms split)
|
||||||
|
wall ms: mean 4847.5 p50 6553.0 min 5.1 max 12800.7
|
||||||
|
small forwards (<30% of p50): 2 (warmup/mixed steps)
|
||||||
|
GPU-covered fraction inside main forwards: mean 99.2%
|
||||||
|
inter-forward gap ms: mean 3827.7 p50 553.0 p90 10924.9 max 10924.9
|
||||||
|
|
||||||
|
== cross-stage bins (1.0ms, total 30.87s) ==
|
||||||
|
both busy : 18.79s 60.9%
|
||||||
|
only PP0 : 0.59s 1.9% <- PP1 bubble
|
||||||
|
only PP1 : 0.56s 1.8% <- PP0 bubble
|
||||||
|
neither : 10.93s 35.4% <- sched/CPU idle
|
||||||
|
pipeline efficiency: both/(any busy) = 94.2%
|
||||||
@ -0,0 +1,5 @@
|
|||||||
|
== 1788710165.9880266-TP-0-PP-0.trace.json.gz ==
|
||||||
|
AllReduce: n=79 total=293.6ms min=3.68 p10=3.69 p50=3.70 p90=3.75 p99=4.18 max=4.18 (ms)
|
||||||
|
span 0.9s; top 2s-buckets (offset_s -> ms): 0s->294ms
|
||||||
|
SendRecv : n=3 total=15.2ms min=4.15 p10=4.15 p50=4.85 p90=6.25 p99=6.25 max=6.25 (ms)
|
||||||
|
span 0.0s; top 2s-buckets (offset_s -> ms): 0s->15ms
|
||||||
@ -0,0 +1,3 @@
|
|||||||
|
== 1788710165.9880266-TP-0-PP-0.trace.json.gz ==
|
||||||
|
n= 79 total= 293.6ms mean= 3.716ms ncclDevKernel_AllReduce_Sum_bf16_RING_LL(ncclDevKernelArgsStorage<4096ul>)
|
||||||
|
n= 3 total= 15.2ms mean= 5.083ms ncclDevKernel_SendRecv(ncclDevKernelArgsStorage<4096ul>)
|
||||||
@ -0,0 +1,71 @@
|
|||||||
|
# 第二轮优化实验记录（2026-09-07，基于 09-06 profile 报告方向）
|
||||||
|
|
||||||
|
平台：174.1.60.5，TP4PP2 原样配置（deploy_par.sh chunk 8192 / mem-frac 0.88）。
|
||||||
|
结论：**配置级快赢全部证伪，现有配置即本栈配置最优**；剩余提升均为开发投入。
|
||||||
|
|
||||||
|
## 1. 干净 128k/90% 基线（补缺口，此前 D 场景 bench 带 profiler 作废）
|
||||||
|
|
||||||
|
cc8，8×131072(90% 共享前缀)/512，run-id 20260907，hit rate 0.8999：
|
||||||
|
|
||||||
|
- wall 39.38s；输出 104.02 tok/s；输入 26.63k tok/s（含缓存命中）
|
||||||
|
- **TTFT p50 14.70s**（min 4.64 / max 20.02）
|
||||||
|
- TPOT 均值 51.7ms（p50 53.8 / max 67.9）
|
||||||
|
- e2e 均值 39.24s；per-req 13.05 tok/s；retract 0
|
||||||
|
|
||||||
|
证据：results/bench/128k_clean_baseline_tp4pp2.log
|
||||||
|
|
||||||
|
## 2. chunk 16384 证伪（方向④关闭）
|
||||||
|
|
||||||
|
变体 `deploy_par.sh "<同参数>" nomtp 16384`：启动成功（日志核对 chunked_prefill_size=16384）、
|
||||||
|
GSM8K 冒烟通过，但首个 16k 请求导致全部 8 个调度器崩溃：
|
||||||
|
|
||||||
|
- **根因：CUDA OOM**。flashinfer cutlass MoE `getWorkspaceInfo` 工作区随 chunk 规模增长，
|
||||||
|
16384-token 前向需 **3.08GB**，mem-frac 0.88 下每卡仅 **2.11GB** 空闲
|
||||||
|
（MemoryError: Tried to allocate 3.08 GiB, GPU 7 ... 2.11 GiB is free）
|
||||||
|
- 要跑通需 mem-frac ≤0.85（KV 池 -5%），但收益主体经物理核算不存在：
|
||||||
|
- AR 纯带宽受限（96MB/3.7ms，每次调用固定开销≈0）→"AR 次数减半"不省带宽部分
|
||||||
|
- indexer 的 q×k 总功与 chunk 切分无关（不变量）
|
||||||
|
- 剩余收益（调度/launch 摊薄）≤3%，在单轮 bench 噪声（±3%）内
|
||||||
|
- 代价：KV -5% + 混合相位 Max ITL 翻倍（单次 prefill 阻塞 1.03s→2.06s）
|
||||||
|
- 附带发现：prefill 分块图（backend=breakable）仅捕获 ≤2048 token 档位，
|
||||||
|
8192/16384 的 extend 都走 eager_runner 逐层重放图块——chunk 增大无图红利亦无额外退化
|
||||||
|
|
||||||
|
证据：results/ops/chunk16k_crash_oom.log（OOM 堆栈）、results/ops/deploy_chunk16k.log、
|
||||||
|
results/bench/chunk16k_cc16.log（bench 侧流中断）
|
||||||
|
|
||||||
|
## 3. AR 重叠/量化现成开关扫描（方向①降级为纯上游开发）
|
||||||
|
|
||||||
|
容器内 `launch_server --help` + 源码核查（sglang nightly 20260828）：
|
||||||
|
|
||||||
|
| 开关 | 判定 |
|
||||||
|
|---|---|
|
||||||
|
| `--enable-quant-communications` | **NPU 专用**：server_args 校验 `device != "npu"` 直接 ValueError（"only supported for NPU device"） |
|
||||||
|
| flashinfer allreduce fusion（no_cli 自动管理） | 自动启用门控 `is_sm90_supported() or is_sm100_supported()` → **SM120 被排除**；后端 trtllm/mnnvl 为 NVLink/小 batch 设计 |
|
||||||
|
| `--enable-symm-mem` / `--enable-torch-symm-mem` | 对称内存 one-shot AR，小消息/decode 导向；96MB prefill AR 在 PCIe 不适用 |
|
||||||
|
| `--enable-nccl-nvls` | NVLink 专用 |
|
||||||
|
| `--pp-async-batch-depth` | PP 气泡仅 ~4%（09-06 对齐数据），非瓶颈 |
|
||||||
|
|
||||||
|
**结论：本栈（TP4+PP2+SM120+PCIe）不存在 prefill AR 重叠/量化的现成路径。**
|
||||||
|
开发落点：① 分块 AR 与后续 GEMM 重叠（async-TP 式 TP linear 改造）；
|
||||||
|
② 将 flashinfer AR 融合（fp8 量化 AR+residual+norm）扩展到 SM120 与大 token 数。
|
||||||
|
|
||||||
|
## 4. indexer 方向关闭（方向③：本栈无杠杆）
|
||||||
|
|
||||||
|
- **修正 09-06 报告的原始设想**："前缀 logits 增量缓存"不成立——每请求后缀 q 独有，
|
||||||
|
其对共享前缀键的 logits 是该请求固有功；前缀 index-K 已由 fused_store_index_k_cache 缓存
|
||||||
|
- 既有 skip（`_should_skip_logits_computation`）仅适用 kv_len ≤ index_topk=2048（全选即正确），
|
||||||
|
128k 不适用且语义上必须选索引；CUDA decode 侧 skip 还被 HIP-only 门控
|
||||||
|
- 内核后端无选择：SM120 CUDA 仅 deepgemm 可用（aiter=ROCm、cutedsl=SM100 门控）；
|
||||||
|
nomtp 走 deepgemm split 变体（native 仅 MTP target-verify）
|
||||||
|
- 模型已做跨层省功：`index_topk_freq=4`（部分层复用上层 topk）
|
||||||
|
- roofline：mqa_logits 9.37ms/次（D 场景 PP0）≈ 8192 新 token × ~122k kv × 2048 FLOP/pair(每 rank,
|
||||||
|
32 头×128 维×2 / TP4) ≈ 2.05 TFLOP/rank → ~219 TFLOPS/rank fp8，paged-gather 内核的合理量级，
|
||||||
|
无 5-10× 调优空间
|
||||||
|
- **结论：28% 是 DSA 在 128k 的算法固有成本（扫全 kv 选 top-2048）。**
|
||||||
|
逃逸路径（分层/粗到细索引）是影响精度的上游研究项，非本栈优化
|
||||||
|
|
||||||
|
## 5. 剩余可行方向（按优先级）
|
||||||
|
|
||||||
|
1. **PP+MTP 上游化**（decode 最大杠杆，TPOT -60%+ 潜力；楔死根因与补丁已在 60.5:/root/）
|
||||||
|
2. **AR 分块重叠 / 量化 AR 内核开发**（prefill -15~25% 潜力；本栈无现成开关）
|
||||||
|
3. MoE 胶水融合 + PP1 AllGather 入图（decode -2~4%）
|
||||||
@ -0,0 +1,233 @@
|
|||||||
|
#!/usr/bin/env python3
|
||||||
|
"""bench_128k_hit90.py + profiler arming for scenario D (TP4PP2 on 60.5).
|
||||||
|
|
||||||
|
Identical scenario to D:/sskj/bench_128k_hit90.py (131072 in = 117968 shared
|
||||||
|
prefix + 13104 unique suffix, 512 out): after the shared-prefix warmup returns,
|
||||||
|
arms the sglang torch profiler (POST /start_profile) so the capture window
|
||||||
|
covers exactly the unhit-suffix prefill phase (+ a few decode steps).
|
||||||
|
|
||||||
|
Usage:
|
||||||
|
python3 bench_128k_hit90_prof.py --concurrency 8 --num-requests 8 \
|
||||||
|
--run-id <unique> --arm-profile /data/hf_models/pp2tp4_profiles/D_128k \
|
||||||
|
--profile-steps 16
|
||||||
|
"""
|
||||||
|
import argparse
|
||||||
|
import datetime
|
||||||
|
import json
|
||||||
|
import random
|
||||||
|
import re
|
||||||
|
import statistics
|
||||||
|
import subprocess
|
||||||
|
import time
|
||||||
|
from concurrent.futures import ThreadPoolExecutor
|
||||||
|
|
||||||
|
import requests
|
||||||
|
|
||||||
|
URL = "http://127.0.0.1:30000/generate"
|
||||||
|
CONTAINER = "glm53-nvfp4"
|
||||||
|
|
||||||
|
INPUT_LEN = 131072
|
||||||
|
SHARED_LEN = 117968
|
||||||
|
UNIQUE_LEN = INPUT_LEN - SHARED_LEN
|
||||||
|
OUTPUT_LEN = 512
|
||||||
|
VOCAB_LO, VOCAB_HI = 1000, 100000
|
||||||
|
SEED = 20260904
|
||||||
|
|
||||||
|
sess = requests.Session()
|
||||||
|
sess.trust_env = False
|
||||||
|
|
||||||
|
|
||||||
|
def build_prompts(num_requests, run_id):
|
||||||
|
rng = random.Random(SEED)
|
||||||
|
shared = [rng.randint(VOCAB_LO, VOCAB_HI) for _ in range(SHARED_LEN)]
|
||||||
|
prompts = []
|
||||||
|
for i in range(num_requests):
|
||||||
|
r = random.Random(SEED * 100003 + run_id * 1000003 + i)
|
||||||
|
prompts.append(shared + [r.randint(VOCAB_LO, VOCAB_HI) for _ in range(UNIQUE_LEN)])
|
||||||
|
return shared, prompts
|
||||||
|
|
||||||
|
|
||||||
|
def warmup(shared):
|
||||||
|
suffix = [random.Random(777).randint(VOCAB_LO, VOCAB_HI) for _ in range(64)]
|
||||||
|
payload = {
|
||||||
|
"input_ids": shared + suffix,
|
||||||
|
"sampling_params": {"max_new_tokens": 8, "temperature": 0.0, "ignore_eos": True},
|
||||||
|
}
|
||||||
|
t0 = time.perf_counter()
|
||||||
|
r = sess.post(URL, json=payload, timeout=900)
|
||||||
|
dt = time.perf_counter() - t0
|
||||||
|
print(f"[warmup] http={r.status_code} wall={dt:.2f}s", flush=True)
|
||||||
|
|
||||||
|
|
||||||
|
def arm_profile(output_dir, num_steps):
|
||||||
|
body = {
|
||||||
|
"output_dir": output_dir,
|
||||||
|
"num_steps": num_steps,
|
||||||
|
"activities": ["CPU", "GPU"],
|
||||||
|
"with_stack": False,
|
||||||
|
"record_shapes": False,
|
||||||
|
"profile_prefix": "pd",
|
||||||
|
}
|
||||||
|
r = sess.post(URL.replace("/generate", "/start_profile"), json=body, timeout=60)
|
||||||
|
print(f"[arm] http={r.status_code} body={r.text.strip()[:80]} steps={num_steps} dir={output_dir}",
|
||||||
|
flush=True)
|
||||||
|
|
||||||
|
|
||||||
|
def bench_one(prompt, idx, results):
|
||||||
|
payload = {
|
||||||
|
"input_ids": prompt,
|
||||||
|
"sampling_params": {"max_new_tokens": OUTPUT_LEN, "temperature": 0.0, "ignore_eos": True},
|
||||||
|
"stream": True,
|
||||||
|
}
|
||||||
|
rec = {"idx": idx}
|
||||||
|
t0 = time.perf_counter()
|
||||||
|
first = last = None
|
||||||
|
first_ct = None
|
||||||
|
final_meta = None
|
||||||
|
max_ct = 0
|
||||||
|
try:
|
||||||
|
with sess.post(URL, json=payload, stream=True, timeout=1800) as resp:
|
||||||
|
for raw in resp.iter_lines():
|
||||||
|
if not raw or not raw.startswith(b"data:"):
|
||||||
|
continue
|
||||||
|
body = raw[5:].strip()
|
||||||
|
if body == b"[DONE]":
|
||||||
|
continue
|
||||||
|
now = time.perf_counter()
|
||||||
|
try:
|
||||||
|
d = json.loads(body)
|
||||||
|
except Exception:
|
||||||
|
continue
|
||||||
|
mi = d.get("meta_info") or {}
|
||||||
|
ct = mi.get("completion_tokens") or 0
|
||||||
|
if ct:
|
||||||
|
max_ct = max(max_ct, ct)
|
||||||
|
if first is None:
|
||||||
|
first = now
|
||||||
|
first_ct = ct
|
||||||
|
last = now
|
||||||
|
if mi.get("finish_reason"):
|
||||||
|
final_meta = mi
|
||||||
|
t_end = time.perf_counter()
|
||||||
|
n_out = max_ct or ((final_meta or {}).get("completion_tokens") or 0)
|
||||||
|
decode_span = (last - first) if (first and last and last > first) else 0.0
|
||||||
|
rec.update(
|
||||||
|
ok=n_out > 0,
|
||||||
|
ttft=(first - t0) if first else None,
|
||||||
|
e2e=t_end - t0,
|
||||||
|
n_out=n_out,
|
||||||
|
first_chunk_tokens=first_ct,
|
||||||
|
decode_span=decode_span,
|
||||||
|
tpot=(decode_span / (n_out - 1)) if n_out > 1 else None,
|
||||||
|
retractions=(final_meta or {}).get("num_retractions"),
|
||||||
|
spec_accept_len=(final_meta or {}).get("spec_accept_length"),
|
||||||
|
)
|
||||||
|
except Exception as e:
|
||||||
|
rec.update(ok=False, error=repr(e))
|
||||||
|
results[idx] = rec
|
||||||
|
|
||||||
|
|
||||||
|
def verify_hit_rate(t_start, t_end):
|
||||||
|
def rfc3339(epoch):
|
||||||
|
return (datetime.datetime.fromtimestamp(epoch, tz=datetime.timezone.utc)
|
||||||
|
.isoformat().replace("+00:00", "Z"))
|
||||||
|
|
||||||
|
try:
|
||||||
|
p = subprocess.run(
|
||||||
|
["docker", "logs", CONTAINER, "--since", rfc3339(t_start), "--until", rfc3339(t_end + 2)],
|
||||||
|
capture_output=True, text=True, timeout=120)
|
||||||
|
text = p.stdout + p.stderr
|
||||||
|
except Exception as e:
|
||||||
|
return {"error": repr(e)}
|
||||||
|
pat = re.compile(r"#new-token: (\d+).*?#cached-token: (\d+)")
|
||||||
|
n_batches = new_tok = cached_tok = 0
|
||||||
|
for line in text.splitlines():
|
||||||
|
if "TP0]" not in line or "Prefill batch" not in line:
|
||||||
|
continue
|
||||||
|
m = pat.search(line)
|
||||||
|
if m:
|
||||||
|
n_batches += 1
|
||||||
|
new_tok += int(m.group(1))
|
||||||
|
cached_tok += int(m.group(2))
|
||||||
|
total = new_tok + cached_tok
|
||||||
|
return {
|
||||||
|
"prefill_batches": n_batches,
|
||||||
|
"new_tokens": new_tok,
|
||||||
|
"cached_tokens": cached_tok,
|
||||||
|
"hit_rate": round(cached_tok / total, 4) if total else None,
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
def stats(vals):
|
||||||
|
vals = [v for v in vals if v is not None]
|
||||||
|
if not vals:
|
||||||
|
return {"mean": None, "p50": None, "max": None}
|
||||||
|
s = sorted(vals)
|
||||||
|
return {
|
||||||
|
"mean": round(statistics.fmean(vals), 4),
|
||||||
|
"p50": round(s[len(s) // 2], 4),
|
||||||
|
"max": round(s[-1], 4),
|
||||||
|
"min": round(s[0], 4),
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
def main():
|
||||||
|
ap = argparse.ArgumentParser()
|
||||||
|
ap.add_argument("--concurrency", type=int, required=True)
|
||||||
|
ap.add_argument("--num-requests", type=int, required=True)
|
||||||
|
ap.add_argument("--run-id", type=int, required=True)
|
||||||
|
ap.add_argument("--arm-profile", default=None, help="output_dir to arm profiler after warmup")
|
||||||
|
ap.add_argument("--profile-steps", type=int, default=16)
|
||||||
|
args = ap.parse_args()
|
||||||
|
|
||||||
|
shared, prompts = build_prompts(args.num_requests, args.run_id)
|
||||||
|
warmup(shared)
|
||||||
|
if args.arm_profile:
|
||||||
|
arm_profile(args.arm_profile, args.profile_steps)
|
||||||
|
|
||||||
|
results = {}
|
||||||
|
t_start = time.time()
|
||||||
|
t0 = time.perf_counter()
|
||||||
|
with ThreadPoolExecutor(max_workers=args.concurrency) as ex:
|
||||||
|
futs = [ex.submit(bench_one, p, i, results) for i, p in enumerate(prompts)]
|
||||||
|
for f in futs:
|
||||||
|
f.result()
|
||||||
|
wall = time.perf_counter() - t0
|
||||||
|
t_end = time.time()
|
||||||
|
|
||||||
|
hit = verify_hit_rate(t_start, t_end)
|
||||||
|
|
||||||
|
ok = [r for r in results.values() if r.get("ok")]
|
||||||
|
n_out_total = sum(r["n_out"] for r in ok)
|
||||||
|
out_tps = [r["n_out"] / r["e2e"] for r in ok if r.get("e2e")]
|
||||||
|
ttft = stats([r.get("ttft") for r in ok])
|
||||||
|
tpot = stats([r.get("tpot") for r in ok])
|
||||||
|
e2e = stats([r.get("e2e") for r in ok])
|
||||||
|
retr = sum(r.get("retractions") or 0 for r in ok)
|
||||||
|
|
||||||
|
summary = {
|
||||||
|
"concurrency": args.concurrency,
|
||||||
|
"num_requests": args.num_requests,
|
||||||
|
"run_id": args.run_id,
|
||||||
|
"ok": len(ok),
|
||||||
|
"failed": args.num_requests - len(ok),
|
||||||
|
"wall_s": round(wall, 2),
|
||||||
|
"input_len": INPUT_LEN,
|
||||||
|
"shared_len": SHARED_LEN,
|
||||||
|
"output_len": OUTPUT_LEN,
|
||||||
|
"output_tokens_total": n_out_total,
|
||||||
|
"output_throughput_tok_s": round(n_out_total / wall, 2) if wall else None,
|
||||||
|
"input_throughput_tok_s": round(INPUT_LEN * len(ok) / wall, 2) if wall else None,
|
||||||
|
"ttft_s": ttft,
|
||||||
|
"tpot_s": tpot,
|
||||||
|
"e2e_s": e2e,
|
||||||
|
"per_req_out_tok_s": stats(out_tps),
|
||||||
|
"retractions_total": retr,
|
||||||
|
"cache_hit_from_logs": hit,
|
||||||
|
}
|
||||||
|
print("\n===== SUMMARY =====")
|
||||||
|
print(json.dumps(summary, indent=2), flush=True)
|
||||||
|
|
||||||
|
|
||||||
|
if __name__ == "__main__":
|
||||||
|
main()
|
||||||
@ -0,0 +1,22 @@
|
|||||||
|
#!/usr/bin/env python3
|
||||||
|
"""Event-category census of torch profiler traces: count + total duration per cat."""
|
||||||
|
import gzip
|
||||||
|
import json
|
||||||
|
import sys
|
||||||
|
from collections import defaultdict
|
||||||
|
|
||||||
|
for path in sys.argv[1:]:
|
||||||
|
with gzip.open(path, "rt") as f:
|
||||||
|
data = json.load(f)
|
||||||
|
events = data.get("traceEvents", []) if isinstance(data, dict) else data
|
||||||
|
cnt = defaultdict(int)
|
||||||
|
dur = defaultdict(float)
|
||||||
|
for e in events:
|
||||||
|
if e.get("ph") != "X":
|
||||||
|
continue
|
||||||
|
c = e.get("cat", "?")
|
||||||
|
cnt[c] += 1
|
||||||
|
dur[c] += e.get("dur", 0) / 1000.0
|
||||||
|
print(f"\n== {path.split('/')[-1]} ==")
|
||||||
|
for c in sorted(cnt, key=lambda k: -dur[k]):
|
||||||
|
print(f" {c:<18} n={cnt[c]:<9} total={dur[c]:>12.1f}ms")
|
||||||
@ -0,0 +1,63 @@
|
|||||||
|
#!/bin/bash
|
||||||
|
# GLM-5.3-NVFP4 并行方案实验部署（60.7/60.8, 8x6000D）
|
||||||
|
# 用法: bash deploy_par.sh "<并行与覆盖参数>" [mtp|nomtp] [chunk] [memfrac]
|
||||||
|
# 例: bash deploy_par.sh "--tp 4 --pp-size 2 --disable-overlap-schedule --max-prefill-tokens 16384" nomtp
|
||||||
|
# 例: bash deploy_par.sh "--tp 8 --dp-size 8 --enable-dp-attention --ep-size 8 --max-prefill-tokens 16384" mtp
|
||||||
|
# 注意: PP 与 overlap/MTP 不兼容（nightly 硬约束），PP 方案必须 nomtp + --disable-overlap-schedule
|
||||||
|
PAR=${1:?usage: deploy_par.sh "<flags>" [mtp|nomtp] [chunk] [memfrac]}
|
||||||
|
MTPMODE=${2:-nomtp}
|
||||||
|
CHUNK=${3:-8192}
|
||||||
|
MEMFRAC=${4:-0.88}
|
||||||
|
|
||||||
|
MTPARGS=""
|
||||||
|
if [ "$MTPMODE" = "mtp" ]; then
|
||||||
|
MTPARGS="--speculative-algorithm EAGLE --speculative-num-steps 3 --speculative-eagle-topk 1 --speculative-num-draft-tokens 4"
|
||||||
|
fi
|
||||||
|
|
||||||
|
# 顽固容器清理（复用 deploy_glm53.sh 逻辑）
|
||||||
|
docker update --restart=no glm53-nvfp4 >/dev/null 2>&1
|
||||||
|
for i in 1 2 3 4 5; do
|
||||||
|
docker rm -f glm53-nvfp4 >/dev/null 2>&1
|
||||||
|
sleep 2
|
||||||
|
docker ps -a --format '{{.Names}}' 2>/dev/null | grep -q '^glm53-nvfp4$' || break
|
||||||
|
done
|
||||||
|
if docker ps -a --format '{{.Names}}' 2>/dev/null | grep -q '^glm53-nvfp4$'; then
|
||||||
|
echo "ERROR: old container cannot be removed"; exit 1
|
||||||
|
fi
|
||||||
|
for i in $(seq 1 15); do ss -ltn 2>/dev/null | grep -q ":30000 " || break; sleep 2; done
|
||||||
|
|
||||||
|
docker run -d --name glm53-nvfp4 --gpus all --shm-size 64g --ipc=host \
|
||||||
|
-v /root/sglang_patch/server_args.py:/sgl-workspace/sglang/python/sglang/srt/server_args.py:ro \
|
||||||
|
-v /root/sglang_patch/layer_setup.py:/sgl-workspace/sglang/python/sglang/srt/model_executor/model_runner_components/layer_setup.py:ro \
|
||||||
|
--restart no -p 30000:30000 \
|
||||||
|
-v /data/hf_models:/data/hf_models \
|
||||||
|
lmsysorg/sglang:nightly-dev-20260828-daf63171 \
|
||||||
|
python3 -m sglang.launch_server \
|
||||||
|
--model-path /data/hf_models/GLM-5.3-NVFP4 \
|
||||||
|
--tp 8 \
|
||||||
|
--mem-fraction-static ${MEMFRAC} \
|
||||||
|
--max-running-requests 16 \
|
||||||
|
--chunked-prefill-size ${CHUNK} \
|
||||||
|
--disable-shared-experts-fusion \
|
||||||
|
--moe-runner-backend flashinfer_cutlass \
|
||||||
|
--disable-flashinfer-autotune \
|
||||||
|
--reasoning-parser glm45 --tool-call-parser glm47 \
|
||||||
|
--enable-hierarchical-cache --hicache-ratio 3 \
|
||||||
|
${MTPARGS} \
|
||||||
|
${PAR} \
|
||||||
|
--host 0.0.0.0 --port 30000
|
||||||
|
|
||||||
|
echo "deployed: par=[${PAR}] mtp=${MTPMODE}; waiting for health..."
|
||||||
|
for i in $(seq 10 10 1800); do
|
||||||
|
code=$(curl -s -o /dev/null -m3 -w '%{http_code}' http://127.0.0.1:30000/health 2>/dev/null)
|
||||||
|
if [ "$code" = "200" ]; then
|
||||||
|
echo "healthy after ${i}s"
|
||||||
|
docker logs glm53-nvfp4 2>&1 | grep -oE "max_total_num_tokens = [0-9]+" | head -1
|
||||||
|
exit 0
|
||||||
|
fi
|
||||||
|
if ! docker ps --format '{{.Names}}' | grep -q '^glm53-nvfp4$'; then
|
||||||
|
echo "CONTAINER DIED after ${i}s"; docker logs --tail 60 glm53-nvfp4 2>&1 | grep -iE "error|assert|not support|incompatible" | tail -8; exit 1
|
||||||
|
fi
|
||||||
|
sleep 10
|
||||||
|
done
|
||||||
|
echo "TIMEOUT waiting for health"; exit 1
|
||||||
@ -0,0 +1,48 @@
|
|||||||
|
#!/usr/bin/env python3
|
||||||
|
"""NCCL event forensics: timeline + duration distribution of AllReduce / SendRecv."""
|
||||||
|
import gzip
|
||||||
|
import json
|
||||||
|
import sys
|
||||||
|
|
||||||
|
def load(path):
|
||||||
|
with gzip.open(path, "rt") as f:
|
||||||
|
data = json.load(f)
|
||||||
|
events = data.get("traceEvents", []) if isinstance(data, dict) else data
|
||||||
|
base = (data.get("baseTimeNanoseconds", 0) / 1000.0) if isinstance(data, dict) else 0.0
|
||||||
|
ar, sr = [], []
|
||||||
|
for e in events:
|
||||||
|
if e.get("ph") != "X" or e.get("cat") != "kernel":
|
||||||
|
continue
|
||||||
|
n = e.get("name", "")
|
||||||
|
if "AllReduce" in n:
|
||||||
|
ar.append((e["ts"] + base, e["dur"]))
|
||||||
|
elif "SendRecv" in n:
|
||||||
|
sr.append((e["ts"] + base, e["dur"]))
|
||||||
|
return ar, sr
|
||||||
|
|
||||||
|
def dist(name, evs):
|
||||||
|
if not evs:
|
||||||
|
print(f" {name}: none")
|
||||||
|
return
|
||||||
|
durs = sorted(d for _, d in evs)
|
||||||
|
n = len(durs)
|
||||||
|
def pct(p): return durs[min(n - 1, int(n * p))]
|
||||||
|
print(f" {name}: n={n} total={sum(durs)/1e3:.1f}ms "
|
||||||
|
f"min={durs[0]/1e3:.2f} p10={pct(0.10)/1e3:.2f} p50={pct(0.50)/1e3:.2f} "
|
||||||
|
f"p90={pct(0.90)/1e3:.2f} p99={pct(0.99)/1e3:.2f} max={durs[-1]/1e3:.2f} (ms)")
|
||||||
|
# time-clustering: bucket into 2s buckets, report top-5 busiest buckets
|
||||||
|
t0 = min(t for t, _ in evs)
|
||||||
|
buckets = {}
|
||||||
|
for t, d in evs:
|
||||||
|
b = int((t - t0) // 2e6)
|
||||||
|
buckets[b] = buckets.get(b, 0.0) + d
|
||||||
|
top = sorted(buckets.items(), key=lambda x: -x[1])[:6]
|
||||||
|
span_s = (max(t for t, _ in evs) - t0) / 1e6
|
||||||
|
print(f" span {span_s:.1f}s; top 2s-buckets (offset_s -> ms): " +
|
||||||
|
", ".join(f"{b*2:.0f}s->{v/1e3:.0f}ms" for b, v in top))
|
||||||
|
|
||||||
|
for path in sys.argv[1:]:
|
||||||
|
print(f"== {path.split('/')[-1]} ==")
|
||||||
|
ar, sr = load(path)
|
||||||
|
dist("AllReduce", ar)
|
||||||
|
dist("SendRecv ", sr)
|
||||||
@ -0,0 +1,237 @@
|
|||||||
|
#!/usr/bin/env python3
|
||||||
|
"""PP stage alignment analysis for sglang torch-profiler traces (TP4PP2).
|
||||||
|
|
||||||
|
Reads two or more per-rank traces (e.g. PP0-TP0 and PP1-TP0), aligns them on the
|
||||||
|
wall clock via baseTimeNanoseconds, and reports:
|
||||||
|
- per-rank: GPU busy/span, kernel count, NCCL breakdown (AllReduce vs Send/Recv),
|
||||||
|
cudaLaunchKernel count (eager vs graph evidence), top kernels
|
||||||
|
- cross-stage 1ms-bin classification: both-busy / only-PP0 / only-PP1 / neither
|
||||||
|
(only-one-stage time = pipeline bubble; neither = scheduler/CPU idle)
|
||||||
|
- forward segmentation (ProfilerStep annotations if present, else gap-based)
|
||||||
|
with per-forward wall, GPU-covered time and inter-forward CPU gap
|
||||||
|
|
||||||
|
Usage:
|
||||||
|
python3 pp_align.py --trace-pp0 <file> --trace-pp1 <file> [--gap-ms 4] [--bin-ms 1]
|
||||||
|
python3 pp_align.py --trace <file> [--gap-ms 4] # single-rank mode
|
||||||
|
"""
|
||||||
|
import argparse
|
||||||
|
import gzip
|
||||||
|
import json
|
||||||
|
import sys
|
||||||
|
from collections import defaultdict
|
||||||
|
|
||||||
|
|
||||||
|
def open_maybe_gz(path):
|
||||||
|
if path.endswith(".gz"):
|
||||||
|
return gzip.open(path, "rt")
|
||||||
|
return open(path, "rt")
|
||||||
|
|
||||||
|
|
||||||
|
def load_trace(path):
|
||||||
|
with open_maybe_gz(path) as f:
|
||||||
|
data = json.load(f)
|
||||||
|
if isinstance(data, list):
|
||||||
|
events, base_ns = data, None
|
||||||
|
else:
|
||||||
|
events = data.get("traceEvents", [])
|
||||||
|
base_ns = data.get("baseTimeNanoseconds")
|
||||||
|
base_us = (base_ns / 1000.0) if base_ns else 0.0
|
||||||
|
kernels = [] # (start_us_abs, dur_us, name)
|
||||||
|
steps = [] # (start_us_abs, dur_us, name) ProfilerStep annotations
|
||||||
|
launch_cnt = 0
|
||||||
|
launch_time = 0.0
|
||||||
|
memcpy_time = 0.0
|
||||||
|
for e in events:
|
||||||
|
if e.get("ph") != "X":
|
||||||
|
continue
|
||||||
|
cat = e.get("cat", "")
|
||||||
|
ts = e.get("ts", 0)
|
||||||
|
dur = e.get("dur", 0)
|
||||||
|
if cat == "kernel":
|
||||||
|
kernels.append((ts + base_us, dur, e.get("name", "?")))
|
||||||
|
elif cat in ("gpu_memcpy", "gpu_memset"):
|
||||||
|
memcpy_time += dur
|
||||||
|
kernels.append((ts + base_us, dur, "::<memcpy/memset>"))
|
||||||
|
elif cat == "user_annotation" and str(e.get("name", "")).startswith("ProfilerStep"):
|
||||||
|
steps.append((ts + base_us, dur, e.get("name", "?")))
|
||||||
|
elif cat == "cuda_runtime":
|
||||||
|
if e.get("name") == "cudaLaunchKernel":
|
||||||
|
launch_cnt += 1
|
||||||
|
launch_time += dur
|
||||||
|
kernels.sort()
|
||||||
|
steps.sort()
|
||||||
|
return {
|
||||||
|
"path": path,
|
||||||
|
"kernels": kernels,
|
||||||
|
"steps": steps,
|
||||||
|
"launch_cnt": launch_cnt,
|
||||||
|
"launch_time": launch_time,
|
||||||
|
"memcpy_time": memcpy_time,
|
||||||
|
"base_ok": base_ns is not None,
|
||||||
|
}
|
||||||
|
|
||||||
|
|
||||||
|
def union_len(intervals):
|
||||||
|
if not intervals:
|
||||||
|
return 0.0
|
||||||
|
intervals = sorted(intervals)
|
||||||
|
tot = 0.0
|
||||||
|
cs, ce = intervals[0]
|
||||||
|
for s, e in intervals[1:]:
|
||||||
|
if s <= ce:
|
||||||
|
ce = max(ce, e)
|
||||||
|
else:
|
||||||
|
tot += ce - cs
|
||||||
|
cs, ce = s, e
|
||||||
|
tot += ce - cs
|
||||||
|
return tot
|
||||||
|
|
||||||
|
|
||||||
|
def classify_nccl(name):
|
||||||
|
if "nccl" not in name.lower():
|
||||||
|
return None
|
||||||
|
if "AllReduce" in name:
|
||||||
|
return "AllReduce"
|
||||||
|
if "SendRecv" in name or "Send" in name or "Recv" in name or "P2P" in name:
|
||||||
|
return "SendRecv"
|
||||||
|
if "Broadcast" in name:
|
||||||
|
return "Broadcast"
|
||||||
|
if "AllGather" in name or "ReduceScatter" in name:
|
||||||
|
return "AG_RS"
|
||||||
|
return "nccl_other"
|
||||||
|
|
||||||
|
|
||||||
|
def rank_report(r, gap_ms):
|
||||||
|
ks = r["kernels"]
|
||||||
|
if not ks:
|
||||||
|
print(f" !! no kernel events in {r['path']}")
|
||||||
|
return
|
||||||
|
span_lo = ks[0][0]
|
||||||
|
span_hi = ks[-1][0] + ks[-1][1]
|
||||||
|
span = span_hi - span_lo
|
||||||
|
busy = union_len([(s, s + d) for s, d, _ in ks])
|
||||||
|
nccl = defaultdict(lambda: [0.0, 0])
|
||||||
|
topk = defaultdict(lambda: [0.0, 0])
|
||||||
|
|
||||||
|
def simp(n):
|
||||||
|
return n.split("<")[0][:100]
|
||||||
|
|
||||||
|
for s, d, n in ks:
|
||||||
|
c = classify_nccl(n)
|
||||||
|
if c:
|
||||||
|
nccl[c][0] += d
|
||||||
|
nccl[c][1] += 1
|
||||||
|
topk[simp(n)][0] += d
|
||||||
|
topk[simp(n)][1] += 1
|
||||||
|
|
||||||
|
print(f" file: {r['path']}")
|
||||||
|
print(f" baseTimeNanoseconds present: {r['base_ok']}")
|
||||||
|
print(f" span {span/1e6:.3f}s GPU-busy {busy/1e6:.3f}s ({100*busy/span:.1f}%) kernels {len(ks)}")
|
||||||
|
print(f" cudaMemcpy/Set GPU time: {r['memcpy_time']/1e3:.1f}ms cudaLaunchKernel calls: {r['launch_cnt']} ({r['launch_time']/1e3:.1f}ms CPU)")
|
||||||
|
if nccl:
|
||||||
|
tot_nccl = sum(v[0] for v in nccl.values())
|
||||||
|
print(f" NCCL total: {tot_nccl/1e3:.1f}ms ({100*tot_nccl/busy:.1f}% of busy, {100*tot_nccl/span:.1f}% of span)")
|
||||||
|
for k, (t, c) in sorted(nccl.items(), key=lambda x: -x[1][0]):
|
||||||
|
print(f" {k:<12} {t/1e3:9.1f}ms n={c}")
|
||||||
|
print(" top kernels (of GPU-busy):")
|
||||||
|
for n, (t, c) in sorted(topk.items(), key=lambda x: -x[1][0])[:15]:
|
||||||
|
print(f" {100*t/busy:5.1f}% {t/1e3:9.1f}ms n={c:<7} {n[:80]}")
|
||||||
|
|
||||||
|
# forward segmentation
|
||||||
|
fwd = []
|
||||||
|
if r["steps"]:
|
||||||
|
fwd = [(s, s + d) for s, d, _ in r["steps"]]
|
||||||
|
src = "ProfilerStep annotations"
|
||||||
|
else:
|
||||||
|
gap = gap_ms * 1000.0
|
||||||
|
cs = ks[0][0]
|
||||||
|
prev_end = ks[0][0] + ks[0][1]
|
||||||
|
for s, d, _ in ks[1:]:
|
||||||
|
e = s + d
|
||||||
|
if s - prev_end > gap:
|
||||||
|
fwd.append((cs, prev_end))
|
||||||
|
cs = s
|
||||||
|
prev_end = max(prev_end, e)
|
||||||
|
fwd.append((cs, prev_end))
|
||||||
|
src = f"gap>{gap_ms}ms split"
|
||||||
|
walls = [e - s for s, e in fwd]
|
||||||
|
covs = [union_len([(s, s + d) for s, d, _ in ks if s >= fs and s + d <= fe]) / (fe - fs)
|
||||||
|
for fs, fe in fwd]
|
||||||
|
gaps = []
|
||||||
|
for i in range(1, len(fwd)):
|
||||||
|
gaps.append(fwd[i][0] - fwd[i - 1][1])
|
||||||
|
print(f" forwards: {len(fwd)} ({src})")
|
||||||
|
if walls:
|
||||||
|
ws = sorted(walls)
|
||||||
|
print(f" wall ms: mean {sum(walls)/len(walls)/1e3:.1f} p50 {ws[len(ws)//2]/1e3:.1f} min {ws[0]/1e3:.1f} max {ws[-1]/1e3:.1f}")
|
||||||
|
small = [w for w in walls if w < 0.3 * ws[len(ws) // 2]]
|
||||||
|
print(f" small forwards (<30% of p50): {len(small)} (warmup/mixed steps)")
|
||||||
|
main_cov = [c for w, c in zip(walls, covs) if w >= 0.3 * ws[len(ws) // 2]]
|
||||||
|
if main_cov:
|
||||||
|
print(f" GPU-covered fraction inside main forwards: mean {sum(main_cov)/len(main_cov)*100:.1f}%")
|
||||||
|
if gaps:
|
||||||
|
gs = sorted(gaps)
|
||||||
|
print(f" inter-forward gap ms: mean {sum(gaps)/len(gaps)/1e3:.1f} p50 {gs[len(gs)//2]/1e3:.1f} p90 {gs[int(len(gs)*0.9)]/1e3:.1f} max {gs[-1]/1e3:.1f}")
|
||||||
|
return fwd
|
||||||
|
|
||||||
|
|
||||||
|
def cross_report(r0, r1, bin_ms):
|
||||||
|
ks0 = [(s, s + d) for s, d, _ in r0["kernels"]]
|
||||||
|
ks1 = [(s, s + d) for s, d, _ in r1["kernels"]]
|
||||||
|
lo = min(ks0[0][0], ks1[0][0])
|
||||||
|
hi = max(ks0[-1][1], ks1[-1][1])
|
||||||
|
nb = int((hi - lo) / (bin_ms * 1000)) + 1
|
||||||
|
b0 = bytearray(nb)
|
||||||
|
b1 = bytearray(nb)
|
||||||
|
for s, e in ks0:
|
||||||
|
i0, i1 = int((s - lo) // (bin_ms * 1000)), int((e - lo) // (bin_ms * 1000))
|
||||||
|
for i in range(max(0, i0), min(nb, i1 + 1)):
|
||||||
|
b0[i] = 1
|
||||||
|
for s, e in ks1:
|
||||||
|
i0, i1 = int((s - lo) // (bin_ms * 1000)), int((e - lo) // (bin_ms * 1000))
|
||||||
|
for i in range(max(0, i0), min(nb, i1 + 1)):
|
||||||
|
b1[i] = 1
|
||||||
|
both = sum(1 for i in range(nb) if b0[i] and b1[i])
|
||||||
|
only0 = sum(1 for i in range(nb) if b0[i] and not b1[i])
|
||||||
|
only1 = sum(1 for i in range(nb) if b1[i] and not b0[i])
|
||||||
|
neither = nb - both - only0 - only1
|
||||||
|
print(f"\n== cross-stage bins ({bin_ms}ms, total {nb*bin_ms/1000:.2f}s) ==")
|
||||||
|
print(f" both busy : {both*bin_ms/1000:8.2f}s {100*both/nb:5.1f}%")
|
||||||
|
print(f" only PP0 : {only0*bin_ms/1000:8.2f}s {100*only0/nb:5.1f}% <- PP1 bubble")
|
||||||
|
print(f" only PP1 : {only1*bin_ms/1000:8.2f}s {100*only1/nb:5.1f}% <- PP0 bubble")
|
||||||
|
print(f" neither : {neither*bin_ms/1000:8.2f}s {100*neither/nb:5.1f}% <- sched/CPU idle")
|
||||||
|
busy_union = both + only0 + only1
|
||||||
|
print(f" pipeline efficiency: both/(any busy) = {100*both/max(1,busy_union):.1f}%")
|
||||||
|
if not (r0["base_ok"] and r1["base_ok"]):
|
||||||
|
print(" !! baseTimeNanoseconds missing on at least one trace: cross-stage alignment unreliable")
|
||||||
|
|
||||||
|
|
||||||
|
def main():
|
||||||
|
ap = argparse.ArgumentParser()
|
||||||
|
ap.add_argument("--trace-pp0")
|
||||||
|
ap.add_argument("--trace-pp1")
|
||||||
|
ap.add_argument("--trace", action="append", default=[], help="single-rank mode, repeatable")
|
||||||
|
ap.add_argument("--gap-ms", type=float, default=4.0)
|
||||||
|
ap.add_argument("--bin-ms", type=float, default=1.0)
|
||||||
|
args = ap.parse_args()
|
||||||
|
|
||||||
|
if args.trace_pp0 and args.trace_pp1:
|
||||||
|
print("== PP0 rank ==")
|
||||||
|
r0 = load_trace(args.trace_pp0)
|
||||||
|
rank_report(r0, args.gap_ms)
|
||||||
|
print("\n== PP1 rank ==")
|
||||||
|
r1 = load_trace(args.trace_pp1)
|
||||||
|
rank_report(r1, args.gap_ms)
|
||||||
|
cross_report(r0, r1, args.bin_ms)
|
||||||
|
elif args.trace:
|
||||||
|
for t in args.trace:
|
||||||
|
print(f"\n== rank: {t} ==")
|
||||||
|
r = load_trace(t)
|
||||||
|
rank_report(r, args.gap_ms)
|
||||||
|
else:
|
||||||
|
ap.print_usage()
|
||||||
|
sys.exit(1)
|
||||||
|
|
||||||
|
|
||||||
|
if __name__ == "__main__":
|
||||||
|
main()
|
||||||
@ -0,0 +1,14 @@
|
|||||||
|
#!/usr/bin/env python3
|
||||||
|
"""Scenario C probe: one 16384-token random request, 1 output token (host-side)."""
|
||||||
|
import random
|
||||||
|
import requests
|
||||||
|
|
||||||
|
rng = random.Random(9090)
|
||||||
|
ids = [rng.randint(1000, 100000) for _ in range(16384)]
|
||||||
|
r = requests.post(
|
||||||
|
"http://127.0.0.1:30000/generate",
|
||||||
|
json={"input_ids": ids,
|
||||||
|
"sampling_params": {"max_new_tokens": 1, "temperature": 0.0, "ignore_eos": True}},
|
||||||
|
timeout=600,
|
||||||
|
)
|
||||||
|
print("C probe:", r.status_code, repr(r.json().get("text", ""))[:80] if r.status_code == 200 else r.text[:300])
|
||||||
@ -0,0 +1,15 @@
|
|||||||
|
#!/usr/bin/env python3
|
||||||
|
"""C-v2 probe: small 2048-token request, 2 output tokens (1 prefill + 2 decode = 3 forwards,
|
||||||
|
auto-stops a num_steps=2 profiler arm with the +1 guard step)."""
|
||||||
|
import random
|
||||||
|
import requests
|
||||||
|
|
||||||
|
rng = random.Random(9091)
|
||||||
|
ids = [rng.randint(1000, 100000) for _ in range(2048)]
|
||||||
|
r = requests.post(
|
||||||
|
"http://127.0.0.1:30000/generate",
|
||||||
|
json={"input_ids": ids,
|
||||||
|
"sampling_params": {"max_new_tokens": 2, "temperature": 0.0, "ignore_eos": True}},
|
||||||
|
timeout=300,
|
||||||
|
)
|
||||||
|
print("C-v2 probe:", r.status_code, repr(r.json().get("text", ""))[:60] if r.status_code == 200 else r.text[:200])
|
||||||
@ -0,0 +1,89 @@
|
|||||||
|
#!/bin/bash
|
||||||
|
# TP4PP2 torch-profiler 采集编排（60.5, GLM-5.3-NVFP4, container glm53-nvfp4:30000）
|
||||||
|
# 用法: bash profile_cap.sh <A|B|C>
|
||||||
|
# A = prefill 主场景: 16x(16384->1) cc16, arm num_steps=36, 输出 A_prefill/
|
||||||
|
# B = decode 主场景: 16x(16384->512) cc16, 等 decode 相位后 arm num_steps=60, 输出 B_decode/
|
||||||
|
# C = mapping 小采集: with_stack 3 步 + 1x(16384->1), 输出 C_map/
|
||||||
|
set -u
|
||||||
|
BASE=/data/hf_models/pp2tp4_profiles
|
||||||
|
URL=http://127.0.0.1:30000
|
||||||
|
mkdir -p /root/bs_results /data/hf_models/bs_results
|
||||||
|
|
||||||
|
health() { [ "$(curl -s -o /dev/null -m3 -w '%{http_code}' $URL/health)" = "200" ]; }
|
||||||
|
|
||||||
|
warm_and_flush() {
|
||||||
|
curl -s -X POST $URL/flush_cache >/dev/null; sleep 3
|
||||||
|
curl -s $URL/generate -H 'Content-Type: application/json' \
|
||||||
|
-d '{"input_ids":[11,22,33,44,55,66,77,88,99,111,222,333,444,555,666,777,888,999,1234,4321,2345,5432,3456,6543,4567,7654,5678,8765,6789,9876,7890,9012,1357,2468,3579,4680,5791,6802,7913,8024,9135,1024,2048,4096,8192,1638,3276,6552,1310,2620,5240,1048,2096,4192,8384,1676,3352,6704,134,268,536,1072,2144,4288],"sampling_params":{"max_new_tokens":8,"temperature":0.0,"ignore_eos":true}}' >/dev/null
|
||||||
|
curl -s -X POST $URL/flush_cache >/dev/null; sleep 3
|
||||||
|
echo "[warm_and_flush] done $(date +%T)"
|
||||||
|
}
|
||||||
|
|
||||||
|
arm() { # dir, num_steps, prefix, with_stack, record_shapes
|
||||||
|
mkdir -p "$1"
|
||||||
|
curl -s -X POST $URL/start_profile -H 'Content-Type: application/json' \
|
||||||
|
-d "{\"output_dir\":\"$1\",\"num_steps\":$2,\"activities\":[\"CPU\",\"GPU\"],\"with_stack\":$4,\"record_shapes\":$5,\"profile_prefix\":\"$3\"}"
|
||||||
|
echo; echo "[arm] prefix=$3 steps=$2 dir=$1 at $(date +%T)"
|
||||||
|
}
|
||||||
|
|
||||||
|
wait_traces() { # dir, min_files
|
||||||
|
local n=0 i
|
||||||
|
for i in $(seq 1 90); do
|
||||||
|
n=$(ls "$1"/*.trace.json.gz 2>/dev/null | wc -l)
|
||||||
|
if [ "$n" -ge "$2" ]; then echo "[traces] $n files ready after ~$((i*5))s"; ls -la "$1" | tail -12; return 0; fi
|
||||||
|
sleep 5
|
||||||
|
done
|
||||||
|
echo "[traces] TIMEOUT: only $n files in $1"; ls -la "$1" 2>/dev/null; return 1
|
||||||
|
}
|
||||||
|
|
||||||
|
run_bench() { # outfile_log, num_prompts, in_len, out_len, seed
|
||||||
|
docker exec glm53-nvfp4 python3 -m sglang.bench_serving \
|
||||||
|
--backend sglang --host 127.0.0.1 --port 30000 \
|
||||||
|
--dataset-name random-ids --tokenizer /data/hf_models/GLM-5.3-NVFP4 \
|
||||||
|
--num-prompts "$2" --random-input-len "$3" --random-output-len "$4" \
|
||||||
|
--random-range-ratio 1.0 --max-concurrency "$2" --temperature 0.0 \
|
||||||
|
--warmup-requests 0 --seed "$5" \
|
||||||
|
--output-file "/data/hf_models/bs_results/$6" \
|
||||||
|
> "$1" 2>&1
|
||||||
|
}
|
||||||
|
|
||||||
|
case $1 in
|
||||||
|
A)
|
||||||
|
health || { echo "server not healthy"; exit 1; }
|
||||||
|
SEED=$(( ($(date +%s) % 900000) + 101 ))
|
||||||
|
warm_and_flush
|
||||||
|
arm $BASE/A_prefill 36 pa false false
|
||||||
|
run_bench /root/bs_results/prof_A_bench.log 16 16384 1 $SEED prof_A_prefill_${SEED}.json
|
||||||
|
echo "[benchA] exit=$?"
|
||||||
|
grep -E "Input token throughput|Mean TTFT|Successful" /root/bs_results/prof_A_bench.log | tail -3
|
||||||
|
wait_traces $BASE/A_prefill 8
|
||||||
|
;;
|
||||||
|
B)
|
||||||
|
health || { echo "server not healthy"; exit 1; }
|
||||||
|
SEED=$(( ($(date +%s) % 900000) + 202 ))
|
||||||
|
warm_and_flush
|
||||||
|
D0=$(docker logs glm53-nvfp4 2>&1 | grep -c "Decode batch")
|
||||||
|
run_bench /root/bs_results/prof_B_bench.log 16 16384 512 $SEED prof_B_decode_${SEED}.json &
|
||||||
|
BPID=$!
|
||||||
|
for i in $(seq 1 120); do
|
||||||
|
D=$(docker logs glm53-nvfp4 2>&1 | grep -c "Decode batch")
|
||||||
|
[ "$D" -gt "$D0" ] && break
|
||||||
|
sleep 2
|
||||||
|
done
|
||||||
|
echo "[B] decode phase after ~$((i*2))s (decode-lines D0=$D0 -> D=$D)"
|
||||||
|
arm $BASE/B_decode 60 pb false false
|
||||||
|
wait $BPID; echo "[benchB] exit=$?"
|
||||||
|
grep -E "Output token throughput|Mean TPOT|Mean TTFT|Successful" /root/bs_results/prof_B_bench.log | tail -4
|
||||||
|
wait_traces $BASE/B_decode 8
|
||||||
|
;;
|
||||||
|
C)
|
||||||
|
health || { echo "server not healthy"; exit 1; }
|
||||||
|
warm_and_flush
|
||||||
|
arm $BASE/C_map 3 pc true true
|
||||||
|
python3 /root/prof_tools/probe_c.py
|
||||||
|
echo "[probeC] exit=$?"
|
||||||
|
wait_traces $BASE/C_map 8
|
||||||
|
;;
|
||||||
|
*)
|
||||||
|
echo "usage: bash profile_cap.sh <A|B|C>"; exit 1 ;;
|
||||||
|
esac
|
||||||
@ -0,0 +1,67 @@
|
|||||||
|
#!/bin/bash
|
||||||
|
# Phase 2 编排：等健康 -> 对照bench(无profiler) -> D采集(128k/90%) -> C-v2(小probe with_stack)
|
||||||
|
set -u
|
||||||
|
URL=http://127.0.0.1:30000
|
||||||
|
BASE=/data/hf_models/pp2tp4_profiles
|
||||||
|
BS=/data/hf_models/bs_results
|
||||||
|
mkdir -p $BS /root/bs_results
|
||||||
|
|
||||||
|
echo "[phase2] waiting for health..."
|
||||||
|
for i in $(seq 1 90); do
|
||||||
|
[ "$(curl -s -o /dev/null -m3 -w '%{http_code}' $URL/health)" = "200" ] && break
|
||||||
|
sleep 10
|
||||||
|
done
|
||||||
|
if [ "$(curl -s -o /dev/null -m3 -w '%{http_code}' $URL/health)" != "200" ]; then
|
||||||
|
echo "[phase2] FATAL: server not healthy after 900s"; exit 1
|
||||||
|
fi
|
||||||
|
echo "[phase2] healthy. smoke gsm8k..."
|
||||||
|
curl -s $URL/generate -H 'Content-Type: application/json' \
|
||||||
|
-d '{"text":"3支铅笔每支2元加1个笔记本8元，付20元找回几元？只答数字。","sampling_params":{"max_new_tokens":64,"temperature":0}}' \
|
||||||
|
| head -c 200; echo
|
||||||
|
|
||||||
|
# 1) clean B bench (no profiler)
|
||||||
|
SEED=$(( ($(date +%s) % 900000) + 303 ))
|
||||||
|
echo "[phase2] clean B bench seed=$SEED start $(date +%T)"
|
||||||
|
docker exec glm53-nvfp4 python3 -m sglang.bench_serving \
|
||||||
|
--backend sglang --host 127.0.0.1 --port 30000 \
|
||||||
|
--dataset-name random-ids --tokenizer /data/hf_models/GLM-5.3-NVFP4 \
|
||||||
|
--num-prompts 16 --random-input-len 16384 --random-output-len 512 \
|
||||||
|
--random-range-ratio 1.0 --max-concurrency 16 --temperature 0.0 \
|
||||||
|
--flush-cache --warmup-requests 1 --seed $SEED \
|
||||||
|
--output-file $BS/cleanB_cc16_${SEED}.json \
|
||||||
|
> /root/bs_results/cleanB.log 2>&1
|
||||||
|
echo "[phase2] clean B exit=$? $(date +%T)"
|
||||||
|
grep -E "Input token throughput|Output token throughput|Mean TTFT|Mean TPOT|Mean ITL|Successful" /root/bs_results/cleanB.log | tail -6
|
||||||
|
|
||||||
|
# 2) D capture (128k / 90% prefix hit, cc8)
|
||||||
|
RUNID=$(( ($(date +%s) % 90000) + 600 ))
|
||||||
|
echo "[phase2] D capture run-id=$RUNID start $(date +%T)"
|
||||||
|
mkdir -p $BASE/D_128k
|
||||||
|
python3 /root/prof_tools/bench_128k_hit90_prof.py \
|
||||||
|
--concurrency 8 --num-requests 8 --run-id $RUNID \
|
||||||
|
--arm-profile $BASE/D_128k --profile-steps 16 \
|
||||||
|
> /root/bs_results/prof_D_bench.log 2>&1
|
||||||
|
echo "[phase2] D bench exit=$? $(date +%T)"
|
||||||
|
grep -E "warmup|arm|ttft|wall_s|output_throughput|hit_rate|new_tokens|cached_tokens" /root/bs_results/prof_D_bench.log | head -12
|
||||||
|
for i in $(seq 1 90); do
|
||||||
|
n=$(ls $BASE/D_128k/*.trace.json.gz 2>/dev/null | wc -l)
|
||||||
|
[ "$n" -ge 8 ] && { echo "[phase2] D traces ready: $n files after ~$((i*5))s"; break; }
|
||||||
|
sleep 5
|
||||||
|
done
|
||||||
|
ls -la $BASE/D_128k | tail -10
|
||||||
|
|
||||||
|
# 3) C-v2 small with_stack mapping probe
|
||||||
|
echo "[phase2] C-v2 start $(date +%T)"
|
||||||
|
mkdir -p $BASE/C_map_v2
|
||||||
|
curl -s -X POST $URL/flush_cache > /dev/null; sleep 3
|
||||||
|
curl -s -X POST $URL/start_profile -H 'Content-Type: application/json' \
|
||||||
|
-d "{\"output_dir\":\"$BASE/C_map_v2\",\"num_steps\":2,\"activities\":[\"CPU\",\"GPU\"],\"with_stack\":true,\"record_shapes\":false,\"profile_prefix\":\"pc2\"}"
|
||||||
|
echo
|
||||||
|
python3 /root/prof_tools/probe_c_v2.py
|
||||||
|
for i in $(seq 1 60); do
|
||||||
|
n=$(ls $BASE/C_map_v2/*.trace.json.gz 2>/dev/null | wc -l)
|
||||||
|
[ "$n" -ge 8 ] && { echo "[phase2] C-v2 traces ready: $n files after ~$((i*5))s"; break; }
|
||||||
|
sleep 5
|
||||||
|
done
|
||||||
|
ls -la $BASE/C_map_v2 | tail -10
|
||||||
|
echo "[phase2] ALL DONE $(date +%T)"
|
||||||
Loading…
x
Reference in New Issue
Block a user