Kernel optimization for LLM training: techniques, trade-offs, and use cases

Why this question

Training a large language model spends almost all of its time inside a small set of GPU kernels — attention, the MLP’s matrix multiplications, normalization, the optimizer step, and, in MoE models, expert dispatch and communication. Because moving data between GPU memory levels costs orders of magnitude more than the arithmetic itself, the same model can train at very different speeds depending on how those kernels are organized: whether operations are fused into one kernel, what precision the GEMMs run at, whether expert routing is block-sparse or padded, and whether communication overlaps computation. These kernel decisions are now the difference between a training run that fits a budget and one that does not — DeepSeek-V3 trained at 180K H800 GPU-hours per trillion tokens with FP8 GEMMs, fused communication, and MoE kernels ³⁵.

The question matters for anyone allocating compute or evaluating “x-times faster training” claims: which kernel techniques are real, how large are the measured gains, what do they cost in engineering effort and numerical risk, and which technique fits which workload? This review maps the 2018-2026 literature on kernel optimization for LLM training, organized by technique family, with the pros, cons, and use cases of each.

Scope and methods

Question. Which kernel-level optimization techniques most improve LLM training efficiency, and what are the trade-offs and use cases of each?

Inclusion criteria. Studies published 2018-2026 (two foundations: 8-bit floating-point training, 2018 ²⁰, and Triton, 2019 ²); kernel-level or operator-level work on attention (forward and backward), GEMM and low-precision training (FP8/FP4/BF16/microscaling), MoE dispatch/combine and expert communication, communication-computation overlap, fused optimizer/normalization kernels, and the tooling used to build them (Triton, CUTLASS, torch.compile, kernel DSLs); peer-reviewed systems/ML venues or widely-used preprints; English.

Exclusion criteria. Inference-only kernels (PagedAttention, FlashDecoding, MARLIN, SageAttention) except where they benchmark the same kernel in training; pure distributed-systems/interconnect papers without a kernel component (covered by the separate LLM-interconnects review) except communication-overlap kernels; algorithmic attention substitutes with no kernel implementation; USENIX-only papers with no DOI; paywalled-only work with no open version; hardware microarchitecture without kernel focus.

Search and screening. Five databases (OpenAlex with API key, arXiv, Crossref, Semantic Scholar, DBLP). 25 concept queries (recorded verbatim in the frontmatter) produced a merged pool of 3,836 records; after title screening (~1,170 topic-gate hits) and full-text screening, 75 sources were included: 66 read in full text, 9 abstract-only (publisher-paywalled: Triton PLDI, PyTorch 2, FasterMoE PPoPP, WASP, PipeMesh, RailS, BF16-FMA, CUTLASS benchmark, low-precision survey, Expert Choice — the last was upgraded via its arXiv version). Approximately 50 exact-title resolution lookups resolved landmark papers; several remembered titles and IDs proved wrong and were corrected or dropped: “MOSS: A Benchmark for Efficient LLM Training Kernels” does not exist (the real MOSS paper is the FP8 microscaling training paper ¹⁷), “SwitchBack” could not be resolved in any database and was dropped, and 1-bit-LAMB was dropped as off-topic (communication compression, not kernels). GitHub-only artifacts without papers (DeepEP, DeepGEMM, DualPipe) are covered through their parent papers and repository READMEs, with their numbers flagged vendor-stated. Every DOI was verified against Crossref or the arXiv abs page; all 75 resolve.

A note on GitHub-only artifacts. DeepEP (MoE dispatch/combine all-to-all kernels), DeepGEMM (FP8 GEMM kernels), and DualPipe (bidirectional pipeline scheduling) have no papers and no DOIs — they are repository-documented ³⁵³⁴³⁶. This review treats their READMEs (retrieved 2026-08-15) as primary documentation and reports their performance as vendor-stated, cross-checked where possible against third-party work ⁶⁶⁶⁷⁶⁸.

The landscape

The literature is young, concentrated, and shaped by one lineage. FlashAttention (2022) established IO-aware tiling as the template ^1112134, and most attention-kernel work since is either a refinement of that template (block-sparse ⁴⁹, INT8 ⁵⁰, vector-relieved ⁴⁶, wavefront-reordered ⁴⁸) or an extension of it to long contexts and distributed memory ²⁹⁵⁶¹. The low-precision training literature runs parallel: FP8 formats were standardized in 2022 ²¹²⁰, production FP8 training arrived with FP8-LM and DeepSeek-V3 ¹⁶³⁵, and 2025 papers pushed to full FP8 ¹⁸¹⁷¹⁹ and FP4 ⁶³⁶⁴. MoE kernel work tracks the model architecture: block-sparse reformulation ²⁴⁹, adaptive parallelism ²⁵, single-kernel fusion ²⁷²⁸, and communication redesign ^69676674.

Three structural facts stand out. First, the field is dominated by a small number of labs and vendors: the FlashAttention lineage is one group’s program ¹¹¹²¹³, the deep-learning-kernel tooling papers come overwhelmingly from NVIDIA-affiliated authors ²³⁵⁵², and DeepSeek’s technical reports carry an outsized share of the production-scale evidence ³⁴³⁵. Second, almost all quantitative evidence is vendor- or single-group-benchmarked; independent head-to-heads are rare, and the best-known independent study found that correctness-based kernel evaluation accepts catastrophically slow kernels ⁵⁴. Third, the era is compressed: the most-cited techniques (FlashAttention, FP8 training, single-kernel MoE) are all from 2022-2025, which means the field is still consolidating — few claims have been independently replicated at the same scale.

Theme 1 — Fused attention kernels: the FlashAttention lineage and its variants

The core result is settled. FlashAttention’s IO-aware tiling — computing attention in blocks that never leave the chip’s SRAM, with online softmax and backward recomputation — trains Transformers 15% faster than the MLPerf 1.1 BERT-large record and 3x faster on GPT-2 at 1K sequence length ¹¹. FlashAttention-2 improved parallelism and work partitioning for roughly 2x over the original, reaching 50-73% of theoretical peak FLOPs/s ¹², and FlashAttention-3 added warp specialization and asynchronous copy/GEMM overlap on Hopper for 1.5-2.0x forward and 1.5-1.75x backward speedups over FlashAttention-2, hitting 740 TFLOPS/s (about 75% of H100 peak) in FP16 ¹³. The I/O complexity analysis proves this tiling is optimal up to constant factors for the attention computation ⁴. A CUTLASS implementation of FlashAttention-2’s forward pass using Hopper TMA bulk copies and WGMMA async matrix operations gains another 20-50% FLOPs/s over the Ampere-optimized reference on H100 ⁵¹.

The training-specific refinements cluster into five groups. First, custom attention variants: FlexAttention provides a PyTorch-native programming model for generating optimized attention kernels, matching FlashAttention-2 on supported variants (1.00-1.22x forward, 0.86-1.05x backward on causal attention) while beating PyTorch’s SDPA by 5.49-8.00x on masked variants, and speeding up torchtune training 2.4x ¹⁴. Second, efficiency for real training shapes: packing multiple sequences with position IDs improves throughput up to 2x over padding with identical loss behavior ⁴³; hiding dropout’s random-number generation inside preceding GEMMs gives 1.26x over sequential execution ⁴⁴; binary block masking dispatches only unmasked blocks for up to 9x on real-world sparse masks ⁴⁵; Block-Sparse FlashAttention prunes roughly half of blocks via calibrated per-layer thresholds for up to 1.10-1.24x with accuracy retention ⁴⁹; sparse QK-dropping and hash-sparse patterns implemented inside the flash kernel speed up transformer training 2.0x at 8K and 3.3x at 16K sequence length ⁴⁰; and vector-relieved variants cut the online-softmax vector-op latency share from ~77% to ~46% for ~2x ⁴⁶. Third, numerical stability in low precision: FlashAttention shows roughly 10x more numeric deviation than baseline attention at BF16 in isolation, but Wasserstein-based analysis bounds its training impact at 2-5x less than feared ⁴²; the first mechanistic account of BF16 flash-attention loss explosion attributes it to low-rank attention representations compounded by biased rounding, and demonstrates a minimal fix ⁴¹. Fourth, long-context and RL-specific kernels: DualKV fuses forward/backward kernels over dual KV regions (shared prompt plus per-sequence response) to eliminate shared-prompt replication in RL training, with 1.63-2.09x policy-update speedups and model FLOPs utilization rising from 36% to 76% ⁴⁷. Fifth, INT8 attention: INT-FlashAttention, the first attention operator with fully INT8 inputs, runs ~72% faster than FP16 FlashAttention with up to 82% memory savings, though the reported benchmarks are inference-time ⁵⁰.

Linear-attention kernels are a separate branch with a different trade-off. FlashLinearAttention-style fused kernels are faster than FlashAttention-2 even at short sequences (~1K) for linear-attention architectures, and Gated Linear Attention transformers train with higher throughput than a same-size Mamba model at matched quality ³⁹. Tiled Flash Linear Attention’s two-level sequence parallelism delivers training (forward+backward) kernels faster than FlashAttention-3 at longer sequences and over 2x faster than Mamba 2 on H100 ⁵⁸. The trade-off: these kernels serve linear-attention architectures, which trade exact softmax attention for recurrence-friendly computation — a model-level choice, not a drop-in kernel swap.

Pros and cons. Pros: fused attention kernels are the single best-documented training optimization, with replicated gains across independent implementations ¹¹¹²¹³, and the IO-optimality proof means the tiling strategy will not be superseded on the same hardware ⁴. Cons: gains are hardware-generation-specific (Hopper’s TMA/WGMMA are what FA3 exploits ¹³⁵¹); the backward pass is where memory and numerical risk concentrate ⁴¹⁴²; and custom-variant flexibility requires either compiler tooling ¹⁴ or expert CUDA work. Use cases: any transformer training run benefits from the stock kernels; the variants pay off for masked/packed/RL workloads ⁴⁵⁴³⁴⁷; linear-attention kernels are for recurrent architectures where exact attention is not required ³⁹⁵⁸.

Theme 2 — GEMM kernels and low-precision training: FP8, FP4, and microscaling

The GEMM kernel layer. Stream-K’s work-centric decomposition — splitting the inner-loop iterations of a matrix product across streaming multiprocessors instead of tiling by output block — reaches peak speedups of up to 14x (FP64) and 6.7x (FP16) over CUTLASS/cuBLAS across 32,824 GEMM shapes ²³. ThunderKittens, a CUDA template library for writing AI kernels, matches cuBLAS and FlashAttention-3 on GEMM and attention (GEMM up to 805 TFLOPS on H100-class hardware) while outperforming the strongest baselines by 10-40% on a range of kernels ⁵⁵. A systematic CUTLASS benchmarking study confirms that tensor-core GEMM kernels require per-shape tuning to reach their potential ⁵⁶. Sparse GPU kernels tuned for deep learning (SpMM/SDDMM) reach 27% of V100 peak and beat cuSPARSE by 3.58x geometric mean, which is what made sparse MoE training viable ⁹.

FP8 training is production-proven at scale, with conditions. The FP8 interchange format (E4M3 for weights/activations, E5M2 for gradients) was defined in 2022 with evidence that FP8 training matches 16-bit quality with unchanged hyperparameters ²¹; the first FP8 training demonstration predates it (2018, chunk-based accumulation with stochastic rounding, ResNet18 test error 42.45% vs 41.96% FP32) ²⁰. FP8-LM showed FP8 gradients, all-reduce, optimizer states, and tensor-parallel communication can train GPT-175B with 39% lower real memory and 75% faster than the Megatron baseline ¹⁶. DeepSeek-V3 trained a 671B-parameter MoE fully in FP8 (with BF16 master weights) at 180K H800 GPU-hours per trillion tokens ³⁵. Subsequent work pushed further: MOSS achieves 34% higher end-to-end training throughput than BF16 on OLMo-7B (45,374 vs 33,805 tokens/s) using microscaling formats and automatic scaling, and 12.3% over a prior FP8 framework at comparable perplexity ¹⁷. μnit Scaling trains 1B-13B models with all hidden linear layers in FP8 using static unit-variance scaling and square-root softmax — no dynamic scale factors — at 25-33% higher throughput than BF16 ¹⁹. Fully-FP8 pushes FP8 into the attention and even the LM head’s GEMMs for up to +43% training throughput at 8B scale (12,764 vs 9,105 tokens/s/GPU), stable to 450B tokens ¹⁸. At trillion-token scale, FP8 training with both Adam moments quantized reaches BF16 parity with ~34% throughput improvement ⁶⁰.

The conditions matter. The counter-evidence is specific and consistent: even standard BF16 mixed precision diverged in ~10% of runs (18 of 188 random seeds) at only 5% of training in one study, and MS-AMP FP8 O1 failed to match BF16 convergence on noisier data — the failures cluster on small runs, low-precision master weights, and noisy gradients ⁶¹. FP4 remains experimental: the first FP4 training framework trains 13B models on 100B tokens with a minor loss gap (2.17 vs 2.07 BF16 for 7B) ⁶³, and mixed-precision FP4 pretraining (FP8 attention, FP4 FFN, FP8 backward) reaches near-FP16 validation loss at ~66% of the theoretical computation cost ⁶⁴. A direct pretraining study found 8-bit per-channel weights and per-token activations match the FP32 baseline with memory savings, while 4-bit activations or gradients destabilize training (activation quantization produced perplexity 418 vs 39.9 baseline) ¹⁰. Microscaling (MX) formats with per-block scaling work as low-friction FP32 drop-ins across 20+ benchmarks, and 6-bit MX enabled the first sub-8-bit training of generative language models ²². On the hardware side, a BF16-only FMA operator family (no FP32 FMA units) trains complex DNNs to FP32-equivalent accuracy at 1.28-1.35x performance ⁶², and the low-precision training survey organizes this landscape by format family and open challenges ⁵⁹.

Pros and cons. Pros: FP8 is the only kernel-level technique with production-scale validation at frontier scale ¹⁶³⁵, delivers memory savings (39%) and throughput gains (25-75%) simultaneously ¹⁶¹⁷¹⁹, and the formats are standardized ²¹. Cons: the numerical risk is real and not fully characterized — divergence is seed-dependent at small scale ⁶¹ and the loss-explosion mechanism is only beginning to be understood ⁴¹; the gains require H100-class hardware with FP8 tensor cores; and each new precision (FP4, MX) restarts the stability engineering ⁶³⁶⁴²². Use cases: FP8 is the default for large-scale pretraining on Hopper/Blackwell ³⁵¹⁸; FP4 and MX are for pushing memory/bandwidth limits where quality loss is acceptable or for frontier labs with stability engineering capacity ⁶³²²; BF16 remains the safe default for small runs and fine-tuning ⁶¹.

Theme 3 — MoE training kernels: from padded to block-sparse to single-kernel

The first generation attacked padding. MoE layers route each token to a subset of experts; naive implementations pad tokens to a uniform expert batch, wasting compute. MegaBlocks reformulated MoE as block-sparse operations with new GPU kernels that never drop tokens, achieving end-to-end training speedups of 1.38x/2.0x/4.35x over Tutel’s padding-based approach ²⁴. Tutel itself contributed adaptive parallelism plus fast SIMT encode/decode kernels and a two-dimensional hierarchical all-to-all, with single-MoE-layer speedups of 4.96x/5.75x on 16/2,048 A100s over Fairseq and 1.14-1.55x end-to-end training on SwinV2-MoE ²⁵. FasterMoE framed the field’s core problem statement — dynamic load imbalance, inefficient synchronous execution, and congested all-to-all communication ²⁶ — and DeepSeekMoE showed that finer-grained expert segmentation plus shared expert isolation reaches LLaMA2-7B quality with a 16B model at ~40% of the compute ³².

The second generation fused the whole layer into one kernel. FlashMoE fuses gate, dispatch, expert GEMMs, and combine into a single persistent GPU-resident kernel with NVSHMEM DMA and in-kernel actor scheduling, delivering up to 6x lower latency and 5.7x higher throughput (17.7M tokens/s on 8 H100s) than state-of-the-art baselines, with SM utilization of 93.17% vs 9.67% for FasterMoE ²⁷. SonicMoE adds IO- and tile-aware optimizations, cutting activation memory by 45% and achieving 1.86x compute throughput over a BF16 MoE kernel on Hopper, sustaining 213B tokens/day on 64 H100s ²⁸. On the tooling side, a pure-Triton fused MoE dispatch (router, permute, grouped expert GEMMs, unpermute) reaches 89-131% of MegaBlocks’ throughput at inference batch sizes — evidence that the DSL path is now viable for MoE dispatch ⁵².

The third generation attacked communication. Expert-parallel training spends 45-67% of its time in all-to-all in some workloads ⁶⁹; LSH-MoE compresses routed activations with cross-polytope hashing and residual error compensation for 1.28-2.2x end-to-end training speedups ⁶⁹. DeepEP provides high-throughput and low-latency all-to-all GPU kernels (MoE dispatch/combine) with FP8 support; per its repository README it reaches 90 GB/s RDMA dispatch bandwidth on CX7 and 726-740 GB/s NVLink bandwidth on SM100 at minimal SM occupancy ³⁵. Third-party work validates the workload and the approach: UBEP re-architects the expert-parallel all-to-all with kernel decomposition and hierarchical token scheduling for up to 52.4% lower all-to-all latency ⁶⁷; NCCL EP builds expert-parallel communication on the NCCL Device API with low-latency (1-128-token) and high-throughput (4096+ token) modes ⁶⁶; RailS exploits rail topologies to turn global load balancing into local scheduling, improving bus bandwidth 20-78% and cutting all-to-all completion time 17-78% ⁷⁴; and Shortcut-connected expert parallelism overlaps dispatch with expert computation for 1.49x training speedups at 100% communication overlap on 8x A30 ⁶⁸. Algorithmic routing work complements the kernels: expert-choice routing trains over 2x faster to the same perplexity than Switch/GShard top-k routing at 8B/64E ³, pre-gating plus preemptive expert migration cuts MoE-block latency 1.7x on average ³³, and PEER (a million small experts with product-key routing) reaches the best compute-optimal perplexity at two FLOP budgets ⁸. The architectural lineage matters for the kernels: Switch Transformers showed 7x pretraining speedups with sparse routing ³⁰, GShard scaled MoE to 600B parameters on 2048 TPU v3s in 4 days ³¹, and DeepSpeed-MoE demonstrated dense-quality training at 5x lower cost with 3.7x model compression ⁷⁰.

Pros and cons. Pros: MoE kernels deliver the largest single-layer gains in the field (up to 5.7x ²⁷²⁵) and directly enable the compute-efficiency economics of modern MoE models ³⁴³⁵. Cons: the kernel work is tightly coupled to routing algorithm and model architecture — changing routing (top-k vs expert-choice vs PEER) changes the kernel problem ³⁸; communication kernels are topology- and hardware-specific ⁶⁷⁷⁴; and the single-kernel approaches are the most engineering-intensive, with FlashMoE’s training/backward fusion still deferred ²⁷. Use cases: block-sparse kernels for dense-scale MoE pretraining ²⁴; single-kernel fusion for latency-sensitive large-batch training and inference ²⁷²⁸; communication redesign for expert-parallel clusters where all-to-all dominates ^69676674; and Triton-based dispatch where portability matters more than peak performance ⁵².

Theme 4 — Communication-overlap and long-context kernels

Attention kernels that span devices. Ring Attention overlaps blockwise KV communication with attention computation, training sequences device-count times longer than prior memory-efficient transformers — 8x on 8 A100s (512K vs 64K tokens) and up to 100M-token contexts at larger scales ²⁹. Striped Attention improves on Ring by striping tokens uniformly across devices instead of contiguous blocks, reaching up to 1.45x end-to-end training throughput at 256K sequence length ⁵. LoongTrain’s 2D attention (head-parallel plus context-parallel) improves model FLOPs utilization by up to 2.88x over Megatron-CP and 1.49-1.53x over DeepSpeed-Ulysses ⁶. DeepSpeed-Ulysses keeps communication volume constant at O(M/P) versus O(M) for baselines, training 2.5x faster with 4x longer sequences than Megatron-LM sequence parallelism and exceeding 175 TFLOPS/GPU at 1M-token training ¹. UPipe refines Ulysses with head-level chunking, cutting attention intermediate-tensor memory by up to 87.5% and enabling 5M-token contexts at matched training speed ⁷¹.

Communication-computation overlap beyond attention. FLUX over-decomposes communication and computation into tiles fused into a single GEMM kernel, overlapping up to 96% of communication for up to 1.24x/1.66x/1.30x training speedups across workloads ⁵⁷. Triton-distributed extends the Triton programming model to generate overlapping distributed kernels, with speedups of 1.09x-44.97x over PyTorch+NCCL baselines ⁵³. PipeMesh uses elastic pipeline scheduling to eliminate most of the 28% all-accelerator idle time caused by communication, improving training throughput 20.1-33.8% over baselines with recomputation adding under 1.9% ⁷³. DualPipe, documented in the DeepSeek-V3 report, achieves full overlap of forward and backward computation-communication phases with reduced pipeline bubbles — the bidirectional schedule that helped DeepSeek-V3 train at 3.7 days per trillion tokens on 2,048 H800s ³⁵. WASP takes a hardware-software co-design route, adding hardware-accelerated automatic warp specialization with warp-level register-file queues for 47% average runtime gains on pipeline-parallel GPU workloads ⁷².

Pros and cons. Pros: these kernels are the only way to train beyond single-device memory, with gains that compound with sequence length (8x-512x context extension ²⁹); overlap techniques are complementary to fusion and precision work ⁵⁷⁵³. Cons: they require multi-GPU clusters and are sensitive to interconnect topology ⁵⁶; the gains are measured against different baselines (Megatron-CP vs Ulysses vs Ring), making cross-paper comparison unreliable; and communication-overlap kernels add scheduling complexity that interacts with pipeline parallelism ⁷³³⁵. Use cases: long-context pretraining and RL with large rollouts ²⁹⁴⁷; frontier-scale training where per-device memory is the binding constraint ¹⁷¹; and MoE/expert-parallel clusters where all-to-all dominates step time ⁶⁹⁶⁷⁶⁶.

Theme 5 — Fused elementwise/optimizer kernels and the tooling layer

Fused training-op kernels. Liger Kernel, a library of Triton kernels for cross-entropy, normalization, RoPE, and other memory-heavy training ops, delivers on average +20% training throughput and -60% GPU memory across popular LLMs (LLaMA3-8B: +42.8% throughput, -54.8% memory at batch size 64) with loss parity ¹⁵. Fusing the optimizer into the forward or backward pass reduces training time by up to 20% by eliminating separate kernel launches and memory round-trips ³⁸. Block-wise dynamic quantization of optimizer states matches 32-bit Adam/AdamW/Momentum performance while cutting state memory 4x (8GB to 2GB for a 1B-parameter model) ⁷. The Muon optimizer achieves roughly 2x computational efficiency over AdamW at compute-optimal training — comparable quality at ~52% of the training FLOPs — and trained Moonlight 3B/16B MoE models on 5.7T tokens ³⁷. torch.compile (TorchDynamo + TorchInductor, which lowers to Triton on GPUs) delivers 1.41x training geometric-mean speedups across 180+ models on A100 with no code changes ⁷⁵.

The tooling layer determines who can build kernels. Triton’s tile-based DSL made custom kernels accessible at “minimal performance cost” without vendor-library dependence ², and it is the substrate of both torch.compile ⁷⁵ and training-kernel libraries like Liger ¹⁵. ThunderKittens shows the CUDA-template alternative — near-hand-tuned performance with higher-level primitives ⁵⁵. But the evaluation gap is real: a 2026 study found correctness-based kernel benchmarks accept catastrophically slow kernels — an idiomatic TileLang LayerNorm passes the KernelBench check while running over 300x slower than PyTorch — and that re-tuning a single kernel at one shape can take minutes ⁵⁴. On the TPU side, JAXBench shows conditioning on curated hardware documentation raises per-sample correctness from 5.8% to 37.3% in autonomous kernel generation, with 1.28x geomean speedup over the XLA baseline ⁶⁵.

Pros and cons. Pros: fused elementwise/optimizer kernels are the lowest-risk optimization — they require no numerical changes, preserve loss curves ¹⁵, and compound with everything else ⁷⁵; DSLs and compilers are what make the whole field accessible to non-vendor teams ²¹⁴. Cons: the per-kernel gains are smaller than attention or FP8 (20% vs 2x) ¹⁵³⁸; compiler-generated kernels can be catastrophically slow on unusual shapes, so benchmarks must be roofline-aware ⁵⁴; and optimizer-level changes (Muon) interact with hyperparameters and have a narrower evidence base than kernel fusion ³⁷. Use cases: fused training-op kernels for any training run, especially memory-bound small-batch settings ¹⁵; torch.compile as the zero-effort default ⁷⁵; DSLs for custom kernels with portability needs ²⁵²; hand-written CUDA/CUTLASS where peak performance is non-negotiable ⁵¹⁵⁵.

Theme 6 — Choosing among techniques: the use-case map

Across the five families, the evidence supports a consistent decision framework. For a standard dense pretraining run on Hopper-class hardware, the default stack is: FlashAttention-3-class attention kernels ¹³, FP8 GEMMs with BF16 master weights ¹⁶³⁵, fused normalization/cross-entropy/optimizer kernels ¹⁵, and torch.compile or equivalent graph compilation ⁷⁵ — together these account for most of the reported 1.4-2x end-to-end gains, with FP8 providing the largest single increment ¹⁷¹⁹. For long-context or RL training, the binding constraint moves to memory and communication, so ring/striped/Ulysses-class kernels and overlap techniques dominate the decision ²⁹⁵¹⁴⁷. For MoE models, the kernel choice is architectural: block-sparse or single-kernel fused dispatch for dense-scale training ²⁴²⁷, and communication redesign (DeepEP-class all-to-all, LSH-MoE compression, rail-aware scheduling) once expert-parallel communication dominates ⁶⁹⁶⁷⁷⁴. For non-NVIDIA hardware or small teams, the DSL/compiler path (Triton, torch.compile) buys most of the gains at a fraction of the engineering cost, with the caveat that unusual shapes need roofline-level validation ²⁵⁴⁵².

Where the evidence disagrees

FP8 training: production-proven or dangerously fragile? DeepSeek-V3, FP8-LM, MOSS, μnit, and Fully-FP8 all report FP8 at matched quality at scale ^1635171918, while the “To FP8 and Back” study reports seed-dependent divergence even in BF16 and FP8 failures on noisier data ⁶¹, and the low-precision attention analysis documents a concrete failure mechanism ⁴¹. The best available resolution is scale and data quality: the failures cluster in small runs and noisy-gradient settings, while the successes are all at 175B-671B scale with careful master-weight and scaling engineering — the disagreement is between “FP8 with production-scale safeguards” and “FP8 as a naive drop-in.”

Single-kernel MoE: universally faster? FlashMoE and SonicMoE report 1.86-5.7x over their chosen baselines ²⁷²⁸, yet the Triton-based dispatch study reports only 89-131% of MegaBlocks at inference batch sizes ⁵², and Tutel’s adaptive approach remains competitive on its own benchmarks ²⁵. The explanation is baseline and shape: single-kernel fusion pays off at large batch sizes and with NVSHMEM-class interconnects, while at small batches the launch overheads are amortized differently and simpler kernels close the gap.

DSL vs hand-written CUDA. The compiler papers report parity or better ¹⁴⁷⁵⁵², the independent evaluation finds order-of-magnitude slowdowns on unusual shapes ⁵⁴, and the CUTLASS/ThunderKittens camp argues for template-level control ⁵⁵⁵¹. This is a real disagreement about the distribution of shapes: DSLs win on the shapes the compiler authors tuned, and lose badly off-distribution; the resolution is workload-specific validation, not a winner.

FlashAttention numerical stability. One study measures ~10x deviation in isolation but low training impact ⁴²; another documents training-time loss explosions in BF16 flash attention with a mechanistic cause ⁴¹. These are consistent once separated by setting — isolation versus full training, and BF16-specific rounding interactions — but the field has not yet produced a unified stability criterion.

Gaps and open questions

No shared benchmark. Every paper benchmarks against a baseline of its own choosing, on hardware of its own generation, at shapes of its own design. MOSS is a step toward standardized FP8 training evaluation ¹⁷ and JAXBench toward standardized kernel-generation evaluation ⁶⁵, but there is no equivalent of a training-side MLPerf for kernels; the “Correct but Slow” study shows why the absence matters ⁵⁴. What would settle it: a community roofline-based training-kernel benchmark across hardware generations with fixed baselines.

The backward pass is under-analyzed. Most kernel papers report forward-pass numbers (Sawtooth is explicitly forward-only ⁴⁸, FlashMoE defers training/backward fusion ²⁷), yet training spends roughly half its time in backward kernels and the numerical risk concentrates there ⁴¹. The backward pass of low-precision attention and single-kernel MoE is the field’s largest unexamined corner.

FP4 and MX are unresolved frontiers. FP4 training works at 13B/100B tokens with a small loss gap ⁶³⁶⁴, and 6-bit MX trains generative LMs ²², but there is no evidence yet at frontier scale, and the stability mechanisms that made FP8 work (master weights, scaling) have not been re-derived for FP4.

Portability is unmeasured. Almost all evidence is NVIDIA Hopper/Ampere; the TPU evidence (GShard, JAXBench) uses different kernel abstractions entirely ³¹⁶⁵, and the Triton cross-platform story is only beginning ⁵². Whether the FP8 and MoE kernel results transfer to AMD, Intel, or TPU training stacks is untested in the retrieved literature.

Confidence and limitations

Confidence in the qualitative map — five technique families, each with real and conditional gains — is high: it is consistent across dozens of independent papers. Confidence in specific numbers is moderate: nearly all performance claims are vendor- or single-group-benchmarked, hardware-specific, and measured against heterogeneous baselines; the two independent evaluations in this corpus (the kernel-evaluation gap study ⁵⁴ and the FP8 stability study ⁶¹) both caution against taking headline speedups at face value.

This review’s own limitations: 9 of 75 sources were abstract-only (publisher paywalled), including several load-bearing IEEE venues (WASP, PipeMesh, RailS, BF16-FMA); three load-bearing artifacts (DeepEP, DeepGEMM, DualPipe) have no papers and are documented from their repositories with vendor-stated numbers; Semantic Scholar was 429-unreachable from this IP all session and contributed nothing; one planned source (HyGIN, IEEE LCA 2026) was dropped for having no retrievable abstract anywhere; and SwitchBack, a remembered low-precision training paper, could not be resolved in any database and was dropped rather than cited from memory. The review covers 2018-2026 with a cutoff of 2026-08-15.