From 7603b226ce675233c4847b36f28eccf72421f0d1 Mon Sep 17 00:00:00 2001
From: Xiaoyu Zhang <35585791+BBuf@users.noreply.github.com>
Date: Thu, 9 Apr 2026 18:23:03 +0800
Subject: [PATCH] Upgrade sglang-torch-profiler-analysis SKILLS (#22440)
---
.../sglang-torch-profiler-analysis/SKILL.md | 206 ++-
.../references/fuse-overlap-catalog.md | 128 +-
.../references/overlap-catalog.md | 77 +-
.../references/trace-workflow.md | 119 --
.../references/validated-workflows.md | 263 ----
.../scripts/analyze_sglang_torch_profile.py | 396 +++---
.../scripts/profile_common.py | 45 -
...ch_profile.py => triage_kernel_helpers.py} | 1211 ++++++++++++-----
...r_overlap.py => triage_overlap_helpers.py} | 220 +--
9 files changed, 1406 insertions(+), 1259 deletions(-)
delete mode 100644 .claude/skills/sglang-torch-profiler-analysis/references/trace-workflow.md
delete mode 100644 .claude/skills/sglang-torch-profiler-analysis/references/validated-workflows.md
rename .claude/skills/sglang-torch-profiler-analysis/scripts/{analyze_sglang_llm_torch_profile.py => triage_kernel_helpers.py} (54%)
rename .claude/skills/sglang-torch-profiler-analysis/scripts/{analyze_sglang_profiler_overlap.py => triage_overlap_helpers.py} (88%)
diff --git a/.claude/skills/sglang-torch-profiler-analysis/SKILL.md b/.claude/skills/sglang-torch-profiler-analysis/SKILL.md
index 1567df1a7..1d7d5d5e9 100644
--- a/.claude/skills/sglang-torch-profiler-analysis/SKILL.md
+++ b/.claude/skills/sglang-torch-profiler-analysis/SKILL.md
@@ -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
diff --git a/.claude/skills/sglang-torch-profiler-analysis/references/fuse-overlap-catalog.md b/.claude/skills/sglang-torch-profiler-analysis/references/fuse-overlap-catalog.md
index 319fc2047..1b2e0c5fe 100644
--- a/.claude/skills/sglang-torch-profiler-analysis/references/fuse-overlap-catalog.md
+++ b/.claude/skills/sglang-torch-profiler-analysis/references/fuse-overlap-catalog.md
@@ -129,14 +129,89 @@ already-known PR family "new".
| PR `#21491` FlashInfer TRTLLM FP8 MoE with fused shared experts | `num_fused_shared_experts`
`trtllm_fp8_block_scale_moe` | `PR #21491`
`python/sglang/srt/layers/moe/fused_moe_triton/fused_moe.py`
`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`
`per_token_quant_fp8` | `PR #22005`
`python/sglang/jit_kernel/csrc/elementwise/fused_add_rmsnorm_per_token_quant.cuh`
`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`
`_gemma_rmsnorm_residual_kernel`
`Gemma4` | `PR #21952`
`python/sglang/srt/layers/gemma4_fused_ops.py`
`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`
`fused_qk_norm_mrope_3d_cache_pts_quant_shuffle`
`rotary_dim` | `PR #20667`
`python/sglang/srt/models/qwen3_5.py`
`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`
`benchmark_combo_kernel`
`q_norm`
`k_norm`
`split_with_sizes` | `PR #21977`
`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`
`fp8_scaled_mm`
`nvjet`
`cudaMemsetAsync` | `PR #22392`
`sgl-kernel/python/sgl_kernel/gemm.py`
`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`
`combine`
`down_gemm` | `PR #21877`
`python/sglang/srt/server_args.py`
`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`
`execute_h2d_async`
`hit-attention`
`merge_state` | `PR #22410`
`python/sglang/srt/layers/attention/nsa_backend.py`
`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`
`gelu_tanh_and_mul`
`gelu_and_mul`
`silu_and_mul_scaled_nvfp4_experts_quantize` | `flashinfer/activation.py`
`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`
`fused_add_rmsnorm`
`fused_add_rmsnorm_quant`
`gemma_rmsnorm`
`gemma_fused_add_rmsnorm`
`fused_rmsnorm_silu`
`rmsnorm_fp4quant`
`add_rmsnorm_fp4quant` | `flashinfer/norm/__init__.py`
`flashinfer/cute_dsl/rmsnorm_fp4quant.py`
`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`
`AllReduceFusionPattern`
`kARResidualRMSNorm`
`kARResidualRMSNormFP8Quant`
`kARResidualRMSNormFP4Quant`
`trtllm_mnnvl_allreduce_fusion` | `flashinfer/comm/allreduce.py`
`flashinfer/comm/trtllm_ar.py`
`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`
`mla_rope_quantize_fp8`
`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`
`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`
`trtllm_bf16_moe`
`trtllm_fp8_per_tensor_scale_moe`
`trtllm_fp8_block_scale_moe`
`trtllm_fp4_block_scale_moe`
`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`
`blockscaled_contiguous_grouped_gemm_finalize_fusion_nvfp4`
`moe_permute`
`moe_unpermute` | `flashinfer/fused_moe/cute_dsl/blockscaled_contiguous_gather_grouped_gemm_swiglu_fusion.py`
`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`
`launch_with_pdl`
`cudaGridDependencySynchronize`
`cudaTriggerProgrammaticLaunchCompletion`
`trigger_completion_at_end=False`
`allreduce_fusion` | `flashinfer/norm/__init__.py`
`flashinfer/activation.py`
`flashinfer/rope.py`
`flashinfer/comm/allreduce.py`
`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`
`main_event`
`memset_event`
`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`
`split_device_green_ctx_by_sm_count`
`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`
`seqlen=0`
`batch_indices < 0` | `PR #2792`
`flashinfer/rope.py`
`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`
`use_prealloc`
`use_cuda_graph` | `PR #2840`
`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`
`cudaTriggerProgrammaticLaunchCompletion`
`inline PTX` | `PR #2720`
`include/flashinfer/comm/trtllm_allreduce_fusion.cuh`
`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`
`non-gated` | `PR #2882`
`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`
`flashinfer_gelu_tanh_and_mul`
`auto_deploy::silu_and_mul`
post-GEMM `silu` + `mul` | `tensorrt_llm/_torch/custom_ops/flashinfer_custom_ops.py`
`tensorrt_llm/_torch/auto_deploy/transform/library/fuse_silu_mul.py`
`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`
`flashinfer_gemma_rmsnorm`
`auto_deploy::flashinfer_rms_norm` | `tensorrt_llm/_torch/custom_ops/flashinfer_custom_ops.py`
`tensorrt_llm/_torch/modules/rms_norm.py`
`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`
`flashinfer_gemma_fused_add_rmsnorm`
`auto_deploy::flashinfer_fused_add_rms_norm_inplace` | `tensorrt_llm/_torch/custom_ops/flashinfer_custom_ops.py`
`tensorrt_llm/_torch/modules/rms_norm.py`
`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`
`flashinfer_rope`
`cos_sin_cache` | `tensorrt_llm/_torch/modules/rotary_embedding.py`
`tensorrt_llm/_torch/auto_deploy/custom_ops/rope/flashinfer_rope.py`
`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`
`BatchPrefillWithPagedKVCacheWrapper`
`BatchDecodeWithPagedKVCacheWrapper`
`auto_deploy::flashinfer_attention_mha_with_cache`
`read_cache_only` | `tensorrt_llm/_torch/attention_backend/flashinfer.py`
`tensorrt_llm/_torch/auto_deploy/custom_ops/attention/flashinfer_attention.py`
`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`
`BatchPrefillWithRaggedKVCacheWrapper`
`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`
`chunked prefill`
`W_kn`
`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`
`BatchMLAPagedAttentionWrapper`
`decode`
`W_kn`
`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`
`trtllm_bf16_moe`
`trtllm_fp8_block_scale_moe`
`trtllm_fp4_block_scale_moe`
`TRTLLM_GEN_FUSED_MOE_USE_FLASHINFER` | `tensorrt_llm/_torch/modules/fused_moe/moe_op_backend.py`
`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`
`selective_state_update`
`flashinfer_ssm` | `tensorrt_llm/_torch/auto_deploy/custom_ops/mamba/flashinfer_backend_mamba.py`
`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`
`record_event_passthrough`
`_aux`
`wait_event` | `tensorrt_llm/_torch/auto_deploy/transform/library/multi_stream_attn.py`
`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`
`begin_aux_stream_passthrough`
`end_aux_stream_passthrough`
`wait_aux_stream_passthrough` | `tensorrt_llm/_torch/auto_deploy/transform/library/multi_stream_moe.py`
`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`
`trtllm_finegrained_fp8_linear`
`record_event_passthrough`
`_aux` | `tensorrt_llm/_torch/auto_deploy/transform/library/multi_stream_gemm.py`
`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`
`residual_add`
`rms_norm`
`fp8 static quant` | `PR #12674`
`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`
`rank 256`
`paged KV-cache`
`gpu append kernel` | `PR #12519`
`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`
`trtllm-gen attention`
`flashinfer`
`kv cache buffer` | `PR #12525`
`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`
`trtllm-gen attention`
`flashinfer` | `PR #12544`
`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`
`flashinfer`
`trtllm_bf16_moe` | `PR #12738`
`PR #12557`
`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`
`mlir_elementwise_fusion`
`piecewise cudagraph`
`caller_stream.synchronize()` | `PR #12847`
`tensorrt_llm/_torch/auto_deploy/transform/library/multi_stream_moe.py`
`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`
`fused_moe_lora_fp8`
`w13_shrink`
`w2_expand` | `vllm/lora/ops/triton_ops/fused_moe_lora_op.py`
`vllm/lora/ops/triton_ops/fused_moe_lora_fp8_op.py`
`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`
`bilinear_pos_embed`
`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`
`shared_experts_stream`
shared expert near router | `vllm/utils/torch_utils.py`
`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`
`all_to_all_single`
`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`
`MiniMax-M2`
`gate kernel` | `PR #38445`
`vllm/model_executor/layers/fused_moe/router/gate_linear.py`
`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`
`QK Norm + RoPE + Cache + Quant` | `PR #38621`
`csrc/fused_qk_norm_rope_cache_quant.cu`
`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`
`MergedColumnParallelLinear`
`weights_proj` | `PR #38684`
`vllm/model_executor/models/deepseek_v2.py`
`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`
`custom_fused_ar_rms`
`RocmAiterAllReduceFusionPass` | `PR #37646`
`vllm/_aiter_ops.py`
`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`
`flashinfer`
`NVFP4`
`rmsnorm + fp4 quant` | `PR #36413`
`vllm/compilation/passes/fusion/rms_quant_fusion.py`
`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`
`MiniMax-M2.5`
`allreduce_rms` | `PR #37045`
`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`
`router_gemm`
`GLM5`
`FI AR RMS fusion` | `PR #39301`
`vllm/model_executor/layers/fused_moe/router/gate_linear.py`
`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`
`flashinfer/activation.py`
`flashinfer/rope.py`
`flashinfer/fused_moe/core.py`
`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`
`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`
`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"
```
diff --git a/.claude/skills/sglang-torch-profiler-analysis/references/overlap-catalog.md b/.claude/skills/sglang-torch-profiler-analysis/references/overlap-catalog.md
index 76d6802f7..6ba1084ab 100644
--- a/.claude/skills/sglang-torch-profiler-analysis/references/overlap-catalog.md
+++ b/.claude/skills/sglang-torch-profiler-analysis/references/overlap-catalog.md
@@ -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`
`combine`
`down_gemm` | `PR #21877`
`python/sglang/srt/server_args.py`
`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`
`execute_h2d_async`
`hit-attention`
`merge_state` | `PR #22410`
`python/sglang/srt/layers/attention/nsa_backend.py`
`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`
`launch_with_pdl`
`cudaGridDependencySynchronize`
`cudaTriggerProgrammaticLaunchCompletion`
`trigger_completion_at_end=False`
`allreduce_fusion` | `flashinfer/norm/__init__.py`
`flashinfer/activation.py`
`flashinfer/rope.py`
`flashinfer/comm/allreduce.py`
`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`
`main_event`
`memset_event`
`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`
`split_device_green_ctx_by_sm_count`
`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`
`use_prealloc`
`use_cuda_graph` | `PR #2840`
`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`
`cudaTriggerProgrammaticLaunchCompletion`
`inline PTX` | `PR #2720`
`include/flashinfer/comm/trtllm_allreduce_fusion.cuh`
`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`
`record_event_passthrough`
`_aux`
`wait_event` | `tensorrt_llm/_torch/auto_deploy/transform/library/multi_stream_attn.py`
`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`
`begin_aux_stream_passthrough`
`end_aux_stream_passthrough`
`wait_aux_stream_passthrough` | `tensorrt_llm/_torch/auto_deploy/transform/library/multi_stream_moe.py`
`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`
`trtllm_finegrained_fp8_linear`
`record_event_passthrough`
`_aux` | `tensorrt_llm/_torch/auto_deploy/transform/library/multi_stream_gemm.py`
`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`
`mlir_elementwise_fusion`
`piecewise cudagraph`
`caller_stream.synchronize()` | `PR #12847`
`tensorrt_llm/_torch/auto_deploy/transform/library/multi_stream_moe.py`
`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`
`shared_experts_stream`
shared expert near router | `vllm/utils/torch_utils.py`
`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`
`all_to_all_single`
`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`
`wk`
`k_norm`
`aux_stream` | `PR #35968`
`vllm/model_executor/models/deepseek_v2.py`
`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`
`router_gemm`
`GLM5`
`FI AR RMS fusion` | `PR #39301`
`vllm/model_executor/layers/fused_moe/router/gate_linear.py`
`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`
`flashinfer/activation.py`
`flashinfer/rope.py`
`flashinfer/fused_moe/core.py`
`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`
`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"
diff --git a/.claude/skills/sglang-torch-profiler-analysis/references/trace-workflow.md b/.claude/skills/sglang-torch-profiler-analysis/references/trace-workflow.md
deleted file mode 100644
index 26740c884..000000000
--- a/.claude/skills/sglang-torch-profiler-analysis/references/trace-workflow.md
+++ /dev/null
@@ -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:
-
-- `-TP-0.trace.json.gz`
-- `-TP-0-DP-0-PP-0-EP-0.trace.json.gz`
-- `merged-.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.
diff --git a/.claude/skills/sglang-torch-profiler-analysis/references/validated-workflows.md b/.claude/skills/sglang-torch-profiler-analysis/references/validated-workflows.md
deleted file mode 100644
index c6219b03c..000000000
--- a/.claude/skills/sglang-torch-profiler-analysis/references/validated-workflows.md
+++ /dev/null
@@ -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.
diff --git a/.claude/skills/sglang-torch-profiler-analysis/scripts/analyze_sglang_torch_profile.py b/.claude/skills/sglang-torch-profiler-analysis/scripts/analyze_sglang_torch_profile.py
index b47693fc8..83b584312 100644
--- a/.claude/skills/sglang-torch-profiler-analysis/scripts/analyze_sglang_torch_profile.py
+++ b/.claude/skills/sglang-torch-profiler-analysis/scripts/analyze_sglang_torch_profile.py
@@ -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__":
diff --git a/.claude/skills/sglang-torch-profiler-analysis/scripts/profile_common.py b/.claude/skills/sglang-torch-profiler-analysis/scripts/profile_common.py
index 7d2ecf7f8..602064d6f 100644
--- a/.claude/skills/sglang-torch-profiler-analysis/scripts/profile_common.py
+++ b/.claude/skills/sglang-torch-profiler-analysis/scripts/profile_common.py
@@ -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
diff --git a/.claude/skills/sglang-torch-profiler-analysis/scripts/analyze_sglang_llm_torch_profile.py b/.claude/skills/sglang-torch-profiler-analysis/scripts/triage_kernel_helpers.py
similarity index 54%
rename from .claude/skills/sglang-torch-profiler-analysis/scripts/analyze_sglang_llm_torch_profile.py
rename to .claude/skills/sglang-torch-profiler-analysis/scripts/triage_kernel_helpers.py
index 40c951f4c..ba43fb2de 100644
--- a/.claude/skills/sglang-torch-profiler-analysis/scripts/analyze_sglang_llm_torch_profile.py
+++ b/.claude/skills/sglang-torch-profiler-analysis/scripts/triage_kernel_helpers.py
@@ -1,11 +1,9 @@
-"""Analyze SGLang LLM torch profiler traces into kernel/category shares."""
+"""Internal kernel attribution helpers for triage-only torch-profiler analysis."""
from __future__ import annotations
-import argparse
import json
import re
-import time
from collections import Counter, defaultdict
from dataclasses import dataclass, field
from pathlib import Path
@@ -14,19 +12,16 @@ from typing import DefaultDict, Dict, Iterable, List, Optional, Sequence, Tuple
from profile_common import (
coerce_optional_int,
contains_any_keyword,
- discover_trace_targets,
extract_trace_events,
has_stream_marker,
is_annotation_event,
is_complete_duration_event,
is_non_kernel_trace_category,
is_trace_metadata_name,
- load_trace_json,
looks_like_python_scope_name,
normalize_repo_relative_path,
normalize_text,
parse_stage,
- run_profiler,
select_heaviest_pid,
)
@@ -197,16 +192,12 @@ NOISE_FRAME_PREFIXES = (
"(",
"python/sglang/srt/distributed/communication_op.py:21 "
- "tensor_model_parallel_fused_allreduce_rmsnorm"
-)
-QWEN3_QK_ROPE_FUSION_PATH = (
- "python/sglang/srt/models/qwen3.py:141 forward_prepare_native"
- "
python/sglang/srt/models/utils.py:230 apply_qk_norm"
- "
python/sglang/jit_kernel/fused_qknorm_rope.py:34 fused_qk_norm_rope_out"
- "
python/sglang/srt/models/qwen3_moe.py:592 apply_qk_norm_rope"
+
+LOW_LEVEL_FRAME_PREFIXES = (
+ "triton/runtime/",
+ "triton/backends/",
+ "torch/_ops.py",
+ "torch/nn/modules/module.py",
)
@@ -295,12 +286,656 @@ class KernelRow:
@dataclass
class FusionOpportunity:
pattern: str
+ status: str
confidence: str
related_us: float
evidence: str
current_locations: str
candidate_path: str
rationale: str
+ covered_row_keys: Tuple[Tuple[str, str, str], ...] = field(
+ default_factory=tuple, repr=False
+ )
+ pattern_span: int = field(default=1, repr=False)
+ has_active_match: bool = field(default=False, repr=False)
+ priority: int = field(default=0, repr=False)
+ subsumes: Tuple[str, ...] = field(default_factory=tuple, repr=False)
+
+
+@dataclass(frozen=True)
+class FusionPatternSpec:
+ pattern: str
+ candidate_path: str
+ active_keywords: Tuple[str, ...] = ()
+ split_groups: Tuple[Tuple[str, ...], ...] = ()
+ rationale_hint: str = ""
+ origin: str = "mainline"
+ model_include: Tuple[str, ...] = ()
+ model_exclude: Tuple[str, ...] = ()
+ min_tp_size: int = 1
+ require_tp: bool = False
+ min_share: float = 0.25
+ likely_share: float = 3.0
+ priority: int = 0
+ subsumes: Tuple[str, ...] = ()
+
+
+FUSION_PATTERN_REGISTRY: Tuple[FusionPatternSpec, ...] = (
+ FusionPatternSpec(
+ pattern="Fused residual add + RMSNorm",
+ candidate_path=(
+ "python/sglang/srt/layers/layernorm.py"
+ "
python/sglang/srt/layers/quantization/modelslim/modelslim.py"
+ ),
+ active_keywords=(
+ "fused_add_rmsnorm",
+ "gemma_fused_add_rmsnorm",
+ "npu_add_rms_norm",
+ "add_rmsnorm_bias",
+ ),
+ rationale_hint=(
+ "Residual add plus RMSNorm already has fused implementations across"
+ " several backends."
+ ),
+ min_share=0.1,
+ likely_share=1.0,
+ ),
+ FusionPatternSpec(
+ pattern="FlashInfer unified allreduce_fusion",
+ candidate_path=(
+ "python/sglang/srt/layers/flashinfer_comm_fusion.py"
+ "
python/sglang/srt/layers/layernorm.py"
+ "
python/sglang/srt/layers/communicator.py"
+ ),
+ active_keywords=(
+ "allreduce_fusion",
+ "fusedaddrmsnormkernel",
+ "flashinfer_comm_fusion.py",
+ ),
+ split_groups=(
+ (
+ "cross_device_reduce",
+ "allreduce",
+ "all_reduce",
+ "custom_all_reduce_ops.py",
+ ),
+ ("rmsnorm", "layernorm", "fused_add_rmsnorm", "layernorm.py"),
+ ),
+ rationale_hint=(
+ "FlashInfer already exposes a TP all-reduce plus residual/RMSNorm"
+ " fusion path."
+ ),
+ require_tp=True,
+ min_tp_size=2,
+ min_share=0.5,
+ likely_share=4.0,
+ ),
+ FusionPatternSpec(
+ pattern="AITER allreduce fusion",
+ candidate_path=(
+ "python/sglang/srt/distributed/communication_op.py"
+ "
python/sglang/srt/layers/communicator.py"
+ "
python/sglang/srt/layers/layernorm.py"
+ ),
+ active_keywords=(
+ "tensor_model_parallel_fused_allreduce_rmsnorm",
+ "apply_aiter_all_reduce_fusion",
+ "custom_fused_ar_rms",
+ ),
+ split_groups=(
+ ("allreduce", "all_reduce", "cross_device_reduce"),
+ ("rmsnorm", "layernorm"),
+ ),
+ rationale_hint=(
+ "ROCm already has an AITER fused all-reduce plus RMSNorm family."
+ ),
+ require_tp=True,
+ min_tp_size=2,
+ min_share=0.5,
+ likely_share=4.0,
+ ),
+ FusionPatternSpec(
+ pattern="Fused activation-and-mul (SwiGLU / GeGLU)",
+ candidate_path="python/sglang/srt/layers/activation.py",
+ active_keywords=("silu_and_mul", "gelu_and_mul", "npu_swiglu"),
+ rationale_hint=(
+ "Packed MLP activation and multiply already has dedicated fused ops."
+ ),
+ min_share=0.1,
+ likely_share=1.0,
+ ),
+ FusionPatternSpec(
+ pattern="In-place QK RMSNorm",
+ candidate_path=(
+ "python/sglang/srt/models/utils.py" "
python/sglang/jit_kernel/norm.py"
+ ),
+ active_keywords=("fused_inplace_qknorm", "minimaxm2rmsnormtp"),
+ split_groups=(("apply_qk_norm", "q_norm", "k_norm", "qknorm"),),
+ rationale_hint=(
+ "Q/K normalization already has in-place or model-specific fused"
+ " implementations."
+ ),
+ min_share=0.3,
+ likely_share=2.0,
+ ),
+ FusionPatternSpec(
+ pattern="Fused QK RMSNorm + RoPE",
+ candidate_path=(
+ "python/sglang/jit_kernel/fused_qknorm_rope.py"
+ "
python/sglang/srt/models/qwen3_moe.py"
+ ),
+ active_keywords=("fused_qknorm_rope", "fused_qk_norm_rope"),
+ split_groups=(
+ ("apply_qk_norm", "q_norm", "k_norm", "qknorm"),
+ ("apply_rope", "rotary", "rope", "mrope"),
+ ),
+ rationale_hint=(
+ "SGLang already ships a fused QK-norm plus RoPE kernel family."
+ ),
+ min_share=0.3,
+ likely_share=2.0,
+ priority=30,
+ ),
+ FusionPatternSpec(
+ pattern="Fused QK RoPE reshape + KV cache write",
+ candidate_path="python/sglang/srt/layers/attention/utils.py",
+ active_keywords=("fused_qk_rope_reshape_and_cache",),
+ split_groups=(
+ ("rotary", "rope", "mrope"),
+ ("reshape", "set_kv", "kv_cache", "cache write", "paged kv"),
+ ),
+ rationale_hint=(
+ "Attention prep already has a fused RoPE plus reshape plus cache"
+ " write path."
+ ),
+ min_share=0.4,
+ likely_share=2.0,
+ priority=40,
+ subsumes=("Fused RoPE + KV cache store",),
+ ),
+ FusionPatternSpec(
+ pattern="Fused RoPE + KV cache store",
+ candidate_path=(
+ "python/sglang/jit_kernel/rope.py" "
python/sglang/srt/models/utils.py"
+ ),
+ active_keywords=("fused_set_kv_buffer",),
+ split_groups=(
+ ("rotary", "rope", "mrope"),
+ ("set_kv_buffer", "kv cache write", "paged kv", "cache write"),
+ ),
+ rationale_hint=(
+ "RoPE application and KV cache storage already have fused fast"
+ " paths in several models."
+ ),
+ min_share=0.3,
+ likely_share=1.5,
+ priority=20,
+ ),
+ FusionPatternSpec(
+ pattern="Fused decode metadata setup",
+ candidate_path=("python/sglang/srt/layers/attention/flashattention_backend.py"),
+ active_keywords=(
+ "normal_decode_set_metadata",
+ "cache_seqlens_int32",
+ "cu_seqlens_k",
+ "swa_page_table",
+ ),
+ rationale_hint=(
+ "Decode metadata setup already has a fused Triton preparation path."
+ ),
+ min_share=0.05,
+ likely_share=0.5,
+ ),
+ FusionPatternSpec(
+ pattern="NSA fused metadata copy for graph replay",
+ candidate_path="python/sglang/jit_kernel/fused_metadata_copy.py",
+ active_keywords=(
+ "fused_metadata_copy",
+ "fused_metadata_copy_multi",
+ "fused_nsa_cache_seqlens",
+ "fused_flashmla_metadata",
+ ),
+ rationale_hint=(
+ "NSA replay metadata copies are already fused into one-kernel" " families."
+ ),
+ min_share=0.02,
+ likely_share=0.2,
+ ),
+ FusionPatternSpec(
+ pattern="DeepSeek MLA fused projection + norm + RoPE",
+ candidate_path=(
+ "python/sglang/srt/models/deepseek_common/attention_forward_methods/"
+ "forward_mla_fused_rope_cpu.py"
+ "
python/sglang/srt/models/deepseek_common/attention_forward_methods/"
+ "forward_mla_fused_rope_rocm.py"
+ ),
+ active_keywords=(
+ "qkv_proj_with_rope_fused_weight",
+ "fused_qkv_a_proj_with_mqa",
+ "forward_absorb_fused_mla_rope",
+ ),
+ split_groups=(
+ ("mla", "qkv_a_proj", "q_a_proj"),
+ ("qknorm", "rmsnorm", "apply_qk_norm"),
+ ("rope", "rotary"),
+ ),
+ rationale_hint=(
+ "DeepSeek MLA has backend-specific fused projection, norm, and"
+ " RoPE prep paths."
+ ),
+ model_include=("deepseek", "glm"),
+ min_share=0.4,
+ likely_share=2.0,
+ priority=80,
+ subsumes=("Fused QK RMSNorm + RoPE",),
+ ),
+ FusionPatternSpec(
+ pattern="Fused QK RoPE concat + MLA cache write",
+ candidate_path=(
+ "python/sglang/srt/layers/rocm_linear_utils.py"
+ "
python/sglang/srt/models/deepseek_common/attention_forward_methods/"
+ "forward_mla.py"
+ ),
+ active_keywords=("fused_qk_rope_cat_and_cache_mla", "set_mla_kv_buffer"),
+ split_groups=(
+ ("mla", "rope", "rotary"),
+ ("cache", "kv_buffer", "concat"),
+ ),
+ rationale_hint=(
+ "MLA RoPE packing and cache write already have fused backend paths."
+ ),
+ model_include=("deepseek", "glm"),
+ min_share=0.3,
+ likely_share=1.5,
+ priority=85,
+ subsumes=("Fused RoPE + KV cache store",),
+ ),
+ FusionPatternSpec(
+ pattern="Qwen3 decode fused QK norm + 3D mRoPE + KV cache write",
+ candidate_path="python/sglang/srt/models/qwen3.py",
+ active_keywords=("fused_qk_norm_mrope_3d_cache_pts_quant_shuffle",),
+ split_groups=(
+ ("apply_qk_norm", "q_norm", "k_norm", "qknorm"),
+ ("mrope", "3d rope", "rotary"),
+ ("cache", "kv_buffer", "paged kv", "cache write"),
+ ),
+ rationale_hint=(
+ "Qwen3-style decode already has a fused QK-norm plus 3D mRoPE plus"
+ " cache-write path."
+ ),
+ model_include=("qwen3",),
+ model_exclude=("qwen3.5", "qwen3_5"),
+ min_share=0.4,
+ likely_share=2.0,
+ priority=90,
+ subsumes=(
+ "Fused QK RMSNorm + RoPE",
+ "Fused QK RoPE reshape + KV cache write",
+ "Fused RoPE + KV cache store",
+ ),
+ ),
+ FusionPatternSpec(
+ pattern="Fused MoE router / top-k / softcapping",
+ candidate_path="python/sglang/srt/layers/moe/router.py",
+ active_keywords=("fusedmoerouter", "fused_moe_router"),
+ split_groups=(
+ ("router", "gate", "router logits"),
+ ("topk", "softmax", "softcap", "tanh"),
+ ),
+ rationale_hint=(
+ "MoE routing already has fused router, softcap, and top-k kernels."
+ ),
+ min_share=0.3,
+ likely_share=1.5,
+ priority=30,
+ ),
+ FusionPatternSpec(
+ pattern="Fused MoE grouped-topk / gate kernels",
+ candidate_path="python/sglang/srt/layers/moe/topk.py",
+ active_keywords=(
+ "fused_topk_deepseek",
+ "moe_fused_gate",
+ "aiter_fused_topk",
+ "kimi_k2_moe_fused_gate",
+ ),
+ split_groups=(
+ ("grouped_topk", "topk", "biased_grouped_topk"),
+ ("gate", "router", "renorm", "routed scaling"),
+ ),
+ rationale_hint=(
+ "Grouped-topk, bias handling, and routed scaling already have fused"
+ " gate kernels."
+ ),
+ min_share=0.3,
+ likely_share=1.5,
+ priority=50,
+ subsumes=("Fused MoE router / top-k / softcapping",),
+ ),
+ FusionPatternSpec(
+ pattern="Fused MoE sum + all-reduce",
+ candidate_path=("python/sglang/srt/layers/moe/fused_moe_triton/fused_moe.py"),
+ active_keywords=("fuse_sum_all_reduce", "enable_fused_moe_sum_all_reduce"),
+ split_groups=(
+ ("fused_moe", "expert", "moe"),
+ ("allreduce", "all_reduce", "cross_device_reduce"),
+ ),
+ rationale_hint=(
+ "The second MoE GEMM already has a fused sum-plus-all-reduce path."
+ ),
+ require_tp=True,
+ min_tp_size=2,
+ min_share=0.4,
+ likely_share=2.0,
+ ),
+ FusionPatternSpec(
+ pattern="Fused MoE activation + quant / re-quant",
+ candidate_path=(
+ "python/sglang/srt/layers/moe/ep_moe/kernels.py"
+ "
python/sglang/jit_kernel/nvfp4.py"
+ "
python/sglang/srt/layers/moe/cutlass_w4a8_moe.py"
+ ),
+ active_keywords=(
+ "silu_and_mul_scaled_fp4",
+ "npu_dequant_swiglu_quant",
+ "swiglu_quant",
+ ),
+ split_groups=(
+ ("silu", "gelu", "act_and_mul"),
+ ("quant", "fp8", "mxfp", "nvfp4", "dequant"),
+ ),
+ rationale_hint=(
+ "Quantized MoE backends already fuse activation with re-quantization."
+ ),
+ min_share=0.3,
+ likely_share=1.5,
+ ),
+ FusionPatternSpec(
+ pattern="DeepSeek comm-prep fused RMSNorm + quant / flatten-quant",
+ candidate_path=(
+ "python/sglang/srt/layers/communicator.py"
+ "
python/sglang/srt/models/deepseek_common/attention_forward_methods/"
+ "forward_mla.py"
+ "
python/sglang/srt/models/deepseek_common/attention_forward_methods/"
+ "forward_mha.py"
+ ),
+ active_keywords=(
+ "fused_rms_fp8_group_quant",
+ "fused_rms_mxfp4_quant",
+ "fused_flatten_fp8_group_quant",
+ "fused_flatten_mxfp4_quant",
+ ),
+ split_groups=(
+ ("rmsnorm", "layernorm", "flatten"),
+ ("fp8", "mxfp4", "quant"),
+ ),
+ rationale_hint=(
+ "DeepSeek comm preparation already fuses norm or flatten work with"
+ " quantization."
+ ),
+ model_include=("deepseek", "glm"),
+ min_share=0.3,
+ likely_share=1.5,
+ ),
+ FusionPatternSpec(
+ pattern="NSA fused top-k transform / page-table build",
+ candidate_path="python/sglang/srt/layers/attention/nsa_backend.py",
+ active_keywords=(
+ "fast_topk_transform_fused",
+ "fast_topk_transform_ragged_fused",
+ ),
+ rationale_hint=(
+ "NSA top-k metadata preparation already has fused transform kernels."
+ ),
+ min_share=0.05,
+ likely_share=0.3,
+ ),
+ FusionPatternSpec(
+ pattern="NSA fused quantize + indexed K-cache store",
+ candidate_path=(
+ "python/sglang/jit_kernel/fused_store_index_cache.py"
+ "
python/sglang/srt/layers/attention/nsa/nsa_indexer.py"
+ ),
+ active_keywords=("fused_store_index_k_cache",),
+ split_groups=(
+ ("act_quant", "quant", "scale_buffer"),
+ ("index_k", "cache", "store"),
+ ),
+ rationale_hint=(
+ "NSA already has a fused quantize-and-indexed-store kernel family."
+ ),
+ min_share=0.2,
+ likely_share=1.0,
+ ),
+ FusionPatternSpec(
+ pattern="Fused sampling temperature + softmax",
+ candidate_path=(
+ "python/sglang/srt/layers/fused_sampling.py"
+ "
python/sglang/srt/layers/sampler.py"
+ ),
+ active_keywords=("fused_temperature_softmax",),
+ split_groups=(
+ ("temperature", "temp_scale"),
+ ("softmax", "sampling"),
+ ),
+ rationale_hint=(
+ "Decode-time sampling already has fused temperature and softmax" " kernels."
+ ),
+ min_share=0.05,
+ likely_share=0.5,
+ ),
+ FusionPatternSpec(
+ pattern="Fused logit softcap",
+ candidate_path=(
+ "python/sglang/srt/layers/elementwise.py"
+ "
python/sglang/srt/layers/logits_processor.py"
+ ),
+ active_keywords=("fused_softcap", "final_logit_softcapping"),
+ rationale_hint=(
+ "Logit softcap math already has dedicated fused elementwise kernels."
+ ),
+ min_share=0.02,
+ likely_share=0.2,
+ ),
+ FusionPatternSpec(
+ pattern="PR #20667 Qwen3.5 fused QK norm + RoPE + KV cache write",
+ candidate_path=(
+ "PR #20667"
+ "
python/sglang/srt/models/qwen3_5.py"
+ "
python/sglang/srt/models/utils.py"
+ ),
+ active_keywords=(
+ "fused_qk_norm_rope_cache_pts_quant_shuffle",
+ "fused_qk_norm_mrope_3d_cache_pts_quant_shuffle",
+ ),
+ split_groups=(
+ ("apply_qk_norm", "qknorm", "q_norm", "k_norm"),
+ ("rotary", "rope", "mrope"),
+ ("cache", "kv_buffer", "cache write"),
+ ),
+ rationale_hint=(
+ "An open SGLang ROCm PR already wires a fused QK-norm plus RoPE"
+ " plus KV-cache family for Qwen3.5."
+ ),
+ origin="inflight",
+ model_include=("qwen3.5", "qwen3_5"),
+ min_share=0.4,
+ likely_share=2.0,
+ priority=100,
+ subsumes=(
+ "Fused QK RMSNorm + RoPE",
+ "Fused QK RoPE reshape + KV cache write",
+ "Fused RoPE + KV cache store",
+ ),
+ ),
+ FusionPatternSpec(
+ pattern="PR #22392 CUTLASS FP8 scaled MM replacing nvjet",
+ candidate_path=(
+ "PR #22392"
+ "
sgl-kernel/python/sgl_kernel/gemm.py"
+ "
python/sglang/srt/layers/quantization/fp8_utils.py"
+ ),
+ active_keywords=("cutlass_scaled_mm", "fp8_scaled_mm"),
+ split_groups=(
+ ("nvjet", "_scaled_mm"),
+ ("memset", "memcpy128"),
+ ),
+ rationale_hint=(
+ "An open SGLang PR already replaces nvjet FP8 GEMM with CUTLASS to"
+ " remove memset bubbles and extra copies."
+ ),
+ origin="inflight",
+ min_share=0.2,
+ likely_share=1.0,
+ priority=90,
+ ),
+ FusionPatternSpec(
+ pattern="vLLM-origin Attention + Quantization",
+ candidate_path=(
+ "vllm/compilation/passes/fusion/attn_quant_fusion.py"
+ "
vllm/docs/design/fusions.md"
+ ),
+ active_keywords=(
+ "merge_attn_states",
+ "attn_quant_fusion",
+ "output_group_scale",
+ ),
+ split_groups=(
+ ("attention", "flash_attn", "flashattention", "mla"),
+ ("quant", "fp8", "nvfp4", "group_scale"),
+ ),
+ rationale_hint=(
+ "vLLM already treats attention-epilogue quantization as a reusable"
+ " fused family."
+ ),
+ origin="upstream",
+ min_share=0.3,
+ likely_share=1.5,
+ ),
+ FusionPatternSpec(
+ pattern="vLLM-origin RMSNorm + Quantization",
+ candidate_path=(
+ "vllm/compilation/passes/fusion/rms_quant_fusion.py"
+ "
vllm/docs/design/fusions.md"
+ ),
+ active_keywords=(
+ "fused_add_rms_norm_static_fp8_quant",
+ "rms_quant_fusion",
+ "norm_quant",
+ ),
+ split_groups=(
+ ("rmsnorm", "layernorm", "fused_add_rms_norm"),
+ ("quant", "fp8", "fp4", "per-group"),
+ ),
+ rationale_hint=(
+ "vLLM already has a compile-time norm-plus-quant fusion family."
+ ),
+ origin="upstream",
+ min_share=0.3,
+ likely_share=1.5,
+ ),
+ FusionPatternSpec(
+ pattern="vLLM-origin SiLU+Mul + Quantization",
+ candidate_path=(
+ "vllm/compilation/passes/fusion/act_quant_fusion.py"
+ "
vllm/docs/design/fusions.md"
+ ),
+ active_keywords=(
+ "silu_mul_quant_fp4",
+ "fused_silu_mul_block_quant",
+ "act_quant_fusion",
+ ),
+ split_groups=(
+ ("silu", "gelu", "act_and_mul"),
+ ("quant", "fp8", "fp4", "block_quant"),
+ ),
+ rationale_hint=(
+ "vLLM already treats activation-plus-quant as a reusable fusion" " family."
+ ),
+ origin="upstream",
+ min_share=0.3,
+ likely_share=1.5,
+ ),
+ FusionPatternSpec(
+ pattern="vLLM-origin DSV3 router GEMM",
+ candidate_path=(
+ "vllm/model_executor/layers/fused_moe/router/gate_linear.py"
+ "
vllm/csrc/moe/dsv3_router_gemm_entry.cu"
+ ),
+ active_keywords=("dsv3_router_gemm", "fp32_router_gemm"),
+ split_groups=(
+ ("router", "gate", "router logits"),
+ ("gemm", "matmul", "cublas", "cutlass"),
+ ),
+ rationale_hint=(
+ "vLLM already has a specialized DeepSeek router GEMM family for"
+ " small decode batches."
+ ),
+ origin="upstream",
+ min_share=0.3,
+ likely_share=1.5,
+ ),
+ FusionPatternSpec(
+ pattern="vLLM-origin DeepSeek min-latency fused QKV-A projection",
+ candidate_path=(
+ "vllm/model_executor/models/deepseek_v2.py"
+ "
vllm/csrc/dsv3_fused_a_gemm.cu"
+ ),
+ active_keywords=("dsv3_fused_a_gemm", "fused_qkv_a_proj"),
+ split_groups=(
+ ("q_a_proj", "kv_a_proj", "weights_proj"),
+ ("gemm", "matmul", "cutlass", "cublas"),
+ ),
+ rationale_hint=(
+ "vLLM already has a fused DeepSeek QKV-A projection family for"
+ " decode-latency reduction."
+ ),
+ origin="upstream",
+ model_include=("deepseek", "glm"),
+ min_share=0.3,
+ likely_share=1.5,
+ ),
+ FusionPatternSpec(
+ pattern="PR #38621 fused QK norm + RoPE + cache + quant",
+ candidate_path=(
+ "PR #38621"
+ "
vllm/csrc/fused_qk_norm_rope_cache_quant.cu"
+ "
vllm/compilation/passes/fusion/qk_norm_rope_cache_quant_fusion.py"
+ ),
+ active_keywords=("fused_qk_norm_rope_cache_quant",),
+ split_groups=(
+ ("qknorm", "q_norm", "k_norm"),
+ ("rope", "rotary", "mrope"),
+ ("cache", "kv_buffer", "cache write"),
+ ("quant", "fp8", "nvfp4"),
+ ),
+ rationale_hint=(
+ "An open vLLM PR already treats QK-norm plus RoPE plus cache plus"
+ " quant as a concrete in-flight fusion family."
+ ),
+ origin="inflight",
+ min_share=0.4,
+ likely_share=2.0,
+ priority=100,
+ subsumes=("vLLM-origin Attention + Quantization",),
+ ),
+ FusionPatternSpec(
+ pattern="PR #37045 MiniMax allreduce_rms kernels",
+ candidate_path=("PR #37045" "
vllm/model_executor/models/minimax_m2.py"),
+ active_keywords=("minimax_allreduce_rms", "minimax_allreduce_rmsnorm"),
+ split_groups=(
+ ("q_norm", "k_norm", "rmsnorm", "minimax"),
+ ("allreduce", "all_reduce", "cross_device_reduce"),
+ ),
+ rationale_hint=(
+ "An open vLLM PR already ports TRTLLM MiniMax allreduce-plus-RMSNorm"
+ " kernels."
+ ),
+ origin="inflight",
+ model_include=("minimax",),
+ min_share=0.3,
+ likely_share=1.5,
+ ),
+)
def short_name(name: str, max_len: int = 96) -> str:
@@ -432,20 +1067,33 @@ def choose_best_location(locations: Dict[str, MappingSiteAggregate]) -> str:
def frame_priority(frame_name: str) -> int:
- text = str(frame_name).strip()
- if text.startswith(NOISE_FRAME_PREFIXES):
+ raw_text = str(frame_name).strip()
+ normalized_text = normalize_source_location(raw_text)
+ if raw_text.startswith(NOISE_FRAME_PREFIXES):
return -20
- if text.startswith("/data/") or text.startswith("/Users/"):
- if "/sglang/" in text:
+ if normalized_text.startswith("python/sglang/"):
+ return 300
+ if normalized_text.startswith("sglang/"):
+ return 290
+ if normalized_text.startswith("sgl_kernel/"):
+ return 260
+ if normalized_text.startswith("triton_kernels/"):
+ return 220
+ if normalized_text.startswith(LOW_LEVEL_FRAME_PREFIXES):
+ return 0
+ if raw_text.startswith("/data/") or raw_text.startswith("/Users/"):
+ if "/sglang/" in raw_text:
return 120
return 100
- if ".py(" in text and "/sglang/" in text:
+ if ".py(" in raw_text and "/sglang/" in raw_text:
return 110
- if ".py(" in text and ("site-packages" in text or text.startswith("torch/")):
+ if ".py:" in normalized_text and (
+ "site-packages" in raw_text or normalized_text.startswith("torch/")
+ ):
return 45
- if ".py(" in text:
+ if ".py:" in normalized_text:
return 35
- if text.startswith(" best_score:
best_key = candidate_key
best_score = score
+ if best_key:
+ return kernels.get(best_key)
+
+ # Long auto-generated kernels such as CUTLASS / FlashAttention templates can
+ # differ in the middle of the symbol while still sharing the same high-level
+ # family. Fall back to a conservative common-prefix match so we can still
+ # recover the higher-level Python callsite from the mapping trace.
+ lowered_compact = normalize_match_text(kernel_name)
+ if len(lowered_compact) < 96:
+ return None
+
+ def common_prefix_len(left: str, right: str) -> int:
+ count = 0
+ for left_ch, right_ch in zip(left, right):
+ if left_ch != right_ch:
+ break
+ count += 1
+ return count
+
+ best_key = None
+ best_score = -1
+ for candidate_key in kernels:
+ candidate_compact = normalize_match_text(candidate_key)
+ if len(candidate_compact) < 96:
+ continue
+ prefix_len = common_prefix_len(lowered_compact, candidate_compact)
+ shorter_len = min(len(lowered_compact), len(candidate_compact))
+ if prefix_len < 64 or prefix_len < int(shorter_len * 0.4):
+ continue
+ score = prefix_len
+ if lowered_compact.startswith(
+ "voidcutlassdevicekernelflash"
+ ) and candidate_compact.startswith("voidcutlassdevicekernelflash"):
+ score += 32
+ if score > best_score:
+ best_key = candidate_key
+ best_score = score
return kernels.get(best_key) if best_key else None
@@ -1036,9 +1721,21 @@ def format_location_for_fusion_display(location: str) -> str:
return f"{match.group('func')} @ {match.group('path')}:{match.group('line')}"
+def normalize_match_text(text: object) -> str:
+ return re.sub(r"[^0-9A-Za-z]+", "", normalize_text(text)).lower()
+
+
def row_matches(row: KernelRow, *needles: str) -> bool:
lowered = " ".join([row.name, row.location, row.cpu_op]).lower()
- return any(needle in lowered for needle in needles)
+ lowered_compact = normalize_match_text(lowered)
+ for needle in needles:
+ needle_lowered = needle.lower()
+ if needle_lowered in lowered:
+ return True
+ needle_compact = normalize_match_text(needle)
+ if needle_compact and needle_compact in lowered_compact:
+ return True
+ return False
def summarize_text(values: Iterable[str], limit: int = 4) -> str:
@@ -1055,11 +1752,19 @@ def summarize_locations(values: Iterable[str], limit: int = 4) -> str:
def summarize_evidence(
- rows: Sequence[KernelRow], total_us: float, limit: int = 3
+ rows: Sequence[KernelRow],
+ total_us: float,
+ limit: int = 3,
+ min_share_pct: float = 1.0,
) -> str:
items = []
- for row in rows[:limit]:
- items.append(f"{row.name} ({pct(row.total_us, total_us):.1f}%)")
+ for row in rows:
+ share = pct(row.total_us, total_us)
+ if share < min_share_pct:
+ continue
+ items.append(f"{row.name} ({share:.1f}%)")
+ if len(items) >= limit:
+ break
return "
".join(items) if items else "-"
@@ -1069,6 +1774,150 @@ def model_path_from_server_args(server_args: Optional[dict]) -> str:
return str(server_args.get("model_path") or server_args.get("model") or "")
+def matching_rows_for_keywords(
+ kernel_rows: Sequence[KernelRow],
+ keywords: Sequence[str],
+) -> List[KernelRow]:
+ if not keywords:
+ return []
+ return [row for row in kernel_rows if row_matches(row, *keywords)]
+
+
+def row_identity(row: KernelRow) -> Tuple[str, str, str]:
+ return (row.name, row.location, row.cpu_op)
+
+
+def merge_kernel_rows(*groups: Sequence[KernelRow]) -> List[KernelRow]:
+ output: List[KernelRow] = []
+ seen = set()
+ for group in groups:
+ for row in group:
+ row_key = row_identity(row)
+ if row_key in seen:
+ continue
+ seen.add(row_key)
+ output.append(row)
+ return output
+
+
+def pattern_model_matches(spec: FusionPatternSpec, model_path: str) -> bool:
+ if spec.model_include and not any(
+ token in model_path for token in spec.model_include
+ ):
+ return False
+ if spec.model_exclude and any(token in model_path for token in spec.model_exclude):
+ return False
+ return True
+
+
+def pattern_status(spec: FusionPatternSpec, has_active_match: bool) -> str:
+ if spec.origin == "mainline":
+ return "active fused path" if has_active_match else "split candidate"
+ if spec.origin == "upstream":
+ return "upstream precedent" if has_active_match else "upstream split precedent"
+ return "in-flight precedent" if has_active_match else "in-flight split precedent"
+
+
+def build_pattern_rationale(
+ spec: FusionPatternSpec,
+ has_active_match: bool,
+ related_us: float,
+ total_us: float,
+) -> str:
+ share = pct(related_us, total_us)
+ if spec.origin == "mainline":
+ if has_active_match:
+ return (
+ f"This trace already hits the `{spec.pattern}` family directly at {share:.1f}% related GPU time. "
+ f"{spec.rationale_hint}"
+ )
+ return (
+ f"Related split kernels occupy {share:.1f}% of cumulative GPU time, and the checked-out SGLang tree "
+ f"already exposes this fusion family. {spec.rationale_hint}"
+ )
+ if spec.origin == "upstream":
+ return (
+ f"This trace matches a reusable upstream vLLM precedent at {share:.1f}% related GPU time. "
+ f"{spec.rationale_hint}"
+ )
+ return (
+ f"This trace matches a PR-backed / in-flight pattern at {share:.1f}% related GPU time. "
+ f"{spec.rationale_hint}"
+ )
+
+
+def pattern_span(spec: FusionPatternSpec) -> int:
+ return max(len(spec.split_groups), 1 if spec.active_keywords else 0)
+
+
+def fusion_priority_key(item: FusionOpportunity) -> Tuple[int, int, int, float]:
+ return (
+ item.priority,
+ item.pattern_span,
+ len(item.covered_row_keys),
+ item.related_us,
+ )
+
+
+def detect_pattern_match(
+ spec: FusionPatternSpec,
+ kernel_rows: Sequence[KernelRow],
+ total_us: float,
+ model_path: str,
+ tp_size: int,
+) -> Optional[FusionOpportunity]:
+ if total_us <= 0:
+ return None
+ if spec.require_tp and tp_size < spec.min_tp_size:
+ return None
+ if not pattern_model_matches(spec, model_path):
+ return None
+
+ active_rows = matching_rows_for_keywords(kernel_rows, spec.active_keywords)
+ split_groups = [
+ matching_rows_for_keywords(kernel_rows, keywords)
+ for keywords in spec.split_groups
+ ]
+ has_active_match = bool(active_rows)
+ has_split_match = bool(split_groups) and all(split_groups)
+ if not has_active_match and not has_split_match:
+ return None
+
+ related_rows = merge_kernel_rows(active_rows, *split_groups)
+ related_us = sum(row.total_us for row in related_rows)
+ if related_us <= 0:
+ return None
+ if not has_active_match and pct(related_us, total_us) < spec.min_share:
+ return None
+
+ return FusionOpportunity(
+ pattern=spec.pattern,
+ status=pattern_status(spec, has_active_match),
+ confidence=(
+ "Likely"
+ if has_active_match or pct(related_us, total_us) >= spec.likely_share
+ else "Conditional"
+ ),
+ related_us=related_us,
+ evidence=summarize_evidence(related_rows, total_us),
+ current_locations=summarize_locations(
+ location for row in related_rows for location in kernel_row_locations(row)
+ ),
+ candidate_path=spec.candidate_path,
+ rationale=build_pattern_rationale(
+ spec=spec,
+ has_active_match=has_active_match,
+ related_us=related_us,
+ total_us=total_us,
+ ),
+ covered_row_keys=tuple(row_identity(row) for row in related_rows),
+ pattern_span=pattern_span(spec),
+ has_active_match=has_active_match,
+ priority=spec.priority,
+ subsumes=spec.subsumes,
+ )
+
+
def detect_fusion_opportunities(
stage: str,
kernel_rows: Sequence[KernelRow],
@@ -1084,106 +1933,31 @@ def detect_fusion_opportunities(
if isinstance(server_args, dict):
tp_size = int(server_args.get("tp_size") or 1)
- comm_rows = [
- row
- for row in kernel_rows
- if row.category == "communication"
- and (
- row_matches(
- row,
- "cross_device_reduce",
- "allreduce",
- "all_reduce",
- "custom_all_reduce_ops.py",
- )
- )
- ]
- comm_us = sum(row.total_us for row in comm_rows)
- if comm_rows and (
- tp_size > 1
- or any(row_matches(row, "custom_all_reduce_ops.py") for row in comm_rows)
- ):
- opportunities.append(
- FusionOpportunity(
- pattern="TP all-reduce + residual/RMSNorm",
- confidence="Likely" if pct(comm_us, total_us) >= 4.0 else "Conditional",
- related_us=comm_us,
- evidence=summarize_evidence(comm_rows, total_us),
- current_locations=summarize_locations(
- location
- for row in comm_rows
- for location in kernel_row_locations(row)
- ),
- candidate_path=ALLREDUCE_FUSION_PATH,
- rationale=(
- f"TP communication already consumes {pct(comm_us, total_us):.1f}% of cumulative GPU kernel "
- "time, and SGLang already exposes a fused allreduce+RMSNorm path for the residual/norm "
- "boundary."
- ),
- )
+ raw_matches: List[FusionOpportunity] = []
+ for spec in FUSION_PATTERN_REGISTRY:
+ opportunity = detect_pattern_match(
+ spec=spec,
+ kernel_rows=kernel_rows,
+ total_us=total_us,
+ model_path=model_path,
+ tp_size=tp_size,
)
+ if opportunity is not None:
+ raw_matches.append(opportunity)
- is_qwen3_dense = (
- "qwen3" in model_path
- and "moe" not in model_path
- and "qwen3-next" not in model_path
- )
- qwen3_qk_rows = [
- row
- for row in kernel_rows
- if row_matches(
- row,
- "apply_qk_norm",
- "fused_inplace_qknorm",
- "qknorm",
- )
- ]
- qwen3_rope_rows = [
- row
- for row in kernel_rows
- if row_matches(
- row,
- "apply_rope",
- "rope.py",
- "rope_inplace",
- "rotary",
- )
- ]
- qwen3_rows: List[KernelRow] = []
- seen_qwen3_keys = set()
- for row in qwen3_qk_rows + qwen3_rope_rows:
- row_key = (row.name, row.location, row.cpu_op)
- if row_key in seen_qwen3_keys:
+ raw_matches.sort(key=fusion_priority_key, reverse=True)
+ consumed_row_keys = set()
+ blocked_patterns = set()
+ for opportunity in raw_matches:
+ if opportunity.pattern in blocked_patterns:
continue
- seen_qwen3_keys.add(row_key)
- qwen3_rows.append(row)
- qwen3_related_us = sum(row.total_us for row in qwen3_rows)
- if (
- is_qwen3_dense
- and qwen3_qk_rows
- and qwen3_rope_rows
- and pct(qwen3_related_us, total_us) >= 1.0
- ):
- opportunities.append(
- FusionOpportunity(
- pattern="Q/K RMSNorm + RoPE before attention",
- confidence="Conditional",
- related_us=qwen3_related_us,
- evidence=summarize_evidence(qwen3_rows, total_us),
- current_locations=summarize_locations(
- location
- for row in qwen3_rows
- for location in kernel_row_locations(row)
- ),
- candidate_path=QWEN3_QK_ROPE_FUSION_PATH,
- rationale=(
- "Dense Qwen3 still prepares Q/K with `apply_qk_norm` and then `rotary_emb` in separate source "
- "steps, while SGLang already ships a fused QK-norm+RoPE kernel path in its JIT/MoE stack."
- ),
- )
- )
-
- opportunities.sort(key=lambda item: item.related_us, reverse=True)
+ if any(
+ row_key in consumed_row_keys for row_key in opportunity.covered_row_keys
+ ):
+ continue
+ opportunities.append(opportunity)
+ consumed_row_keys.update(opportunity.covered_row_keys)
+ blocked_patterns.update(opportunity.subsumes)
return opportunities
@@ -1242,7 +2016,7 @@ def generate_takeaways(
if fusion_opportunities:
top_pattern = fusion_opportunities[0]
takeaways.append(
- f"The strongest source-backed fuse candidate is `{top_pattern.pattern}` with {pct(top_pattern.related_us, total_us):.1f}% related GPU time in this stage."
+ f"The strongest source-backed fuse-pattern match is `{top_pattern.pattern}` ({top_pattern.status}) with {pct(top_pattern.related_us, total_us):.1f}% related GPU time in this stage."
)
return takeaways
@@ -1281,20 +2055,21 @@ def print_fusion_opportunity_table(
opportunities: Sequence[FusionOpportunity],
total_us: float,
) -> None:
- print("\nKernel fuse opportunities (Markdown):")
+ print("\nKernel fuse pattern matches (Markdown):")
print(
- "| Pattern | Confidence | Related GPU time | Share | Evidence kernels | Current kernel Python location | Candidate fused Python path | Rationale |"
+ "| Pattern | Status | Confidence | Related GPU time | Share | Evidence kernels | Current kernel Python location | Reference path | Why it matters |"
)
- print("| --- | --- | ---: | ---: | --- | --- | --- | --- |")
+ print("| --- | --- | --- | ---: | ---: | --- | --- | --- | --- |")
if not opportunities:
print(
- "| No medium-confidence source-backed fusion opportunity matched this trace. | - | - | - | - | - | - | - |"
+ "| No source-backed fuse pattern matched this trace. | - | - | - | - | - | - | - | - |"
)
return
for item in opportunities:
print(
- "| {pattern} | {confidence} | {gpu_time} | {share:.1f}% | {evidence} | {current_locations} | {candidate_path} | {rationale} |".format(
+ "| {pattern} | {status} | {confidence} | {gpu_time} | {share:.1f}% | {evidence} | {current_locations} | {candidate_path} | {rationale} |".format(
pattern=escape_md_cell(item.pattern),
+ status=escape_md_cell(item.status),
confidence=escape_md_cell(item.confidence),
gpu_time=format_ms(item.related_us),
share=pct(item.related_us, total_us),
@@ -1406,177 +2181,3 @@ def print_report(
):
print(f" - {takeaway}")
print()
-
-
-def parse_args(argv: Optional[Sequence[str]] = None) -> argparse.Namespace:
- parser = argparse.ArgumentParser(
- description="Analyze SGLang LLM torch profiler traces into kernel/category shares."
- )
- parser.add_argument("--input", type=str, help="Trace file or profile directory.")
- parser.add_argument("--url", type=str, help="Running SGLang server URL.")
- parser.add_argument(
- "--output-dir", type=str, default=None, help="Output root for live profiling."
- )
- parser.add_argument(
- "--num-steps", type=int, default=5, help="Profiler steps for live mode."
- )
- parser.add_argument(
- "--profile-by-stage", action=argparse.BooleanOptionalAction, default=True
- )
- parser.add_argument(
- "--merge-profiles", action=argparse.BooleanOptionalAction, default=False
- )
- parser.add_argument("--profile-prefix", type=str, default=None)
- parser.add_argument("--probe-requests", type=int, default=1)
- parser.add_argument(
- "--probe-prompt",
- type=str,
- default="Explain what tensor parallelism is in one short paragraph.",
- )
- parser.add_argument("--probe-max-new-tokens", type=int, default=96)
- parser.add_argument("--probe-delay", type=float, default=1.0)
- parser.add_argument(
- "--top-k",
- type=int,
- default=12,
- help="How many top kernels to summarize above the tables.",
- )
- parser.add_argument(
- "--kernel-table-limit",
- type=int,
- default=0,
- help="How many kernels to include in the Markdown kernel table. Use 0 for all kernels.",
- )
- parser.add_argument(
- "--kernel-map",
- type=str,
- default=None,
- help="Existing kernel_map.json from a no-CUDA-graph torch pre-pass.",
- )
- parser.add_argument(
- "--export-kernel-map",
- type=str,
- default=None,
- help="Write kernel-to-Python mapping JSON to this file.",
- )
- parser.add_argument(
- "--all-traces",
- action="store_true",
- help="Analyze every matching trace in the selected run directory.",
- )
- parser.add_argument(
- "--table-only",
- action="store_true",
- help="Print the trace header plus the kernel and fuse-opportunity tables only.",
- )
- return parser.parse_args(argv)
-
-
-def main(argv: Optional[Sequence[str]] = None) -> int:
- args = parse_args(argv)
- if not args.input and not args.url:
- raise SystemExit("Provide either --input or --url.")
-
- input_path = Path(args.input).resolve() if args.input else None
- if args.url:
- profile_dir = run_profiler(
- url=args.url,
- output_dir=args.output_dir,
- num_steps=args.num_steps,
- profile_by_stage=args.profile_by_stage,
- merge_profiles=args.merge_profiles,
- profile_prefix=args.profile_prefix,
- probe_requests=args.probe_requests,
- probe_prompt=args.probe_prompt,
- probe_max_new_tokens=args.probe_max_new_tokens,
- probe_delay=args.probe_delay,
- start_step=None,
- )
- print(f"Generated profile directory: {profile_dir}\n")
- input_path = profile_dir
-
- external_kernel_map = (
- load_kernel_map(Path(args.kernel_map).resolve()) if args.kernel_map else None
- )
- traces, server_args = discover_trace_targets(input_path, all_traces=args.all_traces)
-
- stage_site_stats: DefaultDict[
- str, DefaultDict[str, DefaultDict[str, MappingSiteAggregate]]
- ] = defaultdict(lambda: defaultdict(lambda: defaultdict(MappingSiteAggregate)))
- stage_kernel_categories: DefaultDict[str, Dict[str, str]] = defaultdict(dict)
- global_site_stats: DefaultDict[str, DefaultDict[str, MappingSiteAggregate]] = (
- defaultdict(lambda: defaultdict(MappingSiteAggregate))
- )
- global_kernel_categories: Dict[str, str] = {}
- reports = []
-
- for trace_path in traces:
- trace = load_trace_json(trace_path)
- kernels, cpu_ops, python_frames, launch_events, chosen_pid, window_us = (
- extract_trace_data(trace)
- )
- cpu_ops_by_external_id = build_cpu_op_index(cpu_ops)
- launches_by_correlation = build_launch_index(launch_events)
- local_site_stats = aggregate_kernel_sites(
- kernels,
- cpu_ops_by_external_id,
- python_frames,
- launches_by_correlation=launches_by_correlation,
- )
- stage = parse_stage(trace_path)
- kernel_categories = {
- kernel.canonical_name: kernel.category for kernel in kernels
- }
-
- merge_site_stats(stage_site_stats[stage], local_site_stats)
- merge_site_stats(global_site_stats, local_site_stats)
- stage_kernel_categories[stage].update(kernel_categories)
- global_kernel_categories.update(kernel_categories)
- reports.append(
- (trace_path, kernels, chosen_pid, window_us, stage, kernel_categories)
- )
-
- stage_payloads = {
- stage: build_stage_payload(
- dict(site_stats), stage_kernel_categories.get(stage, {})
- )
- for stage, site_stats in stage_site_stats.items()
- }
- global_payload = build_stage_payload(
- dict(global_site_stats), global_kernel_categories
- )
-
- if args.export_kernel_map:
- export_path = Path(args.export_kernel_map).resolve()
- export_path.parent.mkdir(parents=True, exist_ok=True)
- payload = {
- "version": 2,
- "generated_at": time.strftime("%Y-%m-%dT%H:%M:%S%z"),
- "generated_from": str(input_path),
- "notes": "Kernel-to-Python mapping extracted from torch profiler traces. Prefer generating this from a --disable-cuda-graph --disable-piecewise-cuda-graph pre-pass.",
- "server_args": server_args,
- "stages": stage_payloads,
- "global": global_payload,
- }
- with open(export_path, "w", encoding="utf-8") as handle:
- json.dump(payload, handle, indent=2, ensure_ascii=False)
- print(f"Exported kernel map: {export_path}\n")
-
- for trace_path, kernels, chosen_pid, window_us, stage, _ in reports:
- print_report(
- trace_path=trace_path,
- server_args=server_args,
- kernels=kernels,
- chosen_pid=chosen_pid,
- window_us=window_us,
- local_stage_payload=stage_payloads.get(stage, {"kernels": {}}),
- external_kernel_map=external_kernel_map,
- top_k=args.top_k,
- kernel_table_limit=args.kernel_table_limit,
- table_only=args.table_only,
- )
- return 0
-
-
-if __name__ == "__main__":
- raise SystemExit(main())
diff --git a/.claude/skills/sglang-torch-profiler-analysis/scripts/analyze_sglang_profiler_overlap.py b/.claude/skills/sglang-torch-profiler-analysis/scripts/triage_overlap_helpers.py
similarity index 88%
rename from .claude/skills/sglang-torch-profiler-analysis/scripts/analyze_sglang_profiler_overlap.py
rename to .claude/skills/sglang-torch-profiler-analysis/scripts/triage_overlap_helpers.py
index b2da3d445..98dbd9d3c 100644
--- a/.claude/skills/sglang-torch-profiler-analysis/scripts/analyze_sglang_profiler_overlap.py
+++ b/.claude/skills/sglang-torch-profiler-analysis/scripts/triage_overlap_helpers.py
@@ -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())