[Experiment] GLM-5.3-NVFP4 TP4PP2 torch-profiler profile on 174.1.60.5 (2026-09-06)
- Scenarios: A prefill 16k/1 cc16, B decode 16k/512 cc16, D 128k/90% prefix-hit cc8, E NCCL_PROTO=Simple protocol control
- Key findings: PP bubble disproven (96% both-stage busy, chunked prefill = natural microbatch pipeline);
prefill wall = PP0 chain compute 22.8s(65%) + TP4 AllReduce 10.4s(30%);
AR 96MB/3.7ms => busbw ~39GB/s ~ 75% of practical PCIe Gen5 (near-optimal, RING_LL kernel name is NCCL 2.28 mixed-plan naming);
NCCL_PROTO forcing refuted: docker -e never reaches scheduler ranks (sglang worker spawn drops NCCL env);
decode step 21-26ms (MoE GEMM 35%, SendRecv 18%, PP1 AllGather 273 launches outside cuda graph);
128k/90%: DSA indexer sm120_fp8_mqa_logits 28%/23% = 1.6x the sparse attention it serves;
with_stack pathological on this nightly (2 repros) - flush wedges server
- Full report: D:\sskj\reports\GLM53_TP4PP2_Profile与优化方向_2026-09-06.md
- Raw traces remain on 60.5:/data/hf_models/pp2tp4_profiles/{A_prefill,B_decode,D_128k,E_nccl_simple}/
This commit is contained in:
yy-fighting
parent
8189942353
commit
1e8c36b7d1
@ -0,0 +1,74 @@
|
||||
# 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 指标作废
|
||||
|
||||
## 优化方向（排序见报告 §4）
|
||||
|
||||
① 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,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,61 @@
|
||||
#!/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 \
|
||||
--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