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())