GPU kernel development for LLM inference: Triton, TileLang, FlashMLA, and FlashInfer compared

Why this question

Kernel engineering is now the binding constraint on LLM serving economics: decode is memory-bandwidth-bound, attention dominates latency, and the difference between a generic kernel and a tuned one is routinely 2-5x on the same GPU ⁵³⁴¹. Every serious serving stack — vLLM, SGLang, TensorRT-LLM, DeepSeek’s own deployment — is built on a kernel strategy, and practitioners must choose between writing kernels themselves (Triton, TileLang), integrating a kernel library (FlashInfer), or adopting a vendor’s reference kernels (FlashMLA). The choice is consequential and poorly documented: the four tools are frequently name-dropped together but are not substitutes for each other, and the marketing around each obscures what the evidence actually supports. This review asks which tool is appropriate for which job, and how the performance claims survive contact with independent evaluation.

Scope and methods

Question. How do Triton, TileLang, FlashMLA, and FlashInfer compare as GPU kernel development tools for LLM inference, and which use cases does each serve best?

Criteria. Included: papers on the four tools directly; the techniques they implement (tiled compilation, layout inference, JIT, paged KV cache, MLA, flash-attention variants); evaluations or comparisons involving them; and ecosystem papers documenting adoption (torch.compile, vLLM, SGLang, DeepSeek stack, quantization kernels). Years 2018-2026, English, any venue. Excluded: non-GPU accelerator kernels unless they are explicit portability evidence for these tools; paywalled-only work; USENIX-only papers with no DOI; LLM-kernel-generation benchmarks except where they bear on DSL evaluation quality.

Search and screening. Four databases (OpenAlex, arXiv, Crossref, Semantic Scholar) with 27 concept queries and ~30 exact-title resolution lookups (recorded verbatim in the frontmatter); Semantic Scholar was 429-unreachable from this IP all session and contributed nothing. The merged pool contained 1,886 records; after title screening (~120 candidates) and full-text screening, 55 sources were included: 43 read in full text, 12 abstract-only. A second, training-focused wave was then run at the reviewer’s request: 14 additional concept queries produced a 723-record pool, from which 9 sources were added (all read in full text), bringing the total to 64. Two candidates were dropped for unretrievable metadata (an ACL’26 MLA paper with no abstract anywhere and no arXiv version; an ACL’25 paper rescued at the last minute via its arXiv preprint). One remembered title (“Optimizing GPU Kernels by Reducing Memory Access and Divergence with TileLang”) proved to be a phantom and was dropped rather than forced. Every DOI was verified against Crossref or the arXiv abs page; published versions (PLDI’19 Triton, PPoPP’25 MARLIN, NeurIPS’24 SGLang and FlashAttention-3, NeurIPS’25 FlashMoE and Fully-FP8-GEMM, ASPLOS’25 vAttention) are cited where they exist.

A note on FlashMLA. FlashMLA has no paper, no DOI, and no Zenodo record — it is a GitHub repository ¹¹¹⁰³⁴. This review treats the repository README (retrieved 2026-08-13) as the primary documentation for its capabilities and reports its performance numbers as vendor-stated, cross-checked where possible against third-party kernel studies ¹²⁵⁴. Claims about the tool that cannot be attributed to a retrievable source are flagged as such.

The landscape

The literature divides cleanly into three layers. The foundations are the flash-attention lineage — FlashAttention’s IO-aware tiling ⁸, its parallelization rewrite ⁹, Hopper-era warp specialization ⁴⁴, and the proof that the tiling is I/O-optimal ⁴³ — plus paged KV-cache management for serving ³¹. The tool layer holds the four subjects of this review and their immediate relatives: the compiler lineage from TVM through Ansor and TensorIR to TileLang ⁵⁵⁷¹⁵, the CUDA-template tradition exemplified by CUTLASS ²⁶ and ThunderKittens ⁴⁵, and newer kernel-generating systems such as Axe ²⁴ and ML-Triton ⁴⁸. The system layer shows what the kernels actually power: vLLM ³¹, SGLang ⁵⁹, BatchLLM ⁵⁸, DeepSeek’s serving stack ¹¹¹⁰, and the surveys that frame the efficiency problem ⁶¹⁵³.

Three structural facts stand out. First, the field is young and concentrated: Triton’s paper is from 2019 ⁴⁶, TileLang from 2025 ⁴⁹, FlashInfer from 2025 ⁵¹, FlashMLA from February 2025 (repo only). Second, almost all quantitative evidence is authored by the tool builders themselves or by groups with a stake in the comparison; independent evaluations are rare and only now emerging ³²⁴. Third, the four tools are not competing products but complementary layers: the DSLs are how new kernels get written, and the libraries are how proven kernels get deployed — the overlap is at the boundary where FlashInfer’s JIT and Triton’s compiler both claim the “custom attention variant” use case ⁵¹¹³. A fourth structural fact emerges once training is included: the training-kernel literature (MoE block-sparse kernels ¹⁸, distributed single-kernel MoE ¹, FP8 training GEMMs ³⁸⁵⁶) is written almost entirely in CUDA and CUTLASS, not in the DSLs this review compares — a tooling divide examined in Theme 7.

Theme 1 — What the four tools are

Triton is a tile-based DSL and compiler: programs express computation over statically shaped tiles, and tile-level optimization passes lower them to efficient GPU code, yielding portable matrix-multiplication and convolution kernels “on par with hand-tuned vendor libraries” at higher productivity than CUDA ⁴⁶. Its defining adoption is as the codegen backend of PyTorch’s torch.compile: TorchInductor lowers compiled graphs to Triton kernels, delivering 2.27x inference and 1.41x training geometric-mean speedups across 180+ models on A100 ². Triton is also the default target language for LLM kernel generation ⁶ and the substrate for production training kernels such as Liger ²⁵. Its compiler emits PTX and is JIT-compiled at first call, with third-party backends extending it to AMD, Intel, CPUs, and custom accelerators ³³⁶⁴⁴⁸.

TileLang is a Pythonic DSL on top of the TVM compiler stack, designed so developers “focus on productivity without sacrificing the low-level optimizations necessary for state-of-the-art performance” ⁴⁹. Its differentiators are automatic layout inference and automatic selection of Hopper-specific instructions (TMA bulk copies, WGMMA async matrix operations) with warp specialization, which the authors argue is what lets it beat hand-tuned kernels on attention ⁴⁹. It inherits TVM’s multi-backend infrastructure (CUDA, Metal, CPU paths; the README documents ongoing Blackwell and Metal work), and its paper reports speedups of 1.36x, 1.41x, and 1.70x over FlashAttention-3, Triton, and PyTorch respectively on FlashAttention benchmarks, with the largest gains at short sequence lengths where FlashAttention-3’s fixed tile size is suboptimal ⁴⁹.

FlashInfer is a library and kernel generator for inference: unified APIs for attention, GEMM, and MoE with pluggable backend implementations (FlashAttention-2/3, cuDNN, CUTLASS, TensorRT-LLM) and automatic backend selection ⁵¹. It covers paged and ragged KV cache, decode/prefill/append kernels, MLA, cascade attention, sparse attention, fused prefill+decode (POD), FP8/FP4 quantization, grouped GEMM for LoRA, fused MoE with multiple routing schemes, sorting-free sampling, speculative decoding, and communication primitives — on every NVIDIA architecture from Turing (SM75) through Blackwell ⁵¹. Kernels are JIT-compiled on first use, with pre-built cubin packages to eliminate the compile stall ⁵¹. It is the attention engine of SGLang, vLLM, TensorRT-LLM, TGI, MLC-LLM, and others per the project README, and its paper reports 29-69% inter-token-latency reduction versus compiler (Triton) backends in LLM-serving benchmarks, 28-30% latency reduction for long-context inference, and 13-17% speedup for parallel generation ⁵¹.

FlashMLA is DeepSeek’s library of MLA attention kernels, powering DeepSeek-V3 and V3.2 ¹⁰³⁴. Per the repository README, the dense MLA decoding kernel reaches up to 3,000 GB/s in memory-bound configuration and 660 TFLOPS in compute-bound configuration on H800 SXM5; the September 2025 sparse kernels for DeepSeek Sparse Attention (DSA) reach 640 TFLOPS prefill and 410 TFLOPS decoding with an FP8 KV cache, and 1,460 TFLOPS forward MHA prefill on B200 via NVIDIA’s contributed SM100 kernels. It supports only SM90/SM100, requires CUDA 12.8+, and exposes a minimal Python API (a metadata pre-pass plus one kernel call per decoding step) around a paged KV cache with an optional FP8 format. There is no extension surface: it is a library of specific kernels, not a language. Community ports exist for AMD (via the AITER project), MetaX, Moore Threads, Hygon, Intellifusion, and Iluvatar hardware. The MLA mechanism it implements was introduced in DeepSeek-V2, which reports a 93.3% KV-cache reduction and 5.76x maximum generation throughput versus the prior DeepSeek 67B ¹¹.

Theme 2 — Performance: what the numbers actually support

Attention kernels. On the serving workloads the libraries target, FlashInfer’s advantage over Triton-generated attention is consistent and large: 29-69% lower inter-token latency in serving benchmarks, and up to 20x over FlexAttention’s generated kernels at long sequence lengths, where template block sizes become the bottleneck ⁵¹¹³. FlashAttention-3, the hand-tuned reference, reaches 740 TFLOPS/s FP16 (about 75% of H100 peak) ⁴⁴. A 2025 case study shows a generic Triton paged-attention kernel can be driven from 19.7% to 105.9% of FlashAttention-3’s performance through Q-blocking, parallel tiled softmax, and static launch grids — but only with expert-level optimization, and the 100-300 microsecond Triton launch overhead forced CUDA/HIP-graph integration in deployment ⁴¹. TileLang’s paper claims its generated attention kernels beat FlashAttention-3 by 1.36x on average ⁴⁹, while the independent “Correct but Slow” study found the DSLs’ tuned kernels hold 59.7-97.9% (Triton) and 78.1-94.2% (TileLang) of cuBLAS on GEMM, with both weakest at very large square shapes ³². The honest synthesis: hand-tuned CUDA still defines the ceiling on attention; the DSLs reach or occasionally exceed it on specific shapes with real tuning effort; the libraries ship the tuned result.

MLA kernels. FlashMLA’s 3,000 GB/s decode bandwidth and 660 TFLOPS are vendor-stated ¹⁰. Third-party work both confirms the workload’s importance and disputes the kernel’s optimality. FlashMLA-ETAP (Tencent), a transposed-pipeline MLA kernel for H20, reports 2.78x over FlashMLA at 64K sequence length and 5.24x/4.94x over FlashAttention-3/FlashInfer ¹². TyphoonMLA shows that the “absorbed” MLA formulation FlashMLA uses is not universally best: a hybrid naive+absorb kernel improves attention throughput up to 3.24x over absorb-only baselines (including FlashMLA and FlashInfer) on GPUs with shared-prefix workloads and batch sizes above a threshold of roughly 61 for DeepSeek-V3, at only ~3% HBM overhead, with up to 1.48x end-to-end tokens-per-second over FlashMLA ⁵⁴. The takeaway is that MLA kernel design is still actively contested, and FlashMLA is a strong but not final reference point.

Everything else. The system-level numbers show kernels are one layer of a stack: vLLM’s paged attention delivers 2-4x serving throughput over earlier systems ³¹, SGLang up to 6.4x via radix caching ⁵⁹, BatchLLM 1.3-10.8x via prefix sharing and throughput-oriented batching ⁵⁸, MPK’s persistent mega-kernels up to 1.7x lower end-to-end latency than kernel-per-operator serving ⁷, and FlashDecoding++ up to 4.86x decode speedup over unoptimized baselines ²². Decode kernels are fundamentally bandwidth-bound — a low-batch decode of PaLM 540B on TPU v4 runs at 29 ms per token ³⁹ — which is why quantized attention and GEMM kernels matter: SageAttention is about 2.1x faster than FlashAttention-2 via INT8 attention ⁵⁵, and MARLIN’s ~3.9x on 4-bit weights approaches the theoretical 4x memory-bound ceiling ¹⁷. Triton-based kernels in the ecosystem deliver 65-124% gains from split-K decomposition on GEMM ²³, 89-131% of CUDA Megablocks for MoE dispatch at small batches ³⁶, and roughly 20% training throughput gains at 60% lower memory for Liger ²⁵. One independent characterization of ML compilers on consumer GPUs found TensorRT-LLM’s AOT compilation was the only workflow that consistently beat PyTorch eager end-to-end (up to ~60% on TinyLlama at small batch), while torch.compile-based JIT backends did not — a sobering counterpoint to the compiler papers ⁴.

Theme 3 — Programmability and the evaluation gap

The DSLs’ core promise is productivity: Triton’s pitch is vendor-library performance with custom-kernel flexibility at Python-level effort ⁴⁶, and TileLang’s is concise kernels that “stay close to” hand-tuned performance without manual Hopper programming ⁴⁹. The productivity angle extends up the stack: compiler-native approaches like FlashLight generate attention variants inside the PyTorch compilation pipeline, beating FlexAttention by up to 1.48x on score-modulated variants and improving AlphaFold gated self-attention by over 5x without hand-written CUDA ⁵². The evidence supports the productivity claim more strongly than the performance claim. Triton programs are short enough that LLMs can generate them ⁶, and whole training-kernel libraries are written in it ²⁵. TileLang’s layout inference removes a class of expert decisions entirely ⁴⁹; Axe generalizes layout abstraction to multi-GPU settings and reports 97-100% of cuBLAS on B200 FP16 GEMM, beating FlashInfer by 1.20-1.36x on fused MoE layers ²⁴; ML-Triton’s multi-level lowering reaches at least 95% of Intel’s expert kernels on PVC, including paged attention ⁴⁸.

But the independent “Correct but Slow” study documents a systematic evaluation gap: correctness-based kernel benchmarks accept catastrophically slow kernels — an idiomatic TileLang LayerNorm generated for KernelBench passes correctness while running 323x slower than PyTorch on GH200 (51.5 ms vs 0.16 ms), and re-tuning a single kernel at one shape took 211 seconds ³². The study also found genuine codegen weaknesses: Triton trails on convolution (28.9% of baseline) and large GEMM (59.7%), and TileLang collapses on normalization and reduction kernels (as low as 0.33% of baseline) — with the authors attributing the reduction failures to authoring artifacts rather than the compilers themselves ³². The DSLs also carry runtime costs invisible in microbenchmarks: Triton’s 100-300 microsecond launch overhead ⁴¹ and JIT first-call compilation are real deployment issues that libraries mitigate with CUDA graphs ⁴¹, pre-built cubins ⁵¹, or hybrid JIT plus CUDA-graph runtimes (up to 66% lower time-to-first-token versus TensorRT-LLM in one study) ⁵⁰. A systematic comparison of GPU JIT interfaces finds all surveyed interfaces performance-portable across MI300A and H100 but with materially different overhead and usability trade-offs ¹⁶.

Theme 4 — Portability and hardware coverage

Portability is where the two families diverge most sharply. Triton is the most portable DSL in practice: third-party backends run it on Intel GPUs (ML-Triton reaches ≥95% of Intel’s expert XeTLA kernels) ⁴⁸, on CPUs (a Triton-CPU compilation and runtime co-optimization reaches up to 2.0x on compute-bound GEMM for ARM SVE) ³³, and on custom accelerators (Meta’s MTIA backend with minimal language extensions, deployed across ~60 model types covering 50% of layers) ⁶⁴. TileLang inherits TVM’s multi-backend architecture ⁴⁹⁵. By contrast, FlashInfer and FlashMLA are NVIDIA-only upstream ⁵¹; FlashMLA’s non-NVIDIA story is community ports ¹⁰, and FlashInfer’s README lists no ROCm support. The cost of portability is real on both sides: the ThunderKittens-to-SYCL port (a CUDA-template library, not a DSL) required PTX replacements and register-layout rewrites to come within 5% of Intel’s hand-optimized GEMM ²⁸, and the MTIA Triton backend needed a new compiler backend plus language extensions ⁶⁴. Even Triton’s “write once” promise is qualified: the same Triton source needs per-backend tuning (e.g., split_k differs by GPU) ²³. On the attention side, the only non-NVIDIA production attention-kernel evidence is Ragged Paged Attention on TPU7x, reaching 86% memory-bandwidth utilization in decode — built with the same IO-awareness principles but no code sharing ²⁹.

Theme 5 — Use cases and ecosystems

The four tools occupy different niches, and the ecosystem evidence maps them cleanly:

• Triton is the default choice when you need a custom kernel inside the PyTorch ecosystem: it is what torch.compile generates ², what Liger uses for training ops ²⁵, what quantization kernels are built with ²³¹⁷, what MoE dispatch has been ported to for cross-vendor routing ³⁶, and what distributed kernel overlap is being explored with ⁶⁰. Its profiling ecosystem is maturing (Proton adds multi-level adaptive profiling with SQL-like queries) ⁶³. Its weakness is peak attention performance and launch overhead ⁴¹.
• TileLang fits research and production kernels where per-kernel performance is worth investment: its paper demonstrates attention kernels beating FlashAttention-3 at short-to-medium sequence lengths ⁴⁹, and its TVM lineage makes it a candidate for heterogeneous deployment ⁵¹⁵. It is younger, with a smaller ecosystem and a demonstrated risk of large performance cliffs on kernels outside its strengths ³².
• FlashInfer is the integration choice for serving systems: it provides the attention variants and auxiliary kernels (sampling, speculative decoding, MoE, quantization) that engines need, across the full NVIDIA range, with JIT customization for new attention variants ⁵¹. Its adoption across SGLang, vLLM, TensorRT-LLM, and TGI makes it the de facto attention engine of open-source serving ⁵⁹³¹.
• FlashMLA is for serving DeepSeek-family MLA models at reference performance, and more broadly as the reference implementation for MLA kernel design — its absorb-mode formulation, paged FP8 KV cache, and DSA sparse kernels are now the baseline that third-party MLA kernel papers measure against ¹²⁵⁴³⁴.

Theme 6 — MLA as the shared workload

MLA matters to this comparison because it is the workload FlashMLA exists for and the one where the kernel-design questions are most live. The mechanism compresses KV into a low-rank latent vector, cutting KV cache by 93.3% and boosting maximum generation throughput 5.76x in DeepSeek-V2 ¹¹, and it is retained through V3 ¹⁰ and extended with sparse attention (DSA) in V3.2, which reduces attention complexity from quadratic to O(L·k) with k=2,048 selected tokens per query ³⁴. A hardware-centric analysis argues MLA shifts decode from memory-bound toward compute-bound, with more stable performance on bandwidth-limited hardware ¹⁹. Kernel engineering for MLA is still converging: the absorb formulation ¹¹ versus hybrid naive+absorb with shared prefixes ⁵⁴ versus transposed pipelines for long context ¹², plus FP8 KV caching with bf16 compute ³⁴. The workload is also widening beyond DeepSeek: MHA-to-MLA conversion methods now port the mechanism into existing models with 84-96% KV-cache reductions ¹⁴²⁷, GQA-to-MLA conversions improve accuracy at matched KV budgets ⁶²³⁰, and an analysis of the latent bottleneck finds it globally over-provisioned — mean effective rank 46% of its 128-dimension capacity — implying further compression headroom ⁴². Even the KV-cache transport economics change: routing an absorbed-MLA query row across instances costs ~26x less than moving KV chunks, reshaping disaggregated-serving design ³⁵.

Theme 7 — The training side: where the four tools stand

The primary question of this review is inference-scoped, but the tools are not inference-only, and the training literature is a distinct body of work with a different tooling profile. The boundary within the four: FlashInfer and FlashMLA are forward-only serving libraries by design (no backward passes; everything they ship — paged KV, decode, MLA, sampling — serves generation), while Triton and TileLang are dual-use DSLs whose papers and ecosystems span training and inference ⁴⁹².

Triton’s training footprint is real and growing. torch.compile’s training path lowers to Triton (1.41x geomean on A100) ²; Liger’s fused training kernels deliver ~20% throughput at ~60% lower memory ²⁵; and MOSS — a lossless FP8 training recipe with microscaling — implements its custom GEMM kernels in Triton, matching BF16 accuracy on OLMo-7B and LLaMA-2-7B ⁵⁶. The attention lineage itself is training-born: FlashAttention-1/2/3 were published as training kernels (BERT-large 15% faster, GPT-2 3x ⁸; 225 TFLOPS/s in GPT training ⁹; backward-pass results headline FlashAttention-3 ⁴⁴), and the I/O-optimality proof applies to forward and backward alike ⁴³.

But the training-kernel frontier is written in CUDA and CUTLASS, not the DSLs. The load-bearing training-kernel literature of 2022-2026 is hand-written: MegaBlocks’ block-sparse MoE kernels (never dropping tokens, up to 40% end-to-end training speedups over Tutel and 2.4x over Megatron-LM-trained MoEs) ¹⁸; FlashMoE’s single-kernel distributed MoE (6x latency, 9x GPU utilization, 5.7x throughput versus state-of-the-art baselines despite running FP32 against FP16 baselines, collapsing 33-550 GPU ops in prior implementations to one) ¹; and a ground-up CUTLASS reimplementation of FlashAttention-2 that runs 20-50% faster than the reference FA2 kernel on Hopper ³. The FP8 training line is likewise kernel-driven: FP8-LM cut real memory 39% and trained GPT-175B 75% faster than BF16 Megatron-LM on H100 ³⁸; fully-FP8 architectures now run every transformer GEMM in FP8 in both passes, matching BF16 downstream performance ²¹; the lineage begins with 8-bit FP training at 2-4x speedup without accuracy loss ⁴⁷. Where Triton appears in this literature, it is in the ecosystem layer (Liger, MOSS, torch.compile), not the headline systems — this review found no documented training-system adoption of TileLang, whose attention benchmarks include backward passes but whose training ecosystem is not visible in the literature.

The communication half of training kernels is even less documented. DeepSeek’s MoE training pairs GEMM kernels with all-to-all communication: DeepEP (the MoE communication library, listed at 432 GPU ops in FlashMoE’s profiling of a Megatron-LM+DeepEP stack) and DualPipe (pipeline overlap) are described in the V3 technical report and GitHub repositories rather than papers of their own ¹⁰¹. FlashMoE’s single-kernel design is partly a reaction to this op-count explosion ¹. On the GEMM side, DeepGEMM — DeepSeek’s FP8 GEMM library, also GitHub-only — has become the reference baseline that newer MoE kernel work measures against, with SonicMoE reporting 25% forward and 15% backward speedups over it on Blackwell ²⁰.

What this means for the four tools. For training work, the choice is not Triton versus TileLang versus a library: it is between the expert-CUDA path (MegaBlocks, FlashMoE, CUTLASS — highest ceiling, highest cost) and the Triton-ecosystem path (torch.compile, Liger, MOSS — 20-40% class gains with Python-level effort). FlashInfer and FlashMLA have no training story and show no sign of growing one; their value is that the same tiling techniques they implement are what the training kernels use ⁸⁴³. TileLang’s training applicability is the least-evidenced claim in this review’s scope.

Where the evidence disagrees

Can DSLs match hand-tuned CUDA? TileLang’s paper reports 1.36x over FlashAttention-3 ⁴⁹; the Anatomy study reports Triton reaching 105.9% of FlashAttention-3 after expert optimization ⁴¹; FlashInfer reports compiler backends 29-69% slower on serving workloads ⁵¹; and the independent evaluation finds large, shape-dependent gaps and catastrophic LLM-generated outliers ³². These disagree less than they appear: the DSL papers report best-case tuned kernels on favorable shapes; the serving numbers reflect default-codegen performance under latency constraints; and no study controls for tuning effort. The defensible synthesis is that the DSL ceiling is near hand-tuned on GEMM and attention, but the default is far below it, and the gap is workload- and shape-shaped, not uniform ³²⁴.

Is paged KV cache a win or a loss? vLLM’s paged attention delivers 2-4x serving throughput ³¹, but vAttention quantifies the kernel-level cost: vLLM’s paged decode kernel is up to 2.8x slower than FlashAttention-2’s, paged prefill kernels are up to 37-42% slower, and paged kernels execute 7-13% more instructions ⁴⁰. Both claims are true: paging trades kernel efficiency for system-level memory efficiency, and the disagreement is about where the trade-off should be drawn — FlashInfer and vLLM accept it, vAttention proposes virtual-memory-based management instead ⁴⁰⁵¹.

Which MLA kernel formulation is best? FlashMLA’s absorb mode is the production reference ¹¹, but TyphoonMLA shows hybrid naive+absorb winning by 3.24x on shared-prefix workloads ⁵⁴ and ETAP shows a transposed pipeline winning by 2.78x at 64K context on H20 ¹². The resolution is workload dependence: absorb wins for short-context decode without shared prefixes; the alternatives win where prefix reuse or very long context dominate. No study compares all three on identical hardware and workloads.

JIT versus AOT. FlashInfer mitigates JIT cost with pre-built cubins ⁵¹; Triton’s launch overhead pushes deployments to CUDA graphs ⁴¹; a hybrid JIT-plus-CUDA-graph runtime beats TensorRT-LLM (AOT) by up to 66% time-to-first-token in one setting ⁵⁰; and the independent characterization finds AOT the only consistently winning workflow on consumer GPUs ⁴. The evidence supports AOT for static workloads and JIT+graphs where shape diversity or customization matters; the trade-off is not settled quantitatively.

Training kernels: expert CUDA or DSLs? The training literature contains both a CUDA/CUTLASS frontier (MegaBlocks, FlashMoE, FP8-LM, the CUTLASS-FA2 case study) ¹⁸¹³⁸³ and a working Triton ecosystem (torch.compile, Liger, MOSS) ²²⁵⁵⁶. These do not contradict each other — they operate at different layers (headline systems vs ecosystem kernels) — but the literature gives no direct comparison of the two paths on the same training workload, so the marginal value of the expert-CUDA path over the Triton path is unquantified. A related tension: every FP8 training recipe reviewed claims lossless accuracy at scale ^38563721, but the recipes disagree on why (outlier-reducing architectures ²¹ versus scaling-law interventions ³⁷ versus format innovations ⁵⁶), and none has been independently replicated at frontier scale.

Gaps and open questions

• No independent head-to-head of the four tools. Every direct comparison in this review is authored by a party with a stake (DeepSeek, Tencent, EulerLab, the FlashInfer team, the TileLang team). What would settle the debate: a hardware-pinned benchmark of all four on identical workloads — prefill, decode, paged long-context, MLA with shared prefixes — run by an unaffiliated group.
• FlashMLA has no paper. Its capabilities and numbers rest on a README, a blog, and third-party evaluations. A peer-reviewed specification of the kernel design (absorb formulation, FP8 KV layout, DSA sparse kernels) would make the reference implementable and auditable.
• Serving-level MLA benchmarks are thin. MLA kernel papers measure single-instance throughput; mixed-batching, multi-node, and disaggregated-serving behavior of the different MLA formulations is untested.
• The evaluation gap is one study deep. The “Correct but Slow” finding that correctness benchmarks certify 323x-slow kernels ³² has not been replicated, and it excluded attention — the most important kernel class — because the DSL implementations were algorithmically different.
• Blackwell sparse attention is immature. FlashMLA’s B200 sparse kernels are described in the README as “not really optimized yet”, and no independent sparse-attention kernel evidence exists for SM100.
• Non-NVIDIA coverage is lopsided. FlashInfer has no ROCm support; FlashMLA’s AMD path is a community port; Triton’s AMD path still needs per-kernel tuning ²³³⁶. Whether the serving libraries’ advantages survive off-NVIDIA is untested.
• The MLA rank-overprovisioning result is a small-model finding. The 46% effective-rank utilization result ⁴² comes from a 114M-parameter model; whether it holds at DeepSeek scale is unknown.
• Training-side tooling questions are open. No study compares the expert-CUDA and Triton-ecosystem paths on the same training workload; TileLang has no documented training adoption at all; DeepGEMM, DeepEP, and DualPipe — load-bearing components of the most-cited open training stack — have no peer-reviewed specifications, only repos and third-party mentions ²⁰¹¹⁰; and the four FP8-lossless recipes ^38563721 have not been replicated head-to-head at the same scale.

Confidence and limitations

Confidence is moderate. The qualitative landscape — two families of tools with distinct niches, FlashMLA as a workload-specific reference, the DSL ceiling near hand-tuned with a large default gap, and a training side dominated by expert CUDA with Triton in the ecosystem layer — is consistent across all 64 sources and unlikely to be overturned. The quantitative picture is weaker: most numbers are vendor- or group-reported, benchmark baselines vary, and no study covers all four tools. Limitations of this review: 12 of 64 sources were abstract-only (including the PLDI’19 Triton paper, whose full text is paywalled at ACM; its claims here rest on its abstract and the rescued full abstract); FlashMLA is documented from its repository rather than a paper, and its performance claims are vendor-stated; DeepGEMM, DeepEP, and DualPipe are likewise GitHub-only and are discussed only via the DeepSeek reports and third-party evaluations that name them; one source’s extracted table numbers (MHA2MLA-VLM) could not be fully verified and are cited qualitatively; two candidates with unretrievable metadata were dropped (an ACL’26 MLA paper with no arXiv version, and the phantom “Optimizing GPU Kernels with TileLang” title); Semantic Scholar contributed nothing due to sustained rate-limiting; the review is English-language and NVIDIA-centric by the shape of the literature.