[skill] Remove outdated llm-serving-auto-benchmark skill (#29487)

2026-06-27 14:19:11 +08:00
parent 43435a2f8e
commit c36f166364
60 changed files with 563 additions and 6977 deletions
@@ -1,6 +1,6 @@
 ---
 name: llm-torch-profiler-analysis
-description: "Unified LLM torch-profiler triage skill for `sglang`, `vllm`, and `TensorRT-LLM`. Use it to inspect an existing `trace.json(.gz)` or profile directory, or to drive live profiling against a running server and return one three-table report with kernel, overlap-opportunity, and fuse-pattern tables."
+description: "Unified LLM torch-profiler triage skill for `sglang`, `vllm`, `TensorRT-LLM`, and `TokenSpeed`. Use it to inspect an existing `trace.json(.gz)` or profile directory, or to drive live profiling against a running server when supported and return one three-table report with kernel, overlap-opportunity, and fuse-pattern tables."
 ---

 # Unified LLM Torch Profiler Analysis
@@ -12,6 +12,7 @@ Use this skill for `torch.profiler` analysis across:
 - `sglang`
 - `vllm`
 - `TensorRT-LLM`
+- `TokenSpeed`

 There is only one public workflow:

@@ -50,19 +51,27 @@ add one short note after the tables with exactly one of:

 ## Capability Matrix

-| Capability | SGLang | vLLM | TensorRT-LLM |
-| --- | --- | --- | --- |
-| Existing trace triage | yes | yes | yes |
-| Single-trace live capture | yes | yes, if torch profiler is enabled on server | requires profiler control endpoints |
-| Two-trace mapping+formal triage | yes | yes | yes |
-| Stage-separated live workload | yes | yes | yes, with a writable shared trace dir or per-stage host runner |
-| `--profile-by-stage` capture | yes | no | no |
-| `--profile-prefix` control | yes | usually ignored on HTTP profiler route | usually ignored on HTTP profiler route |
+| Capability | SGLang | vLLM | TensorRT-LLM | TokenSpeed |
+| --- | --- | --- | --- | --- |
+| Existing trace triage | yes | yes | yes | yes |
+| Single-trace live capture | yes | yes, if torch profiler is enabled on server | requires profiler control endpoints | yes, if `/start_profile` and `/stop_profile` are exposed |
+| Two-trace mapping+formal triage | yes | yes | yes | yes |
+| Stage-separated live workload | yes | yes | yes, with a writable shared trace dir or per-stage host runner | yes, via workload-separated HTTP capture |
+| `--profile-by-stage` capture | yes | no | no | no |
+| `--profile-prefix` control | yes | usually ignored on HTTP profiler route | usually ignored on HTTP profiler route | yes, mapped to `profile_id` |

 For TensorRT-LLM, live capture only works when the server exposes `/start_profile` and
 `/stop_profile`, and when the deployment already provides a shared trace path plus the
 required env vars.

+For TokenSpeed, this skill supports both existing trace triage and live capture
+against current servers that expose `/start_profile` and `/stop_profile`.
+The live helper sends `output_dir`, `activities`, `with_stack`,
+`record_shapes`, and `profile_id` in the start payload. TokenSpeed also has its
+own native `profile_by_stage` field for manual capture, but the unified helper
+uses workload-separated `prefill/` and `decode/` directories by default so the
+tables stay comparable across frameworks.
+
 ## Real H100 Validation

 The current reference run is the `4x H100` matrix captured on `2026-04-23` on
@@ -98,6 +107,9 @@ Validated matrix:
 Use this run as the main H100 reference.
 The older `2026-04-22` single-card Qwen3 matrix is still useful for bring-up, but it is
 not the default reference anymore.
+TokenSpeed support was added later and is covered by existing-trace triage and
+HTTP profiler-control support, but it is not part of this older H100 validation
+matrix yet.

 Stage-separated workload validation captured on `2026-05-01` on `h100_sglang`:

@@ -145,12 +157,14 @@ H100 notes:
 - SGLang kernel-site reconstruction keeps sampling disabled in the mapping path so the optimized parser does not perturb SGLang table output; equality rechecks matched for `Mixtral-8x7B-Instruct-v0.1`, `Qwen3-32B`, and `nvidia/NVIDIA-Nemotron-3-Nano-30B-A3B-FP8`
 - vLLM live capture requires `--output-dir` to match the server `torch_profiler_dir`; the validated H100 flow uses `--profiler-config {"profiler":"torch","torch_profiler_dir":"..."}` and then drives `/start_profile` and `/stop_profile`
 - TensorRT-LLM validation stays on `--backend pytorch`; the H100 flow writes the trace with `TLLM_TORCH_PROFILE_TRACE` and then analyzes the saved trace
- the 2026-04-22 TensorRT-LLM 1.0.0 `py_executor.py` profiler setup still needed a `with_stack=True` override for table-quality Python locations, and the matrix runner generated that override under `/data/bbuf/validate/unified_llm_profiler_skill/overrides/trtllm`; re-check this on TensorRT-LLM 1.2.1 or any 1.3.x release-candidate image before assuming the override is still required
+- TensorRT-LLM current mainline was rechecked at `0722c5f47d2cae69ac1a237da51e550dd214532c` on 2026-06-26; the latest delta affects KV eviction / block-offset staging rather than profiler trace controls, so the `b9e1945` profiler evidence still applies: PyTorch profiling uses `record_shapes=True` and `with_modules=True`, but not `with_stack=True`; keep the override path for table-quality Python locations unless the target image proves otherwise
+- TokenSpeed trace analysis has first-class registry rows for native TokenSpeed CuTe DSL MLA, MLA KV pack + FP8 quantize, fused top-k/top-p sampling, persistent lm_head GEMM, and NVFP4 GEMM + SwiGLU + quant; live capture still requires an existing torch-profiler trace until the target TokenSpeed image exposes a supported profiler API
 - on this host, keep all trace roots under `/data/...`, not `/home/...`

 ## When To Use It

- inspect a `torch.profiler` trace or profile directory from `sglang`, `vllm`, or `TensorRT-LLM`
+- inspect a `torch.profiler` trace or profile directory from `sglang`, `vllm`,
+  `TensorRT-LLM`, or `TokenSpeed`
 - profile a live serving endpoint and analyze the result
 - summarize which kernel families dominate prefill or decode
 - map kernels back to Python code paths
@@ -202,7 +216,7 @@ Allowed `--profile-workload` values:
 - `decode`: capture only the one-input / long-output workload
 - `legacy`: keep the old `--probe-prompt` / `--probe-max-new-tokens` behavior

-For `sglang-sota-performance`, do not use the defaults if the slow SGLang
+For `sglang-sota-humanize-loop`, do not use the defaults if the slow SGLang
 benchmark scenario has a known input/output distribution.
 Set the profiler lengths from that slow scenario instead: prefill uses the slow
 input length with output `1`, and decode uses input `1` with the slow output
@@ -283,7 +297,7 @@ and the trace path is shared with the current machine.

 Typical env expectations are:

- `TLLM_PROFILE_START_STOP=1`
+- `TLLM_PROFILE_START_STOP=<start>-<stop>` such as `10-20`
 - `TLLM_TORCH_PROFILE_TRACE=/shared/path/trace.json` or `.json.gz`

 Then run:
@@ -318,11 +332,101 @@ The matrix runner does this automatically on H100 before TensorRT-LLM capture st

 This is the validated TensorRT-LLM flow on `h100_sglang`:

-1. launch `trtllm-serve` with `TLLM_TORCH_PROFILE_TRACE=/data/.../trace.json`
+1. launch `trtllm-serve` with `TLLM_PROFILE_START_STOP=<start>-<stop>` and `TLLM_TORCH_PROFILE_TRACE=/data/.../trace.json`
 2. run a few benchmark requests
 3. analyze the emitted trace with `--input /data/.../trace.json`

-### 5. Two-trace triage from existing profile dirs or traces
+### 5. Single-trace live capture or triage from TokenSpeed
+
+For a running TokenSpeed server that exposes the profiler routes, the unified
+helper can drive live capture:
+
+```bash
+python3 scripts/analyze_llm_torch_profile.py \
+  --framework tokenspeed \
+  --url http://127.0.0.1:8000 \
+  --output-dir /data/bbuf/validate/unified_llm_profiler_skill/runs/example/tokenspeed_profile \
+  --num-steps 5 \
+  --warmup-steps 10 \
+  --no-profile-by-stage \
+  --profile-workload both \
+  --profile-prefix ts-triage
+```
+
+The helper sends `POST /start_profile` with:
+
+- `output_dir`: the `--output-dir` path
+- `activities`: `["CPU", "GPU"]`
+- `with_stack`: `true`
+- `record_shapes`: `false`
+- `profile_id`: `--profile-prefix`, with `-prefill` or `-decode` appended during workload-separated capture
+
+It then sends OpenAI-compatible probe requests and calls `POST /stop_profile`.
+TokenSpeed writes files such as `ts-triage-prefill-TP-0.trace.json.gz` under the
+output directory. If the server was launched with multiple TP ranks, expect one
+trace per rank.
+
+Existing TokenSpeed torch-profiler traces can still be analyzed directly:
+
+```bash
+python3 scripts/analyze_llm_torch_profile.py \
+  --framework tokenspeed \
+  --input /path/to/tokenspeed_profile_dir_or_trace.json.gz
+```
+
+TokenSpeed's own manual profiler control surface can also be used:
+
+```bash
+curl -X POST http://127.0.0.1:8000/start_profile \
+  -H 'Content-Type: application/json' \
+  -d '{"output_dir":"/data/bbuf/profiles/tokenspeed","activities":["CPU","GPU"],"with_stack":true,"record_shapes":false,"profile_id":"ts-manual"}'
+
+# send representative workload here
+
+curl -X POST http://127.0.0.1:8000/stop_profile
+```
+
+For server-side automatic stop, pass `num_steps`. For TokenSpeed-native
+EXTEND/DECODE split, pass `profile_by_stage: true`; this produces files with
+stage suffixes such as `-EXTEND` and `-DECODE`.
+
+TokenSpeed's benchmark driver can capture traces too:
+
+```bash
+tokenspeed bench serve \
+  --base-url http://127.0.0.1:8000 \
+  --model <model> \
+  --dataset-name random \
+  --random-input-len 4090 \
+  --random-output-len 1 \
+  --num-prompts 64 \
+  --profile \
+  --profile-num-steps 5 \
+  --extra-body '{"output_dir":"/data/bbuf/profiles/tokenspeed","activities":["CPU","GPU"],"with_stack":true,"profile_id":"ts-bench"}'
+```
+
+If `output_dir` is omitted, TokenSpeed falls back to `TOKENSPEED_PROFILER_DIR`
+and then `/tmp`.
+
+Use [scripts/probe_llm_server.py](scripts/probe_llm_server.py) with
+`--framework tokenspeed` for a small OpenAI-compatible endpoint probe before or
+after trace collection:
+
+```bash
+python3 scripts/probe_llm_server.py \
+  --framework tokenspeed \
+  --url http://127.0.0.1:8000 \
+  --requests 6 \
+  --max-tokens 48
+```
+
+For `sglang-sota-humanize-loop`, keep TokenSpeed profiler evidence aligned to
+the same slow scenario bucket as the benchmark result. Prefer the unified
+workload-separated live capture when possible; if only a mixed agentic trace is
+available, label that limitation in `analysis/root-cause.md` before comparing
+it to SGLang prefill/decode traces.
+
+### 6. Two-trace triage from existing profile dirs or traces

 ```bash
 python3 scripts/analyze_llm_torch_profile.py \
@@ -332,7 +436,7 @@ python3 scripts/analyze_llm_torch_profile.py \

 Use this when you need stronger overlap attribution and kernel-to-source mapping.

-### 6. Two-trace triage from running servers
+### 7. Two-trace triage from running servers

 ```bash
 python3 scripts/analyze_llm_torch_profile.py \
@@ -350,6 +454,10 @@ For `vllm` or `TensorRT-LLM`, use the same shape but pass:
 - `--formal-output-dir ...`
 - `--no-profile-by-stage`

+For TokenSpeed, either use `--mapping-url` and `--formal-url` against servers
+that expose `/start_profile` and `/stop_profile`, or pass two existing trace
+directories with `--mapping-input` and `--formal-input`.
+
 ## `profile_by_stage`

 `--profile-by-stage` is only meaningful on the SGLang live-capture path.
@@ -363,7 +471,8 @@ For `vllm` or `TensorRT-LLM`, use the same shape but pass:
  bottlenecks.
 - 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`.
- For `vllm` and `TensorRT-LLM`, disable it with `--no-profile-by-stage`.
+- For `vllm`, `TensorRT-LLM`, and `TokenSpeed`, disable it with
+  `--no-profile-by-stage`.

 ## How To Choose The Triage Shape

@@ -28,14 +28,18 @@ overlap opportunity as novel.

 The catalog is grouped by reusable optimization family, not by one specific model.

-Refresh note `2026-05-01`: rescanned current `sglang` and vLLM mainline, then
-rechecked recent merged and open optimization PRs through the GitHub CLI/API.
-The vLLM torch.compile pass inventory is now split out in
+Refresh note `2026-06-26`: rechecked official main heads for SGLang
+`8524678889485801e7a4a12d62015be0c68f7a90`, vLLM
+`abc71548ef029132c3316b902207f254a246d593`, TensorRT-LLM
+`0722c5f47d2cae69ac1a237da51e550dd214532c`, and TokenSpeed
+`5aedf69d6b476baa65571011de6ea60fd5a238a8`. The vLLM torch.compile pass
+inventory is split out in
 [`vllm-torch-compile-fusions.md`](vllm-torch-compile-fusions.md). Stable
-current-code families remain folded into the mainline rows below. New
-status-sensitive rows were added for DeepSeek-V4, GLM5 DSA / PDL, NVFP4 MoE,
-torch.compile decode, vLLM DSV4, vLLM ROCm WMMA, and vLLM GPU/CPU sync-removal
-work. Recheck PR state before treating an in-flight row as shipped.
+current-code families remain folded into the mainline rows below. This refresh
+adds first-class TokenSpeed-origin rows for CuTe DSL MLA, MLA KV pack+FP8
+quantize, sampling, lm_head GEMM, and NVFP4 GEMM+SwiGLU+quant, plus the latest
+SGLang LTX2 Ada-value diffusion fusion. Recheck PR state before treating an
+in-flight row as shipped.

 ## 1. LLM / SRT fused-kernel families

@@ -53,13 +57,13 @@ work. Recheck PR state before treating an in-flight row as shipped.
 | Fused QK RoPE reshape + KV cache write | `fused_qk_rope_reshape_and_cache*`<br>RoPE followed by reshape / cache DtoD | `python/sglang/srt/layers/attention/utils.py::fused_qk_rope_reshape_and_cache` | One Triton kernel applies RoPE to Q / K, reshapes cache layout, and writes K / V directly to paged cache | Treat separate RoPE + reshape + cache-write ladders as an existing attention-prep fusion family. |
 | Fused RoPE + KV cache store | `fused_set_kv_buffer`<br>RoPE followed by KV-store, DtoD, or cache-write kernels | `python/sglang/jit_kernel/rope.py`<br>`python/sglang/srt/models/utils.py::enable_fused_set_kv_buffer` | Shared entrypoints can route to fused RoPE + KV-store or model-side `fused_set_kv_buffer` fast paths | Compare against the fused cache-store path before proposing a new KV rewrite. |
 | Fused decode metadata setup | `normal_decode_set_metadata`<br>`cache_seqlens_int32`<br>`cu_seqlens_k`<br>`page_table`<br>`swa_page_table` | `python/sglang/srt/layers/attention/flashattention_backend.py::normal_decode_set_metadata` | Triton decode path fuses seq-len cast/add, prefix-sum, req-to-token gather, page-table divide, and optional SWA metadata build into 1-2 kernels | If decode exposes multiple tiny metadata kernels before attention, first compare against this existing fused metadata-prep path. |
-| DSA fused metadata copy for graph replay | `fused_metadata_copy`<br>`fused_metadata_copy_multi`<br>`fused_dsa_cache_seqlens`<br>`fused_flashmla_metadata` | `python/sglang/jit_kernel/fused_metadata_copy.py` | CUDA graph replay path fuses multiple metadata copies into one kernel or one multi-destination kernel | Treat bursts of tiny metadata-copy kernels around DSA replay as a missed existing replay fusion. |
+| NSA fused metadata copy for graph replay | `fused_metadata_copy`<br>`fused_metadata_copy_multi`<br>`fused_nsa_cache_seqlens`<br>`fused_flashmla_metadata` | `python/sglang/jit_kernel/fused_metadata_copy.py` | CUDA graph replay path fuses multiple metadata copies into one kernel or one multi-destination kernel | Treat bursts of tiny metadata-copy kernels around NSA replay as a missed existing replay fusion. |
 | DeepSeek MLA fused projection + norm + RoPE | `qkv_proj_with_rope_fused_weight`<br>`fused_qkv_a_proj_with_mqa`<br>`forward_absorb_fused_mla_rope*` | `python/sglang/srt/models/deepseek_common/attention_forward_methods/forward_mla_fused_rope_cpu.py`<br>`python/sglang/srt/models/deepseek_common/attention_forward_methods/forward_mla_fused_rope_rocm.py`<br>`python/sglang/srt/models/deepseek_v2.py` | CPU / ROCm paths fuse DeepSeek MLA projection packing with q / k norm, RoPE, and cache-oriented MLA prep | For DeepSeek MLA, split proj / norm / rope prep is usually an existing backend-specific fuse that did not fire. |
 | Fused QK RoPE concat + MLA cache write | `fused_qk_rope_cat_and_cache_mla`<br>`set_mla_kv_buffer` | `python/sglang/srt/layers/rocm_linear_utils.py`<br>`python/sglang/srt/models/deepseek_common/attention_forward_methods/forward_mla.py` | ROCm MLA path can fuse Q / K RoPE packing, concat, and MLA cache write in one backend-specific op | On DeepSeek / MLA traces, separate RoPE-cat-cache steps are not automatically novel. |
 | Qwen3 decode fused QK norm + 3D mRoPE + KV cache write | `fused_qk_norm_mrope_3d_cache_pts_quant_shuffle`<br>`mrope`<br>decode cache write | `python/sglang/srt/models/qwen3.py` | ROCm / AITER decode path fuses QK norm, 3D mRoPE, and paged KV cache write | On Qwen3-style decode, separate norm + mRoPE + cache-store kernels are not a novel opportunity. |
 | NPU fused split-QKV + RMSNorm + RoPE | `split_qkv_rmsnorm_rope` | `python/sglang/srt/models/llama.py`<br>`python/sglang/srt/models/qwen3.py`<br>`python/sglang/srt/models/qwen3_moe.py`<br>`python/sglang/srt/models/glm4_moe.py` | Ascend path fuses QKV split, Q / K RMSNorm, and RoPE in one op | On NPU traces, separate split / norm / rope kernels usually mean the fused path is unavailable or bypassed. |
 | Fused FP8 quantize + paged KV cache write | `trtllm_fp8_kv_kernel`<br>`fp8 kv cache write`<br>`paged KV cache write` | `python/sglang/srt/layers/attention/triton_ops/trtllm_fp8_kv_kernel.py` | TRTLLM MHA path fuses FP8 quantization, scale computation, and paged K / V cache write | If FP8 KV cache traces show standalone quant plus write kernels, first compare against this existing Triton fuse. |
-| Fused MLA KV cache write + FP8 quant | `set_mla_kv_buffer_fp8_quant*`<br>`set_mla_kv_buffer_triton_fp8_quant` | `python/sglang/srt/mem_cache/utils.py`<br>`python/sglang/srt/mem_cache/memory_pool.py` | MLA / DSA KV pool path can quantize K and write directly into KV storage without a separate concat-and-quant chain | Treat standalone quant + KV-buffer write on MLA paths as missing existing fusion first. |
+| Fused MLA KV cache write + FP8 quant | `set_mla_kv_buffer_fp8_quant*`<br>`set_mla_kv_buffer_triton_fp8_quant` | `python/sglang/srt/mem_cache/utils.py`<br>`python/sglang/srt/mem_cache/memory_pool.py` | MLA / NSA KV pool path can quantize K and write directly into KV storage without a separate concat-and-quant chain | Treat standalone quant + KV-buffer write on MLA paths as missing existing fusion first. |
 | Fused MoE router / top-k / softcapping | `FusedMoeRouter`<br>`fused_moe_router*`<br>router GEMM + `topk` + `tanh` | `python/sglang/srt/layers/moe/router.py` | Single fused router kernel covers router matmul, softcapping, and top-k selection | Treat exposed router matmul + softcap + top-k chains as an existing MoE fusion family. |
 | Fused MoE grouped-topk / gate kernels | `fused_topk_deepseek`<br>`moe_fused_gate`<br>`aiter_fused_topk`<br>`kimi_k2_moe_fused_gate` | `python/sglang/srt/layers/moe/topk.py` | CUDA / ROCm / FlashInfer kernels fuse bias, grouped-topk, renorm, and routed scaling into one gate op | Check backend / model eligibility before proposing a novel router-gate fusion. |
 | Qwen-style shared-expert append into routed top-k output | `_append_shared_to_topk_output`<br>`fused_append_shared_experts_with_weights`<br>`num_fused_shared_experts` | `python/sglang/srt/models/qwen2_moe.py`<br>`python/sglang/srt/layers/moe/moe_runner/triton_utils/fused_moe_triton_kernels.py` | Qwen-style MoE paths can append shared-expert ids and sigmoid gate weights to routed top-k output in one Triton kernel so the shared experts execute inside the fused MoE path | Treat routed top-k plus shared-expert pad / concat ladders as an existing MoE-prep fusion family first. |
@@ -67,8 +71,8 @@ work. Recheck PR state before treating an in-flight row as shipped.
 | Fused MoE sum + all-reduce | routed MoE followed by explicit sum-reduce kernels | `python/sglang/srt/layers/moe/fused_moe_triton/fused_moe.py`<br>`python/sglang/srt/layers/moe/fused_moe_triton/fused_moe_triton_kernels.py` | `fuse_sum_all_reduce=True` path in the second MoE GEMM | Before inventing a new MoE reduction fuse, check whether `enable_fused_moe_sum_all_reduce` is simply off or the quant path is incompatible. |
 | Fused MoE activation + quant / re-quant | `silu_and_mul_*quant*`<br>`npu_dequant_swiglu_quant`<br>`swiglu_quant` | `python/sglang/srt/layers/moe/ep_moe/kernels.py`<br>`python/sglang/jit_kernel/nvfp4.py`<br>`python/sglang/srt/layers/moe/cutlass_w4a8_moe.py`<br>`python/sglang/srt/hardware_backend/npu/quantization/fused_moe_method_npu.py` | Quantized MoE backends fuse SwiGLU / SiLU-and-mul with FP8 / FP4 / NPU re-quant before the second expert GEMM | If MoE traces show standalone activation then quant kernels, first check whether the quantized fused path is missing. |
 | DeepSeek comm-prep fused RMSNorm + quant / flatten-quant | `fused_rms_fp8_group_quant`<br>`fused_rms_mxfp4_quant`<br>`fused_flatten_fp8_group_quant`<br>`fused_flatten_mxfp4_quant` | `python/sglang/srt/layers/communicator.py`<br>`python/sglang/srt/models/deepseek_common/attention_forward_methods/forward_mla.py`<br>`python/sglang/srt/models/deepseek_common/attention_forward_methods/forward_mha.py` | DeepSeek MLA / MHA ROCm paths fuse RMSNorm or flatten with FP8 / MXFP4 quantization for comm / attention prep | On DeepSeek quant traces, split norm + quant or flatten + quant is an existing family, not a new idea. |
-| DSA fused top-k transform / page-table build | `fast_topk_transform_fused`<br>`fast_topk_transform_ragged_fused` | `python/sglang/srt/layers/attention/dsa_backend.py` | DSA can fuse top-k selection with paged / ragged index transform instead of separate top-k plus metadata scatter | If DSA top-k metadata work is split, check `SGLANG_DSA_FUSE_TOPK` and backend support first. |
-| DSA fused quantize + indexed K-cache store | `fused_store_index_k_cache`<br>`act_quant`<br>`index_k_with_scale_buffer` | `python/sglang/jit_kernel/fused_store_index_cache.py`<br>`python/sglang/srt/layers/attention/dsa/dsa_indexer.py` | Single JIT kernel quantizes bf16 K to fp8 + scale and writes directly into DSA index cache | Treat split `act_quant` + buffer-store on CUDA as missing an existing fused store path. |
+| NSA fused top-k transform / page-table build | `fast_topk_transform_fused`<br>`fast_topk_transform_ragged_fused` | `python/sglang/srt/layers/attention/nsa_backend.py` | NSA can fuse top-k selection with paged / ragged index transform instead of separate top-k plus metadata scatter | If NSA top-k metadata work is split, check `SGLANG_NSA_FUSE_TOPK` and backend support first. |
+| NSA fused quantize + indexed K-cache store | `fused_store_index_k_cache`<br>`act_quant`<br>`index_k_with_scale_buffer` | `python/sglang/jit_kernel/fused_store_index_cache.py`<br>`python/sglang/srt/layers/attention/nsa/nsa_indexer.py` | Single JIT kernel quantizes bf16 K to fp8 + scale and writes directly into NSA index cache | Treat split `act_quant` + buffer-store on CUDA as missing an existing fused store path. |
 | Fused sampling temperature + softmax | `fused_temperature_softmax*` | `python/sglang/srt/layers/fused_sampling.py`<br>`python/sglang/srt/layers/sampler.py` | Triton single-pass / multi-pass kernels fuse temperature scaling and softmax during decode | Separate temp-divide + softmax at decode batch sizes is often a missed existing fusion. |
 | Fused logit softcap | `fused_softcap`<br>`final_logit_softcapping` | `python/sglang/srt/layers/elementwise.py`<br>`python/sglang/srt/layers/logits_processor.py` | Triton kernels fuse cast-to-float and softcap / tanh math for logits or generic elementwise softcapping | Treat exposed cast + softcap ladders as an existing Triton fuse family. |
 | Linear-attention packed projection reshuffle | `fused_qkvzba_split_reshape_cat*`<br>`qkvz_proj`<br>`ba_proj`<br>`qkvabz_proj`<br>`fused_qkvbfg_a_proj` | `python/sglang/jit_kernel/triton/gdn_fused_proj.py`<br>`python/sglang/srt/models/qwen3_next.py`<br>`python/sglang/srt/models/qwen3_5.py`<br>`python/sglang/srt/models/kimi_linear.py`<br>`python/sglang/srt/models/jet_nemotron.py` | GDN / Kimi / Jet-style linear-attn models pack multiple projections, then fuse split / reshape / cat into one kernel | Treat split reshape / transpose / cat ladders as an existing linear-attention fusion family. |
@@ -90,7 +94,7 @@ work. Recheck PR state before treating an in-flight row as shipped.
 | Llama4 shared branch vs routed branch overlap | shared expert branch plus routed MoE branch as adjacent windows | `python/sglang/srt/models/llama4.py` | shared expert on current stream, router + topk + routed experts on `alt_stream` | Use Llama4 as the first precedent for branch-level overlap in similar sparse models. |
 | ExaoneMoE shared experts vs router experts overlap | shared expert output and router-expert output form a two-branch window | `python/sglang/srt/models/exaone_moe.py::forward_normal_dual_stream` | shared experts on current stream, router + routed experts on `alt_stream`, explicit join before combine | This is an existing dual-stream MoE overlap family. |
 | Grok residual-MoE branch overlap | dense MLP and block-sparse MoE branches in parallel | `python/sglang/srt/models/grok.py::moe_with_rmoe` | dense MLP on current stream, MoE on `alt_stream`, fused dual residual RMSNorm around boundaries | Treat exposed Grok branch overlap as an existing pattern. |
-| DSA dual-stream overlap | Q-proj, K-proj, RoPE, cache-store, quantization in tight two-stream windows | `python/sglang/srt/layers/attention/dsa/dsa_indexer.py` | Q / K projection split, RoPE split, cache-store vs quantization overlap | DSA already contains several dual-stream overlap precedents. |
+| NSA dual-stream overlap | Q-proj, K-proj, RoPE, cache-store, quantization in tight two-stream windows | `python/sglang/srt/layers/attention/nsa/nsa_indexer.py` | Q / K projection split, RoPE split, cache-store vs quantization overlap | NSA already contains several dual-stream overlap precedents. |
 | MoriEP async dispatch / combine comm stream | `MoriEP`<br>`_comm_stream`<br>`dispatch`<br>`combine`<br>`done_event` | `python/sglang/srt/layers/moe/token_dispatcher/moriep.py` | MoriEP can submit dispatch and combine onto a dedicated communication stream and synchronize only through events | Treat MoriEP comm / compute interleave as an existing MoE overlap family. |
 | Heterogeneous-TP staging scatter overlap | `scatter_stream`<br>`_scatter_stream`<br>`staging` | `python/sglang/srt/disaggregation/common/staging_handler.py`<br>`python/sglang/srt/disaggregation/common/staging_buffer.py` | decode-side staging scatter kernels can run on a dedicated stream while forward continues on the main stream | If decode traces show staging scatter kernels adjacent to forward kernels, classify them against this existing overlap family first. |
 | Generic `alt_stream` overlap families | `alt_stream` plus explicit `wait_stream` / `with torch.cuda.stream(...)` | `qwen2_moe.py`<br>`qwen3_moe.py`<br>`glm4_moe.py`<br>`bailing_moe.py`<br>`llada2.py`<br>`grok.py`<br>`olmo2.py`<br>`step3p5.py`<br>`longcat_flash.py`<br>`falcon_h1.py` | model-specific overlap on attention prep, MoE branches, or cache-store | Search these families before designing a new overlap scheme from scratch. |
@@ -114,6 +118,8 @@ work. Recheck PR state before treating an in-flight row as shipped.
 | Fused diffusion QK norm + RoPE | split QK norm and RoPE in diffusion attention blocks | `python/sglang/jit_kernel/diffusion/qknorm_rope.py`<br>`python/sglang/multimodal_gen/runtime/layers/layernorm.py::apply_qk_norm_rope` | `fused_inplace_qknorm_rope(...)`, with fallback to QK norm plus `apply_flashinfer_rope_qk_inplace(...)` | Distinguish between missing fused qknorm + rope and the existing FlashInfer RoPE fallback. |
 | Z-Image fused `norm(x) * tanh(scale) + shift` | `fused_norm_tanh_mul_add`<br>`tanh(gate) * rmsnorm(x)` | `python/sglang/jit_kernel/diffusion/cutedsl/norm_tanh_mul_add_norm_scale.py`<br>`python/sglang/multimodal_gen/runtime/layers/layernorm.py` | CuTeDSL kernel plus runtime helper for Z-Image residual-form modulation | Treat split Z-Image residual-form modulation as a missing existing diffusion fusion, not a novel idea. |
 | Z-Image fused residual modulation + next norm-scale | `fused_norm_tanh_mul_add_norm_scale`<br>`residual + tanh(gate) * rmsnorm(x)`<br>`ffn_norm1(x) * scale_mlp` | `python/sglang/jit_kernel/diffusion/cutedsl/norm_tanh_mul_add_norm_scale.py`<br>`python/sglang/multimodal_gen/runtime/models/dits/zimage.py` | One CuTeDSL kernel fuses the first residual-form modulation and the next normalization / scale stage | If you see this chain split in Z-Image traces, report it as a missing existing mainline fusion family. |
+| LTX2 fused Ada values | `ltx2_ada_values9`<br>`get_ada_values`<br>`scale_shift_table + timestep.reshape` | `python/sglang/jit_kernel/diffusion/triton/ltx2_ada_values.py`<br>`python/sglang/multimodal_gen/runtime/models/dits/ltx_2.py` | PR `#29390` fuses LTX-2.3 Ada value materialization for video/audio streams and reuses the 9 Ada tensors across self-attention, MLP, and prompt-cross-attention blocks | Treat repeated Ada add/reshape/slice ladders in LTX2 traces as a missing shipped SGLang fusion first. |
+| LTX2 residual-gate add | `diffusion_residual_gate_add`<br>`residual_gate_add`<br>`residual + update * gate` | `python/sglang/jit_kernel/diffusion/residual_gate_add.py`<br>`python/sglang/jit_kernel/csrc/diffusion/residual_gate_add.cuh`<br>`python/sglang/multimodal_gen/runtime/models/dits/ltx_2.py` | PR `#29361` fuses LTX2 `residual + update * gate` sites for attention, cross-attention, and feed-forward updates into one CUDA custom op when dtype, shape, device, and contiguity guards pass | Treat split add/mul gate ladders in LTX2 traces as a missing shipped SGLang fusion first. |
 | Nunchaku fused GELU MLP | `_fused_gelu_mlp`<br>`fused_gelu_mlp` | `python/sglang/multimodal_gen/runtime/models/dits/flux.py` | Nunchaku path fuses `fc1 GEMM + GELU + shift + re-quant + fc2.lora_down` before the second GEMM | Treat split GELU-MLP on Nunchaku checkpoints as an existing fused family, not a new discovery. |

 ## 5. Diffusion kernel-overlap and async-communication families
@@ -134,16 +140,16 @@ Stable entries should be folded into the mainline family rows above.
 | Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
 | --- | --- | --- | --- | --- |
 | PR `#21877` fused grouped down-GEMM + combine | `grouped_gemm_nt_masked`<br>`combine`<br>`fused grouped gemm combine` | `PR #21877`<br>`python/sglang/srt/layers/moe/ep_moe/flashinfer_cutedsl_moe.py`<br>`python/sglang/srt/layers/moe/token_dispatcher/deepep.py` | FlashInfer CuTeDSL kernel fuses the second expert GEMM with DeepEP low-latency combine | Treat this as a concrete upstream MoE fuse / overlap family, not a new thought experiment. |
-| PR `#21889` fused BF16 to FP4 quant + paged KV write | `set_mla_kv_buffer_fp4_quant_kernel`<br>`fp4 kv cache` | `PR #21889`<br>`python/sglang/srt/mem_cache/utils.py` | Triton kernel writes FP4 DSA KV pages directly while quantizing BF16 input | If DSA FP4 KV paths are split into quant plus store, classify them as an in-flight upstream fuse family. |
-| PR `#21889` fused FP4 paged dequant to FP8 + page-table remap | `_dequant_fp4_to_fp8_paged_kernel`<br>`WRITE_PT`<br>`dequant_fp4_paged_decode` | `PR #21889`<br>`python/sglang/srt/layers/attention/dsa/dequant_fp4_to_fp8.py` | Triton kernel reads FP4 pages, writes FP8 directly, and can fuse decode-side page-table remap | Treat this as an upstream in-flight decode-prep fusion family. |
+| PR `#21889` fused BF16 to FP4 quant + paged KV write | `set_mla_kv_buffer_fp4_quant_kernel`<br>`fp4 kv cache` | `PR #21889`<br>`python/sglang/srt/mem_cache/utils.py` | Triton kernel writes FP4 NSA KV pages directly while quantizing BF16 input | If NSA FP4 KV paths are split into quant plus store, classify them as an in-flight upstream fuse family. |
+| PR `#21889` fused FP4 paged dequant to FP8 + page-table remap | `_dequant_fp4_to_fp8_paged_kernel`<br>`WRITE_PT`<br>`dequant_fp4_paged_decode` | `PR #21889`<br>`python/sglang/srt/layers/attention/nsa/dequant_fp4_to_fp8.py` | Triton kernel reads FP4 pages, writes FP8 directly, and can fuse decode-side page-table remap | Treat this as an upstream in-flight decode-prep fusion family. |
 | PR `#21491` FlashInfer TRTLLM FP8 MoE with fused shared experts | `num_fused_shared_experts`<br>`trtllm_fp8_block_scale_moe` | `PR #21491`<br>`python/sglang/srt/layers/moe/fused_moe_triton/fused_moe.py`<br>`python/sglang/srt/models/deepseek_v2.py` | FlashInfer TRTLLM FP8 MoE path can fuse shared experts inside the routed MoE kernel | On FP8 TRTLLM MoE discussions, treat fused shared experts as an upstream pattern that already has a concrete PR. |
 | PR `#22005` fused add + RMSNorm + per-token FP8 quant | `fused_add_rmsnorm_per_token_quant`<br>`per_token_quant_fp8` | `PR #22005`<br>`python/sglang/jit_kernel/csrc/elementwise/fused_add_rmsnorm_per_token_quant.cuh`<br>`python/sglang/jit_kernel/fused_add_rmsnorm_per_token_quant.py` | CUDA JIT kernel keeps normed values in registers and emits BF16 + FP8 outputs plus per-token scales | If FP8 online-quant traces show add+norm followed by per-token quant, treat this as an in-flight upstream CUDA fuse family. |
 | PR `#20667` Qwen3.5 fused QK norm + RoPE + KV cache write | `fused_qk_norm_rope_cache_pts_quant_shuffle`<br>`fused_qk_norm_mrope_3d_cache_pts_quant_shuffle`<br>`rotary_dim` | `PR #20667`<br>`python/sglang/srt/models/qwen3_5.py`<br>`python/sglang/srt/models/utils.py` | ROCm / AITER path fuses Q / K RMSNorm, partial or 3D RoPE, and direct KV cache write for Qwen3.5 attention | Treat split QK-norm + RoPE + cache-store on Qwen3.5 as a concrete in-flight upstream family, not a novel idea. |
 | PR `#22392` CUTLASS FP8 GEMM replacing nvjet | `cutlass_scaled_mm`<br>`fp8_scaled_mm`<br>`nvjet`<br>`cudaMemsetAsync` | `PR #22392`<br>`sgl-kernel/python/sgl_kernel/gemm.py`<br>`python/sglang/srt/layers/quantization/fp8_utils.py` | Runtime replacement swaps nvjet FP8 GEMMs for CUTLASS kernels, removing per-launch memset bubbles and extra output-copy kernels | Treat nvjet GEMM + memset bubble ladders as an in-flight SGLang linear-kernel family before calling them novel. |
 | PR `#18612` NVFP4 CUTLASS MoE fused SiLU+Mul+quant | `silu_and_mul_scaled_nvfp4`<br>`nvfp4 expert quant`<br>`cutlass moe` | `PR #18612`<br>`python/sglang/srt/layers/moe/cutlass_w4a8_moe.py`<br>`python/sglang/jit_kernel/nvfp4.py` | Fuses MoE activation epilogue and NVFP4 expert quantization before the CUTLASS MoE second GEMM | Treat split SiLU+Mul then NVFP4 expert quant in CUTLASS MoE traces as an in-flight upstream SGLang family. |
 | PR `#22918` FlashInfer per-token NVFP4 MoE | `per_token_nvfp4`<br>`trtllm_fp4_block_scale_moe`<br>`FlashInfer MoE` | `PR #22918`<br>`python/sglang/srt/layers/moe/fused_moe_triton/fused_moe.py` | Adds FlashInfer-backed per-token NVFP4 MoE execution so expert quant/dequant work can move into the fused MoE backend | Treat standalone per-token NVFP4 MoE support kernels as a candidate missing backend-selection path, not an automatically novel kernel idea. |
-| PR `#22851` DSA top-k backend and FlashInfer / PyTorch top-k split | `dsa topk`<br>`flashinfer_topk`<br>`pytorch_topk`<br>`fast_topk_transform` | `PR #22851`<br>`python/sglang/srt/layers/attention/dsa_backend.py` | Makes DSA top-k backend selection explicit and aligns fused top-k transform with FlashInfer / PyTorch fallbacks | When DSA top-k dominates decode, first classify it as backend selection or fused-transform eligibility work. |
-| PR `#24125` GLM5 DSA decode CatArrayBatchedCopy removal | `CatArrayBatchedCopy`<br>`GLM-5`<br>`DSA`<br>`TileLang decode` | `PR #24125`<br>`python/sglang/srt/layers/attention/dsa_backend.py` | Skips redundant cat/copy work in the GLM5 DSA TileLang decode path | Treat cat/copy bursts in GLM5 DSA decode as a concrete in-flight cleanup opportunity. |
+| PR `#22851` NSA top-k backend and FlashInfer / PyTorch top-k split | `nsa topk`<br>`flashinfer_topk`<br>`pytorch_topk`<br>`fast_topk_transform` | `PR #22851`<br>`python/sglang/srt/layers/attention/nsa_backend.py` | Makes NSA top-k backend selection explicit and aligns fused top-k transform with FlashInfer / PyTorch fallbacks | When NSA top-k dominates decode, first classify it as backend selection or fused-transform eligibility work. |
+| PR `#24125` GLM5 NSA decode CatArrayBatchedCopy removal | `CatArrayBatchedCopy`<br>`GLM-5`<br>`NSA`<br>`TileLang decode` | `PR #24125`<br>`python/sglang/srt/layers/attention/nsa_backend.py` | Skips redundant cat/copy work in the GLM5 NSA TileLang decode path | Treat cat/copy bursts in GLM5 NSA decode as a concrete in-flight cleanup opportunity. |
 | PR `#24007` MoE LoRA virtual experts for csgmv backend | `csgmv`<br>`virtual experts`<br>`MoE LoRA`<br>`fused_moe_lora` | `PR #24007`<br>`python/sglang/srt/layers/lora_backend.py`<br>`python/sglang/srt/layers/moe` | Routes MoE LoRA adapter work through virtual experts so csgmv-style kernels can batch it instead of launching fragmented adapter work | Treat MoE-LoRA tiny-kernel ladders as an in-flight batching/fusion family. |
 | PR `#24150` torch.compile local decode support | `enable_torch_compile`<br>`local compile`<br>`decode compile`<br>`torchinductor` | `PR #24150`<br>`python/sglang/srt` | Extends SGLang torch.compile coverage to local decode regions, so Inductor-generated fusion may replace hand-authored tiny kernels | When decode traces show compiler-generated kernels or missing named fused kernels, check this in-flight compile path before calling the shape unsupported. |

@@ -237,24 +243,43 @@ contain the same implementation.
 | vLLM-origin RMSNorm (+ residual add) + quant | `RMSNormQuantFusionPass`<br>`fused_add_rms_norm_static_fp8_quant`<br>`per_token_quant`<br>`per_group_quant` | `vllm/compilation/passes/fusion/rms_quant_fusion.py`<br>`vllm/compilation/passes/fusion/rocm_aiter_fusion.py` | Compile-time and ROCm AITER paths fuse RMSNorm or fused-add-RMSNorm with FP8 / FP4 quant output | Treat split norm/add + quant as an upstream fused family, not an unexplored direction. |
 | vLLM-origin SiLU+Mul + quant | `ActivationQuantFusionPass`<br>`SiluMulFp8*`<br>`Nvfp4`<br>`rocm_aiter` | `vllm/compilation/passes/fusion/act_quant_fusion.py`<br>`vllm/compilation/passes/fusion/rocm_aiter_fusion.py` | Activation epilogues fuse `SiLU+Mul` with FP8 / NVFP4 / AITER group quant instead of materializing the BF16 activation first | Treat standalone activation then quant kernels as matching a vLLM-origin precedent. |
 | vLLM-origin add + RMSNorm + pad | `fuse_act_padding`<br>`RocmAiterTritonAddRMSNormPadFusionPass`<br>`add_rmsnorm_pad` | `vllm/compilation/passes/fusion/rocm_aiter_fusion.py`<br>`docs/design/fusions.md` | ROCm / AITER path fuses residual add + RMSNorm directly into the padded layout expected by the next kernel | Treat norm-plus-padding ladders as an existing backend-specific fuse family first. |
-| vLLM-origin attention + output quant | `fuse_attn_quant`<br>`AttnQuantFusionPass`<br>`merge_attn_states`<br>`output_scale`<br>`output_group_scale`<br>`output_block_scale` | `vllm/compilation/passes/fusion/attn_quant_fusion.py`<br>`vllm/v1/attention/ops/merge_attn_states.py`<br>`vllm/csrc/attention/merge_attn_states.cu`<br>`docs/design/fusions.md` | Compile-time fusion pushes FP8 / NVFP4 quantization into the attention epilogue on supported Triton / FlashInfer / ROCm / AITER backends, and mainline `merge_attn_states` kernels already support FP8 output when `output_scale` is provided | Treat attention-output quant and merged-attention quant epilogues as a known upstream family before calling them novel. |
+| vLLM-origin attention + output quant | `fuse_attn_quant`<br>`AttnQuantFusionPass`<br>`merge_attn_states`<br>`output_scale`<br>`output_group_scale`<br>`output_block_scale` | `vllm/compilation/passes/fusion/attn_quant_fusion.py`<br>`vllm/v1/attention/ops/merge_attn_states.py`<br>`csrc/attention/merge_attn_states.cu`<br>`docs/design/fusions.md` | Compile-time fusion pushes FP8 / NVFP4 quantization into the attention epilogue on supported Triton / FlashInfer / ROCm / AITER backends, and mainline `merge_attn_states` kernels already support FP8 output when `output_scale` is provided | Treat attention-output quant and merged-attention quant epilogues as a known upstream family before calling them novel. |
 | vLLM-origin fused QK RMSNorm + RoPE | `fused_qk_norm_rope`<br>`QKNormRoPEFusionPass`<br>`qk norm + rope` | `vllm/compilation/passes/fusion/qk_norm_rope_fusion.py`<br>`vllm/_custom_ops.py`<br>`csrc/fused_qknorm_rope_kernel.cu` | Compile-time and direct custom-op paths fuse per-head Q / K RMSNorm with RoPE | Treat split QK norm + RoPE as a clear vLLM-origin precedent. |
 | vLLM-origin fused reshape + KV cache write | `reshape_and_cache`<br>`triton_reshape_and_cache_flash`<br>`kv cache write` | `vllm/v1/attention/ops/triton_reshape_and_cache_flash.py`<br>`vllm/v1/attention/backends/triton_attn.py` | Triton cache-update kernels reshape K / V into paged-cache layout and can include FP8 KV-cache scale/write logic | Treat reshape / transpose / cache-write ladders as an existing cache-store fusion family. |
 | vLLM-origin fused RoPE + KV cache update | `fuse_rope_kvcache`<br>`RopeKVCacheFusionPass`<br>`triton_rope_and_cache` | `vllm/compilation/passes/fusion/rope_kvcache_fusion.py`<br>`vllm/_aiter_ops.py`<br>`docs/design/fusions.md` | ROCm / AITER compile-time fusion combines RoPE with paged KV cache update instead of launching them separately | Treat split RoPE + cache-store as a known upstream family, especially on ROCm-like paths. |
-| vLLM-origin fused MLA RoPE + concat/cache write | `concat_and_cache_mla_rope_fused`<br>`mla rope cache` | `vllm/_custom_ops.py`<br>`csrc/cache_kernels_fused.cu` | CUDA kernel fuses MLA-oriented RoPE preparation, concat, and cache write into a direct paged-store path | Treat MLA concat + cache-write ladders as a vLLM-origin precedent before calling them novel. |
+| vLLM-origin fused MLA RoPE + unified KV-cache update | `fused_rope_unified_mla_kv_cache_update`<br>`concat_and_cache_mla_rope_fused`<br>`unified_mla_kv_cache_update` | `vllm/compilation/passes/fusion/mla_rope_kvcache_cat_fusion.py`<br>`vllm/_custom_ops.py`<br>`csrc/cache_kernels_fused.cu` | Current vLLM compile pass fuses MLA-oriented RoPE on `q_pe` / `k_pe`, concat, and unified MLA KV-cache update into a direct paged-store path | Treat MLA RoPE + concat + cache-write ladders as a vLLM-origin mainline precedent before calling them novel. |
 | vLLM-origin fused grouped top-k / biased grouped top-k router | `grouped_topk`<br>`biased_grouped_topk`<br>`grouped_topk_fused_kernel` | `vllm/_custom_ops.py`<br>`vllm/_aiter_ops.py`<br>`vllm/model_executor/layers/fused_moe/router/grouped_topk_router.py`<br>`csrc/moe/grouped_topk_kernels.cu` | CUDA / ROCm router kernels fuse grouped score processing, top-k selection, and routed renorm / bias handling | Treat MoE router ladders as matching an upstream grouped-topk family first. |
 | vLLM-origin fused top-k softmax / sigmoid router | `topk_softmax`<br>`topk_sigmoid`<br>`topkGating`<br>`fused_topk` | `vllm/_custom_ops.py`<br>`vllm/_aiter_ops.py`<br>`vllm/model_executor/layers/fused_moe/router/fused_topk_router.py`<br>`vllm/model_executor/layers/fused_moe/router/fused_topk_bias_router.py`<br>`csrc/moe/topk_softmax_kernels.cu` | CUDA and ROCm / AITER router kernels fuse score activation (`softmax` / `sigmoid`), top-k selection, optional bias correction, and routed renorm into one op instead of routing through grouped-topk or eager softmax-plus-topk ladders | Treat standalone score activation -> top-k -> bias / renorm chains as a known upstream fused router family first. |
 | vLLM-origin DSV3 router GEMM | `dsv3_router_gemm`<br>`allow_dsv3_router_gemm`<br>`router logits` | `vllm/_custom_ops.py`<br>`vllm/model_executor/layers/fused_moe/router/gate_linear.py`<br>`csrc/moe/dsv3_router_gemm_entry.cu`<br>`csrc/moe/dsv3_router_gemm_float_out.cu` | Hopper-class CUDA kernel specializes the DeepSeek router linear for small decode batches and can emit FP32 logits directly without a generic GEMM chain | Treat DeepSeek-style router linear paths as an existing upstream specialized fuse, distinct from grouped-topk itself. |
+| vLLM-origin DeepSeek-V4 fused norm + router GEMM | `dsv4_norm_router_gemm`<br>`norm_gate_linear`<br>`router_gemm`<br>`DeepseekV4ForCausalLM` | `vllm/model_executor/layers/fused_moe/router/norm_gate_linear.py`<br>`vllm/model_executor/models/deepseek_v4.py`<br>`csrc/moe/dsv4_norm_router_gemm*` | Current vLLM mainline has a low-latency DeepSeek-V4 path that fuses norm/router-adjacent work into specialized DSV4 router GEMM kernels | Treat DSV4 norm + router ladders as an upstream mainline fused-router family. |
+| vLLM-origin DeepSeek-V4 MHC fused kernels | `mhc_post_pre`<br>`head_compute_mix_kernel`<br>`aiter mhc`<br>`MHC` | `vllm/model_executor/layers/mhc.py`<br>`vllm/model_executor/kernels/mhc/aiter.py`<br>`vllm/model_executor/models/deepseek_v4.py` | Current vLLM mainline includes Tile head-compute kernels, fused `mhc_post_pre`, and ROCm AITER MHC support for DSV4 | Treat DSV4 MHC/head-compute ladders as known upstream kernel families before proposing a new head-compute kernel. |
 | vLLM-origin GPT-OSS router GEMM | `gpt_oss_router_gemm`<br>`router gemm` | `vllm/_custom_ops.py`<br>`vllm/model_executor/layers/fused_moe/router/gate_linear.py`<br>`csrc/moe/gpt_oss_router_gemm.cu` | Model-specific CUDA kernel replaces the router linear plus bias path with one specialized GEMM op | Treat GPT-OSS-style router linear chains as an existing upstream specialized fuse. |
 | vLLM-origin DeepSeek min-latency fused QKV-A projection | `dsv3_fused_a_gemm`<br>`fused_qkv_a_proj`<br>`q_a_proj` | `vllm/model_executor/models/deepseek_v2.py`<br>`vllm/_custom_ops.py`<br>`csrc/dsv3_fused_a_gemm.cu` | Hopper-class CUDA kernel replaces the tiny-batch DeepSeek QKV-A projection path with one specialized min-latency GEMM instead of a generic linear launch | Treat small-batch DeepSeek QKV-A projection ladders as a known upstream fused kernel family first. |
 | vLLM-origin DSV3.2 fused indexer projections | `wk_weights_proj`<br>`MergedColumnParallelLinear`<br>`weights_proj` | `vllm/model_executor/models/deepseek_v2.py`<br>`vllm/model_executor/models/deepseek_mtp.py` | DSV3.2 indexer paths can fuse the `wk` and `weights_proj` projections into one GEMM and carry the matching MTP weight-loading path | Treat paired indexer projection chains as a known upstream fused linear family before calling the opportunity novel. |
+| vLLM-origin ROCm AITER sparse-MLA paged MQA logits | `rocm_aiter_mla_sparse`<br>`paged_mqa_logits`<br>`gluon`<br>`gfx950` | `vllm/v1/attention/ops/rocm_aiter_mla_sparse.py` | Current vLLM enables the AITER/Gluon paged-MQA logits path on gfx950 / MI355X sparse MLA shapes | On AMD sparse-MLA traces, compare against the AITER paged-MQA logits path before proposing a new logits kernel. |
+| vLLM-origin ROCm DSV4 sparse MLA Triton kernels | `rocm_aiter_mla_sparse_dsv4`<br>`flashmla_sparse`<br>`sparse_swa` | `vllm/v1/attention/backends/mla/rocm_aiter_mla_sparse_dsv4.py`<br>`vllm/v1/attention/backends/mla/flashmla_sparse.py`<br>`vllm/v1/attention/backends/mla/sparse_swa.py` | Current vLLM mainline has ROCm DSV4 sparse-MLA Triton backend coverage | On DSV4 ROCm sparse-MLA traces, compare backend selection and sparse-SWA paths before writing a new MLA backend. |
+| vLLM-origin DSV4 dequant gather K cache | `dequant_gather_k_cutedsl`<br>`fused_indexer_q_cutedsl`<br>`cache_utils` | `vllm/v1/attention/ops/deepseek_v4_ops/dequant_gather_k_cutedsl.py`<br>`vllm/v1/attention/ops/deepseek_v4_ops/cache_utils.py` | Current vLLM mainline has a newer CuTe DSL dequant-gather K cache path for DSV4 | Treat K-cache dequant/gather ladders as a known upstream DSV4 cache-kernel family. |
+| vLLM-origin TokenSpeed MLA backend | `TOKENSPEED_MLA`<br>`tokenspeed_mla_decode`<br>`tokenspeed_mla_prefill` | `vllm/v1/attention/backends/mla/tokenspeed_mla.py`<br>`vllm/v1/attention/backends/mla/prefill/tokenspeed_mla.py`<br>`vllm/model_executor/layers/attention/mla_attention.py` | Current vLLM can select TokenSpeed MLA prefill/decode kernels for Blackwell FP8-KV DeepSeek/Kimi MLA shapes when the package and platform gates pass | On Blackwell MLA traces, compare against TokenSpeed MLA backend selection before proposing a new attention kernel. |
 | vLLM-origin MiniMax allreduce_rms kernels | `minimax_allreduce_rms`<br>`minimax_allreduce_rmsnorm`<br>`MiniMax-M2.5`<br>`allreduce_rms` | `vllm/model_executor/models/minimax_m2.py` | TensorRT-LLM-derived MiniMax allreduce-plus-RMSNorm kernels are a concrete upstream TP decode family | Treat MiniMax TP norm + collective ladders as an upstream specialized fusion family. |
 | vLLM-origin CUTLASS scaled MM with scale / bias epilogue | `cutlass_scaled_mm`<br>`cutlass_scaled_mm_azp`<br>`scaled mm` | `vllm/_custom_ops.py`<br>`vllm/model_executor/kernels/linear/scaled_mm/cutlass.py`<br>`csrc/libtorch_stable/quantization/w8a8/cutlass/scaled_mm_entry.cu` | CUTLASS kernels fuse activation scales, weight scales, matmul, and optional bias / AZP epilogues | Treat separate scale-mul + GEMM + bias ladders as a vLLM-origin fused linear family first. |
 | vLLM-origin fused MoE expert execution | `cpu_fused_moe`<br>`rocm_aiter_fused_moe`<br>`FusedMoE` | `vllm/model_executor/layers/fused_moe/layer.py`<br>`vllm/model_executor/layers/fused_moe/cpu_fused_moe.py`<br>`vllm/model_executor/layers/fused_moe/rocm_aiter_fused_moe.py`<br>`vllm/_aiter_ops.py` | MoE backends on CUDA / ROCm / CPU already collapse packed expert execution into fused expert kernels rather than per-expert eager GEMMs | Treat exposed expert-side tiny GEMM ladders as matching an upstream fused-MoE family. |
 | vLLM-origin fused MoE LoRA | `fused_moe_lora`<br>`fused_moe_lora_fp8`<br>`w13_shrink`<br>`w2_expand` | `vllm/lora/ops/triton_ops/fused_moe_lora_op.py`<br>`vllm/lora/ops/triton_ops/fused_moe_lora_fp8_op.py`<br>`vllm/lora/layers/fused_moe.py` | Triton kernels fuse LoRA shrink / expand work into MoE expert execution, including FP8 variants | Treat MoE-LoRA adapter work as an upstream fused family before proposing a brand new kernel. |
 | vLLM-origin ViT fused bilinear position-embedding interpolation | `triton_pos_embed_interpolate`<br>`bilinear_pos_embed`<br>`pos_embed_interpolate_native` | `vllm/model_executor/models/qwen3_vl.py` | Triton kernel fuses bilinear interpolation and spatial-merge reorder for Qwen3-VL ViT position embeddings, replacing many tiny eager kernels | Treat VLM position-embedding ladders as an existing vLLM-origin Triton fusion family. |

-## 15. vLLM-origin kernel-overlap families
+## 15. TokenSpeed-origin fused-kernel families
+
+These rows are direct TokenSpeed families from `lightseekorg/tokenspeed`, not
+only vLLM references to the TokenSpeed package.
+
+| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
+| --- | --- | --- | --- | --- |
+| TokenSpeed CuTe DSL MLA prefill / decode | `tokenspeed_mla_decode`<br>`tokenspeed_mla_prefill`<br>`BlackwellMultiHeadLatentAttentionForward` | `python/tokenspeed/runtime/layers/attention/backends/tokenspeed_mla.py`<br>`tokenspeed-mla/python/tokenspeed_mla/mla_decode.py`<br>`tokenspeed-mla/python/tokenspeed_mla/mla_prefill.py`<br>`tokenspeed-kernel/python/tokenspeed_kernel/ops/attention/tokenspeed_mla/__init__.py` | Blackwell SM100 CuTe DSL MLA kernels cover FP8-KV prefill/decode/verify paths through the `tokenspeed_mla` backend | On TokenSpeed or vLLM+TokenSpeed MLA traces, compare backend selection before proposing a new MLA attention kernel. |
+| TokenSpeed MLA KV pack + FP8 quantize | `_mla_kv_pack_quantize_fp8_kernel`<br>`mla_kv_pack_quantize_fp8`<br>`k_nope` / `k_pe` | `tokenspeed-mla/python/tokenspeed_mla/mla_kv_pack_quantize_fp8.py`<br>`tokenspeed-kernel/python/tokenspeed_kernel/ops/attention/tokenspeed_mla/__init__.py` | One Triton kernel packs `k_nope`, broadcast `k_pe`, and `v`, then writes FP8 K/V for MLA chunked prefill | Treat split K/V concat + FP8 cast ladders as a known TokenSpeed fusion family. |
+| TokenSpeed fused top-k + top-p sampling | `fused_topk_topp`<br>`fused_topk_topp_renorm` | `tokenspeed-kernel/python/tokenspeed_kernel/thirdparty/cuda/fused_topk_topp.py`<br>`tokenspeed-kernel/python/tokenspeed_kernel/thirdparty/cuda/csrc/fused_topk_topp/fused_topk_topp.cu` | CUDA extension fuses top-k, top-p, and renormalization for decode sampling | Treat top-k/top-p/renorm chains in TokenSpeed traces as an existing sampling fusion first. <!-- codespell:ignore thirdparty --> |
+| TokenSpeed persistent lm_head GEMM | `lm_head_gemm`<br>`should_use_fused`<br>`persistent` | `tokenspeed-kernel/python/tokenspeed_kernel/thirdparty/cuda/lm_head_gemm.py`<br>`tokenspeed-kernel/python/tokenspeed_kernel/thirdparty/cuda/csrc/lm_head_gemm.cu` | Shape-gated persistent GEMM replaces `torch.matmul` for selected lm_head / router-like projection shapes | Treat visible lm_head matmul ladders as a candidate for this existing TokenSpeed path before inventing a new logits GEMM. <!-- codespell:ignore thirdparty --> |
+| TokenSpeed NVFP4 GEMM + SwiGLU + quant | `nvfp4_gemm_swiglu_nvfp4_quant`<br>`SwiGLU`<br>`SFC` | `tokenspeed-kernel/python/tokenspeed_kernel/thirdparty/cute_dsl/nvfp4_gemm_swiglu_nvfp4_quant.py` | CuTe DSL kernel fuses block-scaled NVFP4 GEMM, SwiGLU, and optional output quantization | Treat split expert GEMM + activation + FP4 quant chains as matching an upstream TokenSpeed kernel family. <!-- codespell:ignore thirdparty --> |
+
+## 16. vLLM-origin kernel-overlap families

 | Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
 | --- | --- | --- | --- | --- |
@@ -263,7 +288,7 @@ contain the same implementation.
 | vLLM-origin shared-expert aux-stream overlap | `aux_stream`<br>`shared_experts_stream`<br>shared expert near router | `vllm/model_executor/layers/fused_moe/runner/shared_experts.py`<br>`vllm/model_executor/layers/fused_moe/runner/moe_runner_base.py` | MoE shared experts can record the cloned input on `shared_experts_stream`, wait on the caller stream, run in parallel with router-side work, and rejoin before merge | Treat shared-expert vs router overlap as an existing upstream sparse-model family. |
 | vLLM-origin DCP async all-to-all overlap | `dcp_alltoall`<br>`all_to_all_single`<br>`async_op=True` | `vllm/v1/attention/ops/dcp_alltoall.py` | Output / LSE exchange uses async all-to-all handles instead of serializing collective completion on the main path | Treat DCP all-to-all windows as an upstream async-collective family. |

-## 16. vLLM-origin PR-backed / in-flight fused-kernel and kernel-overlap families
+## 17. vLLM-origin PR-backed / in-flight fused-kernel and kernel-overlap families

 | Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
 | --- | --- | --- | --- | --- |
@@ -273,18 +298,18 @@ contain the same implementation.
 | PR `#38621` fused QK norm + RoPE + cache + quant | `fused_qk_norm_rope_cache_quant`<br>`QK Norm + RoPE + Cache + Quant` | `PR #38621`<br>`csrc/fused_qk_norm_rope_cache_quant.cu`<br>`vllm/compilation/passes/fusion/qk_norm_rope_cache_quant_fusion.py` | Draft CUDA kernel and compile-time pass try to fuse QK RMSNorm, RoPE, KV cache write, and optional FP8 quant for small-batch decode | Treat this as an in-flight upstream fusion family before calling a similar idea novel. |
 | PR `#37646` ROCm AITER fused allreduce + RMSNorm | `rocm_aiter_fused_allreduce_rmsnorm`<br>`custom_fused_ar_rms`<br>`RocmAiterAllReduceFusionPass` | `PR #37646`<br>`vllm/_aiter_ops.py`<br>`vllm/compilation/passes/pass_manager.py` | ROCm-specific compile-time path swaps the generic all-reduce fusion pass for an AITER fused allreduce-plus-RMSNorm kernel family | Treat ROCm TP all-reduce + RMSNorm ladders as an in-flight upstream fused-collective family first. |
 | PR `#36413` FlashInfer RMSNorm + FP4 quant fusion | `fuse_norm_quant`<br>`flashinfer`<br>`NVFP4`<br>`rmsnorm + fp4 quant` | `PR #36413`<br>`vllm/compilation/passes/fusion/rms_quant_fusion.py`<br>`vllm/docs/design/fusions.md` | FlashInfer-backed norm-plus-FP4 quant fusion extends the existing RMSNorm+quant family to NVFP4 flows | Treat split RMSNorm + FP4 quant ladders as an upstream in-flight family, not a fresh idea. |
-| PR `#39301` GLM5 router GEMM with PDL overlap | `TRTLLM_ENABLE_PDL`<br>`router_gemm`<br>`GLM5`<br>`FI AR RMS fusion` | `PR #39301`<br>`vllm/model_executor/layers/fused_moe/router/gate_linear.py`<br>`vllm/csrc/moe/dsv3_router_gemm_utils.h` | Extends the specialized router GEMM family to GLM5 hidden size and uses PDL to overlap the router launch with the preceding fused allreduce-plus-RMS block | Treat this as an in-flight upstream router-kernel plus launch-overlap family before calling it novel. |
+| PR `#39301` GLM5 router GEMM with PDL overlap | `TRTLLM_ENABLE_PDL`<br>`router_gemm`<br>`GLM5`<br>`FI AR RMS fusion` | `PR #39301`<br>`vllm/model_executor/layers/fused_moe/router/gate_linear.py`<br>`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. |
 | PR `#41455` ROCm WMMA paged prefill and split-K decode | `wmma`<br>`paged prefill`<br>`split-K decode`<br>`ROCm attention` | `PR #41455`<br>`vllm/v1/attention`<br>`vllm/_aiter_ops.py` | Adds ROCm WMMA attention kernels for paged prefill and split-K decode shapes | Treat split attention support kernels on AMD as an in-flight vLLM attention-kernel family before calling them novel. |
-| PR `#41263` DeepSeek-V4 fused norm / router low-latency path | `DSV4`<br>`fuse norm router`<br>`low latency`<br>`router` | `PR #41263`<br>`vllm/model_executor/models/deepseek_v2.py`<br>`vllm/model_executor/layers/fused_moe/router` | Targets DeepSeek-V4 decode latency by fusing norm / router-adjacent work and low-latency model paths | Treat DSV4 norm-router ladders as a concrete in-flight upstream family. |
-| PR `#41428` DSV4 fused indexer Q quant kernel | `DSV4`<br>`fused Indexer Q quant`<br>`indexer q`<br>`fp4` | `PR #41428`<br>`vllm/model_executor/models/deepseek_v2.py`<br>`vllm/csrc` | Improves the fused DeepSeek-V4 indexer Q quant kernel instead of materializing Q then quantizing separately | Treat DSV4 indexer-Q quant ladders as an in-flight upstream fused quant family. |
-| PR `#41255` DeepSeek-V4 Tile kernels / `head_compute_mix_kernel` | `head_compute_mix_kernel`<br>`Tile kernel`<br>`DSV4`<br>`MLA` | `PR #41255`<br>`vllm/model_executor/models/deepseek_v2.py`<br>`vllm/csrc` | Adds DeepSeek-V4 Tile kernels that mix head compute work in one specialized kernel | Treat DSV4 MLA head-compute ladders as a known in-flight specialized-kernel family. |
+| PR `#41263` DeepSeek-V4 fused norm / router low-latency path | `DSV4`<br>`fuse norm router`<br>`low latency`<br>`router` | `PR #41263`<br>`vllm/model_executor/models/deepseek_v4.py`<br>`vllm/model_executor/layers/fused_moe/router/norm_gate_linear.py` | Merged into current mainline as the DSV4 fused norm + router GEMM family above | Treat this row as provenance for the shipped mainline family, not as merely in-flight. |
+| PR `#41428` DSV4 fused indexer Q quant kernel | `DSV4`<br>`fused Indexer Q quant`<br>`indexer q`<br>`fp4` | `PR #41428`<br>`vllm/model_executor/models/deepseek_v4.py`<br>`vllm/v1/attention/ops/deepseek_v4_ops/fused_indexer_q.py`<br>`vllm/v1/attention/ops/deepseek_v4_ops/fused_indexer_q_cutedsl.py` | Merged current-main improvement to the fused DeepSeek-V4 indexer Q quant kernel instead of materializing Q then quantizing separately | Treat DSV4 indexer-Q quant ladders as an upstream mainline fused quant family. |
+| PR `#41255` DeepSeek-V4 Tile kernels / `head_compute_mix_kernel` | `head_compute_mix_kernel`<br>`Tile kernel`<br>`DSV4`<br>`MLA` | `PR #41255`<br>`vllm/model_executor/models/deepseek_v4.py`<br>`vllm/model_executor/layers/mhc.py`<br>`vllm/model_executor/kernels/mhc/tilelang.py` | Merged current-main DeepSeek-V4 Tile kernels that mix head compute work in one specialized kernel | Treat DSV4 MLA head-compute ladders as a known upstream mainline specialized-kernel family. |
 | PR `#41441` DSV4 all-reduce plus `mhc_post` fusion | `DSV4`<br>`AR+mhc_post`<br>`allreduce`<br>`mhc_post` | `PR #41441`<br>`vllm/model_executor/models/deepseek_v2.py`<br>`vllm/compilation/passes/fusion` | Fuses or overlaps DSV4 all-reduce with post-MLA head-compute work | Treat all-reduce followed by `mhc_post` in DSV4 traces as an in-flight vLLM overlap/fusion family. |
 | PR `#41446` AMD GatedDeltaNet FLA prefill kernels | `GatedDeltaNet`<br>`FLA prefill`<br>`AMD`<br>`Qwen3-Next` | `PR #41446`<br>`vllm/model_executor/models/qwen3_next.py`<br>`vllm/v1/attention` | Optimizes GatedDeltaNet / FLA prefill kernels on AMD linear-attention models | Treat split GDN prefill kernels on ROCm as an in-flight upstream family. |
 | PR `#39748` dual-stream GDN input projection | `dual-stream`<br>`input projection`<br>`GatedDeltaNet`<br>`Qwen3.5` | `PR #39748`<br>`vllm/model_executor/models/qwen3_next.py` | Overlaps sibling input-projection branches for Qwen3 / Qwen3.5 GDN-style blocks | Treat serial GDN input projections as a known in-flight overlap opportunity. |
 | PRs `#41433` / `#41434` / `#41429` / `#40561` GPU/CPU sync removal | `GPU->CPU sync`<br>`cpu sync`<br>`item()`<br>`non_blocking` | `PR #41433`<br>`PR #41434`<br>`PR #41429`<br>`PR #40561` | Removes or gates accidental GPU-to-CPU synchronization points and adds sync-detection coverage | Treat CPU gaps next to small GPU kernels as an upstream vLLM sync-removal family before proposing a kernel-only fix. |
 | PR `#36823` vLLM IR `fused_add_rms_norm` overload | `vllm_ir`<br>`fused_add_rms_norm`<br>`maybe_inplace` | `PR #36823`<br>`vllm/compilation/passes/ir`<br>`vllm/compilation/passes/fusion/rms_quant_fusion.py` | Extends vLLM IR lowering so fused-add-RMSNorm variants remain visible to later compile-time fusions | Treat missing norm/quant compile fusion as potentially an IR-lowering visibility issue. |

-## 17. Important toggles and caveats
+## 18. Important toggles and caveats

 | Toggle / env | Location | Effect on trace interpretation |
 | --- | --- | --- |
@@ -294,7 +319,7 @@ contain the same implementation.
 | `enable_single_batch_overlap` | `python/sglang/srt/server_args.py` | Enables the SBO family. |
 | `enable_fused_moe_sum_all_reduce` | `python/sglang/srt/server_args.py` | Enables fused MoE sum-reduce in the down path. |
 | `SGLANG_BLACKWELL_OVERLAP_SHARED_EXPERTS_OUTSIDE_SBO` | `python/sglang/srt/environ.py` | Alters how DeepSeek-style shared-expert overlap behaves on Blackwell. |
-| `SGLANG_DSA_FUSE_TOPK` | `python/sglang/srt/environ.py` | Gates DSA fused top-k transform / page-table build. |
+| `SGLANG_NSA_FUSE_TOPK` | `python/sglang/srt/environ.py` | Gates NSA fused top-k transform / page-table build. |
 | `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. |
@@ -313,20 +338,27 @@ contain the same implementation.
 | `mlir_elementwise_fusion` | `tensorrt_llm/_torch/auto_deploy/config/default.yaml` | Can absorb merge adds into larger fused kernels, so missing explicit merge nodes in multi-stream traces may be intentional. |
 | `enable_torch_compile` | `python/sglang/srt/server_args.py`<br>`python/sglang/multimodal_gen/runtime/server_args.py` | Compiler-generated fusion / reordering can hide handwritten kernel names; absence of a custom kernel does not always mean absence of fusion. |
 | `enable_fused_grouped_gemm_combine` | `PR #21877` | In-flight path that intentionally disables SBO because combine is folded into down-GEMM. |
-| `PassConfig.fuse_allreduce_rms` | `vllm/config/compilation.py` | Enables vLLM's AllReduce -> RMSNorm (+ residual / quant) compile-time fusion family. |
+| `PassConfig.fuse_allreduce_rms` | `vllm/config/compilation.py` | Enables vLLM's AllReduce -> RMSNorm (+ residual / quant) compile-time fusion family; on ROCm, AITER variants and add-RMSNorm-pad ordering can change the visible kernel split. |
 | `PassConfig.fuse_norm_quant` | `vllm/config/compilation.py` | Enables vLLM's RMSNorm(+residual add) -> FP8 / FP4 quant compile-time fusion family. |
 | `PassConfig.fuse_act_quant` | `vllm/config/compilation.py` | Enables vLLM's `SiLU+Mul -> quant` fusion family, plus ROCm AITER variants where applicable. |
 | `PassConfig.fuse_attn_quant` | `vllm/config/compilation.py` | Enables attention-epilogue quant fusion; requires the right backend / graph visibility, so split kernels may still be expected. |
 | `PassConfig.fuse_mla_dual_rms_norm` | `vllm/config/compilation.py` | Enables the AITER-backed MLA paired-Q/KV RMSNorm fusion family on ROCm. |
 | `PassConfig.enable_qk_norm_rope_fusion` | `vllm/config/compilation.py` | Enables the compile-time QK RMSNorm + RoPE family on CUDA-like backends. |
 | `PassConfig.fuse_rope_kvcache` | `vllm/config/compilation.py` | Enables ROCm / AITER RoPE + KV-cache update fusion and is range-limited by token count. |
+| `PassConfig.fuse_rope_kvcache_cat_mla` | `vllm/config/compilation.py` | Enables the MLA RoPE + unified MLA KV-cache update compile-time fusion family. |
+| `AttentionBackendEnum.TOKENSPEED_MLA` | `vllm/model_executor/layers/attention/mla_attention.py`<br>`vllm/platforms/cuda.py` | Selects the TokenSpeed MLA backend on supported Blackwell FP8-KV MLA shapes when `tokenspeed-mla` is installed. |
+| `rocm_aiter_mla_sparse_dsv4` | `vllm/v1/attention/backends/mla/rocm_aiter_mla_sparse_dsv4.py` | Selects the ROCm DSV4 sparse-MLA backend path on supported AMD shapes. |
 | `PassConfig.fuse_minimax_qk_norm` | `vllm/config/compilation.py` | Enables the MiniMax decode Q/K allreduce-plus-RMSNorm compile-time fusion family. |
 | `PassConfig.fuse_act_padding` | `vllm/config/compilation.py` | Enables the ROCm AITER add-RMSNorm-plus-pad fusion family when AITER is available. |
 | `PassConfig.enable_sp` | `vllm/config/compilation.py` | Rewrites all-reduce into sequence-parallel staging; this is often a prerequisite for the overlap family, not just a pure fuse toggle. |
 | `PassConfig.fuse_gemm_comms` | `vllm/config/compilation.py` | Enables AsyncTP GEMM + collective overlap and auto-enables `enable_sp` when valid. |
-| `TRTLLM_ENABLE_PDL` | `vllm/csrc/dsv3_fused_a_gemm.cu`<br>`vllm/csrc/moe/dsv3_router_gemm_utils.h` | Enables programmatic dependent launch for the DSV3 specialized CUDA kernels, which can change launch grouping and trace shape for router / QKV-A paths. |
+| vLLM PR `#46735` Triton MoE CUDA graph capture fix | `vllm/model_executor/layers/fused_moe/experts/triton_moe.py`<br>`vllm/model_executor/layers/fused_moe/experts/nvfp4_emulation_moe.py` | Latest vLLM mainline fixes CUDA graph capture around Triton / NVFP4-emulation MoE; stale target images may show graph-capture failures or eager fallbacks that are not SGLang kernel wins. |
+| `TRTLLM_ENABLE_PDL` | `csrc/libtorch_stable/dsv3_fused_a_gemm.cu`<br>`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. |
+| TokenSpeed `--attention-backend tokenspeed_mla` | `python/tokenspeed/runtime/layers/attention/backends/tokenspeed_mla.py` | Selects TokenSpeed's native CuTe DSL MLA backend; requires compatible Blackwell FP8-KV MLA shapes, so split MLA support kernels may indicate backend gating rather than a missing kernel. |
+| TokenSpeed `TOKENSPEED_MLA_PREFILL_BACKEND` | `tokenspeed-mla/python/tokenspeed_mla/mla_prefill.py` | Chooses CuTe DSL JIT vs binary prefill backend; trace kernel names can differ even when the same MLA fused family applies. |
+| TokenSpeed `--comm-fusion-max-num-tokens` / `--enable-allreduce-fusion` | `docs/configuration/server.md`<br>`python/tokenspeed/runtime/distributed/comm_backend` | Gates TokenSpeed communication-fusion behavior; inspect these before treating all-reduce + compute separation as a novel overlap gap. |

-## 18. Suggested refresh commands
+## 19. 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.
@@ -26,9 +26,12 @@ necessarily present in the checked-out `sglang` tree, but they should still be
 treated as upstream or analogous kernel-overlap families before labeling an
 overlap opportunity as novel.

-Refresh note `2026-04-22`: rescanned current `sglang`, `flashinfer`,
-`TensorRT-LLM`, and `vllm` mainline overlap paths plus rechecked referenced PR
-state via the GitHub API on `2026-04-22`. Closed-unmerged SGLang
+Refresh note `2026-06-26`: rechecked official main heads for SGLang
+`8524678889485801e7a4a12d62015be0c68f7a90`, vLLM
+`abc71548ef029132c3316b902207f254a246d593`, TensorRT-LLM
+`0722c5f47d2cae69ac1a237da51e550dd214532c`, and TokenSpeed
+`5aedf69d6b476baa65571011de6ea60fd5a238a8`, then added the first
+TokenSpeed-origin communication-fusion row. Closed-unmerged SGLang
 [#22410](https://github.com/sgl-project/sglang/pull/22410) and FlashInfer
 [#2840](https://github.com/flashinfer-ai/flashinfer/pull/2840) were removed
 from the PR-backed sections. SGLang
@@ -48,7 +51,7 @@ upstream overlap references as of this refresh.
 | Llama4 shared branch vs routed branch overlap | shared expert branch plus routed MoE branch as adjacent windows | `python/sglang/srt/models/llama4.py` | shared expert on current stream, router + topk + routed experts on `alt_stream` | Use Llama4 as the first precedent for branch-level overlap in similar sparse models. |
 | ExaoneMoE shared experts vs router experts overlap | shared expert output and router-expert output form a two-branch window | `python/sglang/srt/models/exaone_moe.py::forward_normal_dual_stream` | shared experts on current stream, router + routed experts on `alt_stream`, explicit join before combine | This is an existing dual-stream MoE overlap family. |
 | Grok residual-MoE branch overlap | dense MLP and block-sparse MoE branches in parallel | `python/sglang/srt/models/grok.py::moe_with_rmoe` | dense MLP on current stream, MoE on `alt_stream`, fused dual residual RMSNorm around boundaries | Treat exposed Grok branch overlap as an existing pattern. |
-| DSA dual-stream overlap | Q-proj, K-proj, RoPE, cache-store, quantization in tight two-stream windows | `python/sglang/srt/layers/attention/dsa/dsa_indexer.py` | Q / K projection split, RoPE split, cache-store vs quantization overlap | DSA already contains several dual-stream overlap precedents. |
+| NSA dual-stream overlap | Q-proj, K-proj, RoPE, cache-store, quantization in tight two-stream windows | `python/sglang/srt/layers/attention/nsa/nsa_indexer.py` | Q / K projection split, RoPE split, cache-store vs quantization overlap | NSA already contains several dual-stream overlap precedents. |
 | MoriEP async dispatch / combine comm stream | `MoriEP`<br>`_comm_stream`<br>`dispatch`<br>`combine`<br>`done_event` | `python/sglang/srt/layers/moe/token_dispatcher/moriep.py` | MoriEP can submit dispatch and combine onto a dedicated communication stream and synchronize only through events | Treat MoriEP comm / compute interleave as an existing MoE overlap family. |
 | Generic `alt_stream` overlap families | `alt_stream` plus explicit `wait_stream` / `with torch.cuda.stream(...)` | `qwen2_moe.py`<br>`qwen3_moe.py`<br>`glm4_moe.py`<br>`bailing_moe.py`<br>`llada2.py`<br>`grok.py`<br>`olmo2.py`<br>`step3p5.py`<br>`longcat_flash.py`<br>`falcon_h1.py` | model-specific overlap on attention prep, MoE branches, or cache-store | Search these families before designing a new overlap scheme from scratch. |

@@ -107,7 +110,16 @@ AutoDeploy rather than same-stream PDL windows.
 | TensorRT-LLM multi-stream MoE shared-vs-routed overlap | `multi_stream_moe`<br>`begin_aux_stream_passthrough`<br>`end_aux_stream_passthrough`<br>`wait_aux_stream_passthrough`<br>`mlir_elementwise_fusion`<br>`piecewise cudagraph`<br>`caller_stream.synchronize()` | `tensorrt_llm/_torch/auto_deploy/transform/library/multi_stream_moe.py`<br>`tensorrt_llm/_torch/auto_deploy/utils/multi_stream_utils.py` | Shared-expert work is moved to an auxiliary stream while routed-expert MoE work remains on the main stream and rejoins at the merge node; the same family includes synchronization rules for MLIR-fused kernels and piecewise cudagraph replay | Treat shared-expert vs routed-expert windows, including altered behavior under MLIR / piecewise graph modes, as an existing TensorRT-LLM branch-overlap family. |
 | TensorRT-LLM multi-stream FP8 GEMM fork parallelism | `multi_stream_gemm`<br>`trtllm_finegrained_fp8_linear`<br>`record_event_passthrough`<br>`_aux` | `tensorrt_llm/_torch/auto_deploy/transform/library/multi_stream_gemm.py`<br>`tensorrt_llm/_torch/auto_deploy/utils/multi_stream_utils.py` | Compiler pass identifies fork points with multiple FP8 linears and moves the largest GEMM to the auxiliary stream so sibling GEMMs overlap | Treat sibling FP8 linear branches as an existing TensorRT-LLM overlap family before designing a new stream split. |

-## 8. vLLM-origin kernel-overlap families
+## 8. TokenSpeed-origin kernel-overlap families
+
+These rows are comparative references from `lightseekorg/tokenspeed`. Use them
+when the trace is from TokenSpeed or from a vLLM/TokenSpeed hybrid deployment.
+
+| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
+| --- | --- | --- | --- | --- |
+| TokenSpeed allreduce / communication fusion | `enable_allreduce_fusion`<br>`comm_fusion`<br>`comm_fusion_max_num_tokens`<br>`allreduce` | `docs/configuration/server.md`<br>`python/tokenspeed/runtime/distributed/comm_backend` | TokenSpeed exposes runtime knobs for communication fusion and token-count gating, so all-reduce windows may be a disabled or shape-gated TokenSpeed path | Treat split all-reduce + compute windows in TokenSpeed traces as a comm-fusion eligibility question before calling it a novel overlap opportunity. |
+
+## 9. vLLM-origin kernel-overlap families

 | Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
 | --- | --- | --- | --- | --- |
@@ -116,14 +128,14 @@ AutoDeploy rather than same-stream PDL windows.
 | vLLM-origin shared-expert aux-stream overlap | `aux_stream`<br>`shared_experts_stream`<br>shared expert near router | `vllm/model_executor/layers/fused_moe/runner/shared_experts.py`<br>`vllm/model_executor/layers/fused_moe/runner/moe_runner_base.py` | MoE shared experts can record the cloned input on `shared_experts_stream`, wait on the caller stream, run in parallel with router-side work, and rejoin before merge | Treat shared-expert vs router overlap as an existing upstream sparse-model family. |
 | vLLM-origin DCP async all-to-all overlap | `dcp_alltoall`<br>`all_to_all_single`<br>`async_op=True` | `vllm/v1/attention/ops/dcp_alltoall.py` | Output / LSE exchange uses async all-to-all handles instead of serializing collective completion on the main path | Treat DCP all-to-all windows as an upstream async-collective family. |

-## 9. vLLM-origin PR-backed / in-flight kernel-overlap families
+## 10. vLLM-origin PR-backed / in-flight kernel-overlap families

 | Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
 | --- | --- | --- | --- | --- |
 | PR `#35968` DSV3.2 multi-stream indexer overlap | `weights_proj`<br>`wk`<br>`k_norm`<br>`aux_stream` | `PR #35968`<br>`vllm/model_executor/models/deepseek_v2.py`<br>`vllm/utils/torch_utils.py` | Closed PR explored overlapping the small `weights_proj` GEMM with `wk + k_norm` on a secondary CUDA stream for decode batches instead of serializing both on the default stream | Treat this as a concrete upstream decode-time kernel-overlap family when traces show underutilized projection overlap opportunities. |
 | PR `#39301` GLM5 router GEMM with PDL overlap | `TRTLLM_ENABLE_PDL`<br>`router_gemm`<br>`GLM5`<br>`FI AR RMS fusion` | `PR #39301`<br>`vllm/model_executor/layers/fused_moe/router/gate_linear.py`<br>`vllm/csrc/moe/dsv3_router_gemm_utils.h` | The GLM5 router GEMM path explicitly uses PDL so the router kernel can overlap with the preceding fused allreduce-plus-RMS block on supported GPUs | Treat router-GEMM launch overlap on GLM5-like traces as an in-flight upstream family first. |

-## 10. Important toggles and caveats
+## 11. Important toggles and caveats

 | Toggle / env | Location | Effect on trace interpretation |
 | --- | --- | --- |
@@ -144,8 +156,9 @@ AutoDeploy rather than same-stream PDL 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. |
+| TokenSpeed `--comm-fusion-max-num-tokens` / `--enable-allreduce-fusion` | `docs/configuration/server.md` | Gates TokenSpeed communication fusion; inspect it before treating all-reduce + compute separation as a new overlap gap. |

-## 11. 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.
@@ -155,6 +168,7 @@ the local source trees. They are not used by the triage scripts at runtime.
 FLASHINFER_REPO=${FLASHINFER_REPO:-../flashinfer}
 TRTLLM_REPO=${TRTLLM_REPO:-../TensorRT-LLM}
 VLLM_REPO=${VLLM_REPO:-../vllm}
+TOKENSPEED_REPO=${TOKENSPEED_REPO:-../tokenspeed}

 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
@@ -166,6 +180,8 @@ rg -n "mlir_elementwise_fusion|piecewise|cudagraph|caller_stream.synchronize" "$
 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|maybe_sync_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'
+rg -n "enable_allreduce_fusion|comm_fusion|comm_fusion_max_num_tokens|allreduce|reduce_scatter" "$TOKENSPEED_REPO/python" "$TOKENSPEED_REPO/docs"
+git -C "$TOKENSPEED_REPO" log --all --format='%h %s' | rg -i 'fused|fusion|overlap|allreduce|stream|comm|mla|tokenspeed_mla'
 # 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"
@@ -177,4 +193,5 @@ git -C "$VLLM_REPO" log --all --format='%h %s' | rg -i 'fused|fusion|overlap|all
 #   "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"
+#   "fused OR overlap OR comm_fusion repo:lightseekorg/tokenspeed"
 ```
@@ -1,7 +1,12 @@
 # vLLM Torch Compile Fusion Patterns

-Refresh: `2026-05-01`.
-Source tree: vLLM `origin/main` at `7075df79b`.
+Refresh: `2026-06-26`.
+Source tree: vLLM `origin/main` at
+`abc71548ef029132c3316b902207f254a246d593`; no new LLM compile-fusion pass was
+added after `2317682f9` in this refresh. The mainline `#40392` MLA RoPE +
+KV-cache cat fusion is already included below. Recent post-`#46735` vLLM
+changes include runtime / frontend work such as `#44800` and `#46799`, but they
+do not add a new LLM compile-fusion pass to this inventory.

 Use this file when the fuse-pattern table reports split kernels in a trace and
 you need to decide whether the shape is already covered by vLLM's
@@ -17,7 +22,7 @@ vLLM registers these passes from
 | --- | --- | --- |
 | `enable_sp` | `SequenceParallelismPass` | all-reduce around residual/norm blocks becomes reduce-scatter, local work, and all-gather |
 | `fuse_gemm_comms` | `AsyncTPPass` | GEMM plus reduce-scatter / all-gather overlap through symmetric-memory collectives |
-| `fuse_allreduce_rms` | `AllReduceFusionPass` | all-reduce followed by RMSNorm, optional residual add, optional FP8 / NVFP4 quant |
+| `fuse_allreduce_rms` | `AllReduceFusionPass` or ROCm AITER variant | all-reduce followed by RMSNorm, optional residual add, optional FP8 / NVFP4 quant; current pass ordering runs AITER add-RMSNorm-pad before this fusion when available |
 | `fuse_minimax_qk_norm` | `MiniMaxQKNormPass` | MiniMax Q/K all-reduce plus RMSNorm decode path |
 | `fuse_norm_quant` | `RMSNormQuantFusionPass` | RMSNorm or fused-add-RMSNorm followed by FP8 / FP4 quant |
 | `fuse_norm_quant` + AITER | `RocmAiterRMSNormQuantFusionPass` | ROCm AITER RMSNorm / fused-add-RMSNorm followed by AITER or vLLM quant |
@@ -26,6 +31,7 @@ vLLM registers these passes from
 | `fuse_act_padding` + AITER | `RocmAiterTritonAddRMSNormPadFusionPass` | AITER fused-add-RMSNorm followed by padding into the next layout |
 | `fuse_mla_dual_rms_norm` + AITER | `MLADualRMSNormFusionPass` | MLA paired Q and KV RMSNorms become `fused_mla_dual_rms_norm` |
 | `fuse_rope_kvcache` | `RopeKVCacheFusionPass` | RoPE plus paged KV-cache update, after split cleanup passes |
+| `fuse_rope_kvcache_cat_mla` | `MLARoPEKVCacheCatFusionPass` | MLA RoPE on `q_pe` / `k_pe` plus unified MLA KV-cache update through a fused concat/cache op |
 | `fuse_attn_quant` | `AttnQuantFusionPass` | attention output followed by FP8 / NVFP4 quant |
 | `fuse_attn_quant` | `MLAAttnQuantFusionPass` | MLA attention output followed by FP8 / NVFP4 / FP8 group quant |
 | `enable_qk_norm_rope_fusion` | `QKNormRoPEFusionPass` | Q/K RMSNorm plus RoPE on packed QKV tensors |
@@ -43,6 +49,7 @@ vLLM registers these passes from
 | `fusion/rocm_aiter_fusion.py` | `MLADualRMSNormPattern` | MLA Q branch and KV branch each run RMSNorm | `torch.ops.vllm.fused_mla_dual_rms_norm` backed by AITER fused QK RMSNorm |
 | `fusion/qk_norm_rope_fusion.py` | `QkNormRopePattern` | Q/K RMSNorm, split/getitem reshapes, then RoPE | `_C.fused_qk_norm_rope` |
 | `fusion/rope_kvcache_fusion.py` | `RopeReshapeKVCachePattern` | RoPE output followed by reshape/cache update | `vllm.fused_rope_and_unified_kv_cache_update` |
+| `fusion/mla_rope_kvcache_cat_fusion.py` | `MLARoPEKVCacheCatPattern` | MLA RoPE on `q_pe` and `k_pe` flows into `unified_mla_kv_cache_update` | `vllm.fused_rope_unified_mla_kv_cache_update`, backed by `concat_and_cache_mla_rope_fused` |
 | `fusion/attn_quant_fusion.py` | `AttnFp8StaticQuantPattern`, `AttnNvfp4QuantPattern` | attention output followed by FP8 static quant or NVFP4 quant | backend attention op with fused output quant when supported |
 | `fusion/mla_attn_quant_fusion.py` | `MLAAttnFp8StaticQuantPattern`, `MLAAttnNvfp4QuantPattern`, `MLAAttnFp8GroupQuantPattern` | MLA attention output followed by static FP8, NVFP4, or FP8 group quant | MLA attention op with fused output quant when supported |
 | `fusion/minimax_qk_norm_fusion.py` | `MiniMaxQKNormPattern` | MiniMax `forward_qk`: Q/K variance all-reduce divided by TP world size, then RMS apply | `vllm.minimax_qk_norm_fused` / Lamport fused kernel |
@@ -56,8 +63,9 @@ vLLM registers these passes from
 - If the trace shows attention output followed by quant kernels, compare against
  `AttnQuantFusionPass` or `MLAAttnQuantFusionPass`, not only handwritten
  attention kernels.
- If the trace shows Q/K norm followed by RoPE or cache update, compare both
-  `QKNormRoPEFusionPass` and `RopeKVCacheFusionPass`; they are separate passes.
+- If the trace shows Q/K norm followed by RoPE or cache update, compare
+  `QKNormRoPEFusionPass`, `RopeKVCacheFusionPass`, and the MLA-specific
+  `MLARoPEKVCacheCatFusionPass`; they are separate passes.
 - If the trace is a TP decode trace with visible collectives, check whether
  `enable_sp` and `fuse_gemm_comms` would transform the same region into
  sequence-parallel or AsyncTP overlap.
@@ -34,8 +34,8 @@ def build_triage_parser() -> argparse.ArgumentParser:
    parser = argparse.ArgumentParser(
        prog="analyze_llm_torch_profile.py",
        description=(
-            "Compact LLM torch-profiler triage entrypoint for SGLang, vLLM, and "
-            "TensorRT-LLM. "
+            "Compact LLM torch-profiler triage entrypoint for SGLang, vLLM, "
+            "TensorRT-LLM, and TokenSpeed. "
            "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."
@@ -45,7 +45,17 @@ def build_triage_parser() -> argparse.ArgumentParser:
        "--framework",
        type=str,
        default="auto",
-        choices=["auto", "sglang", "vllm", "trtllm", "tllm", "tensorrt-llm"],
+        choices=[
+            "auto",
+            "sglang",
+            "vllm",
+            "trtllm",
+            "tllm",
+            "tensorrt-llm",
+            "tokenspeed",
+            "token-speed",
+            "ts",
+        ],
        help=(
            "Serving framework. Use auto to detect from trace contents, path hints, "
            "or URL features."
@@ -64,7 +74,9 @@ def build_triage_parser() -> argparse.ArgumentParser:
        help=(
            "Running server URL for single-trace triage. SGLang supports direct "
            "capture via sglang.profiler. vLLM and TensorRT-LLM require a server-side "
-            "torch-profiler output path exposed via --output-dir."
+            "torch-profiler output path exposed via --output-dir. TokenSpeed live "
+            "capture uses the server's /start_profile and /stop_profile endpoints "
+            "when they are available."
        ),
    )
    parser.add_argument(
@@ -74,7 +86,8 @@ def build_triage_parser() -> argparse.ArgumentParser:
        help=(
            "Trace output dir when using --url. For vLLM this should match the "
            "server's torch_profiler_dir. For TensorRT-LLM it should match the "
-            "directory or file path configured by TLLM_TORCH_PROFILE_TRACE."
+            "directory or file path configured by TLLM_TORCH_PROFILE_TRACE. "
+            "For TokenSpeed this is passed as start_profile.output_dir."
        ),
    )
    parser.add_argument(
@@ -83,7 +96,8 @@ def build_triage_parser() -> argparse.ArgumentParser:
        default="triage-trace",
        help=(
            "Profile prefix when generating a trace from --url. SGLang uses it "
-            "directly; vLLM and TensorRT-LLM may ignore it on the HTTP profiler path."
+            "directly; TokenSpeed maps it to profile_id; vLLM and TensorRT-LLM may "
+            "ignore it on the HTTP profiler path."
        ),
    )
    parser.add_argument(
@@ -1,7 +1,7 @@
 """Backwards-compatibility shim for the unified LLM torch-profiler entrypoint.

 The real implementation now lives in ``analyze_llm_torch_profile`` because this
-skill covers SGLang, vLLM, and TensorRT-LLM. Older scripts and runbooks that
+skill covers SGLang, vLLM, TensorRT-LLM, and TokenSpeed. Older scripts and runbooks that
 still invoke ``analyze_sglang_torch_profile.py`` keep working by forwarding to
 that module.
 """
@@ -31,7 +31,7 @@ def parse_args() -> argparse.Namespace:
    parser.add_argument(
        "--framework",
        required=True,
-        choices=("sglang", "vllm", "trtllm"),
+        choices=("sglang", "vllm", "trtllm", "tokenspeed"),
        help="Serving framework.",
    )
    parser.add_argument(
@@ -42,7 +42,7 @@ def parse_args() -> argparse.Namespace:
    parser.add_argument(
        "--model",
        default=None,
-        help="OpenAI model id. Auto-discovered for vLLM and TensorRT-LLM when omitted.",
+        help="OpenAI model id. Auto-discovered for vLLM, TensorRT-LLM, and TokenSpeed when omitted.",
    )
    parser.add_argument(
        "--requests",
@@ -153,7 +153,7 @@ def openai_request(
 def run_probe(args: argparse.Namespace) -> Dict[str, Any]:
    prompts = args.prompt or list(DEFAULT_PROMPTS)
    model = args.model
-    if args.framework in {"vllm", "trtllm"} and not model:
+    if args.framework in {"vllm", "trtllm", "tokenspeed"} and not model:
        model = discover_openai_model(args.url, timeout=args.timeout)

    latencies: List[float] = []
@@ -22,6 +22,7 @@ FRAMEWORK_LABELS = {
    "sglang": "SGLang",
    "vllm": "vLLM",
    "trtllm": "TensorRT-LLM",
+    "tokenspeed": "TokenSpeed",
 }
 TRACE_FILE_PATTERNS = (
    "*.trace.json",
@@ -89,6 +90,9 @@ def canonicalize_framework(value: object) -> str:
        "trtllm": "trtllm",
        "tensorrt-llm": "trtllm",
        "tensorrtllm": "trtllm",
+        "tokenspeed": "tokenspeed",
+        "token-speed": "tokenspeed",
+        "ts": "tokenspeed",
    }
    return aliases.get(lowered, "auto")

@@ -105,6 +109,8 @@ def _normalize_repo_relative_path_cached(text: str) -> str:
        ("python/sglang/", "python/sglang/"),
        ("sgl_kernel/", "sgl_kernel/"),
        ("vllm/", "vllm/"),
+        ("python/tokenspeed/", "python/tokenspeed/"),
+        ("tokenspeed/", "tokenspeed/"),
        ("tensorrt_llm/", "tensorrt_llm/"),
        ("tensorrt-llm/", "tensorrt_llm/"),
    ):
@@ -312,6 +318,8 @@ def detect_framework_from_text(text: object) -> Optional[str]:
    lowered = normalize_text(text).lower()
    if not lowered:
        return None
+    if any(token in lowered for token in ("tokenspeed", "token-speed", "/ts/")):
+        return "tokenspeed"
    if any(
        token in lowered
        for token in (
@@ -333,6 +341,22 @@ def detect_framework_from_server_args(server_args: Optional[dict]) -> Optional[s
    if not isinstance(server_args, dict) or not server_args:
        return None
    lowered_keys = {normalize_text(key).lower() for key in server_args}
+    text = json.dumps(server_args, sort_keys=True)
+    if any(token in text.lower() for token in ("tokenspeed", "token-speed")):
+        return "tokenspeed"
+    if lowered_keys & {
+        "attn_tp_size",
+        "dense_tp_size",
+        "moe_tp_size",
+        "enable_mla_l1_5_cache",
+        "mla_chunk_multiplier",
+        "comm_fusion_max_num_tokens",
+        "enable_allreduce_fusion",
+    }:
+        return "tokenspeed"
+    text_hint = detect_framework_from_text(text)
+    if text_hint:
+        return text_hint
    if lowered_keys & {
        "attention_backend",
        "sampling_backend",
@@ -342,7 +366,7 @@ def detect_framework_from_server_args(server_args: Optional[dict]) -> Optional[s
        "schedule_policy",
    }:
        return "sglang"
-    return detect_framework_from_text(json.dumps(server_args, sort_keys=True))
+    return None


 def detect_framework_from_trace(trace: object) -> Optional[str]:
@@ -402,6 +426,9 @@ def detect_framework_from_url(
        or "decode" in server_info
    ):
        return "sglang"
+    readiness = try_get_json(url.rstrip("/") + "/readiness", timeout=5.0)
+    if readiness is not None:
+        return "tokenspeed"
    models = try_get_json(url.rstrip("/") + "/v1/models")
    if isinstance(models, dict) and isinstance(models.get("data"), list):
        return "vllm"
@@ -794,9 +821,11 @@ def wait_for_profiler_artifact(path: Path, timeout_s: float = 60.0) -> Path:
    return path


-def start_remote_profiler(url: str, framework: str) -> None:
+def start_remote_profiler(
+    url: str, framework: str, payload: Optional[dict] = None
+) -> None:
    try:
-        post_json(url.rstrip("/") + "/start_profile", timeout=60.0)
+        post_json(url.rstrip("/") + "/start_profile", payload=payload, timeout=60.0)
    except Exception as exc:
        if framework == "vllm":
            raise RuntimeError(
@@ -808,11 +837,40 @@ def start_remote_profiler(url: str, framework: str) -> None:
            raise RuntimeError(
                "TensorRT-LLM live torch profiling requires "
                "a server build that exposes POST /start_profile plus the env vars "
-                "TLLM_PROFILE_START_STOP=1 and TLLM_TORCH_PROFILE_TRACE=/shared/path."
+                "TLLM_PROFILE_START_STOP=<start>-<stop> and "
+                "TLLM_TORCH_PROFILE_TRACE=/shared/path."
+            ) from exc
+        if framework == "tokenspeed":
+            raise RuntimeError(
+                "TokenSpeed live torch profiling requires a server build that "
+                "exposes POST /start_profile and POST /stop_profile. The helper "
+                "passes output_dir, activities, and profile_id in the start payload."
            ) from exc
        raise


+def build_remote_profiler_start_payload(
+    framework: str,
+    output_path: Path,
+    profile_prefix: Optional[str],
+    stage: Optional[str],
+) -> Optional[dict]:
+    if framework != "tokenspeed":
+        return None
+
+    profile_id = profile_prefix or "triage-trace"
+    if stage:
+        profile_id = f"{profile_id}-{stage}"
+
+    return {
+        "output_dir": str(output_path),
+        "activities": ["CPU", "GPU"],
+        "with_stack": True,
+        "record_shapes": False,
+        "profile_id": profile_id,
+    }
+
+
 def stop_remote_profiler(url: str, framework: str) -> None:
    try:
        post_json(url.rstrip("/") + "/stop_profile", timeout=300.0)
@@ -829,6 +887,7 @@ def run_remote_profiler(
    framework: str,
    probe_plan: ProbePlan,
    probe_delay: float,
+    profile_prefix: Optional[str] = None,
    stage: Optional[str] = None,
 ) -> Path:
    framework = canonicalize_framework(framework)
@@ -843,7 +902,11 @@ def run_remote_profiler(
        if output_path.exists()
        else set()
    )
-    model = discover_openai_model(url) if framework in {"vllm", "trtllm"} else None
+    model = (
+        discover_openai_model(url)
+        if framework in {"vllm", "trtllm", "tokenspeed"}
+        else None
+    )
    if probe_plan.warmup_requests > 0:
        send_probe_requests(
            url=url,
@@ -854,13 +917,18 @@ def run_remote_profiler(
            model=model,
        )

-    start_remote_profiler(url, framework)
+    start_payload = build_remote_profiler_start_payload(
+        framework=framework,
+        output_path=output_path,
+        profile_prefix=profile_prefix,
+        stage=stage,
+    )
+    start_remote_profiler(url, framework, payload=start_payload)
    stop_error: Optional[BaseException] = None
    try:
        if probe_plan.capture_requests > 0:
-            # `sglang.profiler` performs its own startup work before it reaches
-            # POST /start_profile. A very short delay can send probes too early
-            # and miss the profiling window entirely.
+            # Server-side profilers may do setup work after POST /start_profile.
+            # A very short delay can send probes too early and miss the window.
            time.sleep(max(5.0, probe_delay))
            send_probe_requests(
                url=url,
@@ -1064,14 +1132,14 @@ def run_profiler(
    if profile_by_stage:
        raise ValueError(
            "--profile-by-stage is only supported for SGLang live capture. "
-            "Disable it when profiling vLLM or TensorRT-LLM."
+            "Disable it when profiling vLLM, TensorRT-LLM, or TokenSpeed."
        )
    if merge_profiles:
        raise ValueError(
            "--merge-profiles is only supported for SGLang live capture. "
-            "Disable it when profiling vLLM or TensorRT-LLM."
+            "Disable it when profiling vLLM, TensorRT-LLM, or TokenSpeed."
        )
-    if profile_prefix:
+    if profile_prefix and resolved_framework in {"vllm", "trtllm"}:
        print(
            f"Note: {framework_display_name(resolved_framework)} ignores "
            "--profile-prefix on the HTTP profiler control path.",
@@ -1093,6 +1161,7 @@ def run_profiler(
                warmup_steps=warmup_steps,
            ),
            probe_delay=probe_delay,
+            profile_prefix=profile_prefix,
        )
    output_root = ensure_remote_profiler_output_path(output_dir, resolved_framework)
    for stage in stages:
@@ -1116,6 +1185,7 @@ def run_profiler(
                warmup_steps=warmup_steps,
            ),
            probe_delay=probe_delay,
+            profile_prefix=profile_prefix,
            stage=stage,
        )
    return output_root
@@ -12,9 +12,10 @@ FRAMEWORK_LABELS = {
    "sglang": "SGLang",
    "vllm": "vLLM",
    "trtllm": "TensorRT-LLM",
+    "tokenspeed": "TokenSpeed",
 }

-FRAMEWORK_ORDER = {"sglang": 0, "vllm": 1, "trtllm": 2}
+FRAMEWORK_ORDER = {"sglang": 0, "vllm": 1, "trtllm": 2, "tokenspeed": 3}


 def parse_args(argv: Optional[Sequence[str]] = None) -> argparse.Namespace:
@@ -75,6 +76,8 @@ def framework_key_from_path(path: Path) -> str:
        return "vllm"
    if "trtllm" in lowered or "tensorrt" in lowered:
        return "trtllm"
+    if "tokenspeed" in lowered or "token-speed" in lowered:
+        return "tokenspeed"
    return "other"


@@ -530,16 +530,16 @@ FUSION_PATTERN_REGISTRY: Tuple[FusionPatternSpec, ...] = (
        likely_share=0.5,
    ),
    FusionPatternSpec(
-        pattern="DSA fused metadata copy for graph replay",
+        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_dsa_cache_seqlens",
+            "fused_nsa_cache_seqlens",
            "fused_flashmla_metadata",
        ),
        rationale_hint=(
-            "DSA replay metadata copies are already fused into one-kernel" " families."
+            "NSA replay metadata copies are already fused into one-kernel" " families."
        ),
        min_share=0.02,
        likely_share=0.2,
@@ -744,23 +744,23 @@ FUSION_PATTERN_REGISTRY: Tuple[FusionPatternSpec, ...] = (
        likely_share=1.5,
    ),
    FusionPatternSpec(
-        pattern="DSA fused top-k transform / page-table build",
-        candidate_path="python/sglang/srt/layers/attention/dsa_backend.py",
+        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=(
-            "DSA top-k metadata preparation already has fused transform kernels."
+            "NSA top-k metadata preparation already has fused transform kernels."
        ),
        min_share=0.05,
        likely_share=0.3,
    ),
    FusionPatternSpec(
-        pattern="DSA fused quantize + indexed K-cache store",
+        pattern="NSA fused quantize + indexed K-cache store",
        candidate_path=(
            "python/sglang/jit_kernel/fused_store_index_cache.py"
-            "<br>python/sglang/srt/layers/attention/dsa/dsa_indexer.py"
+            "<br>python/sglang/srt/layers/attention/nsa/nsa_indexer.py"
        ),
        active_keywords=("fused_store_index_k_cache",),
        split_groups=(
@@ -768,7 +768,7 @@ FUSION_PATTERN_REGISTRY: Tuple[FusionPatternSpec, ...] = (
            ("index_k", "cache", "store"),
        ),
        rationale_hint=(
-            "DSA already has a fused quantize-and-indexed-store kernel family."
+            "NSA already has a fused quantize-and-indexed-store kernel family."
        ),
        min_share=0.2,
        likely_share=1.0,
@@ -855,6 +855,180 @@ FUSION_PATTERN_REGISTRY: Tuple[FusionPatternSpec, ...] = (
        likely_share=1.0,
        priority=90,
    ),
+    FusionPatternSpec(
+        pattern="SGLang LTX2 fused Ada values",
+        candidate_path=(
+            "PR #29390"
+            "<br>python/sglang/jit_kernel/diffusion/triton/ltx2_ada_values.py"
+            "<br>python/sglang/multimodal_gen/runtime/models/dits/ltx_2.py"
+        ),
+        active_keywords=(
+            "ltx2_ada_values9",
+            "ltx2_ada_values",
+            "LTX2TransformerBlock",
+        ),
+        split_groups=(
+            ("scale_shift_table", "timestep", "reshape"),
+            ("get_ada_values", "ada", "adaln"),
+            ("slice", "split", "unbind"),
+        ),
+        rationale_hint=(
+            "SGLang mainline fuses LTX-2.3 Ada value materialization for"
+            " video/audio streams; split Ada table add/reshape/slice ladders"
+            " should be checked against this diffusion Triton kernel first."
+        ),
+        origin="upstream",
+        model_include=("ltx", "ltx-2", "ltx2"),
+        min_share=0.2,
+        likely_share=1.0,
+    ),
+    FusionPatternSpec(
+        pattern="SGLang LTX2 residual-gate add CUDA fast path",
+        candidate_path=(
+            "PR #29361"
+            "<br>python/sglang/jit_kernel/diffusion/residual_gate_add.py"
+            "<br>python/sglang/jit_kernel/csrc/diffusion/residual_gate_add.cuh"
+            "<br>python/sglang/multimodal_gen/runtime/models/dits/ltx_2.py"
+        ),
+        active_keywords=(
+            "diffusion_residual_gate_add",
+            "residual_gate_add",
+            "_ltx2_residual_gate_add",
+        ),
+        split_groups=(
+            ("add", "mul", "gate"),
+            ("residual", "update", "gate"),
+            ("hidden_states", "attn_hidden_states", "gate"),
+        ),
+        rationale_hint=(
+            "SGLang mainline fuses LTX2 residual + update * gate sites into"
+            " a CUDA custom op; split add/mul gate ladders should be checked"
+            " against this path before proposing a new diffusion elementwise"
+            " fusion."
+        ),
+        origin="upstream",
+        model_include=("ltx", "ltx-2", "ltx2"),
+        min_share=0.2,
+        likely_share=1.0,
+    ),
+    FusionPatternSpec(
+        pattern="TokenSpeed CuTe DSL MLA prefill / decode",
+        candidate_path=(
+            "python/tokenspeed/runtime/layers/attention/backends/tokenspeed_mla.py"
+            "<br>tokenspeed-mla/python/tokenspeed_mla/mla_decode.py"
+            "<br>tokenspeed-mla/python/tokenspeed_mla/mla_prefill.py"
+            "<br>tokenspeed-kernel/python/tokenspeed_kernel/ops/attention/"
+            "tokenspeed_mla/__init__.py"
+        ),
+        active_keywords=(
+            "tokenspeed_mla_decode",
+            "tokenspeed_mla_prefill",
+            "BlackwellMultiHeadLatentAttentionForward",
+        ),
+        split_groups=(
+            ("mla", "flashmla", "attention", "fmha"),
+            ("prefill", "decode", "verify"),
+            ("fp8", "kv_cache", "page_table"),
+        ),
+        rationale_hint=(
+            "TokenSpeed ships Blackwell CuTe DSL MLA prefill/decode kernels;"
+            " split MLA support kernels should be checked against backend"
+            " selection before being called novel."
+        ),
+        origin="upstream",
+        model_include=("deepseek", "kimi", "qwen3.5", "qwen3_5"),
+        min_share=0.4,
+        likely_share=2.0,
+    ),
+    FusionPatternSpec(
+        pattern="TokenSpeed MLA KV pack + FP8 quantize",
+        candidate_path=(
+            "tokenspeed-mla/python/tokenspeed_mla/mla_kv_pack_quantize_fp8.py"
+            "<br>tokenspeed-kernel/python/tokenspeed_kernel/ops/attention/"
+            "tokenspeed_mla/__init__.py"
+        ),
+        active_keywords=(
+            "_mla_kv_pack_quantize_fp8_kernel",
+            "mla_kv_pack_quantize_fp8",
+        ),
+        split_groups=(
+            ("k_nope", "k_pe", "cat", "concat", "pack"),
+            ("quant", "fp8", "float8"),
+            ("v", "kv", "cache"),
+        ),
+        rationale_hint=(
+            "TokenSpeed fuses MLA K/V pack, concat, and FP8 quantization into"
+            " one Triton kernel for chunked prefill."
+        ),
+        origin="upstream",
+        model_include=("deepseek", "kimi", "qwen3.5", "qwen3_5"),
+        min_share=0.2,
+        likely_share=1.0,
+    ),
+    FusionPatternSpec(
+        pattern="TokenSpeed fused top-k + top-p sampling",
+        candidate_path=(
+            "tokenspeed-kernel/python/tokenspeed_kernel/thirdparty/cuda/"  # codespell:ignore thirdparty
+            "fused_topk_topp.py"
+            "<br>tokenspeed-kernel/python/tokenspeed_kernel/thirdparty/cuda/"  # codespell:ignore thirdparty
+            "csrc/fused_topk_topp/fused_topk_topp.cu"
+        ),
+        active_keywords=("fused_topk_topp", "fused_topk_topp_renorm"),
+        split_groups=(
+            ("topk", "top_k"),
+            ("topp", "top_p"),
+            ("sampling", "renorm", "softmax"),
+        ),
+        rationale_hint=(
+            "TokenSpeed has a fused top-k/top-p renormalization path for"
+            " decode sampling."
+        ),
+        origin="upstream",
+        min_share=0.1,
+        likely_share=0.8,
+    ),
+    FusionPatternSpec(
+        pattern="TokenSpeed persistent lm_head GEMM",
+        candidate_path=(
+            "tokenspeed-kernel/python/tokenspeed_kernel/thirdparty/cuda/"  # codespell:ignore thirdparty
+            "lm_head_gemm.py"
+            "<br>tokenspeed-kernel/python/tokenspeed_kernel/thirdparty/cuda/"  # codespell:ignore thirdparty
+            "csrc/lm_head_gemm.cu"
+        ),
+        active_keywords=("lm_head_gemm",),
+        split_groups=(
+            ("lm_head", "logits", "vocab"),
+            ("gemm", "matmul", "linear"),
+        ),
+        rationale_hint=(
+            "TokenSpeed has a shape-gated persistent lm_head GEMM path; visible"
+            " lm_head matmul ladders should be compared against it."
+        ),
+        origin="upstream",
+        model_include=("kimi", "qwen"),
+        min_share=0.2,
+        likely_share=1.0,
+    ),
+    FusionPatternSpec(
+        pattern="TokenSpeed NVFP4 GEMM + SwiGLU + quant",
+        candidate_path=(
+            "tokenspeed-kernel/python/tokenspeed_kernel/thirdparty/cute_dsl/"  # codespell:ignore thirdparty
+            "nvfp4_gemm_swiglu_nvfp4_quant.py"
+        ),
+        active_keywords=("nvfp4_gemm_swiglu_nvfp4_quant",),
+        split_groups=(
+            ("gemm", "nvfp4", "fp4"),
+            ("swiglu", "silu", "activation", "mul"),
+            ("quant", "scale", "sfc"),
+        ),
+        rationale_hint=(
+            "TokenSpeed's CuTe DSL kernel fuses NVFP4 GEMM, SwiGLU, and"
+            " optional output quantization in one expert-style path."
+        ),
+        origin="upstream",
+        min_share=0.3,
+        likely_share=1.5,
+    ),
    FusionPatternSpec(
        pattern="vLLM-origin Attention + Quantization",
        candidate_path=(
@@ -1235,6 +1409,8 @@ def source_location_priority(location: str) -> int:
        return 290 - penalty
    if text.startswith("vllm/"):
        return 285 - penalty
+    if text.startswith("python/tokenspeed/") or text.startswith("tokenspeed/"):
+        return 283 - penalty
    if text.startswith("tensorrt_llm/"):
        return 280 - penalty
    if text.startswith("sgl_kernel/"):
@@ -1254,6 +1430,8 @@ def is_preferred_source_location(location: str) -> bool:
        text.startswith("python/sglang/")
        or text.startswith("sglang/")
        or text.startswith("vllm/")
+        or text.startswith("python/tokenspeed/")
+        or text.startswith("tokenspeed/")
        or text.startswith("tensorrt_llm/")
        or text.startswith("sgl_kernel/")
    )
@@ -1316,6 +1494,10 @@ def frame_priority(frame_name: str) -> int:
        return 290 - penalty
    if normalized_text.startswith("vllm/"):
        return 285 - penalty
+    if normalized_text.startswith("python/tokenspeed/") or normalized_text.startswith(
+        "tokenspeed/"
+    ):
+        return 283 - penalty
    if normalized_text.startswith("tensorrt_llm/"):
        return 280 - penalty
    if normalized_text.startswith("sgl_kernel/"):
@@ -1329,6 +1511,8 @@ def frame_priority(frame_name: str) -> int:
            return 120
        if "/vllm/" in raw_text:
            return 118
+        if "/tokenspeed/" in raw_text or "/TokenSpeed/" in raw_text:
+            return 117
        if "/TensorRT-LLM/" in raw_text or "/tensorrt_llm/" in raw_text:
            return 116
        return 100
@@ -1336,6 +1520,10 @@ def frame_priority(frame_name: str) -> int:
        return 110
    if ".py(" in raw_text and "/vllm/" in raw_text:
        return 108
+    if ".py(" in raw_text and (
+        "/tokenspeed/" in raw_text or "/TokenSpeed/" in raw_text
+    ):
+        return 107
    if ".py(" in raw_text and (
        "/TensorRT-LLM/" in raw_text or "/tensorrt_llm/" in raw_text
    ):
@@ -2438,6 +2626,8 @@ def fusion_framework_hints(spec: FusionPatternSpec) -> set[str]:
    hints: set[str] = set()
    if "vllm/" in text:
        hints.add("vllm")
+    if any(token in text for token in ("tokenspeed/", "tokenspeed-", "tokenspeed_")):
+        hints.add("tokenspeed")
    if "tensorrt_llm/" in text:
        hints.add("trtllm")
    if any(token in text for token in ("python/sglang/", "sgl-kernel/", "sgl_kernel/")):
@@ -416,6 +416,10 @@ def is_meaningful_python_scope(name: str) -> bool:
        return True
    if normalized.startswith("vllm/"):
        return True
+    if normalized.startswith("python/tokenspeed/") or normalized.startswith(
+        "tokenspeed/"
+    ):
+        return True
    if normalized.startswith("tensorrt_llm/"):
        return True
    if normalized.startswith("sgl_kernel/"):
@@ -697,6 +701,8 @@ def choose_best_scope(scope_chain: Sequence[str]) -> Optional[str]:
            score += 48.0
        elif scope.startswith("vllm/"):
            score += 46.0
+        elif scope.startswith("python/tokenspeed/") or scope.startswith("tokenspeed/"):
+            score += 45.0
        elif scope.startswith("tensorrt_llm/"):
            score += 44.0
        elif scope.startswith("sgl_kernel/"):
@@ -744,6 +750,10 @@ def source_scope_priority(scope: Optional[str]) -> int:
        return 290 - penalty
    if normalized.startswith("vllm/"):
        return 285 - penalty
+    if normalized.startswith("python/tokenspeed/") or normalized.startswith(
+        "tokenspeed/"
+    ):
+        return 283 - penalty
    if normalized.startswith("tensorrt_llm/"):
        return 280 - penalty
    if normalized.startswith("sgl_kernel/"):