minke.yu/sglang
1
0
Fork 0
Code Issues Pull Requests Actions 12 Packages Projects Releases Wiki Activity

Upgrade sglang-torch-profiler-analysis SKILLS (#22440)

Browse Source
This commit is contained in:
Xiaoyu Zhang
2026-04-09 18:23:03 +08:00
committed by GitHub
parent 9fed58805f
commit 7603b226ce
9 changed files with 1406 additions and 1259 deletions
Download Patch File Download Diff File Expand all files Collapse all files
.claude/skills/sglang-torch-profiler-analysis/SKILL.md
+86 -120
View File
@@ -1,58 +1,86 @@
---
name: sglang-torch-profiler-analysis
description: "Unified SGLang torch-profiler skill for trace generation, kernel/category breakdown, two-stage overlap analysis, and small Perfetto trace repair. Use when Codex should inspect an existing `trace.json(.gz)` or profile directory, trigger `sglang.profiler` against a live server, break down prefill/decode GPU time by kernel family, correlate a graph-off mapping trace with a graph-on formal trace to find overlap headroom tied back to Python code, or rewrite a trace so Perfetto renders overlapped events more reliably."
description: "Compact SGLang torch-profiler triage skill. Use when Codex should inspect an existing `trace.json(.gz)` or profile directory, trigger `sglang.profiler` against a live server, and return one compact report with kernel, overlap-opportunity, and fuse-pattern tables. Single-trace triage is enough for quick diagnosis; mapping+formal two-trace triage gives stronger overlap conclusions."
---
# SGLang Torch Profiler Analysis
## Overview
Use this skill for all SGLang torch-profiler work. It replaces the old split between:
Use this skill for SGLang `torch.profiler` analysis.
- kernel/category breakdown
- overlap-specific diagnosis
- small trace post-processing
There is only one public workflow:
Prefer the unified entrypoint:
- `triage`
Use the unified entrypoint:
- [scripts/analyze_sglang_torch_profile.py](scripts/analyze_sglang_torch_profile.py)
This entrypoint exposes four subcommands:
- `triage`: the default compact workflow that prints three main tables
- `breakdown`: one-trace kernel/category share analysis
- `overlap`: required two-trace overlap analysis with source mapping
- `perfetto-fix`: rewrite a trace when Perfetto drops some overlapped lanes
For normal use, prefer `triage`. It already collapses the result into three main tables:
`triage` always prints the same three tables:
- kernel table
- overlap-opportunity table
- fuse-opportunity table
- fuse-pattern table
Internal analyzers live here:
By default, all three tables only render rows at or above `1.0%` cumulative GPU-time share.
Treat anything below that as noise unless the user explicitly asks for a lower cutoff.
- [scripts/analyze_sglang_llm_torch_profile.py](scripts/analyze_sglang_llm_torch_profile.py)
- [scripts/analyze_sglang_profiler_overlap.py](scripts/analyze_sglang_profiler_overlap.py)
- [scripts/profile_common.py](scripts/profile_common.py)
The script-level fuse-pattern table should stay source-backed and deterministic.
Do not build a fuzzy string-matching engine into the script for typo-tolerance.
If exact/source-backed matching is weak but the agent judges that a cluster of kernels
still looks semantically close to a known pattern, add a short AI note after the table
with one of these labels:
- `high`: very likely the same pattern family; naming drift or minor implementation reshaping is the main uncertainty
- `medium`: several signals line up, but one important piece is still ambiguous
- `low`: weak resemblance only; mention it only if it is still worth a human follow-up
## When To Use It
- inspect an SGLang torch profiler trace or profile directory
- profile a live SGLang server and immediately analyze the output
- quantify which kernel families dominate prefill or decode
- compare communication, attention, MoE, quantization, norm, or memory share
- summarize which kernel families dominate prefill or decode
- map kernels back to Python code paths
- judge whether a kernel still has overlap headroom in production shape
- get a text table and a small ASCII timeline without opening Perfetto first
- repair a trace so Perfetto can render overlapping events more faithfully
- judge whether a code path still has overlap headroom
- check whether an already-known fusion or overlap path should have applied
Do not use Nsight Systems as the default path for this workflow. This merged skill is torch-profiler-first.
## Diffusion Backend Gate
## Main Commands
For diffusion benchmark or profiling work, only analyze traces produced by the native
SGLang diffusion backend.
### 1. Compact triage from existing trace directories
If the run that generated the trace logs any of:
- `Falling back to diffusers backend`
- `Using diffusers backend`
- `Loaded diffusers pipeline`
stop the workflow instead of analyzing the trace. Treat it as a backend-selection issue,
not as valid SGLang diffusion profiler evidence.
## Main Flows
### 1. Single-trace triage from an existing profile dir or trace
```bash
python3 scripts/analyze_sglang_torch_profile.py \
--input /path/to/profile_dir_or_trace.json.gz
```
Use this when you want the fastest read on kernel share and likely fused-kernel pattern matches.
The overlap table stays conservative in single-trace mode and will tell you when a mapping/formal pair is needed.
### 2. Single-trace triage from a running server
```bash
python3 scripts/analyze_sglang_torch_profile.py \
--url http://127.0.0.1:30000 \
--num-steps 5 \
--profile-by-stage
```
### 3. Two-trace triage from existing profile dirs or traces
```bash
python3 scripts/analyze_sglang_torch_profile.py triage \
@@ -60,7 +88,9 @@ python3 scripts/analyze_sglang_torch_profile.py triage \
--formal-input /path/to/graph_on_profile_dir
```
### 2. Compact triage from running servers
Use this when you need stronger overlap conclusions and cleaner kernel-to-source attribution.
### 4. Two-trace triage from running servers
```bash
python3 scripts/analyze_sglang_torch_profile.py triage \
@@ -70,50 +100,6 @@ python3 scripts/analyze_sglang_torch_profile.py triage \
--profile-by-stage
```
### 3. Breakdown from an existing trace or profile dir
```bash
python3 scripts/analyze_sglang_torch_profile.py breakdown \
--input /path/to/profile_dir
```
### 4. Breakdown from a running server
```bash
python3 scripts/analyze_sglang_torch_profile.py breakdown \
--url http://127.0.0.1:30000 \
--num-steps 5 \
--profile-by-stage \
--table-only
```
### 5. Two-stage overlap analysis
```bash
python3 scripts/analyze_sglang_torch_profile.py overlap \
--mapping-input /path/to/graph_off_profile_dir \
--formal-input /path/to/graph_on_profile_dir \
--table-only
```
Or profile both servers directly:
```bash
python3 scripts/analyze_sglang_torch_profile.py overlap \
--mapping-url http://127.0.0.1:31025 \
--formal-url http://127.0.0.1:31026 \
--num-steps 5
```
### 6. Perfetto-friendly trace rewrite
```bash
python3 scripts/analyze_sglang_torch_profile.py perfetto-fix \
--input /path/to/trace.json.gz
```
This small repair step is inspired by `torch_utils/src/convert_to_perfetto_compatible/convert_to_perfetto_compatible.py`.
## `profile_by_stage`
`profile_by_stage` is not only for PD disaggregation.
@@ -122,75 +108,59 @@ This small repair step is inspired by `torch_utils/src/convert_to_perfetto_compa
- On the current profile-v2 path inside SGLang, stage-based profiling is effectively the normal path.
- PD-disaggregated serving adds one extra rule: prefill workers and decode workers must be profiled separately. That is stricter than ordinary `profile_by_stage`.
## Which Mode To Choose
## How To Choose The Triage Shape
### `triage`
### Single-trace triage
Use when you want the lowest-friction output:
Use when you want the lowest-friction report:
- one kernel table
- one overlap-opportunity table
- one fuse-opportunity table
- optional stage-aware rows when the trace directory includes both `EXTEND` and `DECODE`
- one trace is already available
- you mainly want kernel share and fusion clues
- you are comparing two runs side by side by running triage once per trace
This is the recommended default for final user-facing reports.
This is the recommended default.
### `breakdown`
### Two-trace triage
Use when you need:
- category share such as attention, communication, MoE, norm, quantize, memory
- top kernels by cumulative GPU time
- stage-aware prefill vs decode summaries
- kernel tables keep full kernel names and full Python locations, already joined with CPU ops
- conservative source-backed fusion opportunities
This mode works with one trace. A graph-off pre-pass plus `--kernel-map` is optional but recommended for the final polished report.
### `overlap`
Use when you need:
- a strong answer about which code paths still have overlap headroom
- a table that says which kernels are already hidden and low ROI, with full kernel names and Python scopes
- dependency-risk hints near adjacent kernels
- an ASCII timeline around the most actionable windows
This mode requires two traces for a final answer:
- a stronger answer about overlap headroom
- graph-off source mapping plus graph-on final behavior
- more trustworthy overlap recommendations in the middle table
1. mapping trace with `--disable-cuda-graph --disable-piecewise-cuda-graph`
2. formal trace with the real serving optimizations enabled
Do not call the mapping pass a "fast profile". It exists to recover `kernel -> cpu_op -> python scope`.
### `perfetto-fix`
Use only when Perfetto fails to render obviously overlapping events cleanly. It is a post-processing utility, not the main analysis flow.
## Workflow
### One-trace breakdown workflow
### Single-trace workflow
1. If the user only wants kernel/category share, one trace is enough.
1. If the user only wants a quick diagnosis, one trace is enough.
2. Prefer rank-local `TP-0` traces over merged traces.
3. For a live server, this skill can call `sglang.profiler` and automatically send a small probe request.
4. Prefer `--profile-by-stage` even on standard serving unless the user explicitly wants an all-stage mixed trace.
### Two-trace overlap workflow
### Two-trace workflow
1. Produce a mapping trace first with graph disabled.
2. Produce a formal trace second with graph enabled and the real serving flags kept on.
3. Run `triage` for the compact three-table report, or `overlap` if you also want source context and ASCII timelines.
3. Run `triage` for the compact three-table report.
4. Read the results in this order:
- kernel table
- overlap-opportunity table
- fuse-opportunity table
- fuse-pattern table
5. Before calling something a "new" optimization idea, compare the top rows against both [references/fuse-overlap-catalog.md](references/fuse-overlap-catalog.md) and [references/overlap-catalog.md](references/overlap-catalog.md). Always check the `PR-backed / in-flight` sections too. Prefer reporting:
- an existing fused or overlap path that should already apply here
- an existing path that appears disabled, unsupported, or regressed in this trace
- an upstream PR-backed pattern that already exists but is not merged into the checked-out tree
- a truly new opportunity only when no catalog entry fits
6. Use the deeper `overlap` report only when you need source context or ASCII timelines beyond the compact three-table artifact.
6. If no exact pattern fully matches but the trace still looks semantically close to a known family, add one flat `AI similarity judgment` note after the tables.
Use `high`, `medium`, or `low` only.
Base that note on the full pattern shape, not on one kernel name alone.
Prefer semantic cues such as producer-consumer chain, source locations, CPU op names, TP context, and model-specific structure.
Do not rewrite the script table itself to include these heuristic judgments.
## References
@@ -198,10 +168,6 @@ Load these only when needed:
- [references/source-map.md](references/source-map.md)
- upstream SGLang profiler entrypoints and trace-writing source paths
- [references/validated-workflows.md](references/validated-workflows.md)
- validated two-pass examples for real SGLang models
- [references/trace-workflow.md](references/trace-workflow.md)
- practical guidance for mapping vs formal traces
- [references/heuristics.md](references/heuristics.md)
- overlap labels, dependency-risk interpretation, and limits
- [references/fuse-overlap-catalog.md](references/fuse-overlap-catalog.md)
@@ -211,13 +177,13 @@ Load these only when needed:
## Output Contract
### For `breakdown`
Return:
- trace path
- trace path or generated profile path
- model/server args when available
- top categories
- top kernels
- kernel table
- overlap-opportunity table
- fuse-pattern table
- optional `AI similarity judgment` note with `high` / `medium` / `low` when exact matching is inconclusive
- one short conclusion about what dominates the run
- any source-backed fusion opportunities worth checking
- whether the overlap conclusion came from single-trace triage or mapping/formal two-trace triage
.claude/skills/sglang-torch-profiler-analysis/references/fuse-overlap-catalog.md
+120 -8
View File
@@ -129,14 +129,89 @@ already-known PR family "new".
| PR `#21491` FlashInfer TRTLLM FP8 MoE with fused shared experts | `num_fused_shared_experts`<br>`trtllm_fp8_block_scale_moe` | `PR #21491`<br>`python/sglang/srt/layers/moe/fused_moe_triton/fused_moe.py`<br>`python/sglang/srt/models/deepseek_v2.py` | FlashInfer TRTLLM FP8 MoE path can fuse shared experts inside the routed MoE kernel | On FP8 TRTLLM MoE discussions, treat fused shared experts as an upstream pattern that already has a concrete PR. |
| PR `#22005` fused add + RMSNorm + per-token FP8 quant | `fused_add_rmsnorm_per_token_quant`<br>`per_token_quant_fp8` | `PR #22005`<br>`python/sglang/jit_kernel/csrc/elementwise/fused_add_rmsnorm_per_token_quant.cuh`<br>`python/sglang/jit_kernel/fused_add_rmsnorm_per_token_quant.py` | CUDA JIT kernel keeps normed values in registers and emits BF16 + FP8 outputs plus per-token scales | If FP8 online-quant traces show add+norm followed by per-token quant, treat this as an in-flight upstream CUDA fuse family. |
| PR `#21952` Gemma4 fused RMSNorm + residual + scalar | `gemma_rmsnorm_residual_scalar`<br>`_gemma_rmsnorm_residual_kernel`<br>`Gemma4` | `PR #21952`<br>`python/sglang/srt/layers/gemma4_fused_ops.py`<br>`python/sglang/srt/models/gemma4_causal.py` | Triton kernel fuses decoder post-FF RMSNorm, residual add, and per-layer scalar multiply into one pass | If Gemma4-style post-FF norm + residual + scalar steps appear split, treat them as an in-flight upstream Triton fuse family. |
| PR `#20667` Qwen3.5 fused QK norm + RoPE + KV cache write | `fused_qk_norm_rope_cache_pts_quant_shuffle`<br>`fused_qk_norm_mrope_3d_cache_pts_quant_shuffle`<br>`rotary_dim` | `PR #20667`<br>`python/sglang/srt/models/qwen3_5.py`<br>`python/sglang/srt/models/utils.py` | ROCm / AITER path fuses Q / K RMSNorm, partial or 3D RoPE, and direct KV cache write for Qwen3.5 attention | Treat split QK-norm + RoPE + cache-store on Qwen3.5 as a concrete in-flight upstream family, not a novel idea. |
| PR `#21977` TorchInductor combo-kernels horizontal Q/K norm fusion | `combo_kernels`<br>`benchmark_combo_kernel`<br>`q_norm`<br>`k_norm`<br>`split_with_sizes` | `PR #21977`<br>`torch._inductor.config.combo_kernels` | TorchInductor horizontally fuses sibling Q-norm and K-norm kernels, often deleting `split_with_sizes` / `clone` ladders in compiled traces | Treat separate Q/K norm ladders in compile-heavy traces as an in-flight compiler-fusion family first. |
| PR `#22392` CUTLASS FP8 GEMM replacing nvjet | `cutlass_scaled_mm`<br>`fp8_scaled_mm`<br>`nvjet`<br>`cudaMemsetAsync` | `PR #22392`<br>`sgl-kernel/python/sgl_kernel/gemm.py`<br>`python/sglang/srt/layers/quantization/fp8_utils.py` | Runtime replacement swaps nvjet FP8 GEMMs for CUTLASS kernels, removing per-launch memset bubbles and extra output-copy kernels | Treat nvjet GEMM + memset bubble ladders as an in-flight SGLang linear-kernel family before calling them novel. |
## 7. PR-backed / in-flight kernel-overlap families
| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
| PR `#21877` fused down-GEMM + combine superseding SBO | `enable_fused_grouped_gemm_combine`<br>`combine`<br>`down_gemm` | `PR #21877`<br>`python/sglang/srt/server_args.py`<br>`python/sglang/srt/layers/moe/token_dispatcher/deepep.py` | Fused combine eliminates the standalone combine window, so SBO is intentionally disabled when this path is on | If the trace discussion is about combine overlap, first classify it as this upstream fused-overlap family. |
| PR `#22410` hiSparse H2D transfer overlap with hit-attention | `transfer_stream`<br>`execute_h2d_async`<br>`hit-attention`<br>`merge_state` | `PR #22410`<br>`python/sglang/srt/layers/attention/nsa_backend.py`<br>`python/sglang/srt/hisparse/hisparse_coordinator.py` | hiSparse decode overlaps host-to-device KV transfer on a transfer stream with hit-attention on the compute stream before merging miss-attention work | Treat hit-attention vs H2D KV transfer windows as a concrete in-flight SGLang overlap family first. |
## 8. vLLM-origin fused-kernel families
## 8. FlashInfer mainline fused-kernel families
These rows are comparative references from `flashinfer`. Use them when a trace
looks like an upstream FlashInfer family even if the current `sglang` checkout
only consumes a subset of that implementation.
| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
| FlashInfer activation / gate epilogues | `silu_and_mul`<br>`gelu_tanh_and_mul`<br>`gelu_and_mul`<br>`silu_and_mul_scaled_nvfp4_experts_quantize` | `flashinfer/activation.py`<br>`flashinfer/quantization/fp4_quantization.py` | FlashInfer covers both the plain activation-plus-mul epilogues and the NVFP4 expert-quantized extension used on MoE expert paths | Treat standalone activation, multiply, and expert-side quant ladders as one existing FlashInfer epilogue family first. |
| FlashInfer norm / residual / quant epilogues | `rmsnorm_quant`<br>`fused_add_rmsnorm`<br>`fused_add_rmsnorm_quant`<br>`gemma_rmsnorm`<br>`gemma_fused_add_rmsnorm`<br>`fused_rmsnorm_silu`<br>`rmsnorm_fp4quant`<br>`add_rmsnorm_fp4quant` | `flashinfer/norm/__init__.py`<br>`flashinfer/cute_dsl/rmsnorm_fp4quant.py`<br>`flashinfer/cute_dsl/add_rmsnorm_fp4quant.py` | The norm family spans plain RMSNorm derivatives, residual-add epilogues, norm+activation, and direct FP8 / NVFP4 output variants instead of materializing each intermediate | Treat split residual add, norm, activation, and quant chains as one existing FlashInfer epilogue family first. |
| FlashInfer allreduce + post-op fusion family | `allreduce_fusion`<br>`AllReduceFusionPattern`<br>`kARResidualRMSNorm`<br>`kARResidualRMSNormFP8Quant`<br>`kARResidualRMSNormFP4Quant`<br>`trtllm_mnnvl_allreduce_fusion` | `flashinfer/comm/allreduce.py`<br>`flashinfer/comm/trtllm_ar.py`<br>`flashinfer/comm/trtllm_mnnvl_ar.py` | TRTLLM and MNNVL backends fuse all-reduce with residual add, RMSNorm, and backend-appropriate quant / norm-output variants | Treat TP collective + norm (+ quant) ladders as an existing FlashInfer fused-collective family first. |
| FlashInfer RoPE + FP8 quant / cache-update family | `rope_quantize_fp8`<br>`mla_rope_quantize_fp8`<br>`rope_quantize_fp8_append_paged_kv_cache` | `flashinfer/rope.py` | The RoPE family covers both RoPE+FP8 output and the larger decode / prefill-prep path that also writes K / V directly into paged KV cache | Treat split RoPE, quant, and cache-write ladders as one existing FlashInfer attention-prep family first. |
| FlashInfer fused DeepSeek grouped-topk routing | `fused_topk_deepseek`<br>`NoAuxTc` | `flashinfer/fused_moe/fused_routing_dsv3.py` | One kernel performs sigmoid+bias, grouped score reduction, group top-k, expert top-k, and routed renorm for DeepSeek-V3-style routing | Treat router score activation -> grouped top-k -> renorm ladders as an existing FlashInfer router family first. |
| FlashInfer fused MoE expert execution | `cutlass_fused_moe`<br>`trtllm_bf16_moe`<br>`trtllm_fp8_per_tensor_scale_moe`<br>`trtllm_fp8_block_scale_moe`<br>`trtllm_fp4_block_scale_moe`<br>`trtllm_mxint4_block_scale_moe` | `flashinfer/fused_moe/core.py` | CUTLASS and TRTLLM backends collapse expert execution, routed combine, and quantized expert variants into fused MoE runners | Treat exposed expert-side tiny GEMM ladders as matching an existing FlashInfer fused-MoE family. |
| FlashInfer CuTeDSL two-stage MoE fusion | `blockscaled_contiguous_gather_grouped_gemm_swiglu_fusion_nvfp4`<br>`blockscaled_contiguous_grouped_gemm_finalize_fusion_nvfp4`<br>`moe_permute`<br>`moe_unpermute` | `flashinfer/fused_moe/cute_dsl/blockscaled_contiguous_gather_grouped_gemm_swiglu_fusion.py`<br>`flashinfer/fused_moe/cute_dsl/blockscaled_contiguous_grouped_gemm_finalize_fusion.py` | The CuTeDSL path fuses gather+GEMM1+SwiGLU in the first stage and finalize+unpermute+scatter-reduce in the second stage, removing standalone `moe_permute` and `moe_unpermute` kernels | Treat multi-kernel MoE ladders around permute / finalize as one existing FlashInfer CuTeDSL family first. |
## 9. FlashInfer mainline kernel-overlap families
| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
| FlashInfer PDL launch-overlap family | `enable_pdl`<br>`launch_with_pdl`<br>`cudaGridDependencySynchronize`<br>`cudaTriggerProgrammaticLaunchCompletion`<br>`trigger_completion_at_end=False`<br>`allreduce_fusion` | `flashinfer/norm/__init__.py`<br>`flashinfer/activation.py`<br>`flashinfer/rope.py`<br>`flashinfer/comm/allreduce.py`<br>`flashinfer/comm/trtllm_ar.py` | FlashInfer uses Programmatic Dependent Launch broadly, and the allreduce path can further advance completion so the next PDL-aware kernel overlaps on the same stream | Treat tight same-stream dependent windows and allreduce-followed-by-kernel windows as one existing FlashInfer launch-overlap family first. |
| FlashInfer CuTeDSL MoE aux-stream async-memset overlap | `aux_stream`<br>`main_event`<br>`memset_event`<br>`use_async_memset` | `flashinfer/fused_moe/cute_dsl/fused_moe.py` | Preallocated MoE output is zeroed on an auxiliary CUDA stream while GEMM1 runs on the main stream, then both streams join before finalize | Treat GEMM1 vs output-zero windows as an existing FlashInfer multi-stream overlap family. |
| FlashInfer green-context SM partition overlap | `split_device_green_ctx`<br>`split_device_green_ctx_by_sm_count`<br>`green_ctx` | `flashinfer/green_ctx.py` | CUDA green contexts partition SMs and create dedicated streams for concurrent kernel families on separate SM slices | Treat full-device two-stream traces and SM-partitioned traces as different manifestations of an existing FlashInfer overlap mechanism. |
## 10. FlashInfer PR-backed / in-flight fused-kernel and kernel-overlap families
| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
| PR `#2792` RoPE + FP8 quant + paged KV append with padding-token support | `rope_quantize_fp8_append_paged_kv_cache`<br>`seqlen=0`<br>`batch_indices < 0` | `PR #2792`<br>`flashinfer/rope.py`<br>`include/flashinfer/pos_enc.cuh` | Extends the existing RoPE+quant+cache-write family to CUDA-graph padding tokens / zero-length sequences instead of introducing a separate kernel ladder | Treat split padding-token handling around RoPE+cache write as an in-flight upstream FlashInfer family first. |
| PR `#2840` CuTeDSL MoE aux-stream overlap race fix | `aux_stream`<br>`use_prealloc`<br>`use_cuda_graph` | `PR #2840`<br>`flashinfer/fused_moe/cute_dsl/fused_moe.py` | Clarifies that async memset overlap is only safe for the preallocated / CUDA-graph case; non-graph mode falls back to main-stream zeroing to avoid races | Treat missing aux-stream overlap in non-graph traces as an intentional safety rule, not a novel opportunity. |
| PR `#2720` PDL runtime-API migration | `cudaGridDependencySynchronize`<br>`cudaTriggerProgrammaticLaunchCompletion`<br>`inline PTX` | `PR #2720`<br>`include/flashinfer/comm/trtllm_allreduce_fusion.cuh`<br>`include/flashinfer/pos_enc.cuh` | Repo-wide migration preserves the existing PDL overlap family while replacing inline PTX with CUDA runtime APIs across norm, RoPE, attention, and MoE codepaths | Treat PDL-looking launch groups as an upstream FlashInfer overlap family even when implementation details differ across revisions. |
| PR `#2882` FP8 per-tensor TRTLLM MoE non-gated activation | `trtllm_fp8_per_tensor_scale_moe`<br>`non-gated` | `PR #2882`<br>`csrc/trtllm_fused_moe_kernel_launcher.cu` | Extends the existing TRTLLM FP8 fused-MoE family to non-gated activations instead of requiring a separate expert path | Treat non-gated FP8 expert ladders as an in-flight upstream FlashInfer extension first. |
## 11. TensorRT-LLM-origin fused-kernel families
These rows are comparative references from `TensorRT-LLM`. Use them when a
trace looks like a TensorRT-LLM or TensorRT-LLM-plus-FlashInfer family even if
the current `sglang` checkout only carries an analogous implementation.
| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
| TensorRT-LLM FlashInfer activation / gate epilogues | `flashinfer_silu_and_mul`<br>`flashinfer_gelu_tanh_and_mul`<br>`auto_deploy::silu_and_mul`<br>post-GEMM `silu` + `mul` | `tensorrt_llm/_torch/custom_ops/flashinfer_custom_ops.py`<br>`tensorrt_llm/_torch/auto_deploy/transform/library/fuse_silu_mul.py`<br>`tensorrt_llm/_torch/models/modeling_gemma3.py` | Runtime custom ops and AutoDeploy rewrite `split/getitem + activation + mul` MLP epilogues into one FlashInfer op, including Gemma3 `gelu_tanh_and_mul` | Treat split gate activation + multiply as an existing TensorRT-LLM/FlashInfer epilogue family first. |
| TensorRT-LLM FlashInfer RMSNorm family | `flashinfer_rmsnorm`<br>`flashinfer_gemma_rmsnorm`<br>`auto_deploy::flashinfer_rms_norm` | `tensorrt_llm/_torch/custom_ops/flashinfer_custom_ops.py`<br>`tensorrt_llm/_torch/modules/rms_norm.py`<br>`tensorrt_llm/_torch/auto_deploy/custom_ops/normalization/rms_norm.py` | Runtime modules and AutoDeploy can lower plain RMSNorm and Gemma RMSNorm directly to FlashInfer kernels | Treat split RMSNorm ladders as an existing TensorRT-LLM norm family before calling them novel. |
| TensorRT-LLM FlashInfer residual add + RMSNorm | `flashinfer_fused_add_rmsnorm`<br>`flashinfer_gemma_fused_add_rmsnorm`<br>`auto_deploy::flashinfer_fused_add_rms_norm_inplace` | `tensorrt_llm/_torch/custom_ops/flashinfer_custom_ops.py`<br>`tensorrt_llm/_torch/modules/rms_norm.py`<br>`tensorrt_llm/_torch/auto_deploy/transform/library/fused_add_rms_norm.py` | Residual add immediately before RMSNorm can collapse to one in-place FlashInfer op, with Gemma variant support | Treat residual add + RMSNorm chains as an existing TensorRT-LLM fused epilogue family first. |
| TensorRT-LLM FlashInfer RoPE with shared cos/sin cache | `flashinfer_apply_rope_with_cos_sin_cache_inplace`<br>`flashinfer_rope`<br>`cos_sin_cache` | `tensorrt_llm/_torch/modules/rotary_embedding.py`<br>`tensorrt_llm/_torch/auto_deploy/custom_ops/rope/flashinfer_rope.py`<br>`tensorrt_llm/_torch/auto_deploy/transform/library/rope.py` | Runtime path applies in-place RoPE from a shared cos/sin cache, while AutoDeploy can prebuild the full cache and lower diverse RoPE graphs to `flashinfer_rope` | Treat separate cos/sin gather + RoPE application ladders as an existing TensorRT-LLM attention-prep family. |
| TensorRT-LLM FlashInfer cached paged attention | `append_paged_kv_cache`<br>`BatchPrefillWithPagedKVCacheWrapper`<br>`BatchDecodeWithPagedKVCacheWrapper`<br>`auto_deploy::flashinfer_attention_mha_with_cache`<br>`read_cache_only` | `tensorrt_llm/_torch/attention_backend/flashinfer.py`<br>`tensorrt_llm/_torch/auto_deploy/custom_ops/attention/flashinfer_attention.py`<br>`docs/source/features/attention.md` | FlashInfer attention backend fuses metadata setup, optional paged-KV append, and prefill/decode wrapper execution, including shared-KV and read-cache-only variants in AutoDeploy | Treat metadata + KV-append + cached-attention ladders as one existing TensorRT-LLM cached-attention family first. |
| TensorRT-LLM FlashInfer MLA regular prefill | `append_paged_mla_kv_cache`<br>`BatchPrefillWithRaggedKVCacheWrapper`<br>`flashinfer_mla` | `tensorrt_llm/_torch/auto_deploy/custom_ops/mla/flashinfer_mla.py` | Regular MLA prefill writes compressed KV pages and runs FlashInfer ragged prefill instead of a split append-plus-prefill ladder | Treat MLA regular-prefill prep as an existing TensorRT-LLM FlashInfer family first. |
| TensorRT-LLM FlashInfer MLA chunked prefill with absorbed `W_kn` | `BatchMLAPagedAttentionWrapper`<br>`chunked prefill`<br>`W_kn`<br>`W_v` | `tensorrt_llm/_torch/auto_deploy/custom_ops/mla/flashinfer_mla.py` | Chunked prefill absorbs `W_kn` into the query-side projection, runs paged MLA attention in compressed space, then projects back with `W_v` | Treat split absorbed-proj + MLA + output-proj ladders as an existing TensorRT-LLM MLA family first. |
| TensorRT-LLM FlashInfer MLA decode with absorbed `W_kn` + `W_v` | `plan_decode`<br>`BatchMLAPagedAttentionWrapper`<br>`decode`<br>`W_kn`<br>`W_v` | `tensorrt_llm/_torch/auto_deploy/custom_ops/mla/flashinfer_mla.py` | Decode path reuses the absorbed-query MLA family and projects the compressed attention output back with `W_v` | Treat similar decode-time absorbed MLA ladders as an existing TensorRT-LLM family, not a new idea. |
| TensorRT-LLM FlashInfer fused MoE backend | `flashinfer.fused_moe`<br>`trtllm_bf16_moe`<br>`trtllm_fp8_block_scale_moe`<br>`trtllm_fp4_block_scale_moe`<br>`TRTLLM_GEN_FUSED_MOE_USE_FLASHINFER` | `tensorrt_llm/_torch/modules/fused_moe/moe_op_backend.py`<br>`tensorrt_llm/_torch/modules/fused_moe/fused_moe_trtllm_gen.py` | TRTLLM-gen MoE can route expert execution and quant helpers through FlashInfer instead of exposing per-expert eager ladders | Treat expert-side tiny GEMM ladders as matching an existing TensorRT-LLM FlashInfer MoE family first. |
| TensorRT-LLM FlashInfer cached SSM / Mamba update | `flashinfer_cached_ssm`<br>`selective_state_update`<br>`flashinfer_ssm` | `tensorrt_llm/_torch/auto_deploy/custom_ops/mamba/flashinfer_backend_mamba.py`<br>`tensorrt_llm/_torch/modules/mamba/mamba2_mixer.py` | Mamba2 paths can lower cached SSM state updates to FlashInfer selective-state-update kernels instead of many smaller state ops | Treat split cached-SSM state update ladders as an existing TensorRT-LLM FlashInfer family first. |
## 12. TensorRT-LLM-origin kernel-overlap families
| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
| TensorRT-LLM multi-stream MLA attention | `multi_stream_mla_attn`<br>`record_event_passthrough`<br>`_aux`<br>`wait_event` | `tensorrt_llm/_torch/auto_deploy/transform/library/multi_stream_attn.py`<br>`tensorrt_llm/_torch/auto_deploy/utils/multi_stream_utils.py` | AutoDeploy rewrites MLA Q/KV forks so the KV projection runs on an auxiliary stream while the Q path stays on the caller stream | Treat exposed Q-branch vs KV-branch overlap as an existing TensorRT-LLM multi-stream family first. |
| TensorRT-LLM multi-stream MoE shared-vs-routed overlap | `multi_stream_moe`<br>`begin_aux_stream_passthrough`<br>`end_aux_stream_passthrough`<br>`wait_aux_stream_passthrough` | `tensorrt_llm/_torch/auto_deploy/transform/library/multi_stream_moe.py`<br>`tensorrt_llm/_torch/auto_deploy/utils/multi_stream_utils.py` | Shared-expert work is moved to an auxiliary stream while routed-expert MoE work remains on the main stream and rejoins at the merge node | Treat shared-expert vs routed-expert windows as an existing TensorRT-LLM branch-overlap family. |
| TensorRT-LLM multi-stream FP8 GEMM fork parallelism | `multi_stream_gemm`<br>`trtllm_finegrained_fp8_linear`<br>`record_event_passthrough`<br>`_aux` | `tensorrt_llm/_torch/auto_deploy/transform/library/multi_stream_gemm.py`<br>`tensorrt_llm/_torch/auto_deploy/utils/multi_stream_utils.py` | Compiler pass identifies fork points with multiple FP8 linears and moves the largest GEMM to the auxiliary stream so sibling GEMMs overlap | Treat sibling FP8 linear branches as an existing TensorRT-LLM overlap family before designing a new stream split. |
## 13. TensorRT-LLM-origin PR-backed / in-flight fused-kernel and kernel-overlap families
| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
| PR `#12674` fused residual add + RMSNorm + FP8 quant | `triton_fused_add_rms_norm_quant_fp8`<br>`residual_add`<br>`rms_norm`<br>`fp8 static quant` | `PR #12674`<br>`tensorrt_llm/_torch/auto_deploy/transform/library/fuse_rmsnorm_quant_fp8.py` | Open PR adds a pattern-matcher pass that replaces residual-add plus RMSNorm plus FP8 static quant with a fused FlashInfer 0.6.7-backed path | Treat split add + norm + FP8 quant ladders as an in-flight TensorRT-LLM family first. |
| PR `#12519` rank-256 `flashinfer_mla` extension | `flashinfer_mla`<br>`rank 256`<br>`paged KV-cache`<br>`gpu append kernel` | `PR #12519`<br>`tensorrt_llm/_torch/auto_deploy/custom_ops/mla/flashinfer_mla.py` | Open PR extends the existing FlashInfer MLA family with a TRTLLM MLA operator, paged-KV support, and a GPU append kernel for rank-256 setups | Treat rank-256 MLA prep / decode ladders as an in-flight TensorRT-LLM MLA family, not a novel direction. |
| PR `#12525` FlashInfer TRTLLM-gen FMHA paged-index / buffer rework | `shared paged index`<br>`trtllm-gen attention`<br>`flashinfer`<br>`kv cache buffer` | `PR #12525`<br>`tensorrt_llm/_torch/auto_deploy/custom_ops/attention/flashinfer_attention.py` | Open PR refines the existing FlashInfer TRTLLM-gen cached-attention family by disabling shared paged index and unifying KV-buffer construction | Treat these attention-prep changes as an in-flight implementation evolution of an existing family first. |
| PR `#12544` NVFP4 KV cache support in TRTLLM-gen attention | `NVFP4 KV cache`<br>`trtllm-gen attention`<br>`flashinfer` | `PR #12544`<br>`tensorrt_llm/_torch/auto_deploy/custom_ops/attention/flashinfer_attention.py` | Open PR extends the cached-attention family so the FlashInfer-backed TRTLLM-gen path can build and consume NVFP4 KV buffers directly | Treat split KV-cache quant + buffer-build ladders as an in-flight TensorRT-LLM attention family first. |
| PR `#12738` / `#12557` BF16 TRTLLM-gen MoE through FlashInfer | `bf16 trtllm-gen moe`<br>`flashinfer`<br>`trtllm_bf16_moe` | `PR #12738`<br>`PR #12557`<br>`tensorrt_llm/_torch/modules/fused_moe/fused_moe_trtllm_gen.py` | Open PRs extend the TRTLLM-gen MoE family so BF16 expert execution can route through FlashInfer instead of only CUTLASS-like paths | Treat BF16 expert ladders as an in-flight TensorRT-LLM FlashInfer MoE family. |
| PR `#12847` `multi_stream_moe` sync fix for MLIR and piecewise cudagraphs | `multi_stream_moe`<br>`mlir_elementwise_fusion`<br>`piecewise cudagraph`<br>`caller_stream.synchronize()` | `PR #12847`<br>`tensorrt_llm/_torch/auto_deploy/transform/library/multi_stream_moe.py`<br>`tensorrt_llm/_torch/auto_deploy/utils/multi_stream_utils.py` | Open PR preserves the existing multi-stream MoE overlap family while tightening synchronization when MLIR-fused kernels or piecewise cudagraph replay are present | Treat missing or altered `multi_stream_moe` overlap under MLIR / piecewise graph modes as an in-flight TensorRT-LLM rule first. |
## 14. vLLM-origin fused-kernel families
These rows are comparative references from `vllm`. Use them when a trace looks
similar to an upstream family even if the current `sglang` checkout does not
@@ -164,7 +239,7 @@ contain the same implementation.
| vLLM-origin fused MoE LoRA | `fused_moe_lora`<br>`fused_moe_lora_fp8`<br>`w13_shrink`<br>`w2_expand` | `vllm/lora/ops/triton_ops/fused_moe_lora_op.py`<br>`vllm/lora/ops/triton_ops/fused_moe_lora_fp8_op.py`<br>`vllm/lora/layers/fused_moe.py` | Triton kernels fuse LoRA shrink / expand work into MoE expert execution, including FP8 variants | Treat MoE-LoRA adapter work as an upstream fused family before proposing a brand new kernel. |
| vLLM-origin ViT fused bilinear position-embedding interpolation | `triton_pos_embed_interpolate`<br>`bilinear_pos_embed`<br>`pos_embed_interpolate_native` | `vllm/model_executor/models/qwen3_vl.py` | Triton kernel fuses bilinear interpolation and spatial-merge reorder for Qwen3-VL ViT position embeddings, replacing many tiny eager kernels | Treat VLM position-embedding ladders as an existing vLLM-origin Triton fusion family. |
## 9. vLLM-origin kernel-overlap families
## 15. vLLM-origin kernel-overlap families
| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
@@ -173,7 +248,7 @@ contain the same implementation.
| vLLM-origin shared-expert aux-stream overlap | `aux_stream`<br>`shared_experts_stream`<br>shared expert near router | `vllm/utils/torch_utils.py`<br>`vllm/model_executor/layers/fused_moe/runner/default_moe_runner.py` | MoE shared experts can run on a dedicated aux stream and overlap with router-side work | Treat shared-expert vs router overlap as an existing upstream sparse-model family. |
| vLLM-origin DCP async all-to-all overlap | `dcp_alltoall`<br>`all_to_all_single`<br>`async_op=True` | `vllm/v1/attention/ops/dcp_alltoall.py` | Output / LSE exchange uses async all-to-all handles instead of serializing collective completion on the main path | Treat DCP all-to-all windows as an upstream async-collective family. |
## 10. vLLM-origin PR-backed / in-flight fused-kernel and kernel-overlap families
## 16. vLLM-origin PR-backed / in-flight fused-kernel and kernel-overlap families
| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
@@ -182,8 +257,12 @@ contain the same implementation.
| PR `#38445` MiniMax-M2 FP32 gate kernel | `fp32_router_gemm`<br>`MiniMax-M2`<br>`gate kernel` | `PR #38445`<br>`vllm/model_executor/layers/fused_moe/router/gate_linear.py`<br>`vllm/model_executor/models/minimax_m2.py` | Draft CUDA kernel fuses BF16->FP32 conversion and low-batch router GEMM for MiniMax-M2, replacing up to three kernels on the gate path | Treat MiniMax-M2 gate ladders as an in-flight upstream fused router family first. |
| PR `#38621` fused QK norm + RoPE + cache + quant | `fused_qk_norm_rope_cache_quant`<br>`QK Norm + RoPE + Cache + Quant` | `PR #38621`<br>`csrc/fused_qk_norm_rope_cache_quant.cu`<br>`vllm/compilation/passes/fusion/qk_norm_rope_cache_quant_fusion.py` | Draft CUDA kernel and compile-time pass try to fuse QK RMSNorm, RoPE, KV cache write, and optional FP8 quant for small-batch decode | Treat this as an in-flight upstream fusion family before calling a similar idea novel. |
| PR `#38684` DSV3.2 fused `wk + weights_proj` | `wk_weights_proj`<br>`MergedColumnParallelLinear`<br>`weights_proj` | `PR #38684`<br>`vllm/model_executor/models/deepseek_v2.py`<br>`vllm/model_executor/models/deepseek_mtp.py` | Merged PR fuses the DSV3.2 indexer `wk` and `weights_proj` projections into one GEMM; FP8 weight-loading caveats are being handled in follow-up `PR #38870` | Treat paired indexer projections as a concrete upstream fused linear family before calling the opportunity novel. |
| PR `#37646` ROCm AITER fused allreduce + RMSNorm | `rocm_aiter_fused_allreduce_rmsnorm`<br>`custom_fused_ar_rms`<br>`RocmAiterAllReduceFusionPass` | `PR #37646`<br>`vllm/_aiter_ops.py`<br>`vllm/compilation/passes/pass_manager.py` | ROCm-specific compile-time path swaps the generic all-reduce fusion pass for an AITER fused allreduce-plus-RMSNorm kernel family | Treat ROCm TP all-reduce + RMSNorm ladders as an in-flight upstream fused-collective family first. |
| PR `#36413` FlashInfer RMSNorm + FP4 quant fusion | `fuse_norm_quant`<br>`flashinfer`<br>`NVFP4`<br>`rmsnorm + fp4 quant` | `PR #36413`<br>`vllm/compilation/passes/fusion/rms_quant_fusion.py`<br>`vllm/docs/design/fusions.md` | FlashInfer-backed norm-plus-FP4 quant fusion extends the existing RMSNorm+quant family to NVFP4 flows | Treat split RMSNorm + FP4 quant ladders as an upstream in-flight family, not a fresh idea. |
| PR `#37045` MiniMax TRTLLM `minimax_allreduce_rms` kernels | `minimax_allreduce_rms`<br>`MiniMax-M2.5`<br>`allreduce_rms` | `PR #37045`<br>`vllm/model_executor/models/minimax_m2.py` | Draft kernel ports TensorRT-LLM MiniMax allreduce-plus-RMSNorm kernels into vLLM for TP MiniMax decode | Treat MiniMax TP norm + collective ladders as an in-flight upstream specialized fusion family. |
| PR `#39301` GLM5 router GEMM with PDL overlap | `TRTLLM_ENABLE_PDL`<br>`router_gemm`<br>`GLM5`<br>`FI AR RMS fusion` | `PR #39301`<br>`vllm/model_executor/layers/fused_moe/router/gate_linear.py`<br>`vllm/csrc/moe/dsv3_router_gemm_utils.h` | Extends the specialized router GEMM family to GLM5 hidden size and uses PDL to overlap the router launch with the preceding fused allreduce-plus-RMS block | Treat this as an in-flight upstream router-kernel plus launch-overlap family before calling it novel. |
## 11. Important toggles and caveats
## 17. Important toggles and caveats
| Toggle / env | Location | Effect on trace interpretation |
| --- | --- | --- |
@@ -198,6 +277,18 @@ contain the same implementation.
| `SGLANG_STAGING_USE_TORCH` | `python/sglang/srt/disaggregation/common/staging_buffer.py` | Forces torch fallback for staging gather / scatter, so Triton staging kernels may disappear by design. |
| `SGLANG_VIT_ENABLE_CUDA_GRAPH` | `python/sglang/srt/environ.py` | Can intentionally disable vision `aux_stream` overlap. |
| `SGLANG_ENABLE_FUSED_QKNORM_ROPE` | `python/sglang/multimodal_gen/runtime/layers/layernorm.py` | Gates the diffusion fused qknorm+rope path. |
| `enable_pdl` / `launch_with_pdl` | `flashinfer/norm/__init__.py`<br>`flashinfer/activation.py`<br>`flashinfer/rope.py`<br>`flashinfer/fused_moe/core.py`<br>`flashinfer/comm/allreduce.py` | Enables FlashInfer PDL across many kernels; launch grouping and same-stream overlap can change substantially when it is on. |
| `trigger_completion_at_end` | `flashinfer/comm/allreduce.py` | `False` enables downstream PDL-aware overlap after FlashInfer allreduce fusion; `True` delays completion to kernel end and removes that overlap window. |
| `use_cuda_graph` | `flashinfer/fused_moe/cute_dsl/fused_moe.py` | Enables the preallocated-buffer path and the safe aux-stream async-memset overlap in FlashInfer CuTeDSL MoE. |
| `split_device_green_ctx*` | `flashinfer/green_ctx.py` | Changes trace shape by partitioning SMs into separate green contexts instead of overlapping full-device streams on the default context. |
| `rmsnorm_backend` | `tensorrt_llm/_torch/auto_deploy/config/default.yaml` | Chooses whether AutoDeploy lowers RMSNorm to FlashInfer, so split norm ladders may reflect backend selection rather than a missing fuse. |
| `insert_cached_attention.backend` | `tensorrt_llm/_torch/auto_deploy/config/default.yaml` | Selects the cached-attention backend; `flashinfer` enables the paged-KV cached-attention family. |
| `insert_cached_mla_attention.backend` | `tensorrt_llm/_torch/auto_deploy/config/default.yaml` | Selects the cached MLA backend; `flashinfer_mla` enables the MLA prefill / decode family. |
| `TRTLLM_GEN_FUSED_MOE_USE_FLASHINFER` | `tensorrt_llm/_torch/modules/fused_moe/fused_moe_trtllm_gen.py` | Forces or guards the FlashInfer-backed TRTLLM-gen MoE family, so expert-kernel shape can change substantially when it is set. |
| `multi_stream_moe` | `tensorrt_llm/_torch/auto_deploy/config/default.yaml` | Enables the TensorRT-LLM shared-expert vs routed-expert overlap family. |
| `multi_stream_mla_attn` | `tensorrt_llm/_torch/auto_deploy/config/default.yaml` | Enables the TensorRT-LLM MLA Q-vs-KV branch overlap family. |
| `multi_stream_gemm` | `tensorrt_llm/_torch/auto_deploy/config/default.yaml` | Enables generalized FP8 GEMM fork overlap in TensorRT-LLM AutoDeploy. |
| `mlir_elementwise_fusion` | `tensorrt_llm/_torch/auto_deploy/config/default.yaml` | Can absorb merge adds into larger fused kernels, so missing explicit merge nodes in multi-stream traces may be intentional. |
| `enable_torch_compile` | `python/sglang/srt/server_args.py`<br>`python/sglang/multimodal_gen/runtime/server_args.py` | Compiler-generated fusion / reordering can hide handwritten kernel names; absence of a custom kernel does not always mean absence of fusion. |
| `enable_fused_grouped_gemm_combine` | `PR #21877` | In-flight path that intentionally disables SBO because combine is folded into down-GEMM. |
| `PassConfig.fuse_allreduce_rms` | `vllm/config/compilation.py` | Enables vLLM's AllReduce -> RMSNorm (+ residual / quant) compile-time fusion family. |
@@ -210,21 +301,42 @@ contain the same implementation.
| `PassConfig.fuse_gemm_comms` | `vllm/config/compilation.py` | Enables AsyncTP GEMM + collective overlap and auto-enables `enable_sp` when valid. |
| `TRTLLM_ENABLE_PDL` | `vllm/csrc/dsv3_fused_a_gemm.cu`<br>`vllm/csrc/moe/dsv3_router_gemm_utils.h` | Enables programmatic dependent launch for the DSV3 specialized CUDA kernels, which can change launch grouping and trace shape for router / QKV-A paths. |
## 12. Suggested refresh commands
## 18. Suggested refresh commands
These commands are only for maintainers refreshing this catalog by rescanning
the local source trees. They are not used by the triage scripts at runtime.
```bash
# Optional sibling checkouts used for comparative scanning:
FLASHINFER_REPO=${FLASHINFER_REPO:-../flashinfer}
TRTLLM_REPO=${TRTLLM_REPO:-../TensorRT-LLM}
VLLM_REPO=${VLLM_REPO:-../vllm}
rg -n "fused_add_rmsnorm|gemma_fused_add_rmsnorm|silu_and_mul|gelu_and_mul|fused_qk_rope_reshape_and_cache|fused_set_kv_buffer|fused_metadata_copy|normal_decode_set_metadata" python/sglang
rg -n "MiniMaxM2RMSNormTP|fused_qknorm_rope|fused_qk_rope_cat_and_cache_mla|fused_qk_norm_mrope_3d_cache_pts_quant_shuffle|split_qkv_rmsnorm_rope|trtllm_fp8_kv_kernel|set_mla_kv_buffer_fp8_quant" python/sglang
rg -n "FusedMoeRouter|fused_topk_deepseek|moe_fused_gate|aiter_fused_topk|fused_rms_fp8_group_quant|fast_topk_transform_fused|fused_store_index_k_cache|fused_temperature_softmax|fused_softcap" python/sglang
rg -n "fused_qkvzba_split_reshape_cat|fused_gdn_gating|rms_norm_gated|layer_norm_gated|chunk_gated_delta_rule_fwd_kkt_solve_kernel|fused_recurrent_gated_delta_rule_update|fused_mamba_state_scatter_with_mask|_fused_gather_to_staging_kernel|_fused_scatter_from_staging_kernel" python/sglang
rg -n "single_batch_overlap|alt_stream|shared_expert|_comm_stream|scatter_stream|triton_mrope_fused|ring_attn|all_to_all_single|reorder_for_compute_comm_overlap|use_dual_stream" python/sglang
git log --all --format='%h %s' | rg -i 'fused|fusion|overlap|cutedsl|triton|cuda|rope|topk|quant|combine|allreduce|all_to_all'
rg -n "fused_add_rms_norm|fused_qk_norm_rope|grouped_topk|topk_softmax|topk_sigmoid|dsv3_router_gemm|dsv3_fused_a_gemm|concat_and_cache_mla_rope_fused|gpt_oss_router_gemm|cutlass_scaled_mm|cpu_fused_moe|fused_moe_lora|triton_pos_embed_interpolate" /Users/bbuf/工作目录/Common/vllm/vllm /Users/bbuf/工作目录/Common/vllm/csrc
rg -n "fuse_allreduce_rms|fuse_norm_quant|fuse_act_quant|fuse_attn_quant|enable_qk_norm_rope_fusion|fuse_rope_kvcache|enable_sp|fuse_gemm_comms|RocmAiter|dcp_alltoall|shared_experts_stream|TRTLLM_ENABLE_PDL|wk_weights_proj" /Users/bbuf/工作目录/Common/vllm/vllm /Users/bbuf/工作目录/Common/vllm/docs/design/fusions.md /Users/bbuf/工作目录/Common/vllm/csrc
git -C /Users/bbuf/工作目录/Common/vllm log --all --format='%h %s' | rg -i 'fused|fusion|overlap|triton|cuda|rope|kv cache|topk|router|allreduce|reduce-scatter|all-gather|all_to_all|quant'
rg -n "silu_and_mul|gelu_tanh_and_mul|gelu_and_mul|silu_and_mul_scaled_nvfp4_experts_quantize|rmsnorm_quant|fused_add_rmsnorm|fused_add_rmsnorm_quant|fused_rmsnorm_silu" "$FLASHINFER_REPO/flashinfer"
rg -n "AllReduceFusionPattern|allreduce_fusion|trigger_completion_at_end|rope_quantize_fp8|rope_quantize_fp8_append_paged_kv_cache|fused_topk_deepseek|cutlass_fused_moe|trtllm_.*_moe" "$FLASHINFER_REPO/flashinfer"
rg -n "aux_stream|use_async_memset|split_device_green_ctx|split_device_green_ctx_by_sm_count|enable_pdl|launch_with_pdl" "$FLASHINFER_REPO/flashinfer" "$FLASHINFER_REPO/include"
git -C "$FLASHINFER_REPO" log --all --format='%h %s' | rg -i 'fused|fusion|overlap|pdl|stream|rope|kv|quant|topk|moe'
rg -n "flashinfer_silu_and_mul|flashinfer_gelu_tanh_and_mul|flashinfer_rmsnorm|flashinfer_gemma_rmsnorm|flashinfer_fused_add_rmsnorm|flashinfer_apply_rope_with_cos_sin_cache_inplace" "$TRTLLM_REPO/tensorrt_llm/_torch"
rg -n "flashinfer_attention_mha_with_cache|append_paged_kv_cache|flashinfer_mla|append_paged_mla_kv_cache|flashinfer_cached_ssm|selective_state_update|flashinfer.fused_moe" "$TRTLLM_REPO/tensorrt_llm/_torch" "$TRTLLM_REPO/docs/source"
rg -n "multi_stream_moe|multi_stream_mla_attn|multi_stream_gemm|record_event_passthrough|begin_aux_stream_passthrough|end_aux_stream_passthrough|wait_aux_stream_passthrough" "$TRTLLM_REPO/tensorrt_llm/_torch"
git -C "$TRTLLM_REPO" log --all --format='%h %s' | rg -i 'fused|fusion|overlap|flashinfer|mla|kv cache|multi-stream|stream|rope|rmsnorm|moe'
rg -n "fused_add_rms_norm|fused_qk_norm_rope|grouped_topk|topk_softmax|topk_sigmoid|dsv3_router_gemm|dsv3_fused_a_gemm|concat_and_cache_mla_rope_fused|gpt_oss_router_gemm|cutlass_scaled_mm|cpu_fused_moe|fused_moe_lora|triton_pos_embed_interpolate" "$VLLM_REPO/vllm" "$VLLM_REPO/csrc"
rg -n "fuse_allreduce_rms|fuse_norm_quant|fuse_act_quant|fuse_attn_quant|enable_qk_norm_rope_fusion|fuse_rope_kvcache|enable_sp|fuse_gemm_comms|RocmAiter|dcp_alltoall|shared_experts_stream|TRTLLM_ENABLE_PDL|wk_weights_proj" "$VLLM_REPO/vllm" "$VLLM_REPO/docs/design/fusions.md" "$VLLM_REPO/csrc"
git -C "$VLLM_REPO" log --all --format='%h %s' | rg -i 'fused|fusion|overlap|triton|cuda|rope|kv cache|topk|router|allreduce|reduce-scatter|all-gather|all_to_all|quant'
# GitHub PR scan terms for the connector or web UI:
# "fused OR overlap repo:sgl-project/sglang"
# "triton OR cutedsl OR cuda fused repo:sgl-project/sglang"
# "fused OR overlap repo:flashinfer-ai/flashinfer"
# "pdl OR aux_stream OR green_ctx repo:flashinfer-ai/flashinfer"
# "fused OR overlap repo:NVIDIA/TensorRT-LLM"
# "flashinfer OR mla OR moe OR rmsnorm repo:NVIDIA/TensorRT-LLM"
# "multi-stream OR aux_stream OR cudagraph repo:NVIDIA/TensorRT-LLM"
# "fused OR overlap repo:vllm-project/vllm"
# "triton OR cuda fused repo:vllm-project/vllm"
```
.claude/skills/sglang-torch-profiler-analysis/references/overlap-catalog.md
+71 -6
View File
@@ -64,8 +64,46 @@ overlap opportunity as novel.
| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
| PR `#21877` fused down-GEMM + combine superseding SBO | `enable_fused_grouped_gemm_combine`<br>`combine`<br>`down_gemm` | `PR #21877`<br>`python/sglang/srt/server_args.py`<br>`python/sglang/srt/layers/moe/token_dispatcher/deepep.py` | Fused combine eliminates the standalone combine window, so SBO is intentionally disabled when this path is on | If the trace discussion is about combine overlap, first classify it as this upstream fused-overlap family. |
| PR `#22410` hiSparse H2D transfer overlap with hit-attention | `transfer_stream`<br>`execute_h2d_async`<br>`hit-attention`<br>`merge_state` | `PR #22410`<br>`python/sglang/srt/layers/attention/nsa_backend.py`<br>`python/sglang/srt/hisparse/hisparse_coordinator.py` | hiSparse decode overlaps host-to-device KV transfer on a transfer stream with hit-attention on the compute stream before running miss-attention and merge | Treat hit-attention vs H2D KV transfer windows as an in-flight SGLang overlap family first. |
## 5. vLLM-origin kernel-overlap families
## 5. FlashInfer kernel-overlap families
These rows are comparative references from `flashinfer`. Use them when a trace
looks like an upstream FlashInfer overlap family even if the current `sglang`
checkout only calls part of that implementation.
| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
| FlashInfer PDL launch-overlap family | `enable_pdl`<br>`launch_with_pdl`<br>`cudaGridDependencySynchronize`<br>`cudaTriggerProgrammaticLaunchCompletion`<br>`trigger_completion_at_end=False`<br>`allreduce_fusion` | `flashinfer/norm/__init__.py`<br>`flashinfer/activation.py`<br>`flashinfer/rope.py`<br>`flashinfer/comm/allreduce.py`<br>`flashinfer/comm/trtllm_ar.py` | FlashInfer uses Programmatic Dependent Launch broadly, and the allreduce path can further advance completion so the next PDL-aware kernel overlaps on the same stream | Treat tight same-stream dependent windows and allreduce-followed-by-kernel windows as one existing FlashInfer launch-overlap family first. |
| FlashInfer CuTeDSL MoE aux-stream async-memset overlap | `aux_stream`<br>`main_event`<br>`memset_event`<br>`use_async_memset` | `flashinfer/fused_moe/cute_dsl/fused_moe.py` | Preallocated MoE output is zeroed on an auxiliary CUDA stream while GEMM1 runs on the main stream, then both streams join before finalize | Treat GEMM1 vs output-zero windows as an existing FlashInfer multi-stream overlap family. |
| FlashInfer green-context SM partition overlap | `split_device_green_ctx`<br>`split_device_green_ctx_by_sm_count`<br>`green_ctx` | `flashinfer/green_ctx.py` | CUDA green contexts partition SMs and create dedicated streams for concurrent kernel families on separate SM slices | Treat SM-partitioned concurrency as an existing FlashInfer overlap mechanism, not a novel scheduler idea. |
## 6. FlashInfer PR-backed / in-flight kernel-overlap families
| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
| PR `#2840` CuTeDSL MoE aux-stream overlap race fix | `aux_stream`<br>`use_prealloc`<br>`use_cuda_graph` | `PR #2840`<br>`flashinfer/fused_moe/cute_dsl/fused_moe.py` | Clarifies that async memset overlap is only safe for the preallocated / CUDA-graph case; non-graph mode falls back to main-stream zeroing to avoid races | Treat missing aux-stream overlap in non-graph traces as an intentional safety rule, not a novel opportunity. |
| PR `#2720` PDL runtime-API migration | `cudaGridDependencySynchronize`<br>`cudaTriggerProgrammaticLaunchCompletion`<br>`inline PTX` | `PR #2720`<br>`include/flashinfer/comm/trtllm_allreduce_fusion.cuh`<br>`include/flashinfer/pos_enc.cuh` | Repo-wide migration preserves the existing PDL overlap family while replacing inline PTX with CUDA runtime APIs across norm, RoPE, attention, and MoE codepaths | Treat PDL-looking launch groups as an upstream FlashInfer overlap family even when implementation details differ across revisions. |
## 7. TensorRT-LLM-origin kernel-overlap families
These rows are comparative references from `TensorRT-LLM`. Current mainline
TensorRT-LLM overlap rows are mostly explicit auxiliary-stream rewrites in
AutoDeploy rather than same-stream PDL windows.
| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
| TensorRT-LLM multi-stream MLA attention | `multi_stream_mla_attn`<br>`record_event_passthrough`<br>`_aux`<br>`wait_event` | `tensorrt_llm/_torch/auto_deploy/transform/library/multi_stream_attn.py`<br>`tensorrt_llm/_torch/auto_deploy/utils/multi_stream_utils.py` | AutoDeploy rewrites MLA Q/KV forks so the KV projection runs on an auxiliary stream while the Q path stays on the caller stream | Treat exposed Q-branch vs KV-branch overlap as an existing TensorRT-LLM multi-stream family first. |
| TensorRT-LLM multi-stream MoE shared-vs-routed overlap | `multi_stream_moe`<br>`begin_aux_stream_passthrough`<br>`end_aux_stream_passthrough`<br>`wait_aux_stream_passthrough` | `tensorrt_llm/_torch/auto_deploy/transform/library/multi_stream_moe.py`<br>`tensorrt_llm/_torch/auto_deploy/utils/multi_stream_utils.py` | Shared-expert work is moved to an auxiliary stream while routed-expert MoE work remains on the main stream and rejoins at the merge node | Treat shared-expert vs routed-expert windows as an existing TensorRT-LLM branch-overlap family. |
| TensorRT-LLM multi-stream FP8 GEMM fork parallelism | `multi_stream_gemm`<br>`trtllm_finegrained_fp8_linear`<br>`record_event_passthrough`<br>`_aux` | `tensorrt_llm/_torch/auto_deploy/transform/library/multi_stream_gemm.py`<br>`tensorrt_llm/_torch/auto_deploy/utils/multi_stream_utils.py` | Compiler pass identifies fork points with multiple FP8 linears and moves the largest GEMM to the auxiliary stream so sibling GEMMs overlap | Treat sibling FP8 linear branches as an existing TensorRT-LLM overlap family before designing a new stream split. |
## 8. TensorRT-LLM-origin PR-backed / in-flight kernel-overlap families
| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
| PR `#12847` `multi_stream_moe` sync fix for MLIR and piecewise cudagraphs | `multi_stream_moe`<br>`mlir_elementwise_fusion`<br>`piecewise cudagraph`<br>`caller_stream.synchronize()` | `PR #12847`<br>`tensorrt_llm/_torch/auto_deploy/transform/library/multi_stream_moe.py`<br>`tensorrt_llm/_torch/auto_deploy/utils/multi_stream_utils.py` | Open PR preserves the existing multi-stream MoE overlap family while tightening synchronization when MLIR-fused kernels or piecewise cudagraph replay are present | Treat missing or altered `multi_stream_moe` overlap under MLIR / piecewise graph modes as an in-flight TensorRT-LLM rule first. |
## 9. vLLM-origin kernel-overlap families
| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
@@ -74,13 +112,14 @@ overlap opportunity as novel.
| vLLM-origin shared-expert aux-stream overlap | `aux_stream`<br>`shared_experts_stream`<br>shared expert near router | `vllm/utils/torch_utils.py`<br>`vllm/model_executor/layers/fused_moe/runner/default_moe_runner.py` | MoE shared experts can run on a dedicated aux stream and overlap with router-side work | Treat shared-expert vs router overlap as an existing upstream sparse-model family. |
| vLLM-origin DCP async all-to-all overlap | `dcp_alltoall`<br>`all_to_all_single`<br>`async_op=True` | `vllm/v1/attention/ops/dcp_alltoall.py` | Output / LSE exchange uses async all-to-all handles instead of serializing collective completion on the main path | Treat DCP all-to-all windows as an upstream async-collective family. |
## 6. vLLM-origin PR-backed / in-flight kernel-overlap families
## 10. vLLM-origin PR-backed / in-flight kernel-overlap families
| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
| PR `#35968` DSV3.2 multi-stream indexer overlap | `weights_proj`<br>`wk`<br>`k_norm`<br>`aux_stream` | `PR #35968`<br>`vllm/model_executor/models/deepseek_v2.py`<br>`vllm/utils/torch_utils.py` | Open PR overlaps the small `weights_proj` GEMM with `wk + k_norm` on a secondary CUDA stream for decode batches instead of serializing both on the default stream | Treat this as a concrete upstream decode-time kernel-overlap family when traces show underutilized projection overlap opportunities. |
| PR `#39301` GLM5 router GEMM with PDL overlap | `TRTLLM_ENABLE_PDL`<br>`router_gemm`<br>`GLM5`<br>`FI AR RMS fusion` | `PR #39301`<br>`vllm/model_executor/layers/fused_moe/router/gate_linear.py`<br>`vllm/csrc/moe/dsv3_router_gemm_utils.h` | The GLM5 router GEMM path explicitly uses PDL so the router kernel can overlap with the preceding fused allreduce-plus-RMS block on supported GPUs | Treat router-GEMM launch overlap on GLM5-like traces as an in-flight upstream family first. |
## 7. Important toggles and caveats
## 11. Important toggles and caveats
| Toggle / env | Location | Effect on trace interpretation |
| --- | --- | --- |
@@ -89,22 +128,48 @@ overlap opportunity as novel.
| `SGLANG_DISAGG_STAGING_BUFFER` | `python/sglang/srt/environ.py` | Enables the heterogeneous-TP staging-buffer family and its overlap windows. |
| `SGLANG_STAGING_USE_TORCH` | `python/sglang/srt/disaggregation/common/staging_buffer.py` | Forces torch fallback for staging gather / scatter, so Triton staging kernels may disappear by design. |
| `SGLANG_VIT_ENABLE_CUDA_GRAPH` | `python/sglang/srt/environ.py` | Can intentionally disable vision `aux_stream` overlap. |
| `enable_pdl` / `launch_with_pdl` | `flashinfer/norm/__init__.py`<br>`flashinfer/activation.py`<br>`flashinfer/rope.py`<br>`flashinfer/fused_moe/core.py`<br>`flashinfer/comm/allreduce.py` | Enables FlashInfer PDL across many kernels; launch grouping and same-stream overlap can change substantially when it is on. |
| `trigger_completion_at_end` | `flashinfer/comm/allreduce.py` | `False` enables downstream PDL-aware overlap after FlashInfer allreduce fusion; `True` delays completion to kernel end and removes that overlap window. |
| `use_cuda_graph` | `flashinfer/fused_moe/cute_dsl/fused_moe.py` | Enables the preallocated-buffer path and the safe aux-stream async-memset overlap in FlashInfer CuTeDSL MoE. |
| `split_device_green_ctx*` | `flashinfer/green_ctx.py` | Changes trace shape by partitioning SMs into separate green contexts instead of overlapping full-device streams on the default context. |
| `multi_stream_moe` | `tensorrt_llm/_torch/auto_deploy/config/default.yaml` | Enables the TensorRT-LLM shared-expert vs routed-expert overlap family. |
| `multi_stream_mla_attn` | `tensorrt_llm/_torch/auto_deploy/config/default.yaml` | Enables the TensorRT-LLM MLA Q-vs-KV branch overlap family. |
| `multi_stream_gemm` | `tensorrt_llm/_torch/auto_deploy/config/default.yaml` | Enables generalized FP8 GEMM fork overlap in TensorRT-LLM AutoDeploy. |
| `mlir_elementwise_fusion` | `tensorrt_llm/_torch/auto_deploy/config/default.yaml` | Can absorb merge adds into larger fused kernels, so missing explicit merge nodes in TensorRT-LLM multi-stream traces may be intentional. |
| `enable_torch_compile` | `python/sglang/srt/server_args.py`<br>`python/sglang/multimodal_gen/runtime/server_args.py` | Compiler-generated reordering can hide or rename overlap windows. |
| `enable_fused_grouped_gemm_combine` | `PR #21877` | In-flight path that intentionally disables SBO because combine is folded into down-GEMM. |
| `PassConfig.enable_sp` | `vllm/config/compilation.py` | Enables vLLM's sequence-parallel staging family that creates RS / AG overlap opportunities. |
| `PassConfig.fuse_gemm_comms` | `vllm/config/compilation.py` | Enables AsyncTP GEMM + collective overlap and auto-enables `enable_sp` when valid. |
## 8. Suggested refresh commands
## 12. Suggested refresh commands
These commands are only for maintainers refreshing this catalog by rescanning
the local source trees. They are not used by the triage scripts at runtime.
```bash
# Optional sibling checkouts used for comparative scanning:
FLASHINFER_REPO=${FLASHINFER_REPO:-../flashinfer}
TRTLLM_REPO=${TRTLLM_REPO:-../TensorRT-LLM}
VLLM_REPO=${VLLM_REPO:-../vllm}
rg -n "single_batch_overlap|alt_stream|shared_expert|scatter_stream|_fused_gather_to_staging_kernel|_fused_scatter_from_staging_kernel|async_op=True" python/sglang
rg -n "apply_qk_norm|vision.py|ring_attn|all_to_all_single|reorder_for_compute_comm_overlap|use_dual_stream" python/sglang/multimodal_gen python/sglang/srt
git log --all --format='%h %s' | rg -i 'fused|fusion|overlap|combine|all_to_all|ring attn|stream|triton|cutedsl|cuda'
rg -n "fuse_gemm_comms|enable_sp|fused_matmul_reduce_scatter|fused_all_gather_matmul|shared_experts_stream|dcp_alltoall|async_op=True|aux_stream|maybe_execute_in_parallel" /Users/bbuf/工作目录/Common/vllm/vllm /Users/bbuf/工作目录/Common/vllm/docs/design/fusions.md
git -C /Users/bbuf/工作目录/Common/vllm log --all --format='%h %s' | rg -i 'fused|fusion|overlap|allreduce|reduce-scatter|all-gather|all_to_all|stream|multi-stream|triton|cuda|router'
rg -n "enable_pdl|launch_with_pdl|trigger_completion_at_end|aux_stream|use_async_memset|split_device_green_ctx|split_device_green_ctx_by_sm_count" "$FLASHINFER_REPO/flashinfer" "$FLASHINFER_REPO/include"
git -C "$FLASHINFER_REPO" log --all --format='%h %s' | rg -i 'fused|fusion|overlap|pdl|stream|rope|kv|quant|topk|moe'
rg -n "multi_stream_moe|multi_stream_mla_attn|multi_stream_gemm|record_event_passthrough|begin_aux_stream_passthrough|end_aux_stream_passthrough|wait_aux_stream_passthrough" "$TRTLLM_REPO/tensorrt_llm/_torch"
rg -n "mlir_elementwise_fusion|piecewise|cudagraph|caller_stream.synchronize" "$TRTLLM_REPO/tensorrt_llm/_torch"
git -C "$TRTLLM_REPO" log --all --format='%h %s' | rg -i 'overlap|multi-stream|aux stream|cudagraph|mlir|stream|flashinfer|moe|mla'
rg -n "fuse_gemm_comms|enable_sp|fused_matmul_reduce_scatter|fused_all_gather_matmul|shared_experts_stream|dcp_alltoall|async_op=True|aux_stream|maybe_execute_in_parallel" "$VLLM_REPO/vllm" "$VLLM_REPO/docs/design/fusions.md"
git -C "$VLLM_REPO" log --all --format='%h %s' | rg -i 'fused|fusion|overlap|allreduce|reduce-scatter|all-gather|all_to_all|stream|multi-stream|triton|cuda|router'
# GitHub PR scan terms for the connector or web UI:
# "fused OR overlap repo:sgl-project/sglang"
# "triton OR cutedsl OR cuda overlap repo:sgl-project/sglang"
# "fused OR overlap repo:flashinfer-ai/flashinfer"
# "pdl OR aux_stream OR green_ctx repo:flashinfer-ai/flashinfer"
# "fused OR overlap repo:NVIDIA/TensorRT-LLM"
# "multi-stream OR aux_stream OR cudagraph repo:NVIDIA/TensorRT-LLM"
# "mlir OR piecewise OR flashinfer repo:NVIDIA/TensorRT-LLM"
# "fused OR overlap repo:vllm-project/vllm"
# "triton OR cuda overlap repo:vllm-project/vllm"
# "multi-stream OR aux_stream overlap repo:vllm-project/vllm"
.claude/skills/sglang-torch-profiler-analysis/references/trace-workflow.md
-119
View File
@@ -1,119 +0,0 @@
# Trace Workflow
This skill is based on SGLang's existing profiling workflow.
Use:
- two traces for `triage`
- one trace for `breakdown`
- two traces for `overlap`
- optional post-processing for `perfetto-fix`
`profile_by_stage` is still useful on normal non-PD serving because it separates prefill and decode. PD disaggregation adds an extra requirement beyond that: prefill workers and decode workers must be profiled separately.
## Existing SGLang Sources
- `sglang/.claude/skills/generate-profile/SKILL.md`
- `sglang/docs/developer_guide/benchmark_and_profiling.md`
- `sglang/docs/diffusion/performance/profiling.md`
- `sglang/python/sglang/profiler.py`
- `sglang/python/sglang/srt/utils/profile_utils.py`
- `sglang/python/sglang/srt/utils/profile_merger.py`
## Required Two-Stage Flow
### Stage 1: Mapping trace
Collect a graph-off trace first.
Recommended properties:
- disable `cuda graph`
- disable `piecewise cuda graph` if it would otherwise hide launch attribution
- keep the same model, parallel shape, backend choices, and request pattern as much as possible
Purpose:
- preserve clean kernel launch attribution
- recover `kernel -> cpu_op -> python scope`
### Stage 2: Formal trace
Collect a second trace with the real serving optimizations enabled.
Recommended properties:
- enable `cuda graph` if the real deployment uses it
- enable `piecewise cuda graph` when the model normally captures it
- keep the production MoE, attention, communication, and quantization backends
Purpose:
- measure real overlap under the real schedule
- decide whether a code path still has overlap headroom
The final compact report should be built from:
- source attribution from stage 1
- overlap conclusions from stage 2
The merged skill's `triage` command turns that into three tables:
- kernel table
- overlap-opportunity table
- fuse-opportunity table
## Ways To Produce A Trace
### Live server
```bash
python3 -m sglang.profiler --url http://127.0.0.1:30000 --num-steps 5
```
### One-shot request plus profile
```bash
python3 -m sglang.test.send_one --profile
```
### Bench serving
```bash
export SGLANG_TORCH_PROFILER_DIR=/tmp/sglang-profile
python3 -m sglang.bench_serving --backend sglang --num-prompts 10 --profile
```
If you only call `python3 -m sglang.profiler`, remember that something still has to drive requests through the server while profiling is active. The merged skill's live URL flows handle this automatically by sending a small probe workload.
## Expected Output
Typical trace outputs are:
- `<profile_id>-TP-0.trace.json.gz`
- `<profile_id>-TP-0-DP-0-PP-0-EP-0.trace.json.gz`
- `merged-<profile_id>.trace.json.gz`
- `server_args.json`
For overlap analysis, prefer a single-rank trace over a merged multi-rank trace.
## Optional Perfetto Repair
When Perfetto fails to render clearly overlapping events on the same logical lane, the unified script exposes:
```bash
python3 scripts/analyze_sglang_torch_profile.py perfetto-fix --input /path/to/trace.json.gz
```
This is intentionally a narrow repair step. Do not make it part of the default profiling workflow unless rendering is actually broken.
## Why The Mapping Trace Must Exist
A single graph-on trace may still tell you that overlap is poor, but it is often not enough to say which Python code path owns the kernel.
That is why this skill requires:
- one trace for readable source attribution
- one trace for real overlap behavior
Do not call the graph-off trace a "fast profile" in the final write-up. Its role is source mapping, not shortcutting the real analysis.
.claude/skills/sglang-torch-profiler-analysis/references/validated-workflows.md
-263
View File
@@ -1,263 +0,0 @@
# Validated Workflows
These were the concrete workflows used to validate the skill design against real remote GPU environments.
## 1. Small single-GPU Qwen
Validated as a two-pass workflow on H100.
### Pass 1: no-CUDA-graph mapping pre-pass
Validated example:
```bash
export SGLANG_TORCH_PROFILER_DIR=/tmp/sglang_torch_profile_qwen25_15b_map
CUDA_VISIBLE_DEVICES=5 FLASHINFER_DISABLE_VERSION_CHECK=1 python3 -m sglang.launch_server \
--model-path Qwen/Qwen2.5-1.5B-Instruct \
--host 127.0.0.1 \
--port 32240 \
--disable-cuda-graph \
--disable-piecewise-cuda-graph
```
Then:
```bash
python3 scripts/analyze_sglang_torch_profile.py breakdown \
--url http://127.0.0.1:32240 \
--num-steps 5 \
--profile-by-stage \
--profile-prefix qwen25_15b_map \
--export-kernel-map /tmp/qwen25_15b_kernel_map.json
```
### Pass 2: final optimized profile
```bash
export SGLANG_TORCH_PROFILER_DIR=/tmp/sglang_torch_profile_qwen25_15b_final
CUDA_VISIBLE_DEVICES=5 FLASHINFER_DISABLE_VERSION_CHECK=1 python3 -m sglang.launch_server \
--model-path Qwen/Qwen2.5-1.5B-Instruct \
--host 127.0.0.1 \
--port 32241
FLASHINFER_DISABLE_VERSION_CHECK=1 python3 -m sglang.profiler \
--url http://127.0.0.1:32241 \
--num-steps 5 \
--profile-by-stage \
--profile-prefix qwen25_15b_final
python3 scripts/analyze_sglang_torch_profile.py breakdown \
--input /tmp/sglang_torch_profile_qwen25_15b_final \
--kernel-map /tmp/qwen25_15b_kernel_map.json
```
## 2. Multi-GPU Qwen
Validated as a two-pass workflow on H100 with `Qwen/Qwen3-32B`.
### Pass 1: no-CUDA-graph mapping pre-pass
Validated target options:
- `Qwen/Qwen3-32B` on H100 with `--tp 2`
- `Qwen/Qwen3-Next-80B-A3B-Instruct` on H200 with `--tp 4`
Example:
```bash
export SGLANG_TORCH_PROFILER_DIR=/tmp/sglang_torch_profile_qwen32b_map
CUDA_VISIBLE_DEVICES=1,2 FLASHINFER_DISABLE_VERSION_CHECK=1 python3 -m sglang.launch_server \
--model-path Qwen/Qwen3-32B \
--tp 2 \
--host 127.0.0.1 \
--port 32040 \
--disable-cuda-graph \
--disable-piecewise-cuda-graph
```
Then:
```bash
python3 scripts/analyze_sglang_torch_profile.py breakdown \
--url http://127.0.0.1:32040 \
--num-steps 5 \
--profile-by-stage \
--profile-prefix qwen32b_map \
--export-kernel-map /tmp/qwen32b_kernel_map.json
```
### Pass 2: final optimized profile
```bash
export SGLANG_TORCH_PROFILER_DIR=/tmp/sglang_torch_profile_qwen32b_final
CUDA_VISIBLE_DEVICES=1,2 FLASHINFER_DISABLE_VERSION_CHECK=1 python3 -m sglang.launch_server \
--model-path Qwen/Qwen3-32B \
--tp 2 \
--host 127.0.0.1 \
--port 32041
FLASHINFER_DISABLE_VERSION_CHECK=1 python3 -m sglang.profiler \
--url http://127.0.0.1:32041 \
--num-steps 5 \
--profile-by-stage \
--profile-prefix qwen32b_final
python3 scripts/analyze_sglang_torch_profile.py breakdown \
--input /tmp/sglang_torch_profile_qwen32b_final \
--kernel-map /tmp/qwen32b_kernel_map.json
```
## Notes
- Prefer `TP-0` traces first for kernel share analysis and for the exported kernel map.
- If the directory only contains merged traces, state that explicitly in the final conclusions.
- For stage-aware comparisons, analyze `EXTEND` and `DECODE` separately before summarizing the overall model behavior.
- On current SGLang builds, add `--disable-piecewise-cuda-graph` together with `--disable-cuda-graph` for the mapping pass, otherwise extend/prefill may still run under piecewise CUDA graph.
- `Qwen/Qwen3-4B-Instruct-2507` was present in H200 cache during validation but did not work on the validated stack because of a `Qwen3Config.rope_parameters` compatibility issue.
## 3. Deliberately broken TP fusion rediscovery
Validated on B200 with `Qwen/Qwen2.5-0.5B-Instruct`, `TP=2`.
The validation intentionally commented out the fused TP all-reduce + RMSNorm path inside:
- `python/sglang/srt/layers/layernorm.py`
and forced the code to fall back to:
- plain `tensor_model_parallel_all_reduce`
- then ordinary `norm_module.forward(...)`
### Mapping pass
Graph-off server:
```bash
CUDA_VISIBLE_DEVICES=6,7 python3 -m sglang.launch_server \
--model-path Qwen/Qwen2.5-0.5B-Instruct \
--tp 2 \
--host 127.0.0.1 \
--port 32260 \
--disable-cuda-graph \
--disable-piecewise-cuda-graph
```
Then:
```bash
python3 scripts/analyze_sglang_torch_profile.py breakdown \
--url http://127.0.0.1:32260 \
--num-steps 5 \
--profile-by-stage \
--profile-prefix qwen25_tp2_map \
--export-kernel-map /tmp/qwen25_tp2_map_kernel_map.json
```
Observed result:
- the fuse table rediscovered `TP all-reduce + residual/RMSNorm`
- it pointed back to `python/sglang/srt/layers/layernorm.py:89 _forward_with_allreduce_fusion`
### Formal pass
Graph-on server with the intentionally broken code still in place:
```bash
CUDA_VISIBLE_DEVICES=6,7 python3 -m sglang.launch_server \
--model-path Qwen/Qwen2.5-0.5B-Instruct \
--tp 2 \
--host 127.0.0.1 \
--port 32261
```
Then:
```bash
python3 scripts/analyze_sglang_torch_profile.py triage \
--mapping-input /tmp/sglang-torch-profile-_7gd033i/1775035260.3725493 \
--formal-input /tmp/sglang-torch-profile-64oszwu2/1775035372.2416728
```
Observed result:
- the kernel table showed `void sglang::cross_device_reduce_1stage<__nv_bfloat16, 2>` at roughly `31%` decode share
- the overlap table surfaced that same communication kernel as a top headroom row
- the fuse table again flagged `TP all-reduce + residual/RMSNorm`
This validation demonstrates that the skill can rediscover a real missing fusion path and also surface the exposed overlap opportunity that appears when the fusion is removed.
## 4. Dense Qwen3 QK-norm + RoPE rediscovery
Validated on B200 with `Qwen/Qwen3-32B`, `TP=2`.
### Mapping pass
Graph-off server:
```bash
CUDA_VISIBLE_DEVICES=1,2 python3 -m sglang.launch_server \
--model-path Qwen/Qwen3-32B \
--tp 2 \
--host 127.0.0.1 \
--port 32320 \
--disable-cuda-graph \
--disable-piecewise-cuda-graph
```
### Formal pass
Graph-on server:
```bash
CUDA_VISIBLE_DEVICES=1,2 python3 -m sglang.launch_server \
--model-path Qwen/Qwen3-32B \
--tp 2 \
--host 127.0.0.1 \
--port 32321
```
Observed result:
- on the current `sm100 + trtllm_mha` stack, `apply_qk_norm` and `RoPE` may collapse into a shared generic `void` kernel row
- the useful source evidence still survives in the mapped Python locations:
- `python/sglang/jit_kernel/rope.py:179 apply_rope_with_cos_sin_cache_inplace`
- `python/sglang/srt/models/utils.py:204 apply_qk_norm`
- the fuse detector therefore needs to treat a shared kernel row containing both `QK norm` and `RoPE` evidence as valid
- after broadening the rule, triage produced:
- `decode | Q/K RMSNorm + RoPE before attention | Conditional | 2.34 ms | 4.5%`
This validation demonstrates that dense-Qwen3 fuse detection should be source-evidence driven, not tied only to `norm` or `rope` kernel categories.
## 5. Single-GPU negative control
Validated on B200 with `Qwen/Qwen2.5-0.5B-Instruct`, single GPU, `TP=1`.
Observed result:
- the three main tables were still produced normally
- no medium-confidence source-backed fusion opportunity was emitted
- the skill did not incorrectly report:
- `TP all-reduce + residual/RMSNorm`
- `Q/K RMSNorm + RoPE before attention`
This validation is the negative control for avoiding TP-specific false positives when no TP communication exists.
## 6. MiniMax overlap-heavy trace validation
Validated on B200 using previously captured `MiniMaxAI/MiniMax-M2.5` mapping and formal traces.
Current note:
- a fresh launch on the current B200 `main` repo failed before profiling with:
- `AttributeError: 'MiniMaxM2Config' object has no attribute 'rope_theta'`
- that load-time issue is separate from the profiler skill itself
Using the existing B200 traces still validated the triage behavior:
- the kernel table was dominated by:
- `void sglang::cross_device_reduce_1stage<__half, 4>`
- `fused_moe_kernel`
- the overlap table stayed compact and preserved only actionable overlap rows plus `low-roi-hidden` deprioritization rows
- the fuse table flagged `TP all-reduce + residual/RMSNorm`
This validation demonstrates that the compact three-table artifact still stays readable on a communication-heavy MoE model.
.claude/skills/sglang-torch-profiler-analysis/scripts/analyze_sglang_torch_profile.py
+219 -177
View File
@@ -1,4 +1,4 @@
"""Unified entrypoint for SGLang torch-profiler analysis workflows."""
"""Compact triage entrypoint for SGLang torch-profiler analysis."""
from __future__ import annotations
@@ -8,58 +8,52 @@ from collections import defaultdict
from pathlib import Path
from typing import Dict, List, Optional, Sequence, Tuple
import analyze_sglang_llm_torch_profile as breakdown_cli
import analyze_sglang_profiler_overlap as overlap_cli
import triage_kernel_helpers as kernel_helpers
import triage_overlap_helpers as overlap_helpers
from profile_common import (
discover_trace_targets,
load_server_args,
load_trace_json,
parse_stage,
run_profiler,
write_perfetto_compatible_trace,
)
MIN_RENDER_SHARE_PCT = 1.0
def build_top_level_parser() -> argparse.ArgumentParser:
def build_triage_parser() -> argparse.ArgumentParser:
parser = argparse.ArgumentParser(
prog="analyze_sglang_torch_profile.py",
description=(
"Unified torch-profiler entrypoint for SGLang. "
"Use `breakdown` for kernel/category share analysis, "
"`overlap` for two-trace overlap analysis, `triage` for the compact "
"three-table workflow, or `perfetto-fix` to rewrite a trace into a "
"more Perfetto-friendly form."
"Compact SGLang torch-profiler triage entrypoint. "
"This prints three tables: kernel mapping, overlap opportunities, "
"and fuse opportunities. "
"Use either a single trace/profile input or a mapping+formal two-trace pair."
),
)
parser.add_argument(
"command",
nargs="?",
choices=("breakdown", "overlap", "triage", "perfetto-fix"),
help="Subcommand to run.",
)
return parser
def parse_perfetto_fix_args(argv: Sequence[str]) -> argparse.Namespace:
parser = argparse.ArgumentParser(
prog="analyze_sglang_torch_profile.py perfetto-fix",
description="Rewrite a trace so overlapping kernel lanes render more reliably in Perfetto.",
"--input",
type=str,
default=None,
help="Single trace file or profile directory to triage.",
)
parser.add_argument(
"--input", required=True, help="Input trace.json or trace.json.gz path."
"--url",
type=str,
default=None,
help="Running SGLang server URL for single-trace triage.",
)
parser.add_argument("--output", default=None, help="Optional output path.")
return parser.parse_args(argv)
def parse_triage_args(argv: Sequence[str]) -> argparse.Namespace:
parser = argparse.ArgumentParser(
prog="analyze_sglang_torch_profile.py triage",
description=(
"Run the compact SGLang torch-profiler triage workflow. "
"This prints three stage-aware tables: kernel mapping, overlap opportunities, "
"and fuse opportunities."
),
parser.add_argument(
"--output-dir",
type=str,
default=None,
help="Trace output dir when using --url.",
)
parser.add_argument(
"--profile-prefix",
type=str,
default="triage-trace",
help="Profile prefix when generating a single trace from --url.",
)
parser.add_argument(
"--mapping-input",
@@ -156,7 +150,34 @@ def parse_triage_args(argv: Sequence[str]) -> argparse.Namespace:
default=0,
help="How many overlap rows to print per stage. Use 0 for all kernels.",
)
return parser
def parse_triage_args(argv: Sequence[str]) -> argparse.Namespace:
parser = build_triage_parser()
args = parser.parse_args(argv)
single_trace_mode = bool(args.input) or bool(args.url)
dual_trace_mode = any(
[
args.mapping_input,
args.mapping_url,
args.formal_input,
args.formal_url,
]
)
if single_trace_mode and dual_trace_mode:
parser.error(
"Use either single-trace mode (--input/--url) or two-trace mode "
"(--mapping-* plus --formal-*), not both."
)
if single_trace_mode:
if bool(args.input) == bool(args.url):
parser.error("Provide exactly one of --input or --url.")
return args
if bool(args.mapping_input) == bool(args.mapping_url):
parser.error("Provide exactly one of --mapping-input or --mapping-url.")
if bool(args.formal_input) == bool(args.formal_url):
@@ -201,26 +222,23 @@ def resolve_profile_targets(
def build_mapping_kernel_map(trace_paths: Sequence[Path]) -> dict:
# The graph-off mapping trace is only used to learn stable
# kernel -> Python/CPU-op attribution. The final percentages still come from
# the formal trace.
stage_site_stats = defaultdict(
lambda: defaultdict(lambda: defaultdict(breakdown_cli.MappingSiteAggregate))
lambda: defaultdict(lambda: defaultdict(kernel_helpers.MappingSiteAggregate))
)
stage_kernel_categories: Dict[str, Dict[str, str]] = defaultdict(dict)
global_site_stats = defaultdict(
lambda: defaultdict(breakdown_cli.MappingSiteAggregate)
lambda: defaultdict(kernel_helpers.MappingSiteAggregate)
)
global_kernel_categories: Dict[str, str] = {}
for trace_path in trace_paths:
trace = load_trace_json(trace_path)
kernels, cpu_ops, python_frames, launch_events, _, _ = (
breakdown_cli.extract_trace_data(trace)
kernel_helpers.extract_trace_data(trace)
)
cpu_ops_by_external_id = breakdown_cli.build_cpu_op_index(cpu_ops)
launches_by_correlation = breakdown_cli.build_launch_index(launch_events)
local_site_stats = breakdown_cli.aggregate_kernel_sites(
cpu_ops_by_external_id = kernel_helpers.build_cpu_op_index(cpu_ops)
launches_by_correlation = kernel_helpers.build_launch_index(launch_events)
local_site_stats = kernel_helpers.aggregate_kernel_sites(
kernels,
cpu_ops_by_external_id,
python_frames,
@@ -230,18 +248,18 @@ def build_mapping_kernel_map(trace_paths: Sequence[Path]) -> dict:
kernel_categories = {
kernel.canonical_name: kernel.category for kernel in kernels
}
breakdown_cli.merge_site_stats(stage_site_stats[stage], local_site_stats)
breakdown_cli.merge_site_stats(global_site_stats, local_site_stats)
kernel_helpers.merge_site_stats(stage_site_stats[stage], local_site_stats)
kernel_helpers.merge_site_stats(global_site_stats, local_site_stats)
stage_kernel_categories[stage].update(kernel_categories)
global_kernel_categories.update(kernel_categories)
stage_payloads = {
stage: breakdown_cli.build_stage_payload(
stage: kernel_helpers.build_stage_payload(
dict(site_stats), stage_kernel_categories.get(stage, {})
)
for stage, site_stats in stage_site_stats.items()
}
global_payload = breakdown_cli.build_stage_payload(
global_payload = kernel_helpers.build_stage_payload(
dict(global_site_stats), global_kernel_categories
)
return {"stages": stage_payloads, "global": global_payload}
@@ -252,7 +270,7 @@ def stage_index(stage: str) -> int:
def stage_display(stage: str) -> str:
return breakdown_cli.stage_label(stage)
return kernel_helpers.stage_label(stage)
def pick_trace_for_stage(stage_to_trace: Dict[str, Path], stage: str) -> Optional[Path]:
@@ -282,14 +300,14 @@ def render_kernel_table(rows: Sequence[dict]) -> List[str]:
for row in rows:
lines.append(
"| {stage} | {kernel} | {category} | {gpu_time} | {share:.1f}% | {launches} | {location} | {cpu_op} |".format(
stage=breakdown_cli.escape_md_cell(stage_display(row["stage"])),
kernel=breakdown_cli.escape_md_cell(row["kernel"]),
category=breakdown_cli.escape_md_cell(row["category"]),
gpu_time=breakdown_cli.format_ms(row["total_us"]),
stage=kernel_helpers.escape_md_cell(stage_display(row["stage"])),
kernel=kernel_helpers.escape_md_cell(row["kernel"]),
category=kernel_helpers.escape_md_cell(row["category"]),
gpu_time=kernel_helpers.format_ms(row["total_us"]),
share=row["share_pct"],
launches=row["launches"],
location=breakdown_cli.escape_md_cell(row["location"]),
cpu_op=breakdown_cli.escape_md_cell(row["cpu_op"]),
location=kernel_helpers.escape_md_cell(row["location"]),
cpu_op=kernel_helpers.escape_md_cell(row["cpu_op"]),
)
)
return lines
@@ -300,6 +318,11 @@ def render_overlap_table(rows: Sequence[dict]) -> List[str]:
"| Stage | Priority | Verdict | Kernel | Python scope | Formal signal | Dep risk | Recommendation |",
"| --- | --- | --- | --- | --- | --- | --- | --- |",
]
if not rows:
lines.append(
"| - | - | - | No actionable overlap rows. Use mapping/formal two-trace triage for stronger overlap conclusions. | - | - | - | - |"
)
return lines
for row in rows:
formal_signal = (
f"{row['total_us']:.1f} us, share {row['share_pct']:.1f}%, "
@@ -309,13 +332,13 @@ def render_overlap_table(rows: Sequence[dict]) -> List[str]:
"| "
+ " | ".join(
[
breakdown_cli.escape_md_cell(stage_display(row["stage"])),
kernel_helpers.escape_md_cell(stage_display(row["stage"])),
row["priority"],
row["verdict"],
breakdown_cli.escape_md_cell(row["kernel"]),
breakdown_cli.escape_md_cell(row["python_scope"]),
breakdown_cli.escape_md_cell(formal_signal),
overlap_cli.dependency_risk_label(row["dependency_signal"]),
kernel_helpers.escape_md_cell(row["kernel"]),
kernel_helpers.escape_md_cell(row["python_scope"]),
kernel_helpers.escape_md_cell(formal_signal),
overlap_helpers.dependency_risk_label(row["dependency_signal"]),
row["recommendation"],
]
)
@@ -337,39 +360,52 @@ def render_fuse_table(rows: Sequence[dict]) -> List[str]:
for row in rows:
lines.append(
"| {stage} | {pattern} | {confidence} | {gpu_time} | {share:.1f}% | {evidence} | {current_locations} | {candidate_path} | {rationale} |".format(
stage=breakdown_cli.escape_md_cell(stage_display(row["stage"])),
pattern=breakdown_cli.escape_md_cell(row["pattern"]),
confidence=breakdown_cli.escape_md_cell(row["confidence"]),
gpu_time=breakdown_cli.format_ms(row["related_us"]),
stage=kernel_helpers.escape_md_cell(stage_display(row["stage"])),
pattern=kernel_helpers.escape_md_cell(row["pattern"]),
confidence=kernel_helpers.escape_md_cell(row["confidence"]),
gpu_time=kernel_helpers.format_ms(row["related_us"]),
share=row["share_pct"],
evidence=breakdown_cli.escape_md_cell(row["evidence"]),
current_locations=breakdown_cli.escape_md_cell(
evidence=kernel_helpers.escape_md_cell(row["evidence"]),
current_locations=kernel_helpers.escape_md_cell(
row["current_locations"]
),
candidate_path=breakdown_cli.escape_md_cell(row["candidate_path"]),
rationale=breakdown_cli.escape_md_cell(row["rationale"]),
candidate_path=kernel_helpers.escape_md_cell(row["candidate_path"]),
rationale=kernel_helpers.escape_md_cell(row["rationale"]),
)
)
return lines
def run_triage(args: argparse.Namespace) -> int:
mapping_traces, mapping_server_args = resolve_profile_targets(
label="mapping",
input_path=args.mapping_input,
url=args.mapping_url,
output_dir=args.mapping_output_dir,
profile_prefix=args.mapping_profile_prefix,
args=args,
)
formal_traces, formal_server_args = resolve_profile_targets(
label="formal",
input_path=args.formal_input,
url=args.formal_url,
output_dir=args.formal_output_dir,
profile_prefix=args.formal_profile_prefix,
args=args,
)
single_trace_mode = bool(args.input) or bool(args.url)
if single_trace_mode:
formal_traces, formal_server_args = resolve_profile_targets(
label="input",
input_path=args.input,
url=args.url,
output_dir=args.output_dir,
profile_prefix=args.profile_prefix,
args=args,
)
mapping_traces = formal_traces
mapping_server_args = formal_server_args
else:
mapping_traces, mapping_server_args = resolve_profile_targets(
label="mapping",
input_path=args.mapping_input,
url=args.mapping_url,
output_dir=args.mapping_output_dir,
profile_prefix=args.mapping_profile_prefix,
args=args,
)
formal_traces, formal_server_args = resolve_profile_targets(
label="formal",
input_path=args.formal_input,
url=args.formal_url,
output_dir=args.formal_output_dir,
profile_prefix=args.formal_profile_prefix,
args=args,
)
mapping_kernel_map = build_mapping_kernel_map(mapping_traces)
@@ -378,18 +414,18 @@ def run_triage(args: argparse.Namespace) -> int:
for formal_trace in formal_traces:
trace = load_trace_json(formal_trace)
kernels, _, _, _, _, _ = breakdown_cli.extract_trace_data(trace)
kernels, _, _, _, _, _ = kernel_helpers.extract_trace_data(trace)
if not kernels:
continue
stage = parse_stage(formal_trace)
total_us = sum(kernel.dur for kernel in kernels)
kernel_stats = breakdown_cli.aggregate(
kernel_stats = kernel_helpers.aggregate(
kernels, key_fn=lambda item: item.canonical_name
)
kernel_categories = {
kernel.canonical_name: kernel.category for kernel in kernels
}
full_kernel_rows = breakdown_cli.build_kernel_rows(
full_kernel_rows = kernel_helpers.build_kernel_rows(
stage=stage,
kernel_stats=kernel_stats,
kernel_categories=kernel_categories,
@@ -398,35 +434,41 @@ def run_triage(args: argparse.Namespace) -> int:
),
external_kernel_map=mapping_kernel_map,
)
visible_kernel_rows = breakdown_cli.limit_kernel_rows(
visible_kernel_rows = kernel_helpers.limit_kernel_rows(
full_kernel_rows, args.kernel_table_limit
)
for row in visible_kernel_rows:
share_pct = kernel_helpers.pct(row.total_us, total_us)
if share_pct < MIN_RENDER_SHARE_PCT:
continue
kernel_rows_rendered.append(
{
"stage": stage,
"kernel": row.name,
"category": row.category,
"total_us": row.total_us,
"share_pct": breakdown_cli.pct(row.total_us, total_us),
"share_pct": share_pct,
"launches": row.aggregate.count,
"location": row.location,
"cpu_op": row.cpu_op,
}
)
for item in breakdown_cli.detect_fusion_opportunities(
for item in kernel_helpers.detect_fusion_opportunities(
stage=stage,
kernel_rows=full_kernel_rows,
total_us=total_us,
server_args=formal_server_args or mapping_server_args,
):
share_pct = kernel_helpers.pct(item.related_us, total_us)
if share_pct < MIN_RENDER_SHARE_PCT:
continue
fuse_rows_rendered.append(
{
"stage": stage,
"pattern": item.pattern,
"confidence": item.confidence,
"related_us": item.related_us,
"share_pct": breakdown_cli.pct(item.related_us, total_us),
"share_pct": share_pct,
"evidence": item.evidence,
"current_locations": item.current_locations,
"candidate_path": item.candidate_path,
@@ -437,76 +479,86 @@ def run_triage(args: argparse.Namespace) -> int:
mapping_stage_map = build_stage_trace_map(mapping_traces)
formal_stage_map = build_stage_trace_map(formal_traces)
overlap_rows_rendered: List[dict] = []
for stage in sorted(formal_stage_map, key=stage_index):
formal_trace = formal_stage_map[stage]
mapping_trace = pick_trace_for_stage(mapping_stage_map, stage)
if mapping_trace is None:
continue
mapping_trace_json = load_trace_json(mapping_trace)
mapping_events, mapping_pid = overlap_cli.extract_kernel_events(
mapping_trace_json, args.pid_substring
)
if not mapping_events:
continue
formal_trace_json = load_trace_json(formal_trace)
formal_events, formal_pid = overlap_cli.extract_kernel_events(
formal_trace_json, args.pid_substring
)
if not formal_events:
continue
mapping_bundle = overlap_cli.TraceBundle(
label=f"mapping-{stage}",
trace_path=mapping_trace,
server_args=mapping_server_args,
raw_events=mapping_trace_json.get(
"traceEvents",
mapping_trace_json if isinstance(mapping_trace_json, list) else [],
),
events=mapping_events,
pid=mapping_pid,
)
formal_bundle = overlap_cli.TraceBundle(
label=f"formal-{stage}",
trace_path=formal_trace,
server_args=formal_server_args,
raw_events=formal_trace_json.get(
"traceEvents",
formal_trace_json if isinstance(formal_trace_json, list) else [],
),
events=formal_events,
pid=formal_pid,
)
formal_bundle.overlap_stats = overlap_cli.analyze_overlap(formal_bundle.events)
aggregates = overlap_cli.aggregate_events(formal_bundle.events)
source_map = overlap_cli.build_kernel_source_map(mapping_bundle)
stage_rows = overlap_cli.build_action_rows(
aggregates,
source_map,
formal_bundle.events,
formal_bundle.overlap_stats["total_busy_us"],
table_limit=max(0, args.overlap_table_limit),
)
for row in stage_rows:
overlap_rows_rendered.append(
{
"stage": stage,
"priority": row.priority,
"verdict": row.verdict,
"kernel": row.kernel,
"python_scope": row.python_scope,
"total_us": row.total_us,
"share_pct": row.share_pct,
"exclusive_ratio": row.exclusive_ratio,
"hidden_ratio": row.hidden_ratio,
"dependency_signal": row.dependency_signal,
"recommendation": row.recommendation,
}
if not single_trace_mode:
for stage in sorted(formal_stage_map, key=stage_index):
formal_trace = formal_stage_map[stage]
mapping_trace = pick_trace_for_stage(mapping_stage_map, stage)
if mapping_trace is None:
continue
mapping_trace_json = load_trace_json(mapping_trace)
mapping_events, mapping_pid = overlap_helpers.extract_kernel_events(
mapping_trace_json, args.pid_substring
)
if not mapping_events:
continue
formal_trace_json = load_trace_json(formal_trace)
formal_events, formal_pid = overlap_helpers.extract_kernel_events(
formal_trace_json, args.pid_substring
)
if not formal_events:
continue
mapping_bundle = overlap_helpers.TraceBundle(
label=f"mapping-{stage}",
trace_path=mapping_trace,
server_args=mapping_server_args,
raw_events=mapping_trace_json.get(
"traceEvents",
mapping_trace_json if isinstance(mapping_trace_json, list) else [],
),
events=mapping_events,
pid=mapping_pid,
)
formal_bundle = overlap_helpers.TraceBundle(
label=f"formal-{stage}",
trace_path=formal_trace,
server_args=formal_server_args,
raw_events=formal_trace_json.get(
"traceEvents",
formal_trace_json if isinstance(formal_trace_json, list) else [],
),
events=formal_events,
pid=formal_pid,
)
formal_bundle.overlap_stats = overlap_helpers.analyze_overlap(
formal_bundle.events
)
aggregates = overlap_helpers.aggregate_events(formal_bundle.events)
source_map = overlap_helpers.build_kernel_source_map(mapping_bundle)
stage_rows = overlap_helpers.build_action_rows(
aggregates,
source_map,
formal_bundle.events,
formal_bundle.overlap_stats["total_busy_us"],
table_limit=max(0, args.overlap_table_limit),
)
for row in stage_rows:
if row.share_pct < MIN_RENDER_SHARE_PCT:
continue
overlap_rows_rendered.append(
{
"stage": stage,
"priority": row.priority,
"verdict": row.verdict,
"kernel": row.kernel,
"python_scope": row.python_scope,
"total_us": row.total_us,
"share_pct": row.share_pct,
"exclusive_ratio": row.exclusive_ratio,
"hidden_ratio": row.hidden_ratio,
"dependency_signal": row.dependency_signal,
"recommendation": row.recommendation,
}
)
lines: List[str] = []
lines.append("Triage View")
lines.append(f"Mapping traces: {', '.join(str(path) for path in mapping_traces)}")
lines.append(f"Formal traces: {', '.join(str(path) for path in formal_traces)}")
if single_trace_mode:
lines.append(f"Input traces: {', '.join(str(path) for path in formal_traces)}")
else:
lines.append(
f"Mapping traces: {', '.join(str(path) for path in mapping_traces)}"
)
lines.append(f"Formal traces: {', '.join(str(path) for path in formal_traces)}")
if formal_server_args or mapping_server_args:
server_args = formal_server_args or mapping_server_args
model = server_args.get("model_path") or server_args.get("model")
@@ -527,32 +579,22 @@ def run_triage(args: argparse.Namespace) -> int:
def main(argv: Optional[Sequence[str]] = None) -> int:
argv = list(argv or sys.argv[1:])
top_parser = build_top_level_parser()
triage_parser = build_triage_parser()
if not argv or argv[0] in {"-h", "--help"}:
top_parser.print_help()
triage_parser.print_help()
return 0
command = argv[0]
remainder = argv[1:]
if command == "breakdown":
return breakdown_cli.main(remainder)
if command == "overlap":
return overlap_cli.main(remainder)
if command == "triage":
return run_triage(parse_triage_args(remainder))
if command == "perfetto-fix":
args = parse_perfetto_fix_args(remainder)
output_path = write_perfetto_compatible_trace(
input_path=Path(args.input),
output_path=Path(args.output).resolve() if args.output else None,
if argv[0] == "triage":
argv = argv[1:]
elif not argv[0].startswith("-"):
triage_parser.error(
"This skill now exposes only the compact triage workflow. "
"Use single-trace mode (--input/--url) or mapping+formal two-trace mode."
)
print(f"Perfetto-friendly trace written to: {output_path}")
return 0
return 2
top_parser.error(f"Unknown command: {command}")
return 2
return run_triage(parse_triage_args(argv))
if __name__ == "__main__":
.claude/skills/sglang-torch-profiler-analysis/scripts/profile_common.py
-45
View File
@@ -331,48 +331,3 @@ def select_heaviest_pid(
if preferred:
return max(preferred, key=lambda pid: durations[pid])
return max(durations, key=lambda pid: durations[pid])
def write_perfetto_compatible_trace(
input_path: Path, output_path: Optional[Path] = None
) -> Path:
resolved_input = input_path.resolve()
if output_path is None:
output_name = f"perfetto-compatible-{resolved_input.name}"
output_path = resolved_input.with_name(output_name)
trace = load_trace_json(resolved_input)
output = {key: value for key, value in trace.items() if key != "traceEvents"}
output["traceEvents"] = _perfetto_fix_events(trace.get("traceEvents", []))
output_path.parent.mkdir(parents=True, exist_ok=True)
if str(output_path).endswith(".gz"):
with gzip.open(output_path, "wt", encoding="utf-8") as handle:
json.dump(output, handle)
else:
with open(output_path, "w", encoding="utf-8") as handle:
json.dump(output, handle)
return output_path
def _perfetto_fix_events(events: Sequence[dict]) -> List[dict]:
fixed_events = [dict(event) for event in events]
last_end_time_of_pid_tid: Dict[Tuple[str, str], float] = defaultdict(lambda: -1.0)
for event in fixed_events:
if event.get("ph") != "X" or not _is_perfetto_overlap_interest_event(event):
continue
pid = str(event.get("pid"))
tid = str(event.get("tid"))
ts = float(event.get("ts", 0.0))
dur = float(event.get("dur", 0.0))
while ts < last_end_time_of_pid_tid[(pid, tid)]:
tid = f"{tid}_hack"
event["tid"] = tid
last_end_time_of_pid_tid[(pid, tid)] = ts + dur
return fixed_events
def _is_perfetto_overlap_interest_event(event: dict) -> bool:
args = event.get("args") or {}
return "registers per thread" in args
.claude/skills/sglang-torch-profiler-analysis/scripts/analyze_sglang_llm_torch_profile.py → .claude/skills/sglang-torch-profiler-analysis/scripts/triage_kernel_helpers.py
+906 -305
View File
File diff suppressed because it is too large Load Diff
.claude/skills/sglang-torch-profiler-analysis/scripts/analyze_sglang_profiler_overlap.py → .claude/skills/sglang-torch-profiler-analysis/scripts/triage_overlap_helpers.py
+4 -216
View File
@@ -1,17 +1,4 @@
"""Analyze SGLang PyTorch profiler traces with a two-stage workflow.
This script correlates two traces:
1. A mapping trace, usually collected with CUDA graph disabled, to keep
`kernel -> cpu_op -> python scope` attribution readable.
2. A formal trace, collected with the real serving optimizations enabled, to
measure actual overlap and critical-path exposure.
It prints:
- mapping/formal trace summaries
- an action table that maps formal overlap findings back to Python code
- a few ASCII timelines from the formal trace
"""
"""Internal overlap helpers for triage-only torch-profiler analysis."""
from __future__ import annotations
@@ -365,9 +352,9 @@ def classify_kernel(name: str) -> str:
def is_kernel_event(event: dict) -> bool:
# The overlap script prefers a slightly broader kernel detector than the
# breakdown script, but it still rejects annotations and Python frames up
# front so the later overlap math only sees real GPU work.
# The overlap helpers prefer a slightly broader kernel detector than the
# kernel-attribution helpers, but still reject annotations and Python
# frames up front so the later overlap math only sees real GPU work.
if not is_complete_duration_event(event):
return False
name = normalize_text(event.get("name", ""))
@@ -1567,202 +1554,3 @@ def resolve_trace_source(
events=events,
pid=pid,
)
def parse_args(argv: Optional[Sequence[str]] = None) -> argparse.Namespace:
parser = argparse.ArgumentParser(
description="Analyze SGLang profiler overlap with mapping-trace correlation."
)
parser.add_argument(
"--mapping-input",
type=str,
default=None,
help="Graph-off mapping trace file or directory.",
)
parser.add_argument(
"--mapping-url",
type=str,
default=None,
help="Running graph-off SGLang server URL for the mapping trace.",
)
parser.add_argument(
"--formal-input",
type=str,
default=None,
help="Formal graph-on trace file or directory.",
)
parser.add_argument(
"--formal-url",
type=str,
default=None,
help="Running graph-on SGLang server URL for the formal trace.",
)
parser.add_argument("--input", type=str, default=None, help=argparse.SUPPRESS)
parser.add_argument("--url", type=str, default=None, help=argparse.SUPPRESS)
parser.add_argument(
"--mapping-output-dir",
type=str,
default=None,
help="Trace output dir when using --mapping-url.",
)
parser.add_argument(
"--formal-output-dir",
type=str,
default=None,
help="Trace output dir when using --formal-url.",
)
parser.add_argument(
"--mapping-profile-prefix",
type=str,
default="mapping-trace",
help="Profile prefix for the mapping trace.",
)
parser.add_argument(
"--formal-profile-prefix",
type=str,
default="formal-trace",
help="Profile prefix for the formal trace.",
)
parser.add_argument(
"--start-step",
type=int,
default=None,
help="Pass through to sglang.profiler when generating traces from URLs.",
)
parser.add_argument(
"--num-steps",
type=int,
default=5,
help="Number of steps to profile when generating traces from URLs.",
)
parser.add_argument(
"--profile-by-stage",
action="store_true",
help="Pass through to sglang.profiler.",
)
parser.add_argument(
"--merge-profiles", action="store_true", help="Pass through to sglang.profiler."
)
parser.add_argument(
"--probe-requests",
type=int,
default=1,
help="When generating traces from a URL, send this many probe requests so the profiler captures real work.",
)
parser.add_argument(
"--probe-max-new-tokens",
type=int,
default=None,
help="Override max_new_tokens for the synthetic probe workload. Defaults to max(64, num_steps * 8).",
)
parser.add_argument(
"--probe-prompt",
type=str,
default=(
"Repeat the word overlap many times with spaces so the server performs several decode steps. "
"Do not add analysis or explanations."
),
help="Prompt used for the synthetic probe request when profiling a live URL.",
)
parser.add_argument(
"--probe-delay",
type=float,
default=0.5,
help="Seconds to wait after starting the profiler before sending the synthetic probe request.",
)
parser.add_argument(
"--pid-substring",
type=str,
default=None,
help="Restrict analysis to PIDs containing this substring.",
)
parser.add_argument(
"--table-limit",
type=int,
default=0,
help="Maximum number of rows in the action table. Use 0 to print all kernels.",
)
parser.add_argument(
"--timeline-count",
type=int,
default=3,
help="Number of ASCII windows to render from the formal trace.",
)
parser.add_argument(
"--timeline-width", type=int, default=96, help="ASCII timeline width."
)
parser.add_argument(
"--window-us",
type=float,
default=None,
help="Fixed timeline window size in microseconds.",
)
parser.add_argument(
"--table-only",
action="store_true",
help="Print trace summaries plus the overlap action table only.",
)
args = parser.parse_args(argv)
if args.formal_input is None and args.input is not None:
args.formal_input = args.input
if args.formal_url is None and args.url is not None:
args.formal_url = args.url
if bool(args.mapping_input) == bool(args.mapping_url):
parser.error("Provide exactly one of --mapping-input or --mapping-url.")
if bool(args.formal_input) == bool(args.formal_url):
parser.error("Provide exactly one of --formal-input or --formal-url.")
return args
def main(argv: Optional[Sequence[str]] = None) -> int:
args = parse_args(argv)
mapping_bundle = resolve_trace_source(
label="mapping",
input_path=args.mapping_input,
url=args.mapping_url,
output_dir=args.mapping_output_dir,
profile_prefix=args.mapping_profile_prefix,
args=args,
)
formal_bundle = resolve_trace_source(
label="formal",
input_path=args.formal_input,
url=args.formal_url,
output_dir=args.formal_output_dir,
profile_prefix=args.formal_profile_prefix,
args=args,
)
formal_bundle.overlap_stats = analyze_overlap(formal_bundle.events)
aggregates = aggregate_events(formal_bundle.events)
source_map = build_kernel_source_map(mapping_bundle)
rows = build_action_rows(
aggregates,
source_map,
formal_bundle.events,
formal_bundle.overlap_stats["total_busy_us"],
table_limit=max(0, args.table_limit),
)
report = build_report(
mapping_bundle=mapping_bundle,
formal_bundle=formal_bundle,
source_map=source_map,
aggregates=aggregates,
rows=rows,
window_us=args.window_us,
timeline_count=max(0, args.timeline_count),
width=max(40, args.timeline_width),
table_only=args.table_only,
)
print(report)
return 0
if __name__ == "__main__":
raise SystemExit(main())