# Fuse And Overlap Catalog

This catalog is the source-backed lookup table that the profiler skill should
consult before labeling a fuse or overlap opportunity as novel.

For overlap-only triage, also load `references/overlap-catalog.md`.

This revision is intentionally kernel-scoped. Keep rows here only when they map
to one fused GPU/NPU kernel family, one fused collective-plus-kernel family, or
one profiler-visible stream overlap among GPU kernels / collective kernels.
Host-only scheduler, event-loop, executor, offload, and load-path patterns are
intentionally excluded.

Use it like this:

1. Start from the three `triage` tables.
2. Match top rows against the `Trace keywords` and `Primary code` columns below.
3. If a finding matches an existing row, report it as:
   - an existing optimization path that is missing, disabled, regressed, or unsupported for the current backend, or
   - an already-known family that should be re-applied to the current model shape.
4. Check the mainline comparison sections and the `PR-backed / in-flight` sections too. If a match exists there, do not call it novel; call it an upstream or in-flight pattern instead.
5. Only call a finding "new" when it does not fit any mainline or PR-backed row in this catalog.

The `vLLM-origin` sections below are comparative references. They are not
necessarily present in the checked-out `sglang` tree, but they should still be
treated as upstream or analogous kernel families before labeling a fuse or
overlap opportunity as novel.

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

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. 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

| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
| Fused residual add + RMSNorm | `fused_add_rmsnorm*`<br>`npu_add_rms_norm`<br>`add_rmsnorm_bias`<br>`gemma_fused_add_rmsnorm`<br>`gemma_rmsnorm_residual_scalar`<br>`_gemma_rmsnorm_residual_kernel`<br>residual add right before norm | `python/sglang/srt/layers/layernorm.py`<br>`python/sglang/srt/layers/gemma4_fused_ops.py`<br>`python/sglang/srt/layers/quantization/modelslim/modelslim.py` | Shared CUDA / ROCm / CPU / NPU fused add-RMSNorm implementations, including Gemma, Gemma4 scalar-residual, and NPU-bias variants | Treat split residual add + RMSNorm as an existing cross-backend fusion first, not a new idea. |
| FlashInfer unified `allreduce_fusion` | `cross_device_reduce_1stage*`<br>`all_reduce`<br>`FusedAddRMSNormKernel`<br>`rmsnorm*` | `python/sglang/srt/layers/flashinfer_comm_fusion.py`<br>`python/sglang/srt/layers/layernorm.py::forward_with_allreduce_fusion`<br>`python/sglang/srt/layers/communicator.py::apply_flashinfer_allreduce_fusion` | FlashInfer workspace creation plus `allreduce_fusion(..., pattern=AllReduceFusionPattern.kARResidualRMSNorm, ...)` | First suspect missing / disabled / unsupported FlashInfer allreduce fusion, not a brand new TP fusion idea. |
| AITER allreduce fusion | ROCm all-reduce plus RMSNorm still split | `python/sglang/srt/layers/layernorm.py::forward_with_allreduce_fusion`<br>`python/sglang/srt/distributed/communication_op.py::tensor_model_parallel_fused_allreduce_rmsnorm`<br>`python/sglang/srt/layers/communicator.py::apply_aiter_all_reduce_fusion` | ROCm-side fused TP all-reduce + RMSNorm with fallback to plain all-reduce plus norm | On AMD, rule out existing AITER fusion before proposing a new communication fusion. |
| Fused activation-and-mul (`SwiGLU` / `GeGLU`) | `silu_and_mul`<br>`gelu_and_mul`<br>`npu_swiglu` | `python/sglang/srt/layers/activation.py` | Single op covers activation plus elementwise multiply across CUDA / CPU / NPU / XPU backends | Treat separate activation + mul on packed MLP outputs as missing existing fusion. |
| Fused dual residual RMSNorm | residual add plus two RMSNorm-like kernels around Grok blocks | `python/sglang/srt/layers/elementwise.py::fused_dual_residual_rmsnorm`<br>`python/sglang/srt/models/grok.py` | One Triton kernel computes intermediate residual update and next RMSNorm output together | On Grok-like residual layouts, treat split residual + norm as missing existing fusion. |
| In-place QK RMSNorm | split `q_norm` / `k_norm` kernels | `python/sglang/srt/models/utils.py::apply_qk_norm`<br>`python/sglang/kernels/ops/layernorm/norm.py::fused_inplace_qknorm` | In-place JIT QK norm plus optional `alt_stream` overlap for K | Check shape, dtype, deterministic mode, and in-place legality before proposing a new QK fuse. |
| TorchInductor horizontal Q/K norm combo-kernels | `combo_kernels`<br>`benchmark_combo_kernel`<br>`q_norm`<br>`k_norm`<br>`split_with_sizes` | `torch._inductor.config.combo_kernels` | TorchInductor can horizontally fuse sibling Q-norm and K-norm kernels in compiled traces, often deleting `split_with_sizes` / `clone` ladders | Treat separate Q/K norm ladders in compile-heavy traces as an existing compiler-fusion family first. |
| MiniMax TP fused QK RMSNorm | `MiniMaxM2RMSNormTP`<br>`rms_sumsq_serial`<br>`rms_apply_serial`<br>`forward_qk` | `python/sglang/srt/models/minimax_m2.py` | Triton kernels compute Q / K sumsq together, TP all-reduces shared stats, then apply both RMSNorms together | On MiniMax traces, separate Q norm and K norm are usually a missed model-specific Triton fusion. |
| Fused QK RMSNorm + RoPE | `qknorm*` + `rope*` + `rotary*` as separate steps | `python/sglang/kernels/ops/attention/fused_qknorm_rope.py`<br>`python/sglang/srt/models/qwen3_moe.py` | One JIT kernel applies QK RMSNorm and RoPE in-place on packed QKV | For compatible LLMs, classify split QK norm + RoPE as a missing existing fusion. |
| 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/kernels/ops/attention/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. |
| 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/kernels/ops/attention/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 / 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. |
| Fused MoE dispatch / permute / combine | token permutation<br>dispatch / combine<br>grouped top-k<br>many small MoE support kernels | `python/sglang/srt/layers/moe/fused_moe_triton/layer.py`<br>`python/sglang/srt/layers/moe/fused_moe_triton/fused_moe.py` | `FusedMoE` plus DeepEP / FlashInfer / FuseEP / standard dispatch backends and `permute_fusion=True` | First ask whether the model is missing an existing `FusedMoE`-style path or backend-specific dispatcher path. |
| 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/kernels/ops/quantization/nvfp4_gemm_swiglu_nvfp4_quant.py`<br>`python/sglang/srt/layers/moe/cutlass_w4a8_moe.py`<br>`python/sglang/srt/hardware_backend/npu/quantization/moe_methods.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. |
| 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/kernels/ops/attention/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/kernels/ops/attention/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. |
| Fused GDN gating prep | `fused_gdn_gating`<br>`softplus`<br>`beta_output` | `python/sglang/srt/layers/attention/fla/fused_gdn_gating.py` | Triton kernel computes GDN gate preparation such as `-exp(A_log) * softplus(...)` and `sigmoid(b)` together | On GDN traces, treat split gate-prep elementwise kernels as missing existing fusion first. |
| Fused RMSNorm-gated linear-attention output | `FusedRMSNormGated`<br>`layer_norm_gated_fwd` | `python/sglang/srt/layers/attention/fla/fused_norm_gate.py`<br>`python/sglang/srt/models/qwen3_next.py`<br>`python/sglang/srt/models/kimi_linear.py` | One Triton op covers residual-aware (RMS)Norm plus sigmoid / swish gating | If norm and output gate appear as separate kernels in GDN / Kimi-like blocks, first suspect a missing existing fusion. |
| Fused gated RMSNorm / LayerNorm | `rms_norm_gated`<br>`layer_norm_gated` | `python/sglang/srt/layers/attention/mamba/ops/layernorm_gated.py` | Mamba-derived kernels can fuse normalization with the gating branch `z * sigmoid(z)` | Treat split norm and gate post-processing on Mamba-style blocks as an existing fusion family. |
| Fused linear-attention chunk KKT + solve_tril | `chunk_gated_delta_rule_fwd_kkt_solve_kernel`<br>`scaled_dot_kkt`<br>`solve_tril`<br>`recompute_w_u` | `python/sglang/srt/layers/attention/fla/chunk_fwd.py`<br>`python/sglang/srt/layers/attention/fla/kda.py` | GDN / KDA chunk forward fuses `scaled_dot_kkt + solve_tril` in the prefill / intra-chunk path, then finishes `recompute_w_u` as the next step | Treat split KKT + triangular-solve ladders as an existing linear-attention fusion family first. |
| Fused linear-attention recurrent / KDA update | `fused_sigmoid_gating_delta_rule_update`<br>`fused_recurrent_gated_delta_rule_update`<br>`fused_kda_gate` | `python/sglang/srt/layers/attention/fla/fused_sigmoid_gating_recurrent.py`<br>`python/sglang/srt/layers/attention/fla/fused_recurrent.py`<br>`python/sglang/srt/models/kimi_linear.py`<br>`python/sglang/srt/models/jet_nemotron.py` | Triton / CuTeDSL kernels fuse gating math, optional QK l2norm, recurrent state update, and output generation | Treat split gating + recurrent-update chains as existing linear-attention fusion, not a novel opportunity. |
| Fused Mamba state gather/scatter with mask | `fused_mamba_state_scatter_with_mask`<br>`index_elementwise_kernel` | `python/sglang/srt/layers/attention/mamba/mamba_state_scatter_triton.py` | Triton kernel replaces multiple masked gather / scatter index kernels with one fused update | If Mamba verify/update shows many tiny index kernels, first compare against this existing fused path. |
| Staging-buffer fused gather / scatter | `_fused_gather_to_staging_kernel`<br>`_fused_scatter_from_staging_kernel` | `python/sglang/srt/disaggregation/common/staging_buffer.py` | Triton kernels gather scattered KV slices into contiguous staging memory and scatter them back into KV cache on decode | Treat ladders of tiny gather/scatter/copy kernels in heterogeneous TP staging as missing an existing Triton fusion. |

## 2. LLM / SRT kernel-overlap families

| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
| Single-batch overlap (SBO) | MoE combine, down-gemm, shared-expert work in nearby two-stream windows | `python/sglang/srt/batch_overlap/single_batch_overlap.py` | combine vs down-gemm overlap, combine vs shared-expert overlap, one-stream dispatch+shared overlap, explicit SM partitioning and events | If exposed MoE combine sits near neighboring compute, classify it against SBO before calling it new overlap. |
| Q and K normalization on different streams | Q-side norm and K-side norm on different streams | `python/sglang/srt/models/utils.py::apply_qk_norm`<br>`python/sglang/srt/models/qwen3.py`<br>`python/sglang/srt/models/qwen3_next.py`<br>`python/sglang/srt/models/qwen3_5.py` | Q stays on current stream, K can run on `alt_stream` in capture mode | Treat split Q / K norm as an existing overlap family when `alt_stream` is already wired. |
| DeepSeek shared-expert / routed-expert overlap | shared-expert GEMMs near DeepEP dispatch / combine | `python/sglang/srt/models/deepseek_v2.py`<br>`python/sglang/srt/batch_overlap/single_batch_overlap.py` | shared experts on `alt_stream`, overlap with dispatch / combine and down-gemm, Blackwell-specific env gating | This is an established routed-vs-shared branch overlap pattern, not a novel idea. |
| 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. |
| 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. |

## 3. VLM-specific kernel families

| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
| Vision QK norm with aux stream | vision-side QK norm or norm-like kernels before attention | `python/sglang/srt/layers/attention/vision.py` | vision QK normalization can call shared `apply_qk_norm(...)`, with K-side work on `aux_stream` | If vision QK prep is split, first check this existing aux-stream path. |
| ViT CUDA graph disables vision aux stream | expected vision overlap is absent under ViT graph | `python/sglang/srt/models/internvl.py`<br>`python/sglang/srt/layers/attention/vision.py`<br>`python/sglang/srt/environ.py::SGLANG_VIT_ENABLE_CUDA_GRAPH` | vision `aux_stream` is intentionally disabled when ViT CUDA graph is on | Missing vision overlap may be intentional, not a regression. |
| Fused multimodal RoPE kernel | `triton_mrope_fused`<br>`multimodal_rotary_embedding_cpu`<br>`npu_mrope`<br>`MRotaryEmbedding` | `python/sglang/srt/layers/rotary_embedding/mrope.py`<br>`python/sglang/srt/layers/rotary_embedding/triton_kernels.py`<br>`python/sglang/srt/models/qwen3.py` | CUDA Triton, CPU `sgl_kernel`, and NPU paths already fuse multimodal t / h / w position lookup plus in-place Q / K rotary application | If VLM traces show separate mRoPE gather / shuffle / apply steps, first classify them as a missing existing mRoPE fusion. |

## 4. Diffusion fused-kernel families

| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
| Fused residual + norm + scale + shift | residual add, norm, scale, shift, gate around DiT blocks | `python/sglang/kernels/ops/diffusion/cutedsl/scale_residual_norm_scale_shift.py`<br>`python/sglang/multimodal_gen/runtime/layers/layernorm.py` | `fused_scale_residual_norm_scale_shift(...)` | Treat split residual + norm + modulation as a missing existing diffusion fusion first. |
| Fused norm + scale + shift | norm followed by scale / shift elementwise kernels | `python/sglang/kernels/ops/diffusion/cutedsl/scale_residual_norm_scale_shift.py`<br>`python/sglang/multimodal_gen/runtime/layers/layernorm.py` | `fused_norm_scale_shift(...)` | Existing modulation fusion already covers this family. |
| Triton scale / shift and gate-select kernels | tiny scale / shift or gate-select kernels dominate modulation blocks | `python/sglang/kernels/ops/diffusion/triton/scale_shift.py`<br>`python/sglang/multimodal_gen/runtime/layers/elementwise.py` | `fuse_scale_shift_kernel(...)` and `fuse_layernorm_scale_shift_gate_select01_kernel(...)` | Check whether the runtime is missing these existing Triton fusions. |
| Fused add-RMSNorm and one-pass RMSNorm | residual add plus RMSNorm still split on short hidden sizes | `python/sglang/multimodal_gen/runtime/layers/layernorm.py`<br>`python/sglang/kernels/ops/diffusion/triton/rmsnorm_onepass.py` | `fused_add_rmsnorm(...)` and `triton_one_pass_rms_norm(...)` | For short hidden-size diffusion blocks, this is already an established fusion family. |
| Fused diffusion QK norm + RoPE | split QK norm and RoPE in diffusion attention blocks | `python/sglang/kernels/ops/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/kernels/ops/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/kernels/ops/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/kernels/ops/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/kernels/ops/diffusion/residual_gate_add.py`<br>`python/sglang/kernels/jit/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

| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
| Ulysses sequence-parallel attention | exposed `all_to_all` around attention blocks | `python/sglang/multimodal_gen/runtime/layers/attention/layer.py`<br>`python/sglang/multimodal_gen/runtime/distributed/communication_op.py` | head / sequence redistribution before and after attention | Treat sequence-parallel all-to-all as an existing distributed attention family. |
| USP attention with all-to-all and ring attention | `all_to_all`, ring-attention comm, head / sequence reshards | `python/sglang/multimodal_gen/runtime/layers/attention/layer.py` | `_usp_input_all_to_all(...)`, `_usp_output_all_to_all(...)`, `ring_attn(...)` | This is the primary existing overlap / comm family for many diffusion models. |
| Turbo-layer async all-to-all pipelining | pipelined A2A windows with explicit waits on a comm stream | `python/sglang/multimodal_gen/runtime/layers/attention/turbo_layer.py` | looped `all_to_all_single(..., async_op=True)` plus staged postprocess on a comm stream | Treat exposed turbo A2A windows as an existing pipelined overlap pattern. |
| TorchInductor compute / communication reorder | compiled traces with compute and comm partially interleaved | `python/sglang/multimodal_gen/runtime/pipelines_core/stages/denoising.py`<br>`python/sglang/multimodal_gen/runtime/pipelines_core/stages/model_specific_stages/mova.py` | `torch._inductor.config.reorder_for_compute_comm_overlap = True` | Existing compile-time reordering may already explain partial overlap in diffusion traces. |
| Dual-stream diffusion models | two nearby compute branches inside one DiT / UNet block | `python/sglang/multimodal_gen/runtime/models/dits/hunyuan3d.py` | `use_dual_stream = True` | Treat dual-branch diffusion execution as an existing overlap family. |

## 6. PR-backed / in-flight fused-kernel families

These rows track still-open upstream work or status-sensitive PR families.
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 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/kernels/jit/csrc/elementwise/fused_add_rmsnorm_per_token_quant.cuh`<br>`python/sglang/kernels/jit/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>`python/sglang/kernels/aot/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/kernels/ops/quantization/nvfp4_gemm_swiglu_nvfp4_quant.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` 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. |

## 7. PR-backed / in-flight kernel-overlap families

| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
| PR `#21877` fused down-GEMM + combine superseding SBO | `enable_fused_grouped_gemm_combine`<br>`combine`<br>`down_gemm` | `PR #21877`<br>`python/sglang/srt/server_args.py`<br>`python/sglang/srt/layers/moe/token_dispatcher/deepep.py` | Fused combine eliminates the standalone combine window, so SBO is intentionally disabled when this path is on | If the trace discussion is about combine overlap, first classify it as this upstream fused-overlap family. |
| PR `#23965` PDL for DSV32 / GLM5 kernels | `enable_pdl`<br>`TRTLLM_ENABLE_PDL`<br>`cudaGridDependencySynchronize`<br>`DSV32`<br>`GLM5` | `PR #23965`<br>`python/sglang/srt/layers`<br>`sgl-kernel` | Enables programmatic dependent launch on selected DeepSeek / GLM kernels so dependent decode kernels can overlap launch-to-start gaps | Treat tight same-stream decode windows around DSV32 / GLM5 as an in-flight PDL overlap family. |
| PR `#21878` TTFT / TPOT torch.compile optimization | `enable_torch_compile`<br>`decode graph`<br>`piecewise cudagraph` | `PR #21878`<br>`python/sglang/srt` | Uses compiler and graph capture changes to shave TTFT / TPOT rather than adding one handwritten kernel | If the trace shows many small compiler-visible decode ops, compare against this compile-overlap / graph-capture family first. |
| PR `#24168` batched GPU-to-CPU sync for logprobs / embeddings | `logprobs`<br>`embeddings`<br>`GPU->CPU sync`<br>`batch sync` | `PR #24168`<br>`python/sglang/srt` | Batches per-request synchronization work that can otherwise serialize decode progress around logprob or embedding outputs | Treat per-request CPU sync stalls in logprob / embedding traces as a concrete in-flight SGLang scheduler/data-movement family. |

## 8. FlashInfer mainline fused-kernel families

These rows are comparative references from `flashinfer`. Use them when a trace
looks like an upstream FlashInfer family even if the current `sglang` checkout
only consumes a subset of that implementation.

| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
| FlashInfer activation / gate epilogues | `silu_and_mul`<br>`gelu_tanh_and_mul`<br>`gelu_and_mul`<br>`silu_and_mul_scaled_nvfp4_experts_quantize` | `flashinfer/activation.py`<br>`flashinfer/quantization/fp4_quantization.py` | FlashInfer covers both the plain activation-plus-mul epilogues and the NVFP4 expert-quantized extension used on MoE expert paths | Treat standalone activation, multiply, and expert-side quant ladders as one existing FlashInfer epilogue family first. |
| FlashInfer norm / residual / quant epilogues | `rmsnorm_quant`<br>`fused_add_rmsnorm`<br>`fused_add_rmsnorm_quant`<br>`gemma_rmsnorm`<br>`gemma_fused_add_rmsnorm`<br>`fused_rmsnorm_silu`<br>`rmsnorm_fp4quant`<br>`add_rmsnorm_fp4quant` | `flashinfer/norm/__init__.py`<br>`flashinfer/cute_dsl/rmsnorm_fp4quant.py`<br>`flashinfer/cute_dsl/add_rmsnorm_fp4quant.py` | The norm family spans plain RMSNorm derivatives, residual-add epilogues, norm+activation, and direct FP8 / NVFP4 output variants instead of materializing each intermediate | Treat split residual add, norm, activation, and quant chains as one existing FlashInfer epilogue family first. |
| FlashInfer allreduce + post-op fusion family | `allreduce_fusion`<br>`AllReduceFusionPattern`<br>`kARResidualRMSNorm`<br>`kARResidualRMSNormFP8Quant`<br>`kARResidualRMSNormFP4Quant`<br>`trtllm_mnnvl_allreduce_fusion` | `flashinfer/comm/allreduce.py`<br>`flashinfer/comm/trtllm_ar.py`<br>`flashinfer/comm/trtllm_mnnvl_ar.py` | TRTLLM and MNNVL backends fuse all-reduce with residual add, RMSNorm, and backend-appropriate quant / norm-output variants | Treat TP collective + norm (+ quant) ladders as an existing FlashInfer fused-collective family first. |
| FlashInfer RoPE + FP8 quant / cache-update family | `rope_quantize_fp8`<br>`mla_rope_quantize_fp8`<br>`rope_quantize_fp8_append_paged_kv_cache`<br>`seqlen=0`<br>`batch_indices < 0` | `flashinfer/rope.py` | The RoPE family covers both RoPE+FP8 output and the larger decode / prefill-prep path that writes K / V directly into paged KV cache, including padding-token / zero-length sequence handling | Treat split RoPE, quant, cache-write, and padding-token ladders as one existing FlashInfer attention-prep family first. |
| FlashInfer fused DeepSeek grouped-topk routing | `fused_topk_deepseek`<br>`NoAuxTc` | `flashinfer/fused_moe/fused_routing_dsv3.py` | One kernel performs sigmoid+bias, grouped score reduction, group top-k, expert top-k, and routed renorm for DeepSeek-V3-style routing | Treat router score activation -> grouped top-k -> renorm ladders as an existing FlashInfer router family first. |
| FlashInfer fused MoE expert execution | `cutlass_fused_moe`<br>`trtllm_bf16_moe`<br>`trtllm_fp8_per_tensor_scale_moe`<br>`trtllm_fp8_block_scale_moe`<br>`trtllm_fp4_block_scale_moe`<br>`trtllm_mxint4_block_scale_moe`<br>`non-gated` | `flashinfer/fused_moe/core.py` | CUTLASS and TRTLLM backends collapse expert execution, routed combine, and quantized expert variants into fused MoE runners, including gated and non-gated FP8 per-tensor cases | Treat exposed expert-side tiny GEMM or non-gated FP8 ladders as matching an existing FlashInfer fused-MoE family. |
| FlashInfer CuTeDSL two-stage MoE fusion | `blockscaled_contiguous_gather_grouped_gemm_swiglu_fusion_nvfp4`<br>`blockscaled_contiguous_grouped_gemm_finalize_fusion_nvfp4`<br>`moe_permute`<br>`moe_unpermute` | `flashinfer/fused_moe/cute_dsl/blockscaled_contiguous_gather_grouped_gemm_swiglu_fusion.py`<br>`flashinfer/fused_moe/cute_dsl/blockscaled_contiguous_grouped_gemm_finalize_fusion.py` | The CuTeDSL path fuses gather+GEMM1+SwiGLU in the first stage and finalize+unpermute+scatter-reduce in the second stage, removing standalone `moe_permute` and `moe_unpermute` kernels | Treat multi-kernel MoE ladders around permute / finalize as one existing FlashInfer CuTeDSL family first. |
| FlashInfer SM120 FP4 / groupwise GEMM heuristics | `cutlass_fp4_gemm_sm120`<br>`CutlassTileConfigSM120`<br>`group_gemm_nvfp4_nt_groupwise`<br>`group_gemm_mxfp4_nt_groupwise` | `flashinfer/gemm/gemm_base.py`<br>`include/flashinfer/gemm/fp4_gemm_cutlass_template_sm120.h`<br>`include/flashinfer/gemm/group_gemm_nvfp4_groupwise_sm120.cuh`<br>`csrc/nv_internal/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cpp` | FlashInfer mainline adds SM120-oriented FP4 GEMM selection and b12x CuTeDSL fused-MoE kernels | Treat SM120 FP4 MoE/GEMM tile selection and Blackwell-lite shape restrictions as an upstream FlashInfer kernel family before inventing a local heuristic. |
| FlashInfer MoE `routing_replay_out` support | `routing_replay_out`<br>`mPtrRoutingReplayOut`<br>`trtllm_fp8_block_scale_moe` | `flashinfer/fused_moe/core.py`<br>`csrc/trtllm_fused_moe_kernel_launcher.cu`<br>`csrc/fused_moe/noAuxTcKernels.cu` | TRTLLM-gen MoE kernels can optionally emit compact routing replay metadata without a separate routing-side reconstruction pass | Treat routing-replay writes in MoE traces as part of the upstream FlashInfer TRTLLM MoE family, not a separate postprocess opportunity. |

## 9. FlashInfer mainline kernel-overlap families

| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
| FlashInfer PDL launch-overlap family | `enable_pdl`<br>`launch_with_pdl`<br>`cudaGridDependencySynchronize`<br>`cudaTriggerProgrammaticLaunchCompletion`<br>`trigger_completion_at_end=False`<br>`allreduce_fusion` | `flashinfer/norm/__init__.py`<br>`flashinfer/activation.py`<br>`flashinfer/rope.py`<br>`flashinfer/comm/allreduce.py`<br>`flashinfer/comm/trtllm_ar.py` | FlashInfer uses Programmatic Dependent Launch broadly, and the allreduce path can further advance completion so the next PDL-aware kernel overlaps on the same stream | Treat tight same-stream dependent windows and allreduce-followed-by-kernel windows as one existing FlashInfer launch-overlap family first. |
| FlashInfer CuTeDSL MoE aux-stream async-memset overlap | `aux_stream`<br>`main_event`<br>`memset_event`<br>`use_async_memset` | `flashinfer/fused_moe/cute_dsl/fused_moe.py` | Preallocated MoE output is zeroed on an auxiliary CUDA stream while GEMM1 runs on the main stream, then both streams join before finalize | Treat GEMM1 vs output-zero windows as an existing FlashInfer multi-stream overlap family. |
| FlashInfer green-context SM partition overlap | `split_device_green_ctx`<br>`split_device_green_ctx_by_sm_count`<br>`green_ctx` | `flashinfer/green_ctx.py` | CUDA green contexts partition SMs and create dedicated streams for concurrent kernel families on separate SM slices | Treat full-device two-stream traces and SM-partitioned traces as different manifestations of an existing FlashInfer overlap mechanism. |

## 10. FlashInfer PR-backed / in-flight fused-kernel and kernel-overlap families

| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
| PR `#2720` PDL runtime-API migration | `cudaGridDependencySynchronize`<br>`cudaTriggerProgrammaticLaunchCompletion`<br>`inline PTX` | `PR #2720`<br>`include/flashinfer/comm/trtllm_allreduce_fusion.cuh`<br>`include/flashinfer/pos_enc.cuh` | Repo-wide migration preserves the existing PDL overlap family while replacing inline PTX with CUDA runtime APIs across norm, RoPE, attention, and MoE codepaths | Treat PDL-looking launch groups as an upstream FlashInfer overlap family even when implementation details differ across revisions. |

## 11. TensorRT-LLM-origin fused-kernel families

These rows are comparative references from `TensorRT-LLM`. Use them when a
trace looks like a TensorRT-LLM or TensorRT-LLM-plus-FlashInfer family even if
the current `sglang` checkout only carries an analogous implementation.

| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
| TensorRT-LLM FlashInfer activation / gate epilogues | `flashinfer_silu_and_mul`<br>`flashinfer_gelu_tanh_and_mul`<br>`auto_deploy::silu_and_mul`<br>post-GEMM `silu` + `mul` | `tensorrt_llm/_torch/custom_ops/flashinfer_custom_ops.py`<br>`tensorrt_llm/_torch/auto_deploy/transform/library/fuse_silu_mul.py`<br>`tensorrt_llm/_torch/models/modeling_gemma3.py` | Runtime custom ops and AutoDeploy rewrite `split/getitem + activation + mul` MLP epilogues into one FlashInfer op, including Gemma3 `gelu_tanh_and_mul` | Treat split gate activation + multiply as an existing TensorRT-LLM/FlashInfer epilogue family first. |
| TensorRT-LLM FlashInfer RMSNorm family | `flashinfer_rmsnorm`<br>`flashinfer_gemma_rmsnorm`<br>`auto_deploy::flashinfer_rms_norm` | `tensorrt_llm/_torch/custom_ops/flashinfer_custom_ops.py`<br>`tensorrt_llm/_torch/modules/rms_norm.py`<br>`tensorrt_llm/_torch/auto_deploy/custom_ops/normalization/rms_norm.py` | Runtime modules and AutoDeploy can lower plain RMSNorm and Gemma RMSNorm directly to FlashInfer kernels | Treat split RMSNorm ladders as an existing TensorRT-LLM norm family before calling them novel. |
| TensorRT-LLM FlashInfer residual add + RMSNorm | `flashinfer_fused_add_rmsnorm`<br>`flashinfer_gemma_fused_add_rmsnorm`<br>`auto_deploy::flashinfer_fused_add_rms_norm_inplace` | `tensorrt_llm/_torch/custom_ops/flashinfer_custom_ops.py`<br>`tensorrt_llm/_torch/modules/rms_norm.py`<br>`tensorrt_llm/_torch/auto_deploy/transform/library/fused_add_rms_norm.py` | Residual add immediately before RMSNorm can collapse to one in-place FlashInfer op, with Gemma variant support | Treat residual add + RMSNorm chains as an existing TensorRT-LLM fused epilogue family first. |
| TensorRT-LLM Triton fused residual add + RMSNorm + FP8 quant | `triton_fused_add_rms_norm_quant_fp8`<br>`fuse_rmsnorm_quant_fp8`<br>`fp8 static quant` | `tensorrt_llm/_torch/auto_deploy/custom_ops/normalization/triton_fused_add_rms_norm_quant_fp8.py`<br>`tensorrt_llm/_torch/auto_deploy/transform/library/fuse_rmsnorm_quant_fp8.py` | Mainline AutoDeploy can rewrite residual-add plus RMSNorm plus FP8 static quant into one Triton op that emits BF16 norm output, FP8 quant output, and residual-add output together | Treat split add + norm + FP8 quant ladders as an existing TensorRT-LLM mainline family first. |
| TensorRT-LLM FlashInfer RoPE with shared cos/sin cache | `flashinfer_apply_rope_with_cos_sin_cache_inplace`<br>`flashinfer_rope`<br>`cos_sin_cache` | `tensorrt_llm/_torch/modules/rotary_embedding.py`<br>`tensorrt_llm/_torch/auto_deploy/custom_ops/rope/flashinfer_rope.py`<br>`tensorrt_llm/_torch/auto_deploy/transform/library/rope.py` | Runtime path applies in-place RoPE from a shared cos/sin cache, while AutoDeploy can prebuild the full cache and lower diverse RoPE graphs to `flashinfer_rope` | Treat separate cos/sin gather + RoPE application ladders as an existing TensorRT-LLM attention-prep family. |
| TensorRT-LLM FlashInfer cached paged attention | `append_paged_kv_cache`<br>`BatchPrefillWithPagedKVCacheWrapper`<br>`BatchDecodeWithPagedKVCacheWrapper`<br>`auto_deploy::flashinfer_attention_mha_with_cache`<br>`read_cache_only` | `tensorrt_llm/_torch/attention_backend/flashinfer.py`<br>`tensorrt_llm/_torch/auto_deploy/custom_ops/attention/flashinfer_attention.py`<br>`docs/source/features/attention.md` | FlashInfer attention backend fuses metadata setup, optional paged-KV append, and prefill/decode wrapper execution, including shared-KV and read-cache-only variants in AutoDeploy | Treat metadata + KV-append + cached-attention ladders as one existing TensorRT-LLM cached-attention family first. |
| TensorRT-LLM FlashInfer MLA regular prefill | `append_paged_mla_kv_cache`<br>`BatchPrefillWithRaggedKVCacheWrapper`<br>`flashinfer_mla`<br>`rank 256`<br>`gpu append kernel` | `tensorrt_llm/_torch/auto_deploy/custom_ops/mla/flashinfer_mla.py` | Regular MLA prefill writes compressed KV pages and runs FlashInfer ragged prefill instead of a split append-plus-prefill ladder, with rank-256 paged-KV setups using the GPU append path | Treat MLA regular-prefill prep as an existing TensorRT-LLM FlashInfer family first. |
| TensorRT-LLM FlashInfer MLA chunked prefill with absorbed `W_kn` | `BatchMLAPagedAttentionWrapper`<br>`chunked prefill`<br>`W_kn`<br>`W_v` | `tensorrt_llm/_torch/auto_deploy/custom_ops/mla/flashinfer_mla.py` | Chunked prefill absorbs `W_kn` into the query-side projection, runs paged MLA attention in compressed space, then projects back with `W_v` | Treat split absorbed-proj + MLA + output-proj ladders as an existing TensorRT-LLM MLA family first. |
| TensorRT-LLM FlashInfer MLA decode with absorbed `W_kn` + `W_v` | `plan_decode`<br>`BatchMLAPagedAttentionWrapper`<br>`decode`<br>`W_kn`<br>`W_v` | `tensorrt_llm/_torch/auto_deploy/custom_ops/mla/flashinfer_mla.py` | Decode path reuses the absorbed-query MLA family and projects the compressed attention output back with `W_v` | Treat similar decode-time absorbed MLA ladders as an existing TensorRT-LLM family, not a new idea. |
| TensorRT-LLM FlashInfer fused MoE backend | `flashinfer.fused_moe`<br>`trtllm_bf16_moe`<br>`trtllm_fp8_block_scale_moe`<br>`trtllm_fp4_block_scale_moe`<br>`TRTLLM_GEN_FUSED_MOE_USE_FLASHINFER` | `tensorrt_llm/_torch/modules/fused_moe/moe_op_backend.py`<br>`tensorrt_llm/_torch/modules/fused_moe/fused_moe_trtllm_gen.py` | TRTLLM-gen MoE can route expert execution and quant helpers through FlashInfer instead of exposing per-expert eager ladders | Treat expert-side tiny GEMM ladders as matching an existing TensorRT-LLM FlashInfer MoE family first. |
| TensorRT-LLM FlashInfer cached SSM / Mamba update | `flashinfer_cached_ssm`<br>`selective_state_update`<br>`flashinfer_ssm` | `tensorrt_llm/_torch/auto_deploy/custom_ops/mamba/flashinfer_backend_mamba.py`<br>`tensorrt_llm/_torch/modules/mamba/mamba2_mixer.py` | Mamba2 paths can lower cached SSM state updates to FlashInfer selective-state-update kernels instead of many smaller state ops | Treat split cached-SSM state update ladders as an existing TensorRT-LLM FlashInfer family first. |

## 12. TensorRT-LLM-origin kernel-overlap families

| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
| TensorRT-LLM multi-stream MLA attention | `multi_stream_mla_attn`<br>`record_event_passthrough`<br>`_aux`<br>`wait_event` | `tensorrt_llm/_torch/auto_deploy/transform/library/multi_stream_attn.py`<br>`tensorrt_llm/_torch/auto_deploy/utils/multi_stream_utils.py` | AutoDeploy rewrites MLA Q/KV forks so the KV projection runs on an auxiliary stream while the Q path stays on the caller stream | Treat exposed Q-branch vs KV-branch overlap as an existing TensorRT-LLM multi-stream family first. |
| TensorRT-LLM multi-stream MoE shared-vs-routed overlap | `multi_stream_moe`<br>`begin_aux_stream_passthrough`<br>`end_aux_stream_passthrough`<br>`wait_aux_stream_passthrough`<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 `multi_stream_moe` 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. |

## 13. TensorRT-LLM-origin PR-backed / in-flight fused-kernel and kernel-overlap families

| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
| PR `#12525` FlashInfer TRTLLM-gen FMHA paged-index / buffer rework | `shared paged index`<br>`trtllm-gen attention`<br>`flashinfer`<br>`kv cache buffer` | `PR #12525`<br>`tensorrt_llm/_torch/auto_deploy/custom_ops/attention/flashinfer_attention.py` | Open PR refines the existing FlashInfer TRTLLM-gen cached-attention family by disabling shared paged index and unifying KV-buffer construction | Treat these attention-prep changes as an in-flight implementation evolution of an existing family first. |
| PR `#12544` NVFP4 KV cache support in TRTLLM-gen attention | `NVFP4 KV cache`<br>`trtllm-gen attention`<br>`flashinfer` | `PR #12544`<br>`tensorrt_llm/_torch/auto_deploy/custom_ops/attention/flashinfer_attention.py` | Open PR extends the cached-attention family so the FlashInfer-backed TRTLLM-gen path can build and consume NVFP4 KV buffers directly | Treat split KV-cache quant + buffer-build ladders as an in-flight TensorRT-LLM attention family first. |
| PR `#12738` / `#12557` BF16 TRTLLM-gen MoE through FlashInfer | `bf16 trtllm-gen moe`<br>`flashinfer`<br>`trtllm_bf16_moe` | `PR #12738`<br>`PR #12557`<br>`tensorrt_llm/_torch/modules/fused_moe/fused_moe_trtllm_gen.py` | Open PRs extend the TRTLLM-gen MoE family so BF16 expert execution can route through FlashInfer instead of only CUTLASS-like paths | Treat BF16 expert ladders as an in-flight TensorRT-LLM FlashInfer MoE family. |

## 14. vLLM-origin fused-kernel families

These rows are comparative references from `vllm`. Use them when a trace looks
similar to an upstream family even if the current `sglang` checkout does not
contain the same implementation.

| Pattern | Trace keywords | Primary code | Existing path | Skill should conclude |
| --- | --- | --- | --- | --- |
| vLLM-origin fused residual add + RMSNorm | `fused_add_rms_norm*`<br>residual add right before RMSNorm | `vllm/model_executor/layers/layernorm.py`<br>`vllm/_custom_ops.py`<br>`csrc/layernorm_kernels.cu`<br>`csrc/cpu/layernorm.cpp` | Custom CUDA / CPU fused add-RMSNorm op reused directly and as a building block for later compile-time fusions | Treat split residual add + RMSNorm as a long-standing vLLM-origin precedent before calling the opportunity novel in sglang. |
| vLLM-origin AllReduce + RMSNorm (+ residual / quant) | `fuse_allreduce_rms`<br>`AllReduceFusionPass`<br>`allreduce + rmsnorm` | `vllm/compilation/passes/fusion/allreduce_rms_fusion.py`<br>`docs/design/fusions.md` | Compile-time patterns cover `AllReduce -> RMSNorm(+residual_add)` and optional FP8 / NVFP4 quant suffixes | Treat TP collective + norm (+ quant) ladders as a known vLLM-origin fusion family first. |
| 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>`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 + 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. 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 |
| --- | --- | --- | --- | --- |
| vLLM-origin AsyncTP GEMM + collective overlap | `fuse_gemm_comms`<br>`fused_matmul_reduce_scatter`<br>`fused_all_gather_matmul` | `vllm/compilation/passes/fusion/collective_fusion.py`<br>`docs/design/fusions.md` | AsyncTP overlaps GEMM with reduce-scatter / all-gather via symmetric-memory collectives | Treat GEMM+comm windows as a clear vLLM-origin overlap precedent first. |
| vLLM-origin Sequence Parallelism staging | `enable_sp`<br>`ReduceScatter`<br>`AllGather`<br>`SequenceParallelismPass` | `vllm/compilation/passes/fusion/sequence_parallelism.py`<br>`docs/design/fusions.md` | Sequence-parallel rewrites all-reduce into RS -> local norm -> AG so later passes can overlap comm and compute | Treat RS / AG staging around norm blocks as an upstream overlap-enabling family. |
| 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. |

## 17. vLLM-origin PR-backed / in-flight fused-kernel and 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 `#37110` Triton attention + per-group FP8 dynamic quant | `group_size=128`<br>`group_size=64`<br>`output_group_scale`<br>`per-group FP8` | `PR #37110`<br>`vllm/compilation/passes/fusion/attn_quant_fusion.py`<br>`vllm/v1/attention/ops/triton_unified_attention.py` | In-flight Triton attention epilogue computes per-group FP8 scales and quantizes output directly instead of launching a separate group-quant kernel | Treat attention + per-group FP8 quant as a concrete upstream vLLM family, not a novel idea. |
| PR `#38445` MiniMax-M2 FP32 gate kernel | `fp32_router_gemm`<br>`MiniMax-M2`<br>`gate kernel` | `PR #38445`<br>`vllm/model_executor/layers/fused_moe/router/gate_linear.py`<br>`vllm/model_executor/models/minimax_m2.py` | Draft CUDA kernel fuses BF16->FP32 conversion and low-batch router GEMM for MiniMax-M2, replacing up to three kernels on the gate path | Treat MiniMax-M2 gate ladders as an in-flight upstream fused router family first. |
| PR `#38621` fused QK norm + RoPE + cache + quant | `fused_qk_norm_rope_cache_quant`<br>`QK Norm + RoPE + Cache + Quant` | `PR #38621`<br>`csrc/fused_qk_norm_rope_cache_quant.cu`<br>`vllm/compilation/passes/fusion/qk_norm_rope_cache_quant_fusion.py` | Draft CUDA kernel and compile-time pass try to fuse QK RMSNorm, RoPE, KV cache write, and optional FP8 quant for small-batch decode | Treat this as an in-flight upstream fusion family before calling a similar idea novel. |
| PR `#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>`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_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. |

## 18. Important toggles and caveats

| Toggle / env | Location | Effect on trace interpretation |
| --- | --- | --- |
| `enable_flashinfer_allreduce_fusion` | `python/sglang/srt/server_args.py` | Enables the FlashInfer TP allreduce fusion family. |
| `enable_aiter_allreduce_fusion` | `python/sglang/srt/server_args.py` | Enables ROCm AITER TP allreduce fusion. |
| `enable_deterministic_inference` | `python/sglang/srt/server_args.py` | Can intentionally disable or change some fast fusion paths, especially AITER allreduce fusion and some sampling / router choices, so split kernels may be expected. |
| `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_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. |
| `SGLANG_ENABLE_FUSED_QKNORM_ROPE` | `python/sglang/multimodal_gen/runtime/layers/layernorm.py` | Gates the diffusion fused qknorm+rope path. |
| `enable_pdl` / `launch_with_pdl` | `flashinfer/norm/__init__.py`<br>`flashinfer/activation.py`<br>`flashinfer/rope.py`<br>`flashinfer/fused_moe/core.py`<br>`flashinfer/comm/allreduce.py` | Enables FlashInfer PDL across many kernels; launch grouping and same-stream overlap can change substantially when it is on. |
| `trigger_completion_at_end` | `flashinfer/comm/allreduce.py` | `False` enables downstream PDL-aware overlap after FlashInfer allreduce fusion; `True` delays completion to kernel end and removes that overlap window. |
| `use_cuda_graph` | `flashinfer/fused_moe/cute_dsl/fused_moe.py` | Enables the preallocated-buffer path and the safe aux-stream async-memset overlap in FlashInfer CuTeDSL MoE. |
| `split_device_green_ctx*` | `flashinfer/green_ctx.py` | Changes trace shape by partitioning SMs into separate green contexts instead of overlapping full-device streams on the default context. |
| `rmsnorm_backend` | `tensorrt_llm/_torch/auto_deploy/config/default.yaml` | Chooses whether AutoDeploy lowers RMSNorm to FlashInfer, so split norm ladders may reflect backend selection rather than a missing fuse. |
| `insert_cached_attention.backend` | `tensorrt_llm/_torch/auto_deploy/config/default.yaml` | Selects the cached-attention backend; `flashinfer` enables the paged-KV cached-attention family. |
| `insert_cached_mla_attention.backend` | `tensorrt_llm/_torch/auto_deploy/config/default.yaml` | Selects the cached MLA backend; `flashinfer_mla` enables the MLA prefill / decode family. |
| `TRTLLM_GEN_FUSED_MOE_USE_FLASHINFER` | `tensorrt_llm/_torch/modules/fused_moe/fused_moe_trtllm_gen.py` | Forces or guards the FlashInfer-backed TRTLLM-gen MoE family, so expert-kernel shape can change substantially when it is set. |
| `multi_stream_moe` | `tensorrt_llm/_torch/auto_deploy/config/default.yaml` | Enables the TensorRT-LLM shared-expert vs routed-expert overlap family. |
| `multi_stream_mla_attn` | `tensorrt_llm/_torch/auto_deploy/config/default.yaml` | Enables the TensorRT-LLM MLA Q-vs-KV branch overlap family. |
| `multi_stream_gemm` | `tensorrt_llm/_torch/auto_deploy/config/default.yaml` | Enables generalized FP8 GEMM fork overlap in TensorRT-LLM AutoDeploy. |
| `mlir_elementwise_fusion` | `tensorrt_llm/_torch/auto_deploy/config/default.yaml` | Can absorb merge adds into larger fused kernels, so missing explicit merge nodes in multi-stream traces may be intentional. |
| `enable_torch_compile` | `python/sglang/srt/server_args.py`<br>`python/sglang/multimodal_gen/runtime/server_args.py` | Compiler-generated fusion / reordering can hide handwritten kernel names; absence of a custom kernel does not always mean absence of fusion. |
| `enable_fused_grouped_gemm_combine` | `PR #21877` | In-flight path that intentionally disables SBO because combine is folded into down-GEMM. |
| `PassConfig.fuse_allreduce_rms` | `vllm/config/compilation.py` | Enables vLLM's AllReduce -> RMSNorm (+ residual / quant) compile-time fusion family; 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. |
| 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. |

## 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.

```bash
# Optional sibling checkouts used for comparative scanning:
FLASHINFER_REPO=${FLASHINFER_REPO:-../flashinfer}
TRTLLM_REPO=${TRTLLM_REPO:-../TensorRT-LLM}
VLLM_REPO=${VLLM_REPO:-../vllm}

rg -n "fused_add_rmsnorm|gemma_fused_add_rmsnorm|silu_and_mul|gelu_and_mul|fused_qk_rope_reshape_and_cache|fused_set_kv_buffer|fused_metadata_copy|normal_decode_set_metadata|_append_shared_to_topk_output|fused_append_shared_experts_with_weights" python/sglang
rg -n "MiniMaxM2RMSNormTP|fused_qknorm_rope|fused_qk_rope_cat_and_cache_mla|fused_qk_norm_mrope_3d_cache_pts_quant_shuffle|split_qkv_rmsnorm_rope|trtllm_fp8_kv_kernel|set_mla_kv_buffer_fp8_quant" python/sglang
rg -n "FusedMoeRouter|fused_topk_deepseek|moe_fused_gate|aiter_fused_topk|fused_rms_fp8_group_quant|fast_topk_transform_fused|fused_store_index_k_cache|fused_temperature_softmax|fused_softcap" python/sglang
rg -n "fused_qkvzba_split_reshape_cat|fused_gdn_gating|rms_norm_gated|layer_norm_gated|chunk_gated_delta_rule_fwd_kkt_solve_kernel|fused_recurrent_gated_delta_rule_update|fused_mamba_state_scatter_with_mask|_fused_gather_to_staging_kernel|_fused_scatter_from_staging_kernel" python/sglang
rg -n "single_batch_overlap|alt_stream|shared_expert|_comm_stream|scatter_stream|triton_mrope_fused|ring_attn|all_to_all_single|reorder_for_compute_comm_overlap|use_dual_stream" python/sglang
git log --all --format='%h %s' | rg -i 'fused|fusion|overlap|cutedsl|triton|cuda|rope|topk|quant|combine|allreduce|all_to_all'
rg -n "silu_and_mul|gelu_tanh_and_mul|gelu_and_mul|silu_and_mul_scaled_nvfp4_experts_quantize|rmsnorm_quant|fused_add_rmsnorm|fused_add_rmsnorm_quant|fused_rmsnorm_silu" "$FLASHINFER_REPO/flashinfer"
rg -n "AllReduceFusionPattern|allreduce_fusion|trigger_completion_at_end|rope_quantize_fp8|rope_quantize_fp8_append_paged_kv_cache|fused_topk_deepseek|cutlass_fused_moe|trtllm_.*_moe" "$FLASHINFER_REPO/flashinfer"
rg -n "aux_stream|use_async_memset|split_device_green_ctx|split_device_green_ctx_by_sm_count|enable_pdl|launch_with_pdl" "$FLASHINFER_REPO/flashinfer" "$FLASHINFER_REPO/include"
git -C "$FLASHINFER_REPO" log --all --format='%h %s' | rg -i 'fused|fusion|overlap|pdl|stream|rope|kv|quant|topk|moe'
rg -n "flashinfer_silu_and_mul|flashinfer_gelu_tanh_and_mul|flashinfer_rmsnorm|flashinfer_gemma_rmsnorm|flashinfer_fused_add_rmsnorm|flashinfer_apply_rope_with_cos_sin_cache_inplace|triton_fused_add_rms_norm_quant_fp8|fuse_rmsnorm_quant_fp8" "$TRTLLM_REPO/tensorrt_llm/_torch"
rg -n "flashinfer_attention_mha_with_cache|append_paged_kv_cache|flashinfer_mla|append_paged_mla_kv_cache|flashinfer_cached_ssm|selective_state_update|flashinfer.fused_moe" "$TRTLLM_REPO/tensorrt_llm/_torch" "$TRTLLM_REPO/docs/source"
rg -n "multi_stream_moe|multi_stream_mla_attn|multi_stream_gemm|record_event_passthrough|begin_aux_stream_passthrough|end_aux_stream_passthrough|wait_aux_stream_passthrough" "$TRTLLM_REPO/tensorrt_llm/_torch"
git -C "$TRTLLM_REPO" log --all --format='%h %s' | rg -i 'fused|fusion|overlap|flashinfer|mla|kv cache|multi-stream|stream|rope|rmsnorm|moe'
rg -n "fused_add_rms_norm|merge_attn_states|fused_qk_norm_rope|grouped_topk|topk_softmax|topk_sigmoid|dsv3_router_gemm|dsv3_fused_a_gemm|concat_and_cache_mla_rope_fused|gpt_oss_router_gemm|cutlass_scaled_mm|cpu_fused_moe|fused_moe_lora|triton_pos_embed_interpolate" "$VLLM_REPO/vllm" "$VLLM_REPO/csrc"
rg -n "fuse_allreduce_rms|fuse_norm_quant|fuse_act_quant|fuse_attn_quant|enable_qk_norm_rope_fusion|fuse_rope_kvcache|enable_sp|fuse_gemm_comms|RocmAiter|dcp_alltoall|shared_experts_stream|TRTLLM_ENABLE_PDL|wk_weights_proj" "$VLLM_REPO/vllm" "$VLLM_REPO/docs/design/fusions.md" "$VLLM_REPO/csrc"
git -C "$VLLM_REPO" log --all --format='%h %s' | rg -i 'fused|fusion|overlap|triton|cuda|rope|kv cache|topk|router|allreduce|reduce-scatter|all-gather|all_to_all|quant'
# GitHub PR scan terms for the connector or web UI:
#   "fused OR overlap repo:sgl-project/sglang"
#   "triton OR cutedsl OR cuda fused repo:sgl-project/sglang"
#   "fused OR overlap repo:flashinfer-ai/flashinfer"
#   "pdl OR aux_stream OR green_ctx repo:flashinfer-ai/flashinfer"
#   "fused OR overlap repo:NVIDIA/TensorRT-LLM"
#   "flashinfer OR mla OR moe OR rmsnorm repo:NVIDIA/TensorRT-LLM"
#   "multi-stream OR aux_stream OR cudagraph repo:NVIDIA/TensorRT-LLM"
#   "fused OR overlap repo:vllm-project/vllm"
#   "triton OR cuda fused repo:vllm-project/vllm"
```
