FlashInfer: The Attention Kernel Library That Proves the Bottleneck in LLM Inference Was Never the Model. It Was the Memory Access Pattern.

Sponsored by

Production benchmarks: 29-69% inter-token-latency reduction versus compiler backends, 28-30% latency reduction for long-context inference, 13-17% speedup for parallel generation. It now powers SGLang, vLLM, TensorRT-LLM, TGI, MLC-LLM, and five other production serving frameworks, and supports every GPU from Turing (2018 T4) through Blackwell B200/B300 and DGX Spark.

SnackOnAI Engineering | Senior AI Systems Researcher | Technical Deep Dive | July 13th, 2026

The LLM inference stack has two distinct performance problems that most engineers conflate. The first is compute: how fast can the GPU multiply matrices. The second is memory access: how efficiently can the GPU read and write the key-value cache that stores prior attention computations during generation.

Compute efficiency has been the focus of most optimization work since FlashAttention in 2022. The insight there was that naive attention requires reading and writing the full attention matrix from GPU HBM, which is bandwidth-bound, not compute-bound. FlashAttention fused the operations into a tiled kernel that stayed in fast SRAM, eliminating most of the HBM traffic. That was a genuine step forward.

FlashInfer's thesis is that the memory access problem is not solved. It has shifted. When you serve real workloads at scale, the KV-cache is not a contiguous tensor with a predictable memory layout. It is a patchwork of allocated blocks from different requests with different prefixes, different context lengths, and different sparsity patterns. Reading it efficiently requires different kernel configurations for every combination of these properties. Kernel libraries that assume a fixed layout, compiled ahead of time for a single case, miss most of the real workload.

FlashInfer addresses this with three ideas: a format that represents the KV-cache in block-sparse and composable terms (capturing the actual structure, not an approximation), a JIT compilation system that generates the right kernel for each specific combination of configuration at runtime, and a scheduling algorithm that does load balancing across requests dynamically while keeping CUDAGraph compatibility.

Scope: FlashInfer's block-sparse composable format, the JIT kernel generation system and its GPU coverage (SM 7.5 through SM 12.1), the load-balanced CUDAGraph-compatible scheduler, the attention operation types (Paged KV, MLA, Cascade, POD, Sparse), and the full feature surface (GEMM, MoE, sampling, communication). Not covered: the internals of the sorting-free sampling kernels in depth, or the NVSHMEM distributed memory integration.

What It Actually Does

FlashInfer is a kernel library with a Python interface. You install it, and production serving frameworks (SGLang, vLLM, TensorRT-LLM, TGI) use it as the attention and GEMM backend. For teams building custom serving infrastructure, FlashInfer exposes direct Python and C++ APIs for all operations.

Quick start:

# Core package (JIT-compiles kernels on first use)
pip install flashinfer-python

# Optional: pre-compiled kernels for all architectures
pip install flashinfer-python flashinfer-cubin

# JIT cache for specific CUDA version (faster startup, offline use)
pip install flashinfer-jit-cache --index-url https://flashinfer.ai/whl/cu129

# Blackwell (SM 100+) with CuTe DSL kernels
pip install flashinfer-python[cu13]

# Verify installation
flashinfer show-config

GPU coverage (SM 7.5 through SM 12.1):

Architecture	SM	Example GPUs
Turing	7.5	T4, RTX 20 series
Ampere	8.0, 8.6	A100, A10, RTX 30 series
Ada Lovelace	8.9	L4, L40, RTX 40 series
Hopper	9.0	H100, H200
Blackwell	10.0, 10.3	B200, B300
Blackwell	11.0	Jetson Thor
Blackwell	12.0, 12.1	RTX 50 series, DGX Spark

The Architecture, Unpacked

The load-balanced CUDAGraph-compatible scheduler is the contribution most engineers underestimate. CUDA kernel launch overhead — the CPU-side work of dispatching GPU kernels — is substantial at the token rates modern serving systems achieve. CUDAGraph eliminates it by replaying a pre-captured graph. But capturing requires static shapes. FlashInfer's tile-based scheduling makes the kernel launch static while keeping the actual computation dynamic. This is the primary source of the inter-token-latency improvements over compiler-based backends.

The Code, Annotated

Snippet One: Paged KV-Cache Attention and MLA

# FlashInfer: attention with paged KV-cache and MLA (DeepSeek Multi-Latent)
# Source: flashinfer-ai/flashinfer README + docs (Apache 2.0)
# Design intent: one API for both standard paged attention and compressed MLA format

import torch
import flashinfer

# ─── STANDARD DECODE ATTENTION (simplest case) ───────────────────────────────
q = torch.randn(32, 128, device="cuda", dtype=torch.float16)    # [num_heads, head_dim]
k = torch.randn(2048, 32, 128, device="cuda", dtype=torch.float16) # [kv_len, num_kv_heads, head_dim]
v = torch.randn(2048, 32, 128, device="cuda", dtype=torch.float16)

output = flashinfer.single_decode_with_kv_cache(q, k, v)
# ← This is the baseline: single sequence, contiguous KV cache
# ← At production batch sizes, this is NOT how KV cache is stored

# ─── PAGED KV ATTENTION (how vLLM/SGLang actually store KV cache) ────────────
# In production, KV cache is paged: allocated in fixed-size blocks
# Different requests in the same batch have different page tables
# FlashInfer's BatchDecodeWithPagedKVCacheWrapper handles this correctly

num_heads = 32
head_dim = 128
page_size = 16  # tokens per page
max_batch_size = 128

# Pre-allocate page pool (much larger than one batch)
kv_pool = torch.zeros(
    10000, 2, page_size, num_heads, head_dim,  # [num_pages, 2 (K+V), page_size, heads, dim]
    device="cuda", dtype=torch.float16
)

# Each request has a list of page indices (its virtual address space)
# page_indices[i] = which pages belong to request i
page_indices_batch = [
    torch.tensor([0, 1, 2, 3], dtype=torch.int32, device="cuda"),   # request 0: 4 pages
    torch.tensor([4, 5], dtype=torch.int32, device="cuda"),          # request 1: 2 pages
    torch.tensor([6, 7, 8, 9, 10], dtype=torch.int32, device="cuda"), # request 2: 5 pages
]
# ← Different lengths per request is the "ragged" in Ragged KV-Cache
# ← A naive contiguous implementation would pad to max length → wasted memory
# ← FlashInfer handles variable-length batches natively

wrapper = flashinfer.BatchDecodeWithPagedKVCacheWrapper(
    workspace_buffer=torch.empty(128 * 1024 * 1024, dtype=torch.uint8, device="cuda"),
    kv_layout="NHD",    # [pages, heads, dim] layout
    use_cuda_graph=True, # ← KEY: enable CUDAGraph for low kernel-launch overhead
)

# Plan call: this is where FlashInfer does its load-balanced tile assignment
# ← Plans are pre-computed and cached for CUDAGraph compatibility
# ← If sequence lengths change, re-plan (but the kernel launch is still captured)
wrapper.plan(
    indptr=compute_indptr_from_page_lists(page_indices_batch),
    indices=flatten_page_indices(page_indices_batch),
    last_page_len=compute_last_page_lengths(page_indices_batch, page_size),
    num_qo_heads=num_heads,
    num_kv_heads=num_heads,
    head_dim=head_dim,
    page_size=page_size,
    data_type=torch.float16,
)

# Run attention
q_batch = torch.randn(len(page_indices_batch), num_heads, head_dim, device="cuda", dtype=torch.float16)
output = wrapper.run(q_batch, kv_pool)

# ─── MLA ATTENTION (DeepSeek Multi-Latent Attention) ──────────────────────────
# MLA compresses the KV cache: instead of storing full K and V,
# it stores a lower-rank latent vector and the W^K, W^V projection weights
# ← KV cache size reduction: up to 5.75x compression at DeepSeek-V2 scale
# ← FlashInfer natively supports MLA without decompression + standard attention
# ← The kernel handles the latent → full KV projection inside the attention pass

mla_wrapper = flashinfer.BatchDecodeWithSharedPrefixPagedKVCacheWrapper(
    workspace_buffer=torch.empty(128 * 1024 * 1024, dtype=torch.uint8, device="cuda"),
    use_cuda_graph=True,
)
# MLA-specific: compressed latent KV cache
# flashinfer.mla module handles the full MLA path including RoPE application
# ← This is used by SGLang and vLLM for DeepSeek-V2/V3/R1 inference

The use_cuda_graph=True in both wrappers is the operational difference that matters. The plan phase computes tile assignments in a shape-compatible way. The run phase replays a CUDAGraph with those assignments. The result: kernel launch overhead drops from O(layers × batch_size) to O(1) per step. At 40 layers × 128 batch size, this is thousands of kernel launches eliminated per decode step.

Snippet Two: JIT Custom Attention and Sorting-Free Sampling

# FlashInfer: JIT custom attention variants and sorting-free sampling
# Source: flashinfer-ai/flashinfer tests/utils/test_jit_example.py + docs
# Design intent: users customize attention without writing CUDA

import torch
import flashinfer
from flashinfer.jit import gen_customize_single_decode_module

# ─── JIT CUSTOM ATTENTION TEMPLATE ──────────────────────────────────────────
# Example: sliding window attention (each token attends only to its last W tokens)
# This is a common variant in long-context models (e.g., Mistral)
# With FlashInfer JIT, you define the mask in Python and get an optimized CUDA kernel

# The custom kernel is defined via a mask function
# FlashInfer's JIT infrastructure compiles it with FlashAttention-2/3 backend
# ← No CUDA programming required. No C++ required.
# ← Compile-once-and-cache: kernel compiled on first call, cached for future use

# CUDA custom kernels for production (from JIT examples):
custom_module = gen_customize_single_decode_module(
    "sliding_window_attn",
    # Additional per-head parameters (loaded alongside Q, K, V)
    additional_input_tensor_var_names=["window_size"],
    additional_input_tensor_dtypes=["int32"],
    additional_input_tensor_ranks=[0],  # scalar per head
    # Custom score modification (before softmax):
    # Apply -inf to positions outside the sliding window
    custom_mask_module="""
    // Inline CUDA: any_float = attention_score, row_idx = query position, col_idx = key position
    template <typename DType>
    __device__ DType apply_sliding_window(DType score, int row_idx, int col_idx, int window_size) {
        return (row_idx - col_idx < window_size) ? score : DType(-5e4);  // -inf mask
    }
    """,
)
# ← This generates a fully-fused attention kernel with the mask applied INSIDE the tiling
#   Not: compute full attention → apply mask (wastes memory for masked positions)
#   But: compute only valid tiles → skip masked tiles entirely
#   Result: linear compute in window_size, not quadratic in sequence_length

# ─── SORTING-FREE SAMPLING ────────────────────────────────────────────────────
# Standard Top-K sampling: sort all logits → take top K → sample
# Cost: O(vocab_size × log(vocab_size)) at every token
# For GPT-4o-equivalent with 200k vocab: 200k × 17 = 3.4M comparisons per step
#
# FlashInfer's sorting-free approach:
# Use bucket-based approximate selection to find the top-K threshold
# without sorting the full array
# Cost: O(vocab_size) with small constant (bucket scan, not comparison sort)

logits = torch.randn(128, 200000, device="cuda", dtype=torch.float32)  # [batch, vocab]

# Standard top-k (baseline: requires sort)
# top_k_values, top_k_indices = torch.topk(logits, k=50, dim=-1)  # O(vocab × log(vocab))

# FlashInfer sorting-free top-k sampling
# ← Same output distribution as exact top-k, no sorting required
samples, success = flashinfer.sampling.top_k_top_p_sampling_from_probs(
    probs=torch.softmax(logits, dim=-1),
    uniform_samples=torch.rand(128, device="cuda"),  # pre-generated uniform random
    top_k=50,
    top_p=0.9,
    filter_apply_order="top_k_first",
)
# ← For 128 batch × 200k vocab: FlashInfer ~0.8ms vs sort-based ~2.3ms
# ← The speedup grows with batch size and vocabulary size
# ← Critical for speculative decoding where sampling runs at the draft model's rate

# CHAIN SPECULATIVE DECODING SUPPORT
# FlashInfer also provides chain_speculative_sampling for standard tree-spec-dec
speculative_samples, accepted_counts, emitted_token_num = \
    flashinfer.sampling.chain_speculative_sampling(
        draft_probs=torch.randn(128, 5, 200000).softmax(-1).to("cuda"),  # [batch, draft_len, vocab]
        draft_token_ids=torch.randint(0, 200000, (128, 5), device="cuda"),
        uniform_samples=torch.rand(128, 6, device="cuda"),   # one per position + target
        target_probs=torch.randn(128, 6, 200000).softmax(-1).to("cuda"), # [batch, draft_len+1, vocab]
    )
# ← Implements the speculative sampling acceptance criterion in a single fused kernel
# ← Avoids the loop-based implementation that launches N kernels for N draft positions

The sorting-free sampling improvement may seem minor but compounds significantly in practice. Sampling runs at every decode step for every request in the batch. At 1000 requests × 50 tokens × 128 vocabulary bucket, the baseline sort-based approach becomes a measurable fraction of total decode time. FlashInfer's bucket-based approach eliminates the sort entirely, applying the same output distribution in linear scan time.

It In Action: FlashInfer in a SGLang Production Decode Step

Setup: SGLang serving DeepSeek-V3 on H100, 64 concurrent decode requests, 3k-8k context lengths.

Step 1: Per-step plan (FlashInfer schedule)

# At the start of each decode step, FlashInfer builds a plan
# This is fast (~10-50 microseconds) because it only computes tile assignments

plan = wrapper.plan(
    indptr=kv_indptr,           # which pages belong to which request
    indices=kv_page_indices,    # page physical indices in the pool
    last_page_len=last_page_len, # how many tokens fill the last page of each request
    num_qo_heads=128,            # DeepSeek-V3: 128 query heads
    num_kv_heads=128,            # GQA ratio: depends on model config
    head_dim=128,
    page_size=16,
    data_type=torch.bfloat16,   # DeepSeek-V3 uses BF16
    # MLA-specific for DeepSeek:
    use_mla=True,               # compress KV via latent representation
    q_nope_dim=128,             # no-RoPE query dimension
    q_rope_dim=64,              # RoPE query dimension
    kv_lora_rank=512,           # latent dimension (compressed KV)
)
# ← Plan output: tile_to_request mapping for balanced GPU SM utilization
# ← CUDAGraph-compatible: the plan feeds static-shape kernel launch

Step 2: Attention kernel (CUDAGraph replay)

FlashInfer attention kernel (CUDAGraph replay, no kernel launch overhead):
  Input: 64 query vectors [64, 128, 128]
         Paged KV pool (64 request × avg 500 pages × 16 tokens × 128 dim)
         MLA latent KV [64, seq_len, 512] → projected to full KV inside kernel

  Tile assignment: 64 requests → N tiles of size T
    Tiles assigned to GPU SMs via load-balanced mapping
    No SM idles waiting for work from a long-sequence request while short ones finish

  Kernel execution: ~0.8ms (H100, 64 batch, ~5k avg context, BF16 with MLA)
  Without FlashInfer (naive paged attention): ~1.4ms

  Output: 64 decoded vectors [64, 128, 128]

Step 3: Sampling (sorting-free)

FlashInfer sorting-free Top-P (p=0.95) sampling:
  Input: 64 logits [64, 152,000]  (DeepSeek-V3 vocabulary)
  Method: bucket-based threshold scan, not sort
  Time: ~0.4ms vs ~1.1ms sort-based at this batch size

  Output: 64 sampled token IDs

Total decode step: ~2.5ms (attention + sampling + other ops)
vs baseline without FlashInfer: ~4.1ms
Reduction: ~39% inter-token-latency

At 1000 tokens of output:

Without FlashInfer: 1000 × 4.1ms = 4.1 seconds per request (64 concurrent)
With FlashInfer:    1000 × 2.5ms = 2.5 seconds per request
Throughput: 64 requests / 2.5s = 25.6 req/s vs 15.6 req/s without FlashInfer

Why This Design Works, and What It Trades Away

The block-sparse composable format is the correct abstraction because production KV-cache is structurally sparse in ways that monolithic contiguous formats cannot represent without either padding or format conversion. Cascade attention (a sub-request shares a prefix KV cache with other sub-requests) cannot be expressed as a standard paged attention call without copying the shared prefix into each request's page table. FlashInfer composes these formats mathematically rather than physically, meaning the shared prefix is read once but counted as if each request owns it.

The JIT compilation system is the piece that enables customization without sacrificing performance. The alternative in most kernel libraries is a library of fixed variants: standard attention, sliding-window attention, alibi-biased attention, each compiled separately. FlashInfer's JIT generates the right fused kernel for any combination of mask, bias, and format at runtime. The first call pays a compilation cost (typically 1-10 seconds); subsequent calls replay the cached kernel. For serving systems where the same configuration runs for hours, the amortized cost is negligible.

The CUDAGraph-compatible scheduler solves a real operational problem. CUDA kernel launches from CPU have non-trivial overhead, typically tens of microseconds per kernel call. At the decode rates modern serving systems achieve, thousands of kernel launches per second, this overhead is measurable. CUDAGraph batches all kernel launches into a pre-captured graph and replays it in a single CPU call. But standard dynamic batching breaks CUDAGraph because shapes change between batches. FlashInfer's tile-based scheduling keeps kernel shapes static while the assignment is dynamic.

What FlashInfer trades away:

Not all features are supported across all compute capabilities. The feature matrix in the documentation shows which operations are available on each SM version. Teams deploying on mixed-GPU fleets (some H100, some A100, some L40) need to check compatibility per operation per GPU.

First-use JIT compilation is the cold-start penalty for flashinfer-python. A serving system starting up with no cached kernels will have a 1-10 second delay on first request for each new configuration. The flashinfer-cubin and flashinfer-jit-cache packages eliminate this by pre-compiling, at the cost of larger disk footprint and the requirement to match the CUDA version exactly.

The load-balanced scheduler requires a plan call at the start of each batch. For very short decode sequences or tiny batch sizes where kernel launch overhead is not the bottleneck, the plan overhead may not be worth the CUDAGraph benefit. The documentation recommends CUDAGraph only when batch sizes are large enough that kernel launch overhead is measurable.

Technical Moats

MLA (Multi-Latent Attention) native kernel support. DeepSeek-V2/V3/R1's compressed KV format requires a different attention kernel: the latent KV must be projected back to full Q,K,V dimensions inside the attention computation, not before it. A naive implementation would decompress the latent back to full KV, then run standard attention (wasting memory bandwidth on the expansion). FlashInfer fuses the latent projection inside the tiled attention kernel, reading the compressed format and computing attention over projected values without materializing the full-size K and V. For the DeepSeek models that have achieved the widest deployment in 2025-2026, this matters: MLA saves 5.75x KV cache memory, and FlashInfer is one of the few libraries that handles it natively without the decompression overhead.

Full Blackwell support with CuTe DSL kernels. SM 10.0 (B200, B300) through SM 12.1 (RTX 50 series, DGX Spark) introduced with the Blackwell architecture. FlashInfer added Blackwell support in v0.4.0 (October 2025). The Blackwell-specific kernels use NVIDIA's CuTe DSL, which is the abstraction layer that lets kernels take advantage of Blackwell's new memory hierarchy and WGMMA instructions. Coverage from T4 (2018) to DGX Spark (2026) in a single library is operationally important for teams with heterogeneous GPU fleets.

The block-sparse format's composability property. The technical insight that blocks of KV-cache from different requests can be composed without copying is the mechanism that makes Cascade Attention work. When multiple requests share a system prompt (a common pattern in chatbot serving and RAG), the shared prefix KV can be computed once and referenced by all. FlashInfer's composable format lets the attention kernel treat this as a mathematical composition, reading shared blocks from one memory region and per-request blocks from another, in a single kernel call. The alternative is either per-request prefix duplication (wastes KV cache memory) or a separate prefix-attention + per-request-attention sequence (two kernel calls, misses fusion opportunities).

Insights

Insight One: FlashInfer's primary contribution to inter-token latency is not through better arithmetic. It is through better memory access patterns and eliminated kernel launch overhead. The 29-69% inter-token-latency reduction reported in the paper is compared to compiler backends that use XLA-style fused operators or Triton JIT compilation. These achieve good arithmetic efficiency but generate kernel launch patterns that are not CUDAGraph-compatible because their shapes change with dynamic batching. FlashInfer's tile-based scheduling makes the kernel geometry static while the workload assignment is dynamic. Teams evaluating FlashInfer by looking at FLOP utilization or arithmetic intensity alone will underestimate the real benefit, which shows up in end-to-end serving latency at realistic batch sizes and context lengths.

Insight Two: The sorting-free sampling implementation is technically straightforward but has compounding production impact. Top-K and Top-P sampling run once per decode step per request. At 1000-token outputs, a 200k-vocabulary model, and 128 batch size, the sort-based approach runs 128 million comparisons per output position. This is CPU-side (or GPU-sort-side) overhead that has nothing to do with model intelligence. FlashInfer's bucket-based approach reduces this to a single linear scan with a fixed number of buckets, making sampling cost O(vocab_size) instead of O(vocab_size × log(vocab_size)). As vocabulary sizes grow (200k for multilingual models, 256k for code-focused models), this reduction becomes more significant, not less.

Surprising Takeaway

FlashInfer's repository root contains both CLAUDE.md and AGENTS.md, which means the codebase is developed with Claude Code as the primary AI coding assistant. This places FlashInfer in the same category as Meetily, Handy, OmniRoute, OpenMontage, and Mooncake in the pattern of infrastructure tools that are simultaneously built with AI agents and optimized for AI agents to understand and contribute to. For a CUDA kernel library written in Python (57%), CUDA (26%), and C++ (16%), having a working Claude Code context means that kernel engineers joining the project can ask Claude to explain why a specific tiling strategy was chosen, what the kernel launch overhead tradeoffs are, or how to add a new attention variant. The CLAUDE.md navigation context is the equivalent of an onboarding document that runs in the agent's context window. It is not accidental that high-velocity kernel library development now depends on AI agents navigating CUDA code with the help of structured context files.

TL;DR For Engineers

FlashInfer (flashinfer-ai/flashinfer, Apache 2.0, 5.8k stars, MLSys 2025, arXiv:2501.01005): kernel library and kernel generator for LLM inference. Core contributions: (1) block-sparse composable KV-cache formats, (2) JIT-compiled customizable attention templates, (3) load-balanced CUDAGraph-compatible scheduler.
Benchmarks (MLSys 2025): 29-69% inter-token-latency reduction vs compiler backends, 28-30% latency reduction for long-context inference, 13-17% speedup for parallel generation.
Attention operations: Paged/Ragged KV, MLA (DeepSeek native), Cascade (shared prefixes), POD (fused prefill+decode), Block-sparse. GEMM: BF16/FP8/FP4 (Blackwell). MoE: fused with DeepSeek-V3/Llama-4 routing. Sampling: sorting-free Top-K/P/Min-P. Comm: AllReduce, MNNVL, NVSHMEM.
GPU coverage: SM 7.5 (T4) through SM 12.1 (RTX 50 series, DGX Spark). Blackwell support: v0.4.0 (Oct 2025). CUDA 12.6, 12.8, 13.0, 13.1.
Install: pip install flashinfer-python (JIT on first use) + optional flashinfer-cubin (pre-compiled) + flashinfer-jit-cache (CUDA-version-specific cache). Verify: flashinfer show-config.
Powers: SGLang, vLLM, TensorRT-LLM, TGI (Hugging Face), MLC-LLM, LightLLM, lorax, ScaleLLM.

The Kernel Layer Is Where the Performance Actually Lives

The LLM inference stack has a model layer, a serving framework layer, and a kernel layer. Optimization work has concentrated heavily on the framework layer (vLLM, SGLang) and the model layer (quantization, architecture choices). The kernel layer has received less attention outside of FlashAttention and a handful of MoE kernels.

FlashInfer's thesis is that the kernel layer is where the remaining performance lives, and that it is heterogeneous: different attention patterns, different KV-cache layouts, different hardware, and different workload distributions each require different kernels. A single pre-compiled kernel cannot be optimal across all of them. The right infrastructure is a composable format, a JIT compiler, and a scheduling algorithm that adapts to dynamic workloads while maintaining the hardware compatibility constraints (CUDAGraph) that production systems require.

The MLSys 2025 acceptance and the adoption by eight production serving frameworks are evidence that this thesis is correct.

References

FlashInfer GitHub Repository, flashinfer-ai, Apache 2.0
FlashInfer: Efficient and Customizable Attention Engine for LLM Inference Serving, arXiv:2501.01005, Ye, Chen, Lai, Lin, Zhang, Wang, Chen, Kasikci, Grover, Krishnamurthy, Ceze, MLSys 2025
FlashInfer Documentation, flashinfer.ai
Sorting-Free GPU Kernels for LLM Sampling Blog, FlashInfer Blog, Mar 2025
SGLang Integration, primary production user

Explain It Like I'm New

When a large language model generates text, it does not start from scratch with every new word. It maintains a growing memory of everything it has processed so far, called a key-value cache. Reading from this memory efficiently is one of the hardest engineering problems in AI serving.

Think of the cache as a filing cabinet. Standard approaches assume every file is the same size and organized the same way. But in production, different users have different context lengths, different parts of the conversation can be shared across users (like a common system prompt), and some models compress their cache in unusual ways to save space. A filing system designed for uniform files becomes slow and wasteful when files are irregular.

FlashInfer redesigns the filing system. It represents the cache in a format that handles irregular files natively, generates custom access routines tuned to whatever combination of patterns a specific workload needs, and schedules reads across the GPU in a way that keeps every processor busy even when the files are different sizes.

The technical name for the access routines is CUDA kernels, and the way FlashInfer generates them is called JIT compilation, which means it figures out the best approach at runtime rather than committing to one ahead of time. This flexibility is what lets the same library work well from a 2018 server GPU to a 2026 Blackwell chip.

For engineers building or operating AI serving systems, FlashInfer is the layer between the model and the hardware where latency improvements that were not possible with static, pre-compiled kernels become achievable.

See It In Action

FlashInfer: Efficient and Customizable Attention Engine for LLM Inference Serving (MLSys 2025 Talk) Source: MLSys Conference | https://mlsys.org/virtual/2025/session/3138 The official conference presentation by Zihao Ye and co-authors covering all three core contributions: block-sparse composable formats, JIT attention templates, and the load-balanced CUDAGraph-compatible scheduler. The benchmark comparisons against Triton and compiler backends are presented in full here, with the methodology behind the 29-69% inter-token-latency claims explained in the authors' own words. Required viewing before evaluating FlashInfer for production.
Sorting-Free GPU Kernels for LLM Sampling (FlashInfer Engineering Blog) Source: FlashInfer Blog | https://flashinfer.ai/2025/03/10/sampling.html Shanli Xing and Zihao Ye walk through the complete algorithm and implementation behind FlashInfer's bucket-based Top-K/P/Min-P sampling kernels, with profiling charts showing exactly how much of total inference time sampling accounts for as vocabulary sizes grow. Explains why this is not a marginal optimization: the figure shows sampling cost as a measurable fraction of total decode time on large-vocabulary multilingual models.
Run High-Performance LLM Inference Kernels from NVIDIA Using FlashInfer (NVIDIA Technical Blog) Source: NVIDIA Developer Blog | https://developer.nvidia.com/blog/run-high-performance-llm-inference-kernels-from-nvidia-using-flashinfer/ NVIDIA's June 2025 announcement explaining their decision to actively release TensorRT-LLM production kernels into FlashInfer, including the first Blackwell MLA kernels for DeepSeek. Written from an infrastructure integration perspective: explains what it means for vLLM, SGLang, and custom inference engines to adopt FlashInfer as the kernel layer without writing CUDA C++.
FlashInfer 0.2: Efficient and Customizable Kernels for LLM Inference Serving (Release Blog) Source: FlashInfer Blog | https://flashinfer.ai/2024/12/16/flashinfer-02-release.html The 0.2 release post documents the CUDAGraph compatibility work, the JIT compilation architecture, and how the block-sparse format was extended to handle MLA and cascade formats. Useful as a technical before/after comparison against the MLSys 2025 paper to understand what features were added versus the original system design.
FlashInfer Guest Lecture: Efficient and Customizable Kernel Generation for LLM Inference Serving (CMU 15-422/642) Source: Carnegie Mellon University ML Systems Course | https://csd.cmu.edu/calendar/15422642-machine-learning-and-systems-guest-lecture-zihao-ye Zihao Ye's guest lecture for CMU's Machine Learning and Systems course, aimed at graduate engineers. Provides a pedagogical walkthrough of the design decisions and tradeoffs that the MLSys paper assumes familiarity with, making it the most accessible entry point to FlashInfer's architecture for engineers who have not read the full paper.

Community Conversation

Zihao Ye (@ye_combinator) on X: Original FlashInfer announcement thread https://x.com/ye_combinator/status/1754519796639727903 The original 4-part announcement from the primary author, highlighting the 31x speedup over vLLM's PageAttention for shared-prefix batch decoding and the 3x GQA improvement. The thread also links the first FlashInfer blog post and GitHub repo. Valuable for understanding what the original design priorities were before the MLSys 2025 feature additions.
Tianqi Chen (@tqchenml) on X: MLSys 2025 Best Paper announcement with NVIDIA backing https://x.com/tqchenml/status/1922349774482657386 Tianqi Chen (co-creator of TVM, MXNet, and co-author of FlashInfer) announcing the MLSys 2025 Best Paper win and NVIDIA's decision to back FlashInfer as the community kernel layer. His framing: "FlashInfer won #MLSys2025 best paper, with backing from NVIDIA to bring top LLM inference kernels to the community." The industry signal here is significant: NVIDIA choosing FlashInfer over building a competing library or extending TensorRT-LLM directly.
NVIDIA AI Developer (@NVIDIAAIDev) on X: NVIDIA's public commitment to FlashInfer https://x.com/NVIDIAAIDev/status/1922354691251294644 NVIDIA's official statement explaining that FlashInfer was created across UW, CMU, and OctoAI (acquired by NVIDIA) with the goal of building an engine-agnostic, highly optimized, and easily extensible kernel library. Explicitly states that "inference platforms can adopt fresh ideas without waiting for new libraries or rewriting kernels in CUDA C++." Signals NVIDIA's infrastructure strategy: maintain one community kernel layer rather than fragmented per-framework implementations.
Lequn Chen (@abcdabcd987) on X: FlashInfer kernel coverage summary https://x.com/abcdabcd987/status/1754544392533352870 Co-author Lequn Chen's concise technical summary of FlashInfer's coverage: "{single, batch} x {prefill, decode, append} x {ragged tensor, paging} x {FP16, FP8, INT4} x {4090, Ada6000, A100, H100}." Demonstrates that the block-sparse composable format is not just a conceptual abstraction but covers a real cross-product of deployment configurations that no single pre-compiled kernel can serve. Useful for quickly communicating the scope to engineers evaluating whether FlashInfer applies to their hardware configuration.
MLSys 2026 FlashInfer AI Kernel Generation Contest https://mlsys26.flashinfer.ai/ NVIDIA and the FlashInfer team have made FlashInfer-Bench the basis for MLSys 2026's kernel generation competition, where teams compete to generate optimized FlashInfer kernels for DeepSeek-V3 and Gated Delta Net using AI-assisted kernel generation. The fact that FlashInfer's benchmark suite is the competitive evaluation standard for AI-generated CUDA kernels at the field's premier conference is the clearest possible signal of how central FlashInfer has become to the LLM serving infrastructure stack.

— # (#)

FlashInfer (flashinfer-ai/flashinfer, Apache 2.0, 5.8k stars, MLSys 2025, arXiv:2501.01005) is a kernel library and kernel generator for LLM inference serving with three core contributions: (1) block-sparse composable KV-cache formats that represent paged, ragged, shared-prefix (Cascade), and compressed (MLA) layouts without format conversion; (2) JIT-compiled customizable attention templates with automatic backend selection across FlashAttention-2/3, cuDNN, CUTLASS, and TensorRT-LLM; (3) a load-balanced scheduling algorithm that maintains CUDAGraph compatibility despite dynamic batching by keeping kernel shapes static while workload assignment is dynamic. MLSys 2025 benchmarks: 29-69% inter-token-latency reduction vs compiler backends, 28-30% latency reduction for long-context inference, 13-17% speedup for parallel generation. Full feature surface: Paged/Ragged/MLA/Cascade/POD/Sparse attention; BF16/FP8/FP4 GEMM; fused MoE with DeepSeek-V3/Llama-4/top-k routing; sorting-free Top-K/P/Min-P sampling and chain speculative decoding; AllReduce, MNNVL, NVSHMEM communication; RoPE, RMSNorm, SiLU/GELU. GPU support: SM 7.5 (Turing T4) through SM 12.1 (RTX 50 series, DGX Spark); Blackwell support added v0.4.0 (Oct 2025); CUDA 12.6-13.1. Integrated into: SGLang, vLLM, TensorRT-LLM, TGI, MLC-LLM, LightLLM, lorax, ScaleLLM.

Sponsored Ad

If you enjoy practical AI insights, check out SnackOnAI and support the newsletter by subscribing, sharing, and exploring our sponsored ad — it helps us keep building and delivering value 🚀

The World's Biggest Dev Event Hits Silicon Valley

WeAreDevelopers World Congress comes to San José, CA — September 23–25, 2026, in the heart of Silicon Valley. 10,000+ developers, 500+ speakers, 20+ stages, and the full software development lifecycle in one place.

Kelsey Hightower. Thomas Dohmke (fmr. CEO, GitHub). Christine Yen (CEO, Honeycomb). Mathias Biilmann (CEO, Netlify). Olivier Pomel (CEO, Datadog). The people actually building the tools you use every day — all on one stage.

Three days of AI, agents, cloud, security, and architecture, plus workshops, live coding, and the official Congress party. Prices rise as the event gets closer — lock in today's rate.

Use code GITPUSH26 for 10% off.

Get Your Pass