diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/README.md b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/README.md new file mode 100644 index 0000000..e742406 --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/README.md @@ -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 --profile-steps 16`（D） +4. `python3 scripts/pp_align.py `；`python3 scripts/nccl_forensics.py `；`python3 scripts/nccl_names.py ` +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 进程并等显存归零）。 diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/config.env b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/config.env new file mode 100644 index 0000000..6768b0a --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/config.env @@ -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" diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/A_census.txt b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/A_census.txt new file mode 100644 index 0000000..195e4b2 --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/A_census.txt @@ -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 diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/A_nccl_forensics.txt b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/A_nccl_forensics.txt new file mode 100644 index 0000000..cb8f5ed --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/A_nccl_forensics.txt @@ -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 diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/A_ppalign_tp0.txt b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/A_ppalign_tp0.txt new file mode 100644 index 0000000..9581892 --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/A_ppalign_tp0.txt @@ -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% diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/A_triage_pp0tp0.txt b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/A_triage_pp0tp0.txt new file mode 100644 index 0000000..cbceb1f --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/A_triage_pp0tp0.txt @@ -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 | gemm | 3380.00 ms | 6.0% | 3180 | unresolved | cuLaunchKernel | +| void cutlass::Kernel2 | 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 > | 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
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
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. | diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/A_triage_pp1tp0.txt b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/A_triage_pp1tp0.txt new file mode 100644 index 0000000..f43fc63 --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/A_triage_pp1tp0.txt @@ -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 | 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 | 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 > | gemm | 668.88 ms | 1.9% | 1326 | unresolved | cudaLaunchKernelExC | +| void cutlass::Kernel2 | 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
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
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. | diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/B_ppalign_tp0.txt b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/B_ppalign_tp0.txt new file mode 100644 index 0000000..e6ea67e --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/B_ppalign_tp0.txt @@ -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% diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/B_triage_pp0tp0.txt b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/B_triage_pp0tp0.txt new file mode 100644 index 0000000..1aad1b9 --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/B_triage_pp0tp0.txt @@ -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 | 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 | 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 | 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
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
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. | diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/B_triage_pp1tp0.txt b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/B_triage_pp1tp0.txt new file mode 100644 index 0000000..04e7079 --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/B_triage_pp1tp0.txt @@ -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 | 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 | 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 | gemm | 21.88 ms | 1.6% | 60 | unresolved | cudaGraphLaunch | +| void cutlass::Kernel2 | 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
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
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. | diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/D_ppalign_tp0.txt b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/D_ppalign_tp0.txt new file mode 100644 index 0000000..64178d7 --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/D_ppalign_tp0.txt @@ -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% diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/E_nccl_forensics_tp0.txt b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/E_nccl_forensics_tp0.txt new file mode 100644 index 0000000..7b2f288 --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/E_nccl_forensics_tp0.txt @@ -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 diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/E_nccl_names_tp0.txt b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/E_nccl_names_tp0.txt new file mode 100644 index 0000000..f108a80 --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/results/analysis/E_nccl_names_tp0.txt @@ -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>) diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/scripts/bench_128k_hit90_prof.py b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/scripts/bench_128k_hit90_prof.py new file mode 100644 index 0000000..73b9b6e --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/scripts/bench_128k_hit90_prof.py @@ -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 --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() diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/scripts/census.py b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/scripts/census.py new file mode 100644 index 0000000..dc7ad60 --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/scripts/census.py @@ -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") diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/scripts/deploy_par_605.sh b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/scripts/deploy_par_605.sh new file mode 100644 index 0000000..aeb4f38 --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/scripts/deploy_par_605.sh @@ -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 "" [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 diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/scripts/nccl_forensics.py b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/scripts/nccl_forensics.py new file mode 100644 index 0000000..92909e2 --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/scripts/nccl_forensics.py @@ -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) diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/scripts/pp_align.py b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/scripts/pp_align.py new file mode 100644 index 0000000..3bffded --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/scripts/pp_align.py @@ -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 --trace-pp1 [--gap-ms 4] [--bin-ms 1] + python3 pp_align.py --trace [--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, "::")) + 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() diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/scripts/probe_c.py b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/scripts/probe_c.py new file mode 100644 index 0000000..a57be1d --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/scripts/probe_c.py @@ -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]) diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/scripts/probe_c_v2.py b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/scripts/probe_c_v2.py new file mode 100644 index 0000000..823b930 --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/scripts/probe_c_v2.py @@ -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]) diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/scripts/profile_cap.sh b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/scripts/profile_cap.sh new file mode 100644 index 0000000..674f64a --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/scripts/profile_cap.sh @@ -0,0 +1,89 @@ +#!/bin/bash +# TP4PP2 torch-profiler 采集编排（60.5, GLM-5.3-NVFP4, container glm53-nvfp4:30000） +# 用法: bash profile_cap.sh +# 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 "; exit 1 ;; +esac diff --git a/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/scripts/run_phase2.sh b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/scripts/run_phase2.sh new file mode 100644 index 0000000..74be2bb --- /dev/null +++ b/experiments/pro6000/glm53_nvfp4_pro6000d_sglang_tp4pp2_profile/scripts/run_phase2.sh @@ -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)"