llm-torch-profiler-analysis

Unified LLM torch-profiler triage skill for `sglang`, `vllm`, `TensorRT-LLM`, and `TokenSpeed`. Use it to inspect an existing `trace.json(.gz)` or profile directory, or to drive live profiling against a running server when supported and return one three-table report with kernel, overlap-opportunity, and fuse-pattern tables.

Install
npx skills add 'https://github.com/sgl-project/sglang/tree/main/.claude/skills/llm-torch-profiler-analysis'
Download bundle ↓
main · a9fb1c3Scanned 2026-09-17

Contributors

GitHub-linked commit authors for this SKILL.md at the saved revision. Co-authors and history before file renames are not included.

File history ↗
View on GitHub
← Back to SKILL.md

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

PatternTrace keywordsPrimary codeExisting pathSkill should conclude
Fused residual add + RMSNormfused_add_rmsnorm*npu_add_rms_normadd_rmsnorm_biasgemma_fused_add_rmsnormgemma_rmsnorm_residual_scalar_gemma_rmsnorm_residual_kernelresidual add right before normpython/sglang/srt/layers/layernorm.pypython/sglang/srt/layers/gemma4_fused_ops.pypython/sglang/srt/layers/quantization/modelslim/modelslim.pyShared CUDA / ROCm / CPU / NPU fused add-RMSNorm implementations, including Gemma, Gemma4 scalar-residual, and NPU-bias variantsTreat split residual add + RMSNorm as an existing cross-backend fusion first, not a new idea.
FlashInfer unified allreduce_fusioncross_device_reduce_1stage*all_reduceFusedAddRMSNormKernelrmsnorm*python/sglang/srt/layers/flashinfer_comm_fusion.pypython/sglang/srt/layers/layernorm.py::forward_with_allreduce_fusionpython/sglang/srt/layers/communicator.py::apply_flashinfer_allreduce_fusionFlashInfer 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 fusionROCm all-reduce plus RMSNorm still splitpython/sglang/srt/layers/layernorm.py::forward_with_allreduce_fusionpython/sglang/srt/distributed/communication_op.py::tensor_model_parallel_fused_allreduce_rmsnormpython/sglang/srt/layers/communicator.py::apply_aiter_all_reduce_fusionROCm-side fused TP all-reduce + RMSNorm with fallback to plain all-reduce plus normOn AMD, rule out existing AITER fusion before proposing a new communication fusion.
Fused activation-and-mul (SwiGLU / GeGLU)silu_and_mulgelu_and_mulnpu_swiglupython/sglang/srt/layers/activation.pySingle op covers activation plus elementwise multiply across CUDA / CPU / NPU / XPU backendsTreat separate activation + mul on packed MLP outputs as missing existing fusion.
Fused dual residual RMSNormresidual add plus two RMSNorm-like kernels around Grok blockspython/sglang/srt/layers/elementwise.py::fused_dual_residual_rmsnormpython/sglang/srt/models/grok.pyOne Triton kernel computes intermediate residual update and next RMSNorm output togetherOn Grok-like residual layouts, treat split residual + norm as missing existing fusion.
In-place QK RMSNormsplit q_norm / k_norm kernelspython/sglang/srt/models/utils.py::apply_qk_normpython/sglang/kernels/ops/layernorm/norm.py::fused_inplace_qknormIn-place JIT QK norm plus optional alt_stream overlap for KCheck shape, dtype, deterministic mode, and in-place legality before proposing a new QK fuse.
TorchInductor horizontal Q/K norm combo-kernelscombo_kernelsbenchmark_combo_kernelq_normk_normsplit_with_sizestorch._inductor.config.combo_kernelsTorchInductor can horizontally fuse sibling Q-norm and K-norm kernels in compiled traces, often deleting split_with_sizes / clone laddersTreat separate Q/K norm ladders in compile-heavy traces as an existing compiler-fusion family first.
MiniMax TP fused QK RMSNormMiniMaxM2RMSNormTPrms_sumsq_serialrms_apply_serialforward_qkpython/sglang/srt/models/minimax_m2.pyTriton kernels compute Q / K sumsq together, TP all-reduces shared stats, then apply both RMSNorms togetherOn MiniMax traces, separate Q norm and K norm are usually a missed model-specific Triton fusion.
Fused QK RMSNorm + RoPEqknorm* + rope* + rotary* as separate stepspython/sglang/kernels/ops/attention/fused_qknorm_rope.pypython/sglang/srt/models/qwen3_moe.pyOne JIT kernel applies QK RMSNorm and RoPE in-place on packed QKVFor compatible LLMs, classify split QK norm + RoPE as a missing existing fusion.
Fused QK RoPE reshape + KV cache writefused_qk_rope_reshape_and_cache*RoPE followed by reshape / cache DtoDpython/sglang/srt/layers/attention/utils.py::fused_qk_rope_reshape_and_cacheOne Triton kernel applies RoPE to Q / K, reshapes cache layout, and writes K / V directly to paged cacheTreat separate RoPE + reshape + cache-write ladders as an existing attention-prep fusion family.
Fused RoPE + KV cache storefused_set_kv_bufferRoPE followed by KV-store, DtoD, or cache-write kernelspython/sglang/kernels/ops/attention/rope.pypython/sglang/srt/models/utils.py::enable_fused_set_kv_bufferShared entrypoints can route to fused RoPE + KV-store or model-side fused_set_kv_buffer fast pathsCompare against the fused cache-store path before proposing a new KV rewrite.
Fused decode metadata setupnormal_decode_set_metadatacache_seqlens_int32cu_seqlens_kpage_tableswa_page_tablepython/sglang/srt/layers/attention/flashattention_backend.py::normal_decode_set_metadataTriton decode path fuses seq-len cast/add, prefix-sum, req-to-token gather, page-table divide, and optional SWA metadata build into 1-2 kernelsIf decode exposes multiple tiny metadata kernels before attention, first compare against this existing fused metadata-prep path.
NSA fused metadata copy for graph replayfused_metadata_copyfused_metadata_copy_multifused_nsa_cache_seqlensfused_flashmla_metadatapython/sglang/kernels/ops/attention/fused_metadata_copy.pyCUDA graph replay path fuses multiple metadata copies into one kernel or one multi-destination kernelTreat bursts of tiny metadata-copy kernels around NSA replay as a missed existing replay fusion.
DeepSeek MLA fused projection + norm + RoPEqkv_proj_with_rope_fused_weightfused_qkv_a_proj_with_mqaforward_absorb_fused_mla_rope*python/sglang/srt/models/deepseek_common/attention_forward_methods/forward_mla_fused_rope_cpu.pypython/sglang/srt/models/deepseek_common/attention_forward_methods/forward_mla_fused_rope_rocm.pypython/sglang/srt/models/deepseek_v2.pyCPU / ROCm paths fuse DeepSeek MLA projection packing with q / k norm, RoPE, and cache-oriented MLA prepFor DeepSeek MLA, split proj / norm / rope prep is usually an existing backend-specific fuse that did not fire.
Fused QK RoPE concat + MLA cache writefused_qk_rope_cat_and_cache_mlaset_mla_kv_bufferpython/sglang/srt/layers/rocm_linear_utils.pypython/sglang/srt/models/deepseek_common/attention_forward_methods/forward_mla.pyROCm MLA path can fuse Q / K RoPE packing, concat, and MLA cache write in one backend-specific opOn DeepSeek / MLA traces, separate RoPE-cat-cache steps are not automatically novel.
Qwen3 decode fused QK norm + 3D mRoPE + KV cache writefused_qk_norm_mrope_3d_cache_pts_quant_shufflemropedecode cache writepython/sglang/srt/models/qwen3.pyROCm / AITER decode path fuses QK norm, 3D mRoPE, and paged KV cache writeOn Qwen3-style decode, separate norm + mRoPE + cache-store kernels are not a novel opportunity.
NPU fused split-QKV + RMSNorm + RoPEsplit_qkv_rmsnorm_ropepython/sglang/srt/models/llama.pypython/sglang/srt/models/qwen3.pypython/sglang/srt/models/qwen3_moe.pypython/sglang/srt/models/glm4_moe.pyAscend path fuses QKV split, Q / K RMSNorm, and RoPE in one opOn NPU traces, separate split / norm / rope kernels usually mean the fused path is unavailable or bypassed.
Fused FP8 quantize + paged KV cache writetrtllm_fp8_kv_kernelfp8 kv cache writepaged KV cache writepython/sglang/srt/layers/attention/triton_ops/trtllm_fp8_kv_kernel.pyTRTLLM MHA path fuses FP8 quantization, scale computation, and paged K / V cache writeIf FP8 KV cache traces show standalone quant plus write kernels, first compare against this existing Triton fuse.
Fused MLA KV cache write + FP8 quantset_mla_kv_buffer_fp8_quant*set_mla_kv_buffer_triton_fp8_quantpython/sglang/srt/mem_cache/utils.pypython/sglang/srt/mem_cache/memory_pool.pyMLA / NSA KV pool path can quantize K and write directly into KV storage without a separate concat-and-quant chainTreat standalone quant + KV-buffer write on MLA paths as missing existing fusion first.
Fused MoE router / top-k / softcappingFusedMoeRouterfused_moe_router*router GEMM + topk + tanhpython/sglang/srt/layers/moe/router.pySingle fused router kernel covers router matmul, softcapping, and top-k selectionTreat exposed router matmul + softcap + top-k chains as an existing MoE fusion family.
Fused MoE grouped-topk / gate kernelsfused_topk_deepseekmoe_fused_gateaiter_fused_topkkimi_k2_moe_fused_gatepython/sglang/srt/layers/moe/topk.pyCUDA / ROCm / FlashInfer kernels fuse bias, grouped-topk, renorm, and routed scaling into one gate opCheck backend / model eligibility before proposing a novel router-gate fusion.
Qwen-style shared-expert append into routed top-k output_append_shared_to_topk_outputfused_append_shared_experts_with_weightsnum_fused_shared_expertspython/sglang/srt/models/qwen2_moe.pypython/sglang/srt/layers/moe/moe_runner/triton_utils/fused_moe_triton_kernels.pyQwen-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 pathTreat routed top-k plus shared-expert pad / concat ladders as an existing MoE-prep fusion family first.
Fused MoE dispatch / permute / combinetoken permutationdispatch / combinegrouped top-kmany small MoE support kernelspython/sglang/srt/layers/moe/fused_moe_triton/layer.pypython/sglang/srt/layers/moe/fused_moe_triton/fused_moe.pyFusedMoE plus DeepEP / FlashInfer / FuseEP / standard dispatch backends and permute_fusion=TrueFirst ask whether the model is missing an existing FusedMoE-style path or backend-specific dispatcher path.
Fused MoE sum + all-reducerouted MoE followed by explicit sum-reduce kernelspython/sglang/srt/layers/moe/fused_moe_triton/fused_moe.pypython/sglang/srt/layers/moe/fused_moe_triton/fused_moe_triton_kernels.pyfuse_sum_all_reduce=True path in the second MoE GEMMBefore 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-quantsilu_and_mul_*quant*npu_dequant_swiglu_quantswiglu_quantpython/sglang/srt/layers/moe/ep_moe/kernels.pypython/sglang/kernels/ops/quantization/nvfp4_gemm_swiglu_nvfp4_quant.pypython/sglang/srt/layers/moe/cutlass_w4a8_moe.pypython/sglang/srt/hardware_backend/npu/quantization/moe_methods.pyQuantized MoE backends fuse SwiGLU / SiLU-and-mul with FP8 / FP4 / NPU re-quant before the second expert GEMMIf MoE traces show standalone activation then quant kernels, first check whether the quantized fused path is missing.
DeepSeek comm-prep fused RMSNorm + quant / flatten-quantfused_rms_fp8_group_quantfused_rms_mxfp4_quantfused_flatten_fp8_group_quantfused_flatten_mxfp4_quantpython/sglang/srt/layers/communicator.pypython/sglang/srt/models/deepseek_common/attention_forward_methods/forward_mla.pypython/sglang/srt/models/deepseek_common/attention_forward_methods/forward_mha.pyDeepSeek MLA / MHA ROCm paths fuse RMSNorm or flatten with FP8 / MXFP4 quantization for comm / attention prepOn DeepSeek quant traces, split norm + quant or flatten + quant is an existing family, not a new idea.
NSA fused top-k transform / page-table buildfast_topk_transform_fusedfast_topk_transform_ragged_fusedpython/sglang/srt/layers/attention/nsa_backend.pyNSA can fuse top-k selection with paged / ragged index transform instead of separate top-k plus metadata scatterIf NSA top-k metadata work is split, check SGLANG_NSA_FUSE_TOPK and backend support first.
NSA fused quantize + indexed K-cache storefused_store_index_k_cacheact_quantindex_k_with_scale_bufferpython/sglang/kernels/ops/attention/fused_store_index_cache.pypython/sglang/srt/layers/attention/nsa/nsa_indexer.pySingle JIT kernel quantizes bf16 K to fp8 + scale and writes directly into NSA index cacheTreat split act_quant + buffer-store on CUDA as missing an existing fused store path.
Fused sampling temperature + softmaxfused_temperature_softmax*python/sglang/srt/layers/fused_sampling.pypython/sglang/srt/layers/sampler.pyTriton single-pass / multi-pass kernels fuse temperature scaling and softmax during decodeSeparate temp-divide + softmax at decode batch sizes is often a missed existing fusion.
Fused logit softcapfused_softcapfinal_logit_softcappingpython/sglang/srt/layers/elementwise.pypython/sglang/srt/layers/logits_processor.pyTriton kernels fuse cast-to-float and softcap / tanh math for logits or generic elementwise softcappingTreat exposed cast + softcap ladders as an existing Triton fuse family.
Linear-attention packed projection reshufflefused_qkvzba_split_reshape_cat*qkvz_projba_projqkvabz_projfused_qkvbfg_a_projpython/sglang/kernels/ops/attention/triton_gdn_fused_proj.pypython/sglang/srt/models/qwen3_next.pypython/sglang/srt/models/qwen3_5.pypython/sglang/srt/models/kimi_linear.pypython/sglang/srt/models/jet_nemotron.pyGDN / Kimi / Jet-style linear-attn models pack multiple projections, then fuse split / reshape / cat into one kernelTreat split reshape / transpose / cat ladders as an existing linear-attention fusion family.
Fused GDN gating prepfused_gdn_gatingsoftplusbeta_outputpython/sglang/srt/layers/attention/fla/fused_gdn_gating.pyTriton kernel computes GDN gate preparation such as -exp(A_log) * softplus(...) and sigmoid(b) togetherOn GDN traces, treat split gate-prep elementwise kernels as missing existing fusion first.
Fused RMSNorm-gated linear-attention outputFusedRMSNormGatedlayer_norm_gated_fwdpython/sglang/srt/layers/attention/fla/fused_norm_gate.pypython/sglang/srt/models/qwen3_next.pypython/sglang/srt/models/kimi_linear.pyOne Triton op covers residual-aware (RMS)Norm plus sigmoid / swish gatingIf norm and output gate appear as separate kernels in GDN / Kimi-like blocks, first suspect a missing existing fusion.
Fused gated RMSNorm / LayerNormrms_norm_gatedlayer_norm_gatedpython/sglang/srt/layers/attention/mamba/ops/layernorm_gated.pyMamba-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_trilchunk_gated_delta_rule_fwd_kkt_solve_kernelscaled_dot_kktsolve_trilrecompute_w_upython/sglang/srt/layers/attention/fla/chunk_fwd.pypython/sglang/srt/layers/attention/fla/kda.pyGDN / KDA chunk forward fuses scaled_dot_kkt + solve_tril in the prefill / intra-chunk path, then finishes recompute_w_u as the next stepTreat split KKT + triangular-solve ladders as an existing linear-attention fusion family first.
Fused linear-attention recurrent / KDA updatefused_sigmoid_gating_delta_rule_updatefused_recurrent_gated_delta_rule_updatefused_kda_gatepython/sglang/srt/layers/attention/fla/fused_sigmoid_gating_recurrent.pypython/sglang/srt/layers/attention/fla/fused_recurrent.pypython/sglang/srt/models/kimi_linear.pypython/sglang/srt/models/jet_nemotron.pyTriton / CuTeDSL kernels fuse gating math, optional QK l2norm, recurrent state update, and output generationTreat split gating + recurrent-update chains as existing linear-attention fusion, not a novel opportunity.
Fused Mamba state gather/scatter with maskfused_mamba_state_scatter_with_maskindex_elementwise_kernelpython/sglang/srt/layers/attention/mamba/mamba_state_scatter_triton.pyTriton kernel replaces multiple masked gather / scatter index kernels with one fused updateIf 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_fused_scatter_from_staging_kernelpython/sglang/srt/disaggregation/common/staging_buffer.pyTriton kernels gather scattered KV slices into contiguous staging memory and scatter them back into KV cache on decodeTreat ladders of tiny gather/scatter/copy kernels in heterogeneous TP staging as missing an existing Triton fusion.

2. LLM / SRT kernel-overlap families

PatternTrace keywordsPrimary codeExisting pathSkill should conclude
Single-batch overlap (SBO)MoE combine, down-gemm, shared-expert work in nearby two-stream windowspython/sglang/srt/batch_overlap/single_batch_overlap.pycombine vs down-gemm overlap, combine vs shared-expert overlap, one-stream dispatch+shared overlap, explicit SM partitioning and eventsIf exposed MoE combine sits near neighboring compute, classify it against SBO before calling it new overlap.
Q and K normalization on different streamsQ-side norm and K-side norm on different streamspython/sglang/srt/models/utils.py::apply_qk_normpython/sglang/srt/models/qwen3.pypython/sglang/srt/models/qwen3_next.pypython/sglang/srt/models/qwen3_5.pyQ stays on current stream, K can run on alt_stream in capture modeTreat split Q / K norm as an existing overlap family when alt_stream is already wired.
DeepSeek shared-expert / routed-expert overlapshared-expert GEMMs near DeepEP dispatch / combinepython/sglang/srt/models/deepseek_v2.pypython/sglang/srt/batch_overlap/single_batch_overlap.pyshared experts on alt_stream, overlap with dispatch / combine and down-gemm, Blackwell-specific env gatingThis is an established routed-vs-shared branch overlap pattern, not a novel idea.
Llama4 shared branch vs routed branch overlapshared expert branch plus routed MoE branch as adjacent windowspython/sglang/srt/models/llama4.pyshared expert on current stream, router + topk + routed experts on alt_streamUse Llama4 as the first precedent for branch-level overlap in similar sparse models.
ExaoneMoE shared experts vs router experts overlapshared expert output and router-expert output form a two-branch windowpython/sglang/srt/models/exaone_moe.py::forward_normal_dual_streamshared experts on current stream, router + routed experts on alt_stream, explicit join before combineThis is an existing dual-stream MoE overlap family.
Grok residual-MoE branch overlapdense MLP and block-sparse MoE branches in parallelpython/sglang/srt/models/grok.py::moe_with_rmoedense MLP on current stream, MoE on alt_stream, fused dual residual RMSNorm around boundariesTreat exposed Grok branch overlap as an existing pattern.
NSA dual-stream overlapQ-proj, K-proj, RoPE, cache-store, quantization in tight two-stream windowspython/sglang/srt/layers/attention/nsa/nsa_indexer.pyQ / K projection split, RoPE split, cache-store vs quantization overlapNSA already contains several dual-stream overlap precedents.
MoriEP async dispatch / combine comm streamMoriEP_comm_streamdispatchcombinedone_eventpython/sglang/srt/layers/moe/token_dispatcher/moriep.pyMoriEP can submit dispatch and combine onto a dedicated communication stream and synchronize only through eventsTreat MoriEP comm / compute interleave as an existing MoE overlap family.
Heterogeneous-TP staging scatter overlapscatter_stream_scatter_streamstagingpython/sglang/srt/disaggregation/common/staging_handler.pypython/sglang/srt/disaggregation/common/staging_buffer.pydecode-side staging scatter kernels can run on a dedicated stream while forward continues on the main streamIf decode traces show staging scatter kernels adjacent to forward kernels, classify them against this existing overlap family first.
Generic alt_stream overlap familiesalt_stream plus explicit wait_stream / with torch.cuda.stream(...)qwen2_moe.pyqwen3_moe.pyglm4_moe.pybailing_moe.pyllada2.pygrok.pyolmo2.pystep3p5.pylongcat_flash.pyfalcon_h1.pymodel-specific overlap on attention prep, MoE branches, or cache-storeSearch these families before designing a new overlap scheme from scratch.

3. VLM-specific kernel families

PatternTrace keywordsPrimary codeExisting pathSkill should conclude
Vision QK norm with aux streamvision-side QK norm or norm-like kernels before attentionpython/sglang/srt/layers/attention/vision.pyvision QK normalization can call shared apply_qk_norm(...), with K-side work on aux_streamIf vision QK prep is split, first check this existing aux-stream path.
ViT CUDA graph disables vision aux streamexpected vision overlap is absent under ViT graphpython/sglang/srt/models/internvl.pypython/sglang/srt/layers/attention/vision.pypython/sglang/srt/environ.py::SGLANG_VIT_ENABLE_CUDA_GRAPHvision aux_stream is intentionally disabled when ViT CUDA graph is onMissing vision overlap may be intentional, not a regression.
Fused multimodal RoPE kerneltriton_mrope_fusedmultimodal_rotary_embedding_cpunpu_mropeMRotaryEmbeddingpython/sglang/srt/layers/rotary_embedding/mrope.pypython/sglang/srt/layers/rotary_embedding/triton_kernels.pypython/sglang/srt/models/qwen3.pyCUDA Triton, CPU sgl_kernel, and NPU paths already fuse multimodal t / h / w position lookup plus in-place Q / K rotary applicationIf VLM traces show separate mRoPE gather / shuffle / apply steps, first classify them as a missing existing mRoPE fusion.

4. Diffusion fused-kernel families

PatternTrace keywordsPrimary codeExisting pathSkill should conclude
Fused residual + norm + scale + shiftresidual add, norm, scale, shift, gate around DiT blockspython/sglang/kernels/ops/diffusion/norm/scale_residual_norm_cutedsl.pypython/sglang/multimodal_gen/runtime/layers/layernorm.pyfused_scale_residual_norm_scale_shift(...)Treat split residual + norm + modulation as a missing existing diffusion fusion first.
Fused norm + scale + shiftnorm followed by scale / shift elementwise kernelspython/sglang/kernels/ops/diffusion/norm/scale_residual_norm_cutedsl.pypython/sglang/multimodal_gen/runtime/layers/layernorm.pyfused_norm_scale_shift(...)Existing modulation fusion already covers this family.
Triton scale / shift and gate-select kernelstiny scale / shift or gate-select kernels dominate modulation blockspython/sglang/kernels/ops/diffusion/modulate/scale_shift_triton.pypython/sglang/multimodal_gen/runtime/layers/elementwise.pyfuse_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 RMSNormresidual add plus RMSNorm still split on short hidden sizespython/sglang/multimodal_gen/runtime/layers/layernorm.pypython/sglang/kernels/ops/diffusion/norm/rmsnorm_onepass_triton.pyfused_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 + RoPEsplit QK norm and RoPE in diffusion attention blockspython/sglang/kernels/ops/diffusion/rope/qknorm_rope_jit.pypython/sglang/multimodal_gen/runtime/layers/layernorm.py::apply_qk_norm_ropefused_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) + shiftfused_norm_tanh_mul_addtanh(gate) * rmsnorm(x)python/sglang/kernels/ops/diffusion/cutedsl/norm_tanh_mul_add_norm_scale.pypython/sglang/multimodal_gen/runtime/layers/layernorm.pyCuTeDSL kernel plus runtime helper for Z-Image residual-form modulationTreat split Z-Image residual-form modulation as a missing existing diffusion fusion, not a novel idea.
Z-Image fused residual modulation + next norm-scalefused_norm_tanh_mul_add_norm_scaleresidual + tanh(gate) * rmsnorm(x)ffn_norm1(x) * scale_mlppython/sglang/kernels/ops/diffusion/cutedsl/norm_tanh_mul_add_norm_scale.pypython/sglang/multimodal_gen/runtime/models/dits/zimage.pyOne CuTeDSL kernel fuses the first residual-form modulation and the next normalization / scale stageIf you see this chain split in Z-Image traces, report it as a missing existing mainline fusion family.
LTX2 fused Ada valuesltx2_ada_values9get_ada_valuesscale_shift_table + timestep.reshapepython/sglang/kernels/ops/diffusion/modulate/ltx2_ada_values_triton.pypython/sglang/multimodal_gen/runtime/models/dits/ltx_2.pyPR #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 blocksTreat repeated Ada add/reshape/slice ladders in LTX2 traces as a missing shipped SGLang fusion first.
LTX2 residual-gate adddiffusion_residual_gate_addresidual_gate_addresidual + update * gatepython/sglang/kernels/ops/diffusion/modulate/residual_gate_add_jit.pypython/sglang/kernels/jit/csrc/diffusion/residual_gate_add.cuhpython/sglang/multimodal_gen/runtime/models/dits/ltx_2.pyPR #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 passTreat split add/mul gate ladders in LTX2 traces as a missing shipped SGLang fusion first.
Nunchaku fused GELU MLP_fused_gelu_mlpfused_gelu_mlppython/sglang/multimodal_gen/runtime/models/dits/flux.pyNunchaku path fuses fc1 GEMM + GELU + shift + re-quant + fc2.lora_down before the second GEMMTreat split GELU-MLP on Nunchaku checkpoints as an existing fused family, not a new discovery.

5. Diffusion kernel-overlap and async-communication families

PatternTrace keywordsPrimary codeExisting pathSkill should conclude
Ulysses sequence-parallel attentionexposed all_to_all around attention blockspython/sglang/multimodal_gen/runtime/layers/attention/layer.pypython/sglang/multimodal_gen/runtime/distributed/communication_op.pyhead / sequence redistribution before and after attentionTreat sequence-parallel all-to-all as an existing distributed attention family.
USP attention with all-to-all and ring attentionall_to_all, ring-attention comm, head / sequence reshardspython/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 pipeliningpipelined A2A windows with explicit waits on a comm streampython/sglang/multimodal_gen/runtime/layers/attention/turbo_layer.pylooped all_to_all_single(..., async_op=True) plus staged postprocess on a comm streamTreat exposed turbo A2A windows as an existing pipelined overlap pattern.
TorchInductor compute / communication reordercompiled traces with compute and comm partially interleavedpython/sglang/multimodal_gen/runtime/pipelines_core/stages/denoising.pypython/sglang/multimodal_gen/runtime/pipelines_core/stages/model_specific_stages/mova.pytorch._inductor.config.reorder_for_compute_comm_overlap = TrueExisting compile-time reordering may already explain partial overlap in diffusion traces.
Dual-stream diffusion modelstwo nearby compute branches inside one DiT / UNet blockpython/sglang/multimodal_gen/runtime/models/dits/hunyuan3d.pyuse_dual_stream = TrueTreat 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.

PatternTrace keywordsPrimary codeExisting pathSkill should conclude
PR #21877 fused grouped down-GEMM + combinegrouped_gemm_nt_maskedcombinefused grouped gemm combinePR #21877python/sglang/srt/layers/moe/ep_moe/flashinfer_cutedsl_moe.pypython/sglang/srt/layers/moe/token_dispatcher/deepep.pyFlashInfer CuTeDSL kernel fuses the second expert GEMM with DeepEP low-latency combineTreat this as a concrete upstream MoE fuse / overlap family, not a new thought experiment.
PR #21889 fused BF16 to FP4 quant + paged KV writeset_mla_kv_buffer_fp4_quant_kernelfp4 kv cachePR #21889python/sglang/srt/mem_cache/utils.pyTriton kernel writes FP4 NSA KV pages directly while quantizing BF16 inputIf 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_kernelWRITE_PTdequant_fp4_paged_decodePR #21889python/sglang/srt/layers/attention/nsa/dequant_fp4_to_fp8.pyTriton kernel reads FP4 pages, writes FP8 directly, and can fuse decode-side page-table remapTreat this as an upstream in-flight decode-prep fusion family.
PR #21491 FlashInfer TRTLLM FP8 MoE with fused shared expertsnum_fused_shared_expertstrtllm_fp8_block_scale_moePR #21491python/sglang/srt/layers/moe/fused_moe_triton/fused_moe.pypython/sglang/srt/models/deepseek_v2.pyFlashInfer TRTLLM FP8 MoE path can fuse shared experts inside the routed MoE kernelOn 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 quantfused_add_rmsnorm_per_token_quantper_token_quant_fp8PR #22005python/sglang/kernels/jit/csrc/elementwise/fused_add_rmsnorm_per_token_quant.cuhpython/sglang/kernels/jit/fused_add_rmsnorm_per_token_quant.pyCUDA JIT kernel keeps normed values in registers and emits BF16 + FP8 outputs plus per-token scalesIf 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 writefused_qk_norm_rope_cache_pts_quant_shufflefused_qk_norm_mrope_3d_cache_pts_quant_shufflerotary_dimPR #20667python/sglang/srt/models/qwen3_5.pypython/sglang/srt/models/utils.pyROCm / AITER path fuses Q / K RMSNorm, partial or 3D RoPE, and direct KV cache write for Qwen3.5 attentionTreat 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 nvjetcutlass_scaled_mmfp8_scaled_mmnvjetcudaMemsetAsyncPR #22392python/sglang/kernels/aot/python/sgl_kernel/gemm.pypython/sglang/srt/layers/quantization/fp8_utils.pyRuntime replacement swaps nvjet FP8 GEMMs for CUTLASS kernels, removing per-launch memset bubbles and extra output-copy kernelsTreat 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+quantsilu_and_mul_scaled_nvfp4nvfp4 expert quantcutlass moePR #18612python/sglang/srt/layers/moe/cutlass_w4a8_moe.pypython/sglang/kernels/ops/quantization/nvfp4_gemm_swiglu_nvfp4_quant.pyFuses MoE activation epilogue and NVFP4 expert quantization before the CUTLASS MoE second GEMMTreat split SiLU+Mul then NVFP4 expert quant in CUTLASS MoE traces as an in-flight upstream SGLang family.
PR #22918 FlashInfer per-token NVFP4 MoEper_token_nvfp4trtllm_fp4_block_scale_moeFlashInfer MoEPR #22918python/sglang/srt/layers/moe/fused_moe_triton/fused_moe.pyAdds FlashInfer-backed per-token NVFP4 MoE execution so expert quant/dequant work can move into the fused MoE backendTreat 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 splitnsa topkflashinfer_topkpytorch_topkfast_topk_transformPR #22851python/sglang/srt/layers/attention/nsa_backend.pyMakes NSA top-k backend selection explicit and aligns fused top-k transform with FlashInfer / PyTorch fallbacksWhen NSA top-k dominates decode, first classify it as backend selection or fused-transform eligibility work.
PR #24125 GLM5 NSA decode CatArrayBatchedCopy removalCatArrayBatchedCopyGLM-5NSATileLang decodePR #24125python/sglang/srt/layers/attention/nsa_backend.pySkips redundant cat/copy work in the GLM5 NSA TileLang decode pathTreat cat/copy bursts in GLM5 NSA decode as a concrete in-flight cleanup opportunity.
PR #24007 MoE LoRA virtual experts for csgmv backendcsgmvvirtual expertsMoE LoRAfused_moe_loraPR #24007python/sglang/srt/layers/lora_backend.pypython/sglang/srt/layers/moeRoutes MoE LoRA adapter work through virtual experts so csgmv-style kernels can batch it instead of launching fragmented adapter workTreat MoE-LoRA tiny-kernel ladders as an in-flight batching/fusion family.
PR #24150 torch.compile local decode supportenable_torch_compilelocal compiledecode compiletorchinductorPR #24150python/sglang/srtExtends SGLang torch.compile coverage to local decode regions, so Inductor-generated fusion may replace hand-authored tiny kernelsWhen 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

PatternTrace keywordsPrimary codeExisting pathSkill should conclude
PR #21877 fused down-GEMM + combine superseding SBOenable_fused_grouped_gemm_combinecombinedown_gemmPR #21877python/sglang/srt/server_args.pypython/sglang/srt/layers/moe/token_dispatcher/deepep.pyFused combine eliminates the standalone combine window, so SBO is intentionally disabled when this path is onIf the trace discussion is about combine overlap, first classify it as this upstream fused-overlap family.
PR #23965 PDL for DSV32 / GLM5 kernelsenable_pdlTRTLLM_ENABLE_PDLcudaGridDependencySynchronizeDSV32GLM5PR #23965python/sglang/srt/layerssgl-kernelEnables programmatic dependent launch on selected DeepSeek / GLM kernels so dependent decode kernels can overlap launch-to-start gapsTreat tight same-stream decode windows around DSV32 / GLM5 as an in-flight PDL overlap family.
PR #21878 TTFT / TPOT torch.compile optimizationenable_torch_compiledecode graphpiecewise cudagraphPR #21878python/sglang/srtUses compiler and graph capture changes to shave TTFT / TPOT rather than adding one handwritten kernelIf 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 / embeddingslogprobsembeddingsGPU->CPU syncbatch syncPR #24168python/sglang/srtBatches per-request synchronization work that can otherwise serialize decode progress around logprob or embedding outputsTreat 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.

PatternTrace keywordsPrimary codeExisting pathSkill should conclude
FlashInfer activation / gate epiloguessilu_and_mulgelu_tanh_and_mulgelu_and_mulsilu_and_mul_scaled_nvfp4_experts_quantizeflashinfer/activation.pyflashinfer/quantization/fp4_quantization.pyFlashInfer covers both the plain activation-plus-mul epilogues and the NVFP4 expert-quantized extension used on MoE expert pathsTreat standalone activation, multiply, and expert-side quant ladders as one existing FlashInfer epilogue family first.
FlashInfer norm / residual / quant epiloguesrmsnorm_quantfused_add_rmsnormfused_add_rmsnorm_quantgemma_rmsnormgemma_fused_add_rmsnormfused_rmsnorm_silurmsnorm_fp4quantadd_rmsnorm_fp4quantflashinfer/norm/__init__.pyflashinfer/cute_dsl/rmsnorm_fp4quant.pyflashinfer/cute_dsl/add_rmsnorm_fp4quant.pyThe norm family spans plain RMSNorm derivatives, residual-add epilogues, norm+activation, and direct FP8 / NVFP4 output variants instead of materializing each intermediateTreat split residual add, norm, activation, and quant chains as one existing FlashInfer epilogue family first.
FlashInfer allreduce + post-op fusion familyallreduce_fusionAllReduceFusionPatternkARResidualRMSNormkARResidualRMSNormFP8QuantkARResidualRMSNormFP4Quanttrtllm_mnnvl_allreduce_fusionflashinfer/comm/allreduce.pyflashinfer/comm/trtllm_ar.pyflashinfer/comm/trtllm_mnnvl_ar.pyTRTLLM and MNNVL backends fuse all-reduce with residual add, RMSNorm, and backend-appropriate quant / norm-output variantsTreat TP collective + norm (+ quant) ladders as an existing FlashInfer fused-collective family first.
FlashInfer RoPE + FP8 quant / cache-update familyrope_quantize_fp8mla_rope_quantize_fp8rope_quantize_fp8_append_paged_kv_cacheseqlen=0batch_indices < 0flashinfer/rope.pyThe 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 handlingTreat split RoPE, quant, cache-write, and padding-token ladders as one existing FlashInfer attention-prep family first.
FlashInfer fused DeepSeek grouped-topk routingfused_topk_deepseekNoAuxTcflashinfer/fused_moe/fused_routing_dsv3.pyOne kernel performs sigmoid+bias, grouped score reduction, group top-k, expert top-k, and routed renorm for DeepSeek-V3-style routingTreat router score activation -> grouped top-k -> renorm ladders as an existing FlashInfer router family first.
FlashInfer fused MoE expert executioncutlass_fused_moetrtllm_bf16_moetrtllm_fp8_per_tensor_scale_moetrtllm_fp8_block_scale_moetrtllm_fp4_block_scale_moetrtllm_mxint4_block_scale_moenon-gatedflashinfer/fused_moe/core.pyCUTLASS and TRTLLM backends collapse expert execution, routed combine, and quantized expert variants into fused MoE runners, including gated and non-gated FP8 per-tensor casesTreat exposed expert-side tiny GEMM or non-gated FP8 ladders as matching an existing FlashInfer fused-MoE family.
FlashInfer CuTeDSL two-stage MoE fusionblockscaled_contiguous_gather_grouped_gemm_swiglu_fusion_nvfp4blockscaled_contiguous_grouped_gemm_finalize_fusion_nvfp4moe_permutemoe_unpermuteflashinfer/fused_moe/cute_dsl/blockscaled_contiguous_gather_grouped_gemm_swiglu_fusion.pyflashinfer/fused_moe/cute_dsl/blockscaled_contiguous_grouped_gemm_finalize_fusion.pyThe 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 kernelsTreat multi-kernel MoE ladders around permute / finalize as one existing FlashInfer CuTeDSL family first.
FlashInfer SM120 FP4 / groupwise GEMM heuristicscutlass_fp4_gemm_sm120CutlassTileConfigSM120group_gemm_nvfp4_nt_groupwisegroup_gemm_mxfp4_nt_groupwiseflashinfer/gemm/gemm_base.pyinclude/flashinfer/gemm/fp4_gemm_cutlass_template_sm120.hinclude/flashinfer/gemm/group_gemm_nvfp4_groupwise_sm120.cuhcsrc/nv_internal/tensorrt_llm/kernels/cutlass_kernels/cutlass_heuristic.cppFlashInfer mainline adds SM120-oriented FP4 GEMM selection and b12x CuTeDSL fused-MoE kernelsTreat 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 supportrouting_replay_outmPtrRoutingReplayOuttrtllm_fp8_block_scale_moeflashinfer/fused_moe/core.pycsrc/trtllm_fused_moe_kernel_launcher.cucsrc/fused_moe/noAuxTcKernels.cuTRTLLM-gen MoE kernels can optionally emit compact routing replay metadata without a separate routing-side reconstruction passTreat 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

PatternTrace keywordsPrimary codeExisting pathSkill should conclude
FlashInfer PDL launch-overlap familyenable_pdllaunch_with_pdlcudaGridDependencySynchronizecudaTriggerProgrammaticLaunchCompletiontrigger_completion_at_end=Falseallreduce_fusionflashinfer/norm/__init__.pyflashinfer/activation.pyflashinfer/rope.pyflashinfer/comm/allreduce.pyflashinfer/comm/trtllm_ar.pyFlashInfer uses Programmatic Dependent Launch broadly, and the allreduce path can further advance completion so the next PDL-aware kernel overlaps on the same streamTreat 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 overlapaux_streammain_eventmemset_eventuse_async_memsetflashinfer/fused_moe/cute_dsl/fused_moe.pyPreallocated MoE output is zeroed on an auxiliary CUDA stream while GEMM1 runs on the main stream, then both streams join before finalizeTreat GEMM1 vs output-zero windows as an existing FlashInfer multi-stream overlap family.
FlashInfer green-context SM partition overlapsplit_device_green_ctxsplit_device_green_ctx_by_sm_countgreen_ctxflashinfer/green_ctx.pyCUDA green contexts partition SMs and create dedicated streams for concurrent kernel families on separate SM slicesTreat 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

PatternTrace keywordsPrimary codeExisting pathSkill should conclude
PR #2720 PDL runtime-API migrationcudaGridDependencySynchronizecudaTriggerProgrammaticLaunchCompletioninline PTXPR #2720include/flashinfer/comm/trtllm_allreduce_fusion.cuhinclude/flashinfer/pos_enc.cuhRepo-wide migration preserves the existing PDL overlap family while replacing inline PTX with CUDA runtime APIs across norm, RoPE, attention, and MoE codepathsTreat 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.

PatternTrace keywordsPrimary codeExisting pathSkill should conclude
TensorRT-LLM FlashInfer activation / gate epiloguesflashinfer_silu_and_mulflashinfer_gelu_tanh_and_mulauto_deploy::silu_and_mulpost-GEMM silu + multensorrt_llm/_torch/custom_ops/flashinfer_custom_ops.pytensorrt_llm/_torch/auto_deploy/transform/library/fuse_silu_mul.pytensorrt_llm/_torch/models/modeling_gemma3.pyRuntime custom ops and AutoDeploy rewrite split/getitem + activation + mul MLP epilogues into one FlashInfer op, including Gemma3 gelu_tanh_and_mulTreat split gate activation + multiply as an existing TensorRT-LLM/FlashInfer epilogue family first.
TensorRT-LLM FlashInfer RMSNorm familyflashinfer_rmsnormflashinfer_gemma_rmsnormauto_deploy::flashinfer_rms_normtensorrt_llm/_torch/custom_ops/flashinfer_custom_ops.pytensorrt_llm/_torch/modules/rms_norm.pytensorrt_llm/_torch/auto_deploy/custom_ops/normalization/rms_norm.pyRuntime modules and AutoDeploy can lower plain RMSNorm and Gemma RMSNorm directly to FlashInfer kernelsTreat split RMSNorm ladders as an existing TensorRT-LLM norm family before calling them novel.
TensorRT-LLM FlashInfer residual add + RMSNormflashinfer_fused_add_rmsnormflashinfer_gemma_fused_add_rmsnormauto_deploy::flashinfer_fused_add_rms_norm_inplacetensorrt_llm/_torch/custom_ops/flashinfer_custom_ops.pytensorrt_llm/_torch/modules/rms_norm.pytensorrt_llm/_torch/auto_deploy/transform/library/fused_add_rms_norm.pyResidual add immediately before RMSNorm can collapse to one in-place FlashInfer op, with Gemma variant supportTreat residual add + RMSNorm chains as an existing TensorRT-LLM fused epilogue family first.
TensorRT-LLM Triton fused residual add + RMSNorm + FP8 quanttriton_fused_add_rms_norm_quant_fp8fuse_rmsnorm_quant_fp8fp8 static quanttensorrt_llm/_torch/auto_deploy/custom_ops/normalization/triton_fused_add_rms_norm_quant_fp8.pytensorrt_llm/_torch/auto_deploy/transform/library/fuse_rmsnorm_quant_fp8.pyMainline 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 togetherTreat split add + norm + FP8 quant ladders as an existing TensorRT-LLM mainline family first.
TensorRT-LLM FlashInfer RoPE with shared cos/sin cacheflashinfer_apply_rope_with_cos_sin_cache_inplaceflashinfer_ropecos_sin_cachetensorrt_llm/_torch/modules/rotary_embedding.pytensorrt_llm/_torch/auto_deploy/custom_ops/rope/flashinfer_rope.pytensorrt_llm/_torch/auto_deploy/transform/library/rope.pyRuntime 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_ropeTreat separate cos/sin gather + RoPE application ladders as an existing TensorRT-LLM attention-prep family.
TensorRT-LLM FlashInfer cached paged attentionappend_paged_kv_cacheBatchPrefillWithPagedKVCacheWrapperBatchDecodeWithPagedKVCacheWrapperauto_deploy::flashinfer_attention_mha_with_cacheread_cache_onlytensorrt_llm/_torch/attention_backend/flashinfer.pytensorrt_llm/_torch/auto_deploy/custom_ops/attention/flashinfer_attention.pydocs/source/features/attention.mdFlashInfer attention backend fuses metadata setup, optional paged-KV append, and prefill/decode wrapper execution, including shared-KV and read-cache-only variants in AutoDeployTreat metadata + KV-append + cached-attention ladders as one existing TensorRT-LLM cached-attention family first.
TensorRT-LLM FlashInfer MLA regular prefillappend_paged_mla_kv_cacheBatchPrefillWithRaggedKVCacheWrapperflashinfer_mlarank 256gpu append kerneltensorrt_llm/_torch/auto_deploy/custom_ops/mla/flashinfer_mla.pyRegular 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 pathTreat MLA regular-prefill prep as an existing TensorRT-LLM FlashInfer family first.
TensorRT-LLM FlashInfer MLA chunked prefill with absorbed W_knBatchMLAPagedAttentionWrapperchunked prefillW_knW_vtensorrt_llm/_torch/auto_deploy/custom_ops/mla/flashinfer_mla.pyChunked prefill absorbs W_kn into the query-side projection, runs paged MLA attention in compressed space, then projects back with W_vTreat 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_vplan_decodeBatchMLAPagedAttentionWrapperdecodeW_knW_vtensorrt_llm/_torch/auto_deploy/custom_ops/mla/flashinfer_mla.pyDecode path reuses the absorbed-query MLA family and projects the compressed attention output back with W_vTreat similar decode-time absorbed MLA ladders as an existing TensorRT-LLM family, not a new idea.
TensorRT-LLM FlashInfer fused MoE backendflashinfer.fused_moetrtllm_bf16_moetrtllm_fp8_block_scale_moetrtllm_fp4_block_scale_moeTRTLLM_GEN_FUSED_MOE_USE_FLASHINFERtensorrt_llm/_torch/modules/fused_moe/moe_op_backend.pytensorrt_llm/_torch/modules/fused_moe/fused_moe_trtllm_gen.pyTRTLLM-gen MoE can route expert execution and quant helpers through FlashInfer instead of exposing per-expert eager laddersTreat expert-side tiny GEMM ladders as matching an existing TensorRT-LLM FlashInfer MoE family first.
TensorRT-LLM FlashInfer cached SSM / Mamba updateflashinfer_cached_ssmselective_state_updateflashinfer_ssmtensorrt_llm/_torch/auto_deploy/custom_ops/mamba/flashinfer_backend_mamba.pytensorrt_llm/_torch/modules/mamba/mamba2_mixer.pyMamba2 paths can lower cached SSM state updates to FlashInfer selective-state-update kernels instead of many smaller state opsTreat split cached-SSM state update ladders as an existing TensorRT-LLM FlashInfer family first.

12. TensorRT-LLM-origin kernel-overlap families

PatternTrace keywordsPrimary codeExisting pathSkill should conclude
TensorRT-LLM multi-stream MLA attentionmulti_stream_mla_attnrecord_event_passthrough_auxwait_eventtensorrt_llm/_torch/auto_deploy/transform/library/multi_stream_attn.pytensorrt_llm/_torch/auto_deploy/utils/multi_stream_utils.pyAutoDeploy rewrites MLA Q/KV forks so the KV projection runs on an auxiliary stream while the Q path stays on the caller streamTreat exposed Q-branch vs KV-branch overlap as an existing TensorRT-LLM multi-stream family first.
TensorRT-LLM multi-stream MoE shared-vs-routed overlapmulti_stream_moebegin_aux_stream_passthroughend_aux_stream_passthroughwait_aux_stream_passthroughmlir_elementwise_fusionpiecewise cudagraphcaller_stream.synchronize()tensorrt_llm/_torch/auto_deploy/transform/library/multi_stream_moe.pytensorrt_llm/_torch/auto_deploy/utils/multi_stream_utils.pyShared-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 replayTreat 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 parallelismmulti_stream_gemmtrtllm_finegrained_fp8_linearrecord_event_passthrough_auxtensorrt_llm/_torch/auto_deploy/transform/library/multi_stream_gemm.pytensorrt_llm/_torch/auto_deploy/utils/multi_stream_utils.pyCompiler pass identifies fork points with multiple FP8 linears and moves the largest GEMM to the auxiliary stream so sibling GEMMs overlapTreat 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

PatternTrace keywordsPrimary codeExisting pathSkill should conclude
PR #12525 FlashInfer TRTLLM-gen FMHA paged-index / buffer reworkshared paged indextrtllm-gen attentionflashinferkv cache bufferPR #12525tensorrt_llm/_torch/auto_deploy/custom_ops/attention/flashinfer_attention.pyOpen PR refines the existing FlashInfer TRTLLM-gen cached-attention family by disabling shared paged index and unifying KV-buffer constructionTreat these attention-prep changes as an in-flight implementation evolution of an existing family first.
PR #12544 NVFP4 KV cache support in TRTLLM-gen attentionNVFP4 KV cachetrtllm-gen attentionflashinferPR #12544tensorrt_llm/_torch/auto_deploy/custom_ops/attention/flashinfer_attention.pyOpen PR extends the cached-attention family so the FlashInfer-backed TRTLLM-gen path can build and consume NVFP4 KV buffers directlyTreat split KV-cache quant + buffer-build ladders as an in-flight TensorRT-LLM attention family first.
PR #12738 / #12557 BF16 TRTLLM-gen MoE through FlashInferbf16 trtllm-gen moeflashinfertrtllm_bf16_moePR #12738PR #12557tensorrt_llm/_torch/modules/fused_moe/fused_moe_trtllm_gen.pyOpen PRs extend the TRTLLM-gen MoE family so BF16 expert execution can route through FlashInfer instead of only CUTLASS-like pathsTreat 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.

PatternTrace keywordsPrimary codeExisting pathSkill should conclude
vLLM-origin fused residual add + RMSNormfused_add_rms_norm*residual add right before RMSNormvllm/model_executor/layers/layernorm.pyvllm/_custom_ops.pycsrc/layernorm_kernels.cucsrc/cpu/layernorm.cppCustom CUDA / CPU fused add-RMSNorm op reused directly and as a building block for later compile-time fusionsTreat 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_rmsAllReduceFusionPassallreduce + rmsnormvllm/compilation/passes/fusion/allreduce_rms_fusion.pydocs/design/fusions.mdCompile-time patterns cover AllReduce -> RMSNorm(+residual_add) and optional FP8 / NVFP4 quant suffixesTreat TP collective + norm (+ quant) ladders as a known vLLM-origin fusion family first.
vLLM-origin RMSNorm (+ residual add) + quantRMSNormQuantFusionPassfused_add_rms_norm_static_fp8_quantper_token_quantper_group_quantvllm/compilation/passes/fusion/rms_quant_fusion.pyvllm/compilation/passes/fusion/rocm_aiter_fusion.pyCompile-time and ROCm AITER paths fuse RMSNorm or fused-add-RMSNorm with FP8 / FP4 quant outputTreat split norm/add + quant as an upstream fused family, not an unexplored direction.
vLLM-origin SiLU+Mul + quantActivationQuantFusionPassSiluMulFp8*Nvfp4rocm_aitervllm/compilation/passes/fusion/act_quant_fusion.pyvllm/compilation/passes/fusion/rocm_aiter_fusion.pyActivation epilogues fuse SiLU+Mul with FP8 / NVFP4 / AITER group quant instead of materializing the BF16 activation firstTreat standalone activation then quant kernels as matching a vLLM-origin precedent.
vLLM-origin add + RMSNorm + padfuse_act_paddingRocmAiterTritonAddRMSNormPadFusionPassadd_rmsnorm_padvllm/compilation/passes/fusion/rocm_aiter_fusion.pydocs/design/fusions.mdROCm / AITER path fuses residual add + RMSNorm directly into the padded layout expected by the next kernelTreat norm-plus-padding ladders as an existing backend-specific fuse family first.
vLLM-origin attention + output quantfuse_attn_quantAttnQuantFusionPassmerge_attn_statesoutput_scaleoutput_group_scaleoutput_block_scalevllm/compilation/passes/fusion/attn_quant_fusion.pyvllm/v1/attention/ops/merge_attn_states.pycsrc/attention/merge_attn_states.cudocs/design/fusions.mdCompile-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 providedTreat attention-output quant and merged-attention quant epilogues as a known upstream family before calling them novel.
vLLM-origin fused QK RMSNorm + RoPEfused_qk_norm_ropeQKNormRoPEFusionPassqk norm + ropevllm/compilation/passes/fusion/qk_norm_rope_fusion.pyvllm/_custom_ops.pycsrc/fused_qknorm_rope_kernel.cuCompile-time and direct custom-op paths fuse per-head Q / K RMSNorm with RoPETreat split QK norm + RoPE as a clear vLLM-origin precedent.
vLLM-origin fused reshape + KV cache writereshape_and_cachetriton_reshape_and_cache_flashkv cache writevllm/v1/attention/ops/triton_reshape_and_cache_flash.pyvllm/v1/attention/backends/triton_attn.pyTriton cache-update kernels reshape K / V into paged-cache layout and can include FP8 KV-cache scale/write logicTreat reshape / transpose / cache-write ladders as an existing cache-store fusion family.
vLLM-origin fused RoPE + KV cache updatefuse_rope_kvcacheRopeKVCacheFusionPasstriton_rope_and_cachevllm/compilation/passes/fusion/rope_kvcache_fusion.pyvllm/_aiter_ops.pydocs/design/fusions.mdROCm / AITER compile-time fusion combines RoPE with paged KV cache update instead of launching them separatelyTreat split RoPE + cache-store as a known upstream family, especially on ROCm-like paths.
vLLM-origin fused MLA RoPE + unified KV-cache updatefused_rope_unified_mla_kv_cache_updateconcat_and_cache_mla_rope_fusedunified_mla_kv_cache_updatevllm/compilation/passes/fusion/mla_rope_kvcache_cat_fusion.pyvllm/_custom_ops.pycsrc/cache_kernels_fused.cuCurrent vLLM compile pass fuses MLA-oriented RoPE on q_pe / k_pe, concat, and unified MLA KV-cache update into a direct paged-store pathTreat 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 routergrouped_topkbiased_grouped_topkgrouped_topk_fused_kernelvllm/_custom_ops.pyvllm/_aiter_ops.pyvllm/model_executor/layers/fused_moe/router/grouped_topk_router.pycsrc/moe/grouped_topk_kernels.cuCUDA / ROCm router kernels fuse grouped score processing, top-k selection, and routed renorm / bias handlingTreat MoE router ladders as matching an upstream grouped-topk family first.
vLLM-origin fused top-k softmax / sigmoid routertopk_softmaxtopk_sigmoidtopkGatingfused_topkvllm/_custom_ops.pyvllm/_aiter_ops.pyvllm/model_executor/layers/fused_moe/router/fused_topk_router.pyvllm/model_executor/layers/fused_moe/router/fused_topk_bias_router.pycsrc/moe/topk_softmax_kernels.cuCUDA 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 laddersTreat standalone score activation -> top-k -> bias / renorm chains as a known upstream fused router family first.
vLLM-origin DSV3 router GEMMdsv3_router_gemmallow_dsv3_router_gemmrouter logitsvllm/_custom_ops.pyvllm/model_executor/layers/fused_moe/router/gate_linear.pycsrc/moe/dsv3_router_gemm_entry.cucsrc/moe/dsv3_router_gemm_float_out.cuHopper-class CUDA kernel specializes the DeepSeek router linear for small decode batches and can emit FP32 logits directly without a generic GEMM chainTreat DeepSeek-style router linear paths as an existing upstream specialized fuse, distinct from grouped-topk itself.
vLLM-origin DeepSeek-V4 fused norm + router GEMMdsv4_norm_router_gemmnorm_gate_linearrouter_gemmDeepseekV4ForCausalLMvllm/model_executor/layers/fused_moe/router/norm_gate_linear.pyvllm/model_executor/models/deepseek_v4.pycsrc/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 kernelsTreat DSV4 norm + router ladders as an upstream mainline fused-router family.
vLLM-origin DeepSeek-V4 MHC fused kernelsmhc_post_prehead_compute_mix_kernelaiter mhcMHCvllm/model_executor/layers/mhc.pyvllm/model_executor/kernels/mhc/aiter.pyvllm/model_executor/models/deepseek_v4.pyCurrent vLLM mainline includes Tile head-compute kernels, fused mhc_post_pre, and ROCm AITER MHC support for DSV4Treat DSV4 MHC/head-compute ladders as known upstream kernel families before proposing a new head-compute kernel.
vLLM-origin GPT-OSS router GEMMgpt_oss_router_gemmrouter gemmvllm/_custom_ops.pyvllm/model_executor/layers/fused_moe/router/gate_linear.pycsrc/moe/gpt_oss_router_gemm.cuModel-specific CUDA kernel replaces the router linear plus bias path with one specialized GEMM opTreat GPT-OSS-style router linear chains as an existing upstream specialized fuse.
vLLM-origin DeepSeek min-latency fused QKV-A projectiondsv3_fused_a_gemmfused_qkv_a_projq_a_projvllm/model_executor/models/deepseek_v2.pyvllm/_custom_ops.pycsrc/dsv3_fused_a_gemm.cuHopper-class CUDA kernel replaces the tiny-batch DeepSeek QKV-A projection path with one specialized min-latency GEMM instead of a generic linear launchTreat small-batch DeepSeek QKV-A projection ladders as a known upstream fused kernel family first.
vLLM-origin DSV3.2 fused indexer projectionswk_weights_projMergedColumnParallelLinearweights_projvllm/model_executor/models/deepseek_v2.pyvllm/model_executor/models/deepseek_mtp.pyDSV3.2 indexer paths can fuse the wk and weights_proj projections into one GEMM and carry the matching MTP weight-loading pathTreat paired indexer projection chains as a known upstream fused linear family before calling the opportunity novel.
vLLM-origin ROCm AITER sparse-MLA paged MQA logitsrocm_aiter_mla_sparsepaged_mqa_logitsgluongfx950vllm/v1/attention/ops/rocm_aiter_mla_sparse.pyCurrent vLLM enables the AITER/Gluon paged-MQA logits path on gfx950 / MI355X sparse MLA shapesOn 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 kernelsrocm_aiter_mla_sparse_dsv4flashmla_sparsesparse_swavllm/v1/attention/backends/mla/rocm_aiter_mla_sparse_dsv4.pyvllm/v1/attention/backends/mla/flashmla_sparse.pyvllm/v1/attention/backends/mla/sparse_swa.pyCurrent vLLM mainline has ROCm DSV4 sparse-MLA Triton backend coverageOn DSV4 ROCm sparse-MLA traces, compare backend selection and sparse-SWA paths before writing a new MLA backend.
vLLM-origin DSV4 dequant gather K cachedequant_gather_k_cutedslfused_indexer_q_cutedslcache_utilsvllm/v1/attention/ops/deepseek_v4_ops/dequant_gather_k_cutedsl.pyvllm/v1/attention/ops/deepseek_v4_ops/cache_utils.pyCurrent vLLM mainline has a newer CuTe DSL dequant-gather K cache path for DSV4Treat K-cache dequant/gather ladders as a known upstream DSV4 cache-kernel family.
vLLM-origin TokenSpeed MLA backendTOKENSPEED_MLAtokenspeed_mla_decodetokenspeed_mla_prefillvllm/v1/attention/backends/mla/tokenspeed_mla.pyvllm/v1/attention/backends/mla/prefill/tokenspeed_mla.pyvllm/model_executor/layers/attention/mla_attention.pyCurrent vLLM can select TokenSpeed MLA prefill/decode kernels for Blackwell FP8-KV DeepSeek/Kimi MLA shapes when the package and platform gates passOn Blackwell MLA traces, compare against TokenSpeed MLA backend selection before proposing a new attention kernel.
vLLM-origin MiniMax allreduce_rms kernelsminimax_allreduce_rmsminimax_allreduce_rmsnormMiniMax-M2.5allreduce_rmsvllm/model_executor/models/minimax_m2.pyTensorRT-LLM-derived MiniMax allreduce-plus-RMSNorm kernels are a concrete upstream TP decode familyTreat MiniMax TP norm + collective ladders as an upstream specialized fusion family.
vLLM-origin CUTLASS scaled MM with scale / bias epiloguecutlass_scaled_mmcutlass_scaled_mm_azpscaled mmvllm/_custom_ops.pyvllm/model_executor/kernels/linear/scaled_mm/cutlass.pycsrc/libtorch_stable/quantization/w8a8/cutlass/scaled_mm_entry.cuCUTLASS kernels fuse activation scales, weight scales, matmul, and optional bias / AZP epiloguesTreat separate scale-mul + GEMM + bias ladders as a vLLM-origin fused linear family first.
vLLM-origin fused MoE expert executioncpu_fused_moerocm_aiter_fused_moeFusedMoEvllm/model_executor/layers/fused_moe/layer.pyvllm/model_executor/layers/fused_moe/cpu_fused_moe.pyvllm/model_executor/layers/fused_moe/rocm_aiter_fused_moe.pyvllm/_aiter_ops.pyMoE backends on CUDA / ROCm / CPU already collapse packed expert execution into fused expert kernels rather than per-expert eager GEMMsTreat exposed expert-side tiny GEMM ladders as matching an upstream fused-MoE family.
vLLM-origin fused MoE LoRAfused_moe_lorafused_moe_lora_fp8w13_shrinkw2_expandvllm/lora/ops/triton_ops/fused_moe_lora_op.pyvllm/lora/ops/triton_ops/fused_moe_lora_fp8_op.pyvllm/lora/layers/fused_moe.pyTriton kernels fuse LoRA shrink / expand work into MoE expert execution, including FP8 variantsTreat MoE-LoRA adapter work as an upstream fused family before proposing a brand new kernel.
vLLM-origin ViT fused bilinear position-embedding interpolationtriton_pos_embed_interpolatebilinear_pos_embedpos_embed_interpolate_nativevllm/model_executor/models/qwen3_vl.pyTriton kernel fuses bilinear interpolation and spatial-merge reorder for Qwen3-VL ViT position embeddings, replacing many tiny eager kernelsTreat 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.

PatternTrace keywordsPrimary codeExisting pathSkill should conclude
TokenSpeed CuTe DSL MLA prefill / decodetokenspeed_mla_decodetokenspeed_mla_prefillBlackwellMultiHeadLatentAttentionForwardpython/tokenspeed/runtime/layers/attention/backends/tokenspeed_mla.pytokenspeed-mla/python/tokenspeed_mla/mla_decode.pytokenspeed-mla/python/tokenspeed_mla/mla_prefill.pytokenspeed-kernel/python/tokenspeed_kernel/ops/attention/tokenspeed_mla/__init__.pyBlackwell SM100 CuTe DSL MLA kernels cover FP8-KV prefill/decode/verify paths through the tokenspeed_mla backendOn 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_kernelmla_kv_pack_quantize_fp8k_nope / k_petokenspeed-mla/python/tokenspeed_mla/mla_kv_pack_quantize_fp8.pytokenspeed-kernel/python/tokenspeed_kernel/ops/attention/tokenspeed_mla/__init__.pyOne Triton kernel packs k_nope, broadcast k_pe, and v, then writes FP8 K/V for MLA chunked prefillTreat split K/V concat + FP8 cast ladders as a known TokenSpeed fusion family.
TokenSpeed fused top-k + top-p samplingfused_topk_toppfused_topk_topp_renormtokenspeed-kernel/python/tokenspeed_kernel/thirdparty/cuda/fused_topk_topp.pytokenspeed-kernel/python/tokenspeed_kernel/thirdparty/cuda/csrc/fused_topk_topp/fused_topk_topp.cuCUDA extension fuses top-k, top-p, and renormalization for decode samplingTreat top-k/top-p/renorm chains in TokenSpeed traces as an existing sampling fusion first.
TokenSpeed persistent lm_head GEMMlm_head_gemmshould_use_fusedpersistenttokenspeed-kernel/python/tokenspeed_kernel/thirdparty/cuda/lm_head_gemm.pytokenspeed-kernel/python/tokenspeed_kernel/thirdparty/cuda/csrc/lm_head_gemm.cuShape-gated persistent GEMM replaces torch.matmul for selected lm_head / router-like projection shapesTreat visible lm_head matmul ladders as a candidate for this existing TokenSpeed path before inventing a new logits GEMM.
TokenSpeed NVFP4 GEMM + SwiGLU + quantnvfp4_gemm_swiglu_nvfp4_quantSwiGLUSFCtokenspeed-kernel/python/tokenspeed_kernel/thirdparty/cute_dsl/nvfp4_gemm_swiglu_nvfp4_quant.pyCuTe DSL kernel fuses block-scaled NVFP4 GEMM, SwiGLU, and optional output quantizationTreat split expert GEMM + activation + FP4 quant chains as matching an upstream TokenSpeed kernel family.

16. vLLM-origin kernel-overlap families

PatternTrace keywordsPrimary codeExisting pathSkill should conclude
vLLM-origin AsyncTP GEMM + collective overlapfuse_gemm_commsfused_matmul_reduce_scatterfused_all_gather_matmulvllm/compilation/passes/fusion/collective_fusion.pydocs/design/fusions.mdAsyncTP overlaps GEMM with reduce-scatter / all-gather via symmetric-memory collectivesTreat GEMM+comm windows as a clear vLLM-origin overlap precedent first.
vLLM-origin Sequence Parallelism stagingenable_spReduceScatterAllGatherSequenceParallelismPassvllm/compilation/passes/fusion/sequence_parallelism.pydocs/design/fusions.mdSequence-parallel rewrites all-reduce into RS -> local norm -> AG so later passes can overlap comm and computeTreat RS / AG staging around norm blocks as an upstream overlap-enabling family.
vLLM-origin shared-expert aux-stream overlapaux_streamshared_experts_streamshared expert near routervllm/model_executor/layers/fused_moe/runner/shared_experts.pyvllm/model_executor/layers/fused_moe/runner/moe_runner_base.pyMoE 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 mergeTreat shared-expert vs router overlap as an existing upstream sparse-model family.
vLLM-origin DCP async all-to-all overlapdcp_alltoallall_to_all_singleasync_op=Truevllm/v1/attention/ops/dcp_alltoall.pyOutput / LSE exchange uses async all-to-all handles instead of serializing collective completion on the main pathTreat DCP all-to-all windows as an upstream async-collective family.

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

PatternTrace keywordsPrimary codeExisting pathSkill should conclude
PR #35968 DSV3.2 multi-stream indexer overlapweights_projwkk_normaux_streamPR #35968vllm/model_executor/models/deepseek_v2.pyvllm/utils/torch_utils.pyClosed 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 streamTreat 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 quantgroup_size=128group_size=64output_group_scaleper-group FP8PR #37110vllm/compilation/passes/fusion/attn_quant_fusion.pyvllm/v1/attention/ops/triton_unified_attention.pyIn-flight Triton attention epilogue computes per-group FP8 scales and quantizes output directly instead of launching a separate group-quant kernelTreat attention + per-group FP8 quant as a concrete upstream vLLM family, not a novel idea.
PR #38445 MiniMax-M2 FP32 gate kernelfp32_router_gemmMiniMax-M2gate kernelPR #38445vllm/model_executor/layers/fused_moe/router/gate_linear.pyvllm/model_executor/models/minimax_m2.pyDraft CUDA kernel fuses BF16->FP32 conversion and low-batch router GEMM for MiniMax-M2, replacing up to three kernels on the gate pathTreat MiniMax-M2 gate ladders as an in-flight upstream fused router family first.
PR #38621 fused QK norm + RoPE + cache + quantfused_qk_norm_rope_cache_quantQK Norm + RoPE + Cache + QuantPR #38621csrc/fused_qk_norm_rope_cache_quant.cuvllm/compilation/passes/fusion/qk_norm_rope_cache_quant_fusion.pyDraft CUDA kernel and compile-time pass try to fuse QK RMSNorm, RoPE, KV cache write, and optional FP8 quant for small-batch decodeTreat this as an in-flight upstream fusion family before calling a similar idea novel.
PR #37646 ROCm AITER fused allreduce + RMSNormrocm_aiter_fused_allreduce_rmsnormcustom_fused_ar_rmsRocmAiterAllReduceFusionPassPR #37646vllm/_aiter_ops.pyvllm/compilation/passes/pass_manager.pyROCm-specific compile-time path swaps the generic all-reduce fusion pass for an AITER fused allreduce-plus-RMSNorm kernel familyTreat ROCm TP all-reduce + RMSNorm ladders as an in-flight upstream fused-collective family first.
PR #36413 FlashInfer RMSNorm + FP4 quant fusionfuse_norm_quantflashinferNVFP4rmsnorm + fp4 quantPR #36413vllm/compilation/passes/fusion/rms_quant_fusion.pyvllm/docs/design/fusions.mdFlashInfer-backed norm-plus-FP4 quant fusion extends the existing RMSNorm+quant family to NVFP4 flowsTreat split RMSNorm + FP4 quant ladders as an upstream in-flight family, not a fresh idea.
PR #39301 GLM5 router GEMM with PDL overlapTRTLLM_ENABLE_PDLrouter_gemmGLM5FI AR RMS fusionPR #39301vllm/model_executor/layers/fused_moe/router/gate_linear.pycsrc/moe/dsv3_router_gemm_utils.hExtends the specialized router GEMM family to GLM5 hidden size and uses PDL to overlap the router launch with the preceding fused allreduce-plus-RMS blockTreat 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 decodewmmapaged prefillsplit-K decodeROCm attentionPR #41455vllm/v1/attentionvllm/_aiter_ops.pyAdds ROCm WMMA attention kernels for paged prefill and split-K decode shapesTreat 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 pathDSV4fuse norm routerlow latencyrouterPR #41263vllm/model_executor/models/deepseek_v4.pyvllm/model_executor/layers/fused_moe/router/norm_gate_linear.pyMerged into current mainline as the DSV4 fused norm + router GEMM family aboveTreat this row as provenance for the shipped mainline family, not as merely in-flight.
PR #41428 DSV4 fused indexer Q quant kernelDSV4fused Indexer Q quantindexer qfp4PR #41428vllm/model_executor/models/deepseek_v4.pyvllm/v1/attention/ops/deepseek_v4_ops/fused_indexer_q.pyvllm/v1/attention/ops/deepseek_v4_ops/fused_indexer_q_cutedsl.pyMerged current-main improvement to the fused DeepSeek-V4 indexer Q quant kernel instead of materializing Q then quantizing separatelyTreat DSV4 indexer-Q quant ladders as an upstream mainline fused quant family.
PR #41255 DeepSeek-V4 Tile kernels / head_compute_mix_kernelhead_compute_mix_kernelTile kernelDSV4MLAPR #41255vllm/model_executor/models/deepseek_v4.pyvllm/model_executor/layers/mhc.pyvllm/model_executor/kernels/mhc/tilelang.pyMerged current-main DeepSeek-V4 Tile kernels that mix head compute work in one specialized kernelTreat DSV4 MLA head-compute ladders as a known upstream mainline specialized-kernel family.
PR #41441 DSV4 all-reduce plus mhc_post fusionDSV4AR+mhc_postallreducemhc_postPR #41441vllm/model_executor/models/deepseek_v2.pyvllm/compilation/passes/fusionFuses or overlaps DSV4 all-reduce with post-MLA head-compute workTreat all-reduce followed by mhc_post in DSV4 traces as an in-flight vLLM overlap/fusion family.
PR #41446 AMD GatedDeltaNet FLA prefill kernelsGatedDeltaNetFLA prefillAMDQwen3-NextPR #41446vllm/model_executor/models/qwen3_next.pyvllm/v1/attentionOptimizes GatedDeltaNet / FLA prefill kernels on AMD linear-attention modelsTreat split GDN prefill kernels on ROCm as an in-flight upstream family.
PR #39748 dual-stream GDN input projectiondual-streaminput projectionGatedDeltaNetQwen3.5PR #39748vllm/model_executor/models/qwen3_next.pyOverlaps sibling input-projection branches for Qwen3 / Qwen3.5 GDN-style blocksTreat serial GDN input projections as a known in-flight overlap opportunity.
PRs #41433 / #41434 / #41429 / #40561 GPU/CPU sync removalGPU->CPU synccpu syncitem()non_blockingPR #41433PR #41434PR #41429PR #40561Removes or gates accidental GPU-to-CPU synchronization points and adds sync-detection coverageTreat 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 overloadvllm_irfused_add_rms_normmaybe_inplacePR #36823vllm/compilation/passes/irvllm/compilation/passes/fusion/rms_quant_fusion.pyExtends vLLM IR lowering so fused-add-RMSNorm variants remain visible to later compile-time fusionsTreat missing norm/quant compile fusion as potentially an IR-lowering visibility issue.

18. Important toggles and caveats

Toggle / envLocationEffect on trace interpretation
enable_flashinfer_allreduce_fusionpython/sglang/srt/server_args.pyEnables the FlashInfer TP allreduce fusion family.
enable_aiter_allreduce_fusionpython/sglang/srt/server_args.pyEnables ROCm AITER TP allreduce fusion.
enable_deterministic_inferencepython/sglang/srt/server_args.pyCan 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_overlappython/sglang/srt/server_args.pyEnables the SBO family.
enable_fused_moe_sum_all_reducepython/sglang/srt/server_args.pyEnables fused MoE sum-reduce in the down path.
SGLANG_BLACKWELL_OVERLAP_SHARED_EXPERTS_OUTSIDE_SBOpython/sglang/srt/environ.pyAlters how DeepSeek-style shared-expert overlap behaves on Blackwell.
SGLANG_NSA_FUSE_TOPKpython/sglang/srt/environ.pyGates NSA fused top-k transform / page-table build.
SGLANG_DISAGG_STAGING_BUFFERpython/sglang/srt/environ.pyEnables the heterogeneous-TP staging-buffer family and its overlap windows.
SGLANG_STAGING_USE_TORCHpython/sglang/srt/disaggregation/common/staging_buffer.pyForces torch fallback for staging gather / scatter, so Triton staging kernels may disappear by design.
SGLANG_VIT_ENABLE_CUDA_GRAPHpython/sglang/srt/environ.pyCan intentionally disable vision aux_stream overlap.
SGLANG_ENABLE_FUSED_QKNORM_ROPEpython/sglang/multimodal_gen/runtime/layers/layernorm.pyGates the diffusion fused qknorm+rope path.
enable_pdl / launch_with_pdlflashinfer/norm/__init__.pyflashinfer/activation.pyflashinfer/rope.pyflashinfer/fused_moe/core.pyflashinfer/comm/allreduce.pyEnables FlashInfer PDL across many kernels; launch grouping and same-stream overlap can change substantially when it is on.
trigger_completion_at_endflashinfer/comm/allreduce.pyFalse enables downstream PDL-aware overlap after FlashInfer allreduce fusion; True delays completion to kernel end and removes that overlap window.
use_cuda_graphflashinfer/fused_moe/cute_dsl/fused_moe.pyEnables the preallocated-buffer path and the safe aux-stream async-memset overlap in FlashInfer CuTeDSL MoE.
split_device_green_ctx*flashinfer/green_ctx.pyChanges trace shape by partitioning SMs into separate green contexts instead of overlapping full-device streams on the default context.
rmsnorm_backendtensorrt_llm/_torch/auto_deploy/config/default.yamlChooses whether AutoDeploy lowers RMSNorm to FlashInfer, so split norm ladders may reflect backend selection rather than a missing fuse.
insert_cached_attention.backendtensorrt_llm/_torch/auto_deploy/config/default.yamlSelects the cached-attention backend; flashinfer enables the paged-KV cached-attention family.
insert_cached_mla_attention.backendtensorrt_llm/_torch/auto_deploy/config/default.yamlSelects the cached MLA backend; flashinfer_mla enables the MLA prefill / decode family.
TRTLLM_GEN_FUSED_MOE_USE_FLASHINFERtensorrt_llm/_torch/modules/fused_moe/fused_moe_trtllm_gen.pyForces or guards the FlashInfer-backed TRTLLM-gen MoE family, so expert-kernel shape can change substantially when it is set.
multi_stream_moetensorrt_llm/_torch/auto_deploy/config/default.yamlEnables the TensorRT-LLM shared-expert vs routed-expert overlap family.
multi_stream_mla_attntensorrt_llm/_torch/auto_deploy/config/default.yamlEnables the TensorRT-LLM MLA Q-vs-KV branch overlap family.
multi_stream_gemmtensorrt_llm/_torch/auto_deploy/config/default.yamlEnables generalized FP8 GEMM fork overlap in TensorRT-LLM AutoDeploy.
mlir_elementwise_fusiontensorrt_llm/_torch/auto_deploy/config/default.yamlCan absorb merge adds into larger fused kernels, so missing explicit merge nodes in multi-stream traces may be intentional.
enable_torch_compilepython/sglang/srt/server_args.pypython/sglang/multimodal_gen/runtime/server_args.pyCompiler-generated fusion / reordering can hide handwritten kernel names; absence of a custom kernel does not always mean absence of fusion.
enable_fused_grouped_gemm_combinePR #21877In-flight path that intentionally disables SBO because combine is folded into down-GEMM.
PassConfig.fuse_allreduce_rmsvllm/config/compilation.pyEnables 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_quantvllm/config/compilation.pyEnables vLLM's RMSNorm(+residual add) -> FP8 / FP4 quant compile-time fusion family.
PassConfig.fuse_act_quantvllm/config/compilation.pyEnables vLLM's SiLU+Mul -> quant fusion family, plus ROCm AITER variants where applicable.
PassConfig.fuse_attn_quantvllm/config/compilation.pyEnables attention-epilogue quant fusion; requires the right backend / graph visibility, so split kernels may still be expected.
PassConfig.fuse_mla_dual_rms_normvllm/config/compilation.pyEnables the AITER-backed MLA paired-Q/KV RMSNorm fusion family on ROCm.
PassConfig.enable_qk_norm_rope_fusionvllm/config/compilation.pyEnables the compile-time QK RMSNorm + RoPE family on CUDA-like backends.
PassConfig.fuse_rope_kvcachevllm/config/compilation.pyEnables ROCm / AITER RoPE + KV-cache update fusion and is range-limited by token count.
PassConfig.fuse_rope_kvcache_cat_mlavllm/config/compilation.pyEnables the MLA RoPE + unified MLA KV-cache update compile-time fusion family.
AttentionBackendEnum.TOKENSPEED_MLAvllm/model_executor/layers/attention/mla_attention.pyvllm/platforms/cuda.pySelects the TokenSpeed MLA backend on supported Blackwell FP8-KV MLA shapes when tokenspeed-mla is installed.
rocm_aiter_mla_sparse_dsv4vllm/v1/attention/backends/mla/rocm_aiter_mla_sparse_dsv4.pySelects the ROCm DSV4 sparse-MLA backend path on supported AMD shapes.
PassConfig.fuse_minimax_qk_normvllm/config/compilation.pyEnables the MiniMax decode Q/K allreduce-plus-RMSNorm compile-time fusion family.
PassConfig.fuse_act_paddingvllm/config/compilation.pyEnables the ROCm AITER add-RMSNorm-plus-pad fusion family when AITER is available.
PassConfig.enable_spvllm/config/compilation.pyRewrites all-reduce into sequence-parallel staging; this is often a prerequisite for the overlap family, not just a pure fuse toggle.
PassConfig.fuse_gemm_commsvllm/config/compilation.pyEnables AsyncTP GEMM + collective overlap and auto-enables enable_sp when valid.
vLLM PR #46735 Triton MoE CUDA graph capture fixvllm/model_executor/layers/fused_moe/experts/triton_moe.pyvllm/model_executor/layers/fused_moe/experts/nvfp4_emulation_moe.pyLatest 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_PDLcsrc/libtorch_stable/dsv3_fused_a_gemm.cucsrc/moe/dsv3_router_gemm_utils.hEnables 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_mlapython/tokenspeed/runtime/layers/attention/backends/tokenspeed_mla.pySelects 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_BACKENDtokenspeed-mla/python/tokenspeed_mla/mla_prefill.pyChooses 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-fusiondocs/configuration/server.mdpython/tokenspeed/runtime/distributed/comm_backendGates 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.

# 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"
Referenced from SKILL.md