Kernel fusion for efficient LLM training and inference

Why this question

LLM training and inference costs are dominated by a small set of operations — attention, the MLP block, normalisation, embedding lookups, and the optimiser step — executed as GPU kernels. Because data movement between GPU memory levels costs orders of magnitude more energy and time than floating-point arithmetic, the same algorithm can run at wildly different speeds depending on how its kernels are organised. Kernel fusion is the practice of organising those kernels so that intermediate results never leave the chip. It is the concrete mechanism behind most of the “tokens per watt” and “tokens per second” improvements reported over the past half-decade, and it is the layer at which hardware generations (Ampere, Hopper, Blackwell) are exploited.

The question matters for anyone allocating compute, choosing frameworks, or evaluating claims about “x-times faster training/inference”: which fusion techniques are real, how large are the measured gains, and where does the evidence stop? This review maps the literature on kernel fusion for LLMs between 2020 and 2026, with FlashAttention and Liger Kernel as the two canonical examples.

Scope and methods

Question. Which kernel-fusion methods improve the efficiency of LLM training and inference, and what gains do they deliver?

Inclusion criteria. Studies published 2020–2026 (one foundational exception: the Triton compiler, 2019, which nearly all later fusion work builds on ²⁸); work on kernel-level or operator-level fusion, fused attention, fused training ops, fused quantised kernels, MoE kernels, or kernel-generation tooling for LLMs; peer-reviewed systems/ML venues or widely-used arXiv preprints; English.

Exclusion criteria. Purely algorithmic attention alternatives without kernel implementation; distributed-training and interconnect work (covered in a separate review); KV-cache eviction/compression without kernel contribution; hardware architecture without kernel focus; USENIX-only papers with no DOI (one dropped: Rammer [OSDI’20] has no DOI and no arXiv version, which the site validator cannot accept); library-only works with no citable paper (xFormers; the “vLLM: Easy, Fast, and Cheap” paper does not exist — that title is the project’s GitHub tagline, and the kernel contribution is PagedAttention ⁶).

Search and screening. 20 concept queries × 4 databases (OpenAlex, arXiv, Crossref, Semantic Scholar), 47 curated landmark titles resolved via exact-title lookups (OpenAlex title.search, arXiv ti:, Crossref query.bibliographic), and citation snowballing from FlashAttention, FlashAttention-2, PagedAttention, Liger Kernel, and Triton. Merged pool: 2,826 unique works after DOI/title deduplication. Screen: venue whitelist + topic gate + noise gate → 368 candidates, of which 92 were selected after manual curation and full verification. Every DOI was verified (Crossref per-DOI check for publisher DOIs; abs-page title match for arXiv DOIs); the known wrong-publisher-DOI pattern on arXiv-first papers was handled by swapping to 10.48550/arxiv.* DOIs. Phantom titles from the curator’s memory (e.g., “Split-K Attention: Optimizing Memory-Efficient Attention”, “Unified Paged Attention…”) were dropped when no retrievable record existed anywhere.

Access. 75 of 92 sources were read in full text (arXiv HTML/PDF or publisher copy); 17 were abstract-only (paywalled NeurIPS proceedings pages and a few publisher pages that block bots). Screening counts: 2,826 retrieved → 368 screened candidates → 92 included.

The landscape

The literature has a clear shape: a pre-2022 phase of general operator-fusion research for deep learning, then an explosion centred on attention after FlashAttention (2022), then a widening into training ops, FP8, MoE, and serving kernels, and finally (2025–2026) a turn toward automated and LLM-generated kernels.

Three structural features matter. First, the field is young and preprint-dominated: of the 92 sources, 70 are arXiv preprints, and the median publication year is 2024. Second, it is hardware-concentrated: nearly all measurements are on NVIDIA GPUs (A100, H100, and recently GB10/Blackwell), with a handful of TPU and Ascend results — generalisation to other hardware is largely untested. Third, benchmarks are not standardised: each paper benchmarks against a different baseline (cuDNN, PyTorch eager, xFormers, FlashAttention-2), so “2x” in one paper is not “2x” in another; a 2026 benchmark study found that even correct LLM-generated kernels are often slower than eager execution, underscoring how baseline choice shapes every claim ⁸⁰. The 2019–2021 compiler work (Ansor, TensorIR, DNNFusion, horizontal fusion) frames fusion as a compiler problem; the 2022–2025 work reframes it as a hardware-aware kernel authorship problem; the 2026 frontier frames it as a problem for LLMs themselves.

Theme 1 — I/O-aware attention kernels: the FlashAttention lineage

The central result of the period is that attention can be made I/O-aware: instead of materialising the full N×N attention matrix in slow HBM, tile the computation so that all intermediate reads and writes happen in on-chip SRAM ¹. The mathematical prerequisite — computing softmax with O(1) memory via online rescaling — was established by Rabe and Staats, who showed self-attention can run in O(1)–O(log n) memory with 59x less memory overhead at sequence length 16,384 ⁴. FlashAttention combined that trick with tiling and reported up to 3x wall-clock speedup on GPT-2-scale training and 2.4x on long-range benchmarks, with memory no longer quadratic ¹.

FlashAttention-2 showed the first version was still only 25–40% of peak FLOPs/s — not because of memory traffic but because of poor work partitioning across thread blocks and warps — and fixed it with better parallelism, reaching 2–4x over optimised baselines ². A 2024 I/O-complexity analysis proved the tiling strategy is optimal for a wide range of SRAM sizes — FlashAttention is not a heuristic that happens to work but the minimiser of a proven lower bound ²⁰. FlashAttention-3 moved to the Hopper generation, exploiting Tensor-Core/TMA asynchrony and warp specialisation to reach 740 TFLOPs/s FP16 (75% utilisation) and ~1.2 PFLOPs/s with FP8 ³.

The decode side needed different tricks. Because generation is memory-bound with tiny matrices, FlashDecoding++ introduced asynchronous softmax with a unified max value (removing a ~20% synchronisation overhead), flat-GEMM double buffering, and heuristic dataflow selection, reporting ~4.9x decode speedups over Hugging Face baselines in its evaluation ⁵; its journal successor adds latency and memory optimisation ⁹². PagedAttention attacked the KV-cache memory problem from the memory-management side — paging the KV cache into fixed-size blocks, inspired by OS virtual memory — which reduced fragmentation and allowed 2–4x serving throughput gains at the time ⁶. vAttention later argued that PagedAttention’s non-contiguous virtual layout itself imposes kernel complexity, and showed that keeping virtual contiguity via CUDA virtual-memory APIs works with unmodified attention kernels and improves throughput by up to 1.23x over PagedAttention-style kernels ⁷⁴.

The lineage then specialised in every direction. FlashInfer provides a JIT-compiled, customisable attention engine covering block-sparse KV formats and load-balancing for serving ⁷. Flex Attention turned mask/score customisation into a programming model — users write mask_mod/score_mod in PyTorch and the compiler emits a fused Triton kernel, closing the gap between research attention variants and production kernels ⁸; IBM’s deployment work fused PagedAttention-style paging into FlexAttention kernels for long-context serving ¹⁶. A 2025 case study shows a full paged attention kernel can be built in Triton alone with competitive performance ¹⁵. ByteTransformer fused multi-head attention with architecture-aware padding-free execution for variable-length inputs, 6.13x over stock PyTorch MHA on A10 ¹⁰.

Quantisation and sparsity entered the kernel. SageAttention quantised Q/K/P/V to INT8 within a fused Triton kernel, reaching 2.1–2.7x over FlashAttention-2/xFormers at 340 TOPS ⁹; SageAttention2 pushed to per-thread INT4 with outlier smoothing, ~3x over FlashAttention-2 ⁸⁸. INT-FlashAttention made INT8 quantisation compatible with the fused forward workflow, 72% faster with lower error than baseline INT8 attention ²⁵. Block-Sparse FlashAttention computes exact query-key scores while skipping ~50% of computation on sparse patterns at 99%+ accuracy retention ²⁶; sparse causal FlashAttention (Triton) extended the algorithm to arbitrary causal sparsity with 2.0–3.3x training speedups ¹³. Binary Block Masking made FlashAttention mask-aware for padding-free training ²³.

Recent work pushes the same principles further: VFA pre-computes key-block maxima to relieve vector operations, projecting ~2x (up to 6x projected) gains ¹⁴; Sawtooth Wavefront Reordering reschedules CTA wavefronts on Blackwell so concurrent CTAs reuse L2-resident K/V sectors, cutting L2 misses ≥50% and raising throughput up to 60% on GB10 ¹⁸; DualKV fuses forward and backward attention kernels to eliminate shared-prompt replication in RL training, with 1.63–2.09x GRPO speedups ¹⁷. Two training-side refinements matter for production: packing multiple sequences with position IDs into one fused attention call (substantially higher SFT throughput) ²², and hiding dropout’s RNG generation behind GEMM rather than inside the kernel, which recovers 1.22–1.26x ²⁴.

Blockwise attention across devices — Ring Attention and Striped Attention — extended the fused-kernel idea to the distributed setting, trading communication overlap for near-infinite context ¹¹¹².

Theme 2 — Fusing the rest of the model: training-op kernels and fusion compilers

Attention is only one layer. The training loop’s other ops — RMSNorm, LayerNorm, RoPE, SwiGLU, cross-entropy, the AdamW step — are memory-bound elementwise/reduction operations where each kernel launch round-trips through HBM. Fusion here means computing chains of these ops in one pass.

The enabling technology is Triton, a tile-based DSL and compiler that lets researchers write custom fused CUDA-level kernels in a Python-like language without hand-tuning PTX ²⁸; PyTorch 2’s torch.compile made this automatic for ordinary PyTorch code, lowering to Triton via TorchInductor with 2.27x geomean speedups on training workloads at the time ²⁹. On top of these, Liger Kernel — the second canonical example in this review’s question — ships fused Triton kernels for RMSNorm, LayerNorm, RoPE, SwiGLU, cross-entropy, and the fused linear cross-entropy, reporting ~20% higher training throughput and up to 60% memory savings across models from 1B to 8B ²⁷. ThunderKittens provides a lower-level tile-based abstraction with async-overlap templates, reporting 10–40% gains on attention backward and large wins on SSM/linear-attention kernels ³⁰.

The compiler literature supplies the theory of what to fuse. DNNFusion classified operators and their mathematical properties to expose fusion opportunities 8.8x more often than prior frameworks, with up to 9.3x speedups on DNNs ⁸⁹. FusionStitching fuses memory-intensive ops with complex dependencies into large kernels (up to 2.21x) ⁴². Ansor auto-tuned tensor programs with up to 3.8x CPU / 1.7x GPU gains ⁴⁴; TensorIR generalised loop-nest IR to make tensor primitives first-class ⁴³. A source-level study of XLA’s fusion passes found conservative criteria leave optimisation on the table ⁴⁶. Horizontal fusion — fusing kernels that would otherwise run in sequence across different resources — complements the more common vertical (dataflow) fusion, with 2.5–60.8% gains on fused kernels ⁴⁵.

Two 2025–2026 developments reframe the granularity question. Megakernels: MPK compiles multi-GPU inference into a single persistent mega-kernel with SM-level scheduling (1.7x lower latency) ³⁹; Blockbuster’s data-movement-aware fusion automatically rediscovers FlashAttention and fuses LayerNorm+Matmul, arguing block-level fusion is the right granularity on modern GPUs ⁴⁷; ComFuse fuses memory-intensive subgraphs with compute kernels via stage-stream execution ⁴⁰; RedFuser fuses cascaded reductions (softmax→GEMM) with 2–5x speedups ⁴¹. DeepFusionKernel fuses the SwiGLU chain into one kernel for 13.2% gains on H100 ⁵⁰. The Fused Kernel Library generates a single compile-time-fused kernel from arbitrary function sequences ⁵¹. The other direction — cluster-level fusion — keeps intermediates on-chip across GPUs: ClusterFusion’s ClusterReduce/ClusterGather primitives give 1.61x ⁴⁸, and FLUX fuses communication with computation, overlapping up to 96% of communication time ⁴⁹. Twill formalises software pipelining + warp specialisation as a constraint problem with guaranteed-optimal schedules ⁵².

DeepSeek-V2 is the production-scale existence proof for training-op fusion: its report attributes 42.5% training-cost savings and a 93.3% KV-cache reduction to MLA plus fused kernels for attention and MoE dispatch ³².

Theme 3 — FP8 and low-precision training kernels

Fusion and precision are complementary: FP8 halves the bytes moved by every fused kernel, so the two multiply. NVIDIA’s E4M3/E5M2 interchange format paper established that FP8 can match 16-bit training quality across model families ⁷¹. FP8-LM built the first full FP8 training stack (gradients, optimiser states, communication in FP8) with 75% faster training and 39% less memory than BF16 ³¹. Scaling to trillion-token runs showed the failure mode: late-training divergence attributed to SwiGLU activation outliers, fixable with careful scaling ³⁶. But a critical evaluation found current FP8 recipes “not robust enough” as drop-in BF16 replacements ³⁷, and a 2025 mechanistic analysis showed why low-precision attention training can collapse — similar low-rank representations of Q/K make softmax unstable under quantisation ¹⁹.

The 2025–2026 generation claims to close the gap: µnit Scaling trains 1B–13B models in FP8 without dynamic scaling or special hyperparameters ³⁸; MOSS (microscaling formats + automatic scaling) reports 1.34x end-to-end speedup with lossless accuracy on OLMo-7B ³⁴; FOG runs all GEMMs including attention projections in FP8 for the first time, with up to 43% throughput gains ³⁵; COAT compresses optimiser states and activations to FP8, cutting memory 1.54x ⁵⁴. The largest production data point is DeepSeek-V3: 14.8 trillion tokens trained in 2.788M H800 GPU-hours (~180K GPU-hours per trillion) with a validated FP8 pipeline and custom kernels ³³.

Theme 4 — Hardware-aware kernels for linear attention and state-space models

Sub-quadratic architectures only beat Transformers if their kernels are as good as FlashAttention’s. Mamba’s hardware-aware selective scan — fusing the scan, avoiding materialisation, and recomputing in the backward pass — delivers 5x higher generation throughput than a Transformer baseline at comparable quality ⁵⁵. Lightning Attention-2 combined tiled intra-block softmax with an inter-block linear-attention trick in Triton, keeping speed constant across sequence lengths ⁵⁶. Gated Linear Attention’s FlashLinearAttention kernel is faster than FlashAttention-2 as a standalone layer even at 1K sequence length ⁵⁷, and Tiled Flash Linear Attention adds intra-chunk sequence parallelisation for arbitrary chunk sizes, with the mLSTM variants winning on recall-style tasks ⁵⁸. Google’s Efficiently Scaling Transformer Inference remains the reference for the decoding-side kernel stack: 29 ms/token for PaLM 540B at 76% MFU using partitioning-selection models plus low-level kernel optimisations ⁵⁹.

Theme 5 — Fused MoE kernels

Mixture-of-Experts models move the bottleneck to the routing/dispatch path: tokens must be permuted, gathered, and scattered around expert GEMMs, and naive implementations waste most of the gain. Tutel’s adaptive parallelism and pipelining (zero-cost runtime switching, flexible all-to-all) gave 4.96x/5.75x speedups over state-of-the-art MoE systems ⁶⁰. MegaBlocks reformulated MoE as block-sparse operations with custom kernels (never dropping tokens), enabling 2.4x faster training than dense baselines ⁶¹. DeepSpeed-MoE showed 5x training-cost savings versus quality-equivalent dense models ⁶². FasterMoE attacked the serving side with expert prefetching, adaptive GPU memory, and fused kernels for the dispatch path ⁹¹; dynamic gating alone improved maximum throughput 6.21–11.23x in one study ⁶⁶. A 2024 survey taxonomises the now-large MoE inference-optimisation literature ⁶⁷.

The 2026 wave fuses the entire MoE layer: TritonMoE is a pure-Triton fused dispatch kernel (router, permutation, expert GEMMs, combination) reaching 89–131% of MegaBlocks’ kernels on A100 while cutting dispatch traffic 35% ⁶³; RaMP selects among 134–268 polymorphic fused-MoE kernel configurations with a wave cost model, 1.22x kernel / 1.30x end-to-end ⁶⁴; DA-MoE matches the live routing histogram to offline-tuned kernel dispatch, 1.16–1.29x geomean latency improvements ⁶⁵.

Theme 6 — Quantised and mixed-precision inference kernels

Weight-only quantisation (W4A16, FP8) is only as fast as its GEMM kernel. LLM.int8() established the mixed-precision decomposition that keeps outlier dimensions in FP16, enabling 175B-parameter inference without degradation ⁷⁰. MARLIN is the canonical 4-bit kernel: it sustains close to the theoretical 4x quantisation speedup at batch sizes 16–32 (where prior kernels dropped to 1–2x) via asynchronous loads, L2-resident activations, and striped SM partitioning ⁹⁰. FP6-LLM’s TC-FPx tensor-core kernel (ahead-of-time bit packing, SIMT-efficient dequantisation) reached 1.69–2.65x over FP16 baselines ⁸⁷. Atom pushed fine-grained 4-bit weight-and-activation quantisation with 7.73x over FP16 serving throughput ⁶⁸. Flash-LLM fused unstructured-sparse weight loading with dense compute (Load-as-Sparse/Compute-as-Dense), 2.9x over Sputnik/SparTA ⁶⁹. Fast NF4 dequantisation kernels reach 2.0–2.2x ⁷², and a Triton SplitK-decomposed fused dequant+GEMM kernel speeds skinny W4A16 matmuls 65–124% ⁷³.

Theme 7 — Serving systems built on kernels

Several serving systems are, at their core, a kernel idea plus a scheduler. vLLM’s PagedAttention is the reference case ⁶; vAttention’s virtual-memory alternative ⁷⁴; SGLang contributes RadixAttention (KV reuse across requests) plus a high-performance runtime where fused kernels for attention/MoE are first-class ⁷⁵. For multi-tenant fine-tuning, Punica’s Segmented Gather Matrix-Vector CUDA kernel batches GEMMs across different LoRA adapters on one GPU ⁷⁶, and S-LoRA serves thousands of adapters via unified paging and heterogeneous-batching kernels ⁷⁷; AdaFuse adds token-level pre-gating with a fused switching kernel, >2.4x decode-latency reduction ⁷⁸. A hybrid JIT-CUDA-graph runtime cuts TTFT up to 66% versus TensorRT-LLM on short sequences ⁵³.

Theme 8 — The frontier: LLM-generated kernels

The newest development is using LLMs to write and optimise the fused kernels themselves. The benchmarks came first: TritonBench showed state-of-the-art code LLMs struggle to generate efficient Triton operators (184 curated operators, two evaluation channels) ⁷⁹; KernelBenchX found that across five kernel-generation methods and 176 tasks, 72% of fusion tasks fail for all methods and 46.6% of correct kernels are slower than eager execution ⁸⁰. The optimisers then improved: AutoKernel’s agent loop with a five-stage correctness harness beats PyTorch defaults by 5.29x on RMSNorm ⁸¹; DRTriton, trained with curriculum RL on synthesised PyTorch–Triton pairs, reaches 92% of expert-level kernels versus 19–23% for frontier general LLMs ⁸²; EvoEngineer formalises CUDA kernel optimisation as code evolution with 2.72x median speedups ⁸³; compiler-grounded hierarchical diagnosis (pattern triage → profiling → IR attribution → rewrite) achieves 4.35x geomean on Ascend ⁸⁴; TritonForge’s profiling-guided loop reaches up to 5x ⁸⁵; KernelBrain’s budget-aware search reaches 0.88–6.72x ⁸⁶. The pattern is consistent: LLMs can now write correct kernels and, with feedback loops, match or beat hand-tuned ones on targeted ops — but general fusion remains unsolved.

Where the evidence disagrees

FlashAttention: heuristic or optimum? The I/O-complexity proof says the tiling strategy is optimal for a wide range of SRAM sizes ²⁰, yet the sparse/quantised attention literature claims further large gains ²⁶⁹¹³. The apparent conflict dissolves once separated by what is being optimised: the lower bound governs exact attention’s HBM traffic; sparse and quantised variants change the problem (approximate attention), so both claims stand. What does not dissolve is the numerical-stability tension: FlashAttention shows roughly an order of magnitude more numeric deviation than baseline attention at BF16 in isolation ²¹ — the field largely treats this as acceptable, but it is unresolved for long-context and low-precision training ¹⁹.

FP8 training: ready or not? Production reports (DeepSeek-V3, 14.8T tokens) ³³ and new recipes (µnit, MOSS, FOG) ³⁸³⁴³⁵ claim lossless FP8 training; a careful evaluation found existing methods not robust enough as BF16 replacements ³⁷. The best available resolution is temporal and architectural: the failure mode (outlier-driven divergence, especially in SwiGLU) was identified ³⁶, and the 2025–2026 recipes explicitly target it — the disagreement is between generations of methods measured at different scales, not between camps.

Megakernels vs per-op fusion. MPK/Blockbuster/ComFuse argue for fusing at block granularity into persistent mega-kernels ³⁹⁴⁷⁴⁰, while the Liger-style approach fuses per-operation chains ²⁷ and FlashAttention-3 shows hand-tuned warp-specialised kernels still win at the attention layer ³. The disagreement is about which GPU resources are the bottleneck (launch/occupancy vs memory traffic) and likely depends on hardware generation; no study compares the paradigms head-to-head on the same workload.

How fast is “fast”? Reported speedups for overlapping claims differ by up to 2x across papers (e.g., attention kernels vs FlashAttention-2 baselines of different vintages; MoE dispatch gains of 1.16x vs 6–11x ⁶⁵⁶⁶ from different bottleneck definitions). Benchmark heterogeneity — not contradictory evidence — explains most of this scatter ⁸⁰.

Gaps and open questions

• No shared benchmark for kernel efficiency. KernelBenchX/TritonBench are first steps ⁸⁰⁷⁹, but there is no standard suite of fusion tasks with fixed baselines and hardware — every “x-times” claim in this review is only meaningful against its own baseline. What would settle it: a public, hardware-pinned fusion benchmark adopted across venues.
• Fusion beyond NVIDIA. Almost all evidence is CUDA/Triton on NVIDIA GPUs; TPU attention kernels (e.g., a 2026 TPU paged-attention kernel) and Ascend results are rare ⁸⁴. Whether fusion principles transfer is untested.
• Megakernel vs fused-per-op vs hand-tuned, on identical workloads — the central engineering question of 2026 has no direct comparison study.
• The stability question for exact attention — an order-of-magnitude numeric deviation in isolation ²¹ has no systematic study of downstream effects at scale.
• LLM-generated kernels currently fail most fusion tasks ⁸⁰; whether this closes with better training data (DRTriton-style) or needs different abstractions is open.
• Production verification: outside DeepSeek’s reports ³²³³, almost no fused-kernel claims are verified at datacenter scale by independent parties.

Confidence and limitations

Confidence is moderate. The core direction of evidence is unambiguous and replicated: I/O-aware attention kernels, fused training ops, and FP8 pipelines all deliver large, repeatable gains, and the theoretical underpinning for attention tiling is proven ²⁰. But most sources are preprints; speedup numbers are hardware- and baseline-specific; the 2025–2026 literature (17 of 92 sources) is largely un-peer-reviewed; and 17 sources were abstract-only, including the NeurIPS proceedings versions of FlashAttention and FlashAttention-3 (their arXiv versions were read in full, and their headline numbers are quoted from the abstracts). Three sources were extracted from stubs then rescued (vAttention, TritonBench, FasterMoE) — their entries are more limited than the rest. Papers dropped for retrievability (Rammer; xFormers; two phantom titles) do not affect the conclusions. The review covers English-language, GPU-focused literature; FPGA and non-English work is out of scope.