2026-08-02·48 min read

LLM Inference in 2026: A Field Guide

llm-inferencegpusystemsperformanceml-infrastructure

Almost everything interesting about serving large language models follows from a single fact: reading a prompt and writing a response are two completely different workloads that happen to share the same weights. Reading is compute-bound. Writing is memory-bandwidth-bound. Once you internalise that asymmetry, the rest of the field stops looking like a grab-bag of tricks and starts looking like a set of forced moves.

This is a map of where LLM inference actually stands in mid-2026 — what runs in production, what is still a paper, what got quietly disproved, and where the remaining headroom is. It is written for engineers who are comfortable with systems but have never had to care what an SM is. I have tried to give you the arithmetic rather than the conclusions, because the arithmetic is what lets you evaluate the next claim you read.

Where numbers appear, they come from primary sources — vendor specs, arXiv papers, or engineering blogs from the teams who run these systems. Where the field disagrees with itself, I have said so rather than picking the tidier story. A few well-known results turn out not to survive contact with production, and those are the most useful parts.

Core idea Stack Techniques Hardware Economics Frontier Learn Read all

1. The asymmetry

A transformer generating text does two distinguishable things.

Prefill processes the prompt. All the input tokens go through the network together, in parallel. If your prompt is 2,000 tokens, you load each weight matrix once and use it for 2,000 tokens' worth of arithmetic. The chip's floating-point units are the constraint — this phase is compute-bound. It determines how long the user waits before the first token appears, a metric called TTFT (time to first token).

Decode generates the response, one token at a time, because token N+1 depends on token N. Each step reads every weight in the model from memory to produce a single token per sequence. The chip's memory bandwidth is the constraint — this phase is memory-bandwidth-bound. It determines how fast text streams out, a metric called TPOT (time per output token).

The gap between these two is not small. On an NVIDIA H100, single-stream decode uses roughly 0.3% of the chip's peak floating-point throughput. You have bought one of the most powerful compute devices ever manufactured and you are using it as an expensive memory controller.

The generative idea

Once you see that decode is bandwidth-bound and prefill is compute-bound, most of the field's techniques become predictable. Every decode optimisation is a scheme to move fewer bytes or to amortise a byte across more tokens. Every prefill optimisation is a scheme to keep the tensor cores fed. And the largest structural wins come from noticing that the two phases want physically different machines.

Because decode is bandwidth-bound…	…the field produced
bytes are the currency, not FLOPs	quantization (FP8, FP4), mixture-of-experts, latent attention, KV compression
the floating-point units sit idle	speculative decoding — spend free FLOPs verifying guesses
many sequences can share one weight read	continuous batching, and the throughput-versus-latency tradeoff curve
prefill and decode want opposite machines	chunked prefill, prefill/decode disaggregation
the KV cache is the thing being streamed	a whole 2025–26 literature treating KV as distributed storage

And one corollary that governs all the economics: throughput and per-user latency are in direct tension. Adding a second concurrent request is nearly free in bandwidth terms, because both requests read the same weights. But once you have added enough requests to become compute-bound, every additional one slows everybody down. Your latency target picks your operating point; your operating point picks your cost per token. There is no single "fastest" configuration — there is a Pareto frontier, and any benchmark quoting one number is hiding where on that curve it sits.

2. Doing the arithmetic

This section is the one worth actually working through. If you can do this accounting, you can sanity-check any claim in the rest of the article — and any claim in any vendor benchmark — without running anything.

Two rules of thumb

For a dense transformer with P parameters, ignoring attention itself (which is small for short contexts and large for long ones):

FLOPs per token  ≈  2P        (one multiply + one add per parameter)
Bytes per token  ≈  P × bytes_per_parameter

The factor of 2 is because a multiply-accumulate is two floating-point operations. The bytes figure is what you must stream from memory for every single decode step.

Worked example: Llama-3.1-70B on an H100

Take a 70-billion-parameter model in BF16 (2 bytes per parameter) on an H100 SXM: 989 TFLOP/s of BF16 compute, 3.35 TB/s of HBM bandwidth.

Quantity	Value	Where it comes from
Weight bytes	140 GB	70e9 × 2
Decode floor, one GPU	—	140 GB doesn't fit in 80 GB. You need at least 2 GPUs.
Decode floor, 2× H100	20.9 ms/token	140 GB ÷ (2 × 3.35 TB/s)
→ ceiling on speed	~48 tok/s	1 ÷ 20.9 ms. Nothing you do in software beats this at batch 1.
FLOPs per decode token	140 GFLOP	2 × 70e9
Time if compute-bound	0.07 ms	140 GFLOP ÷ (2 × 989 TFLOP/s)
Ratio	~300×	You are memory-bound by a factor of three hundred.

That last row is the whole argument. There is 300× of headroom in compute sitting unused during decode, which is exactly why speculative decoding — burning spare FLOPs to guess ahead — was invented, and exactly why quantization (which cuts bytes) helps decode far more than it helps prefill.

Now prefill on the same model with a 2,048-token prompt:

FLOPs = 2 × 70e9 × 2048  ≈  287 TFLOP
Time  = 287 TFLOP ÷ (2 × 989 TFLOP/s × 0.5 efficiency)  ≈  290 ms

Same weights, same hardware, and now you are firmly compute-bound. One phase wants bandwidth; the other wants FLOPs. This is not a subtlety — it is a factor-of-hundreds difference in what the bottleneck is.

The KV cache, which is where the memory actually goes

Attention needs the keys and values of every previous token. Recomputing them each step would be quadratic, so they are cached. The cache size is:

KV bytes = 2 × layers × kv_heads × head_dim × seq_len × batch × dtype_bytes
           ↑
           one for K, one for V

For Llama-3.1-70B (80 layers, 8 KV heads after grouped-query attention, 128 head dim) in BF16:

Scenario	KV cache
1 sequence, 4K context	1.3 GB
1 sequence, 128K context	41.9 GB
32 sequences, 8K context	83.9 GB
128 sequences, 8K context	335 GB

Why this table matters

Look at the last two rows against the 140 GB of weights and the 160 GB you get from two H100s. At realistic batch sizes and context lengths, the KV cache is larger than the model. It is also dynamic, unpredictably sized, and different for every request. That is why memory management — not matrix multiplication — became the central engineering problem of LLM serving, and why so much of the last three years of research is really distributed-systems work wearing an ML hat.

Why grouped-query attention exists, in one calculation

The original transformer gives every attention head its own K and V. Llama-3.1-70B has 64 attention heads. If all 64 had their own KV, that 128K-context cache would be 8× larger — 335 GB for a single sequence.

Grouped-query attention (GQA) shares one K/V pair across a group of query heads — here, 8 query heads per KV head. You lose a little quality and cut KV memory 8×. Multi-query attention (MQA) is the extreme case: one KV head for all queries.

DeepSeek's multi-head latent attention (MLA) goes further, projecting KV into a shared low-rank latent of roughly 576 dimensions per token instead of storing per-head keys and values at all. It trades extra compute at decode time — which, per the arithmetic above, you have in enormous surplus — for dramatically fewer bytes. That is a textbook example of reading the roofline and designing against it.

Arithmetic intensity, and why batching works

Arithmetic intensity is the ratio that decides which bottleneck you hit: FLOPs performed per byte moved.

Multiply a batch of B activation vectors by a d × d weight matrix. You do 2Bd² FLOPs and move 2d² bytes of weights. Intensity ≈ B. The batch size is the arithmetic intensity.

Every chip has a corresponding machine balance — peak FLOP/s ÷ peak bytes/s. Below it you are memory-bound; above it you are compute-bound.

Accelerator	Peak BF16	Bandwidth	Machine balance
NVIDIA H100 SXM	989 TFLOP/s	3.35 TB/s	~295
NVIDIA B200	2,250 TFLOP/s	8 TB/s	~281
Google TPU v5e	197 TFLOP/s	819 GB/s	~240
NVIDIA DGX Spark (GB10)	~100 TFLOP/s	273 GB/s	~366
Mac Studio M3 Ultra	~26 TFLOP/s	819 GB/s	~32

So: batch 1 decode has intensity ~1 against a machine balance of ~295. You need roughly 300 concurrent sequences before an H100 stops being memory-bound during decode. That single number explains why every production serving system is built around aggressive batching, and why a chatbot serving one user is an appalling use of a datacentre GPU.

Note the last two rows. Apple's unified memory gives a machine balance of ~32 — it is comparatively bandwidth-rich and compute-poor, which is why Macs decode surprisingly well and prefill badly. The DGX Spark is the opposite extreme at ~366. Same workload, opposite failure modes.

3. The roofline

The roofline model puts all of this on one chart. Horizontal axis: arithmetic intensity. Vertical axis: achievable throughput. The diagonal is the bandwidth limit, the flat ceiling is the compute limit, and every workload sits somewhere underneath.

H100 roofline Decode, batch 1 Decode, batch 256 Prefill, 2048 tokens

Batching slides you rightward along the diagonal — one weight read now serves many sequences. Prefill starts out on the flat ceiling. The distance between the orange dot and the blue line is the entire economic argument for continuous batching.

Three things are worth extracting from this picture.

First, batching moves you along the roofline, it does not lift it. Going from batch 1 to batch 256 buys you two orders of magnitude of aggregate throughput at nearly constant per-token bandwidth cost. That is why serving economics are dominated by concurrency.

Second, the knee is where the tradeoff bites. Past the machine balance, adding requests no longer improves aggregate throughput much but does degrade each user's latency. Good serving systems operate near the knee and know exactly where it is for their model and hardware.

Third, a chip's position on this chart is a design choice with consequences. Compare the H100's balance of 295 with the Mac's 32. On the Mac, batch 1 decode is only about 30× below the knee rather than 300×, which is why a laptop can produce a decent tokens-per-second on a large model while being hopeless at processing a long prompt.

4. Vocabulary

You need these to read anything else in the field.

Term	What it is	Bounded by
TTFT	Time to first token	Prefill and queueing
TPOT / ITL	Time per output token / inter-token latency	Decode
Interactivity	tokens/sec/user — the axis a human actually feels	Decode
Throughput	Total tokens/sec, usually normalised per GPU	Both
Goodput	Requests/sec completed while meeting latency SLOs	The only metric that really matters
MFU	Model FLOPs utilisation — fraction of peak compute achieved	Prefill health
MBU	Model bandwidth utilisation — fraction of peak bandwidth achieved	Decode health
Cost per Mtok	($/hour ÷ tokens/hour) × 10⁶	The bottom line
J/token	Joules per token	The emerging metric — power is now the binding constraint

The important one is goodput. A system can post spectacular throughput numbers while violating every latency commitment it made, simply by batching enormously and making everyone wait. Goodput — introduced properly by the DistServe paper — counts only the requests that finished within their SLO, and it is the metric serious teams optimise.

5. The stack

The field has stratified into five layers. The most consequential change since 2023 is that the inference engine is no longer the whole system — an entire orchestration tier appeared above it.

Layer 5

Gateway / router

LiteLLM · Envoy AI Gateway · Gateway API Inference Extension · OpenRouter — auth, quotas, multi-tenancy, cross-provider fanout, routing requests to different models by difficulty.

Layer 4 — did not exist in 2023

Orchestration

NVIDIA Dynamo · llm-d (CNCF) · KServe · Ray Serve — prefill/decode disaggregation, KV-aware routing, tiered KV offload, SLO-driven autoscaling, topology awareness.

Layer 3

Engine

vLLM · SGLang · TensorRT-LLM · llama.cpp · MLX · LMDeploy — scheduling, paged KV, continuous batching, prefix caching, speculative decoding, parallelism.

Layer 2

Kernels

FlashInfer · CUTLASS / CuTe DSL · Triton · Helion · ThunderKittens — attention, GEMM, grouped-GEMM for MoE, quantized math.

Layer 1

Hardware

NVIDIA Blackwell/Rubin · Google TPU · AMD Instinct · AWS Trainium · Apple Silicon · Groq/Cerebras/Etched — bandwidth, tensor cores, interconnect, numeric formats.

The part people underestimate

Layers 3 through 5 are a distributed systems problem, not a machine-learning one. Scheduling, admission control, cache coherence, load balancing, failure recovery, tail latency. As we will see in the section on open problems, that is also where most of the remaining performance is — and it is far more accessible than kernel engineering to anyone with a backend or infrastructure background.

12. Engines and libraries

vLLM is the default. SGLang is the specialist that eats frontier-scale MoE. TensorRT-LLM is NVIDIA's performance ceiling for NVIDIA-only shops. Above all three, Dynamo and llm-d are converging on the same orchestration design and treat the engine as swappable.

	vLLM	SGLang	TensorRT-LLM
Peak throughput, tuned	Very good	Best on large MoE and high prefix reuse	Best on NVIDIA-tuned dense models, especially Blackwell FP4
Prefix-heavy / agentic	Good (hash prefix cache)	Best (radix tree)	Adequate
Model coverage	Widest — 200+ architectures	Broad, frontier-first	Narrowest
Hardware breadth	Widest — CUDA, ROCm, XPU, TPU, CPU, Neuron, Ascend	Broad	NVIDIA only
Operational complexity	Lowest	Medium	Highest

The most useful thing in this section

Real differences between these engines at equal tuning effort are usually 10–30%, and which direction it goes is workload-dependent. Anyone quoting a 2–3× gap is comparing an untuned configuration against a tuned one.

Before you consider switching engines, turn on prefix caching, chunked prefill, CUDA graphs, and FP8. Configuration beats engine choice most of the time.

Things worth knowing about each

vLLM dominates on model coverage and ecosystem gravity rather than peak throughput. Day-zero support for new architectures, a transformers fallback backend that runs essentially any Hugging Face architecture at near-native speed, PyTorch Foundation governance, 2,000+ contributors, and a release roughly every two weeks. Every major cloud's managed offering is built on it.

SGLang claims trillions of tokens per day across 400,000+ GPUs, with xAI, Cursor, Oracle and LinkedIn named publicly. It is the engine inside DeepSeek-scale deployments and the preferred backend for reinforcement-learning rollouts.

TensorRT-LLM abandoned TensorRT. Release 1.2 made PyTorch the sole execution backend and removed the ahead-of-time compiled-engine path entirely. That is a significant philosophical concession: the compiled-engine model lost to eager PyTorch plus good kernels. TRT-LLM is now best understood as NVIDIA's opinionated PyTorch runtime with privileged access to NVIDIA kernels.

Hugging Face TGI is in maintenance mode, with HF explicitly redirecting users to vLLM, SGLang, llama.cpp and MLX. This was an elegant retreat rather than a defeat: having won the argument that engines should build on transformers model definitions, HF stopped competing on serving. transformers became the reference definition; the engines became the runtime.

Dynamo versus llm-d — the same four problems (KV-aware routing, disaggregation, tiered KV offload, SLO autoscaling), different politics. Dynamo is NVIDIA-aligned with rack-scale topology awareness. llm-d is the vendor-neutral CNCF Sandbox counterpart built on upstream Kubernetes primitives, targeting any model on any accelerator — including TPU and Intel XPU disaggregation, which Dynamo will never do. Choose the orchestrator on your hardware and your politics, not your engine.

The kernel-authoring ladder

CUDA C++ / PTX        ← the floor; you rarely need this
      ▲
CUTLASS 4.x + CuTe    ← Python-native, layouts and atoms exposed.
      ▲                 How Blackwell GEMM and MoE kernels are written now.
Gluon (inside Triton) ← explicit layouts, warp specialisation
      ▲
Triton                ← the workhorse. ~85% of inference kernels.
      ▲
Helion                ← Python DSL that COMPILES TO TRITON and autotunes.
                        ~30 lines for attention vs ~120 in raw Triton.

And above all of it, FlashInfer has become the consolidation point — a shared attention, GEMM and MoE kernel layer sitting under vLLM, SGLang, TensorRT-LLM and MLC-LLM, with a JIT generator that specialises for paged KV layouts, latent attention, sparse patterns and FP8/FP4. NVIDIA now ships kernels through it. In effect, the engines stopped writing their own attention kernels. Before you write anything, check whether FlashInfer already has it.

What to actually use

Situation	Use
Single Apple Silicon Mac	MLX for scripting; LM Studio for a GUI plus a local OpenAI-compatible server; Ollama if you want one command and can lose the last 30%
Single consumer GPU, 24–32 GB	llama.cpp for flexibility; ExLlamaV3 for maximum quality per VRAM byte; vLLM above about 4 concurrent users
One 8-GPU node	vLLM by default; SGLang for heavy shared prefixes, frontier MoE, or RL rollouts; TRT-LLM only after measuring
Multi-node cluster	You need an orchestrator — llm-d or Dynamo — plus a gateway at the edge
Offline / batch scoring	Ray Data LLM over vLLM or SGLang

Maturity All Production Contested Research

Six families of technique, sorted by what actually runs in production. Filter by maturity above.

6. Batching and phase management

All of this is production-standard. If a serving system does not do these three things, it is not a serious serving system.

Continuous batching production

The naive approach batches requests together, pads them to the length of the longest, and waits for all of them to finish. If one request generates 2,000 tokens and the rest generate 50, the whole batch is held hostage.

Continuous batching — from the Orca paper (OSDI 2022) — schedules at iteration granularity instead. After every single forward pass, finished sequences are evicted and queued ones admitted. vLLM implements this by flattening every sequence into one long "super sequence" with position indices and attention masks, so there is no padding at all.

The gain over static batching is routinely 2–4× in real traffic, and it is entirely free — no accuracy cost, no tuning. Universal.

Chunked prefill production

Here is a problem you get for free once you have continuous batching. A user submits a 100,000-token prompt. Prefilling it takes seconds of solid compute. During those seconds, every other user's token stream stops, because the GPU is busy. Classic head-of-line blocking, and it shows up as ugly spikes in inter-token latency.

The fix, from Sarathi-Serve (OSDI 2024): split the prefill into chunks sized by a per-step token budget, and co-schedule those chunks alongside ongoing decodes in the same forward pass. Each step does a bit of prefill work and a bit of decode work. Nobody stalls.

In vLLM's V1 engine this is the scheduler. The scheduler's output is literally a map of {request_id: number_of_tokens_to_process}, and chunked prefill, prefix caching, and speculative decoding all fall out as special cases rather than being bolted on as modes. It is a genuinely elegant piece of design worth reading.

Prefill/decode disaggregation production

The logical conclusion of the asymmetry. If prefill wants FLOPs and decode wants bandwidth, stop running them on the same machine. Run prefill on one pool of workers, decode on another, and ship the KV cache between them over the interconnect.

This eliminates interference entirely and lets you tune parallelism strategy — and even hardware type — per phase. Introduced by DistServe and Splitwise in 2024, it is now the default topology for frontier-scale mixture-of-experts serving. Meta has reported 15–25% total-cost-of-ownership improvement purely from running the two phases on different accelerator types.

A concrete production configuration, and the unglamorous thing that actually helped

A mid-2026 deployment of GLM-5.2 in NVFP4 across 24 B300 GPUs: four prefill nodes running 4-way data parallelism plus expert parallelism, one decode node running 8-way data parallelism plus expert parallelism. Result: mean TTFT at or below 2.5 seconds, and TPOT down from roughly 40 ms to 17 ms.

The single largest contributor to that TPOT improvement was not an algorithm. It was speculative padding on the decode side — pre-padding batch shapes so newly arriving requests could join without triggering a shape change and a kernel recompile. Production wins are frequently plumbing.

There is also healthy skepticism now. MLSys 2026 carried a paper titled "Beyond the Buzz: A Pragmatic Take on Inference Disaggregation." Below some scale threshold, the cost of transferring KV between pools exceeds the benefit of specialising them. Disaggregation is not free and not always right.

7. The KV cache became a distributed storage system

This is the most active systems area in the field, and it has the distinct flavour of we accidentally built a distributed storage system and are now discovering it has a literature.

PagedAttention production

The KV cache grows unpredictably — you do not know how long a response will be when you start it. Allocating a contiguous buffer for the worst case wastes enormous memory; allocating incrementally fragments it.

vLLM's answer, borrowed wholesale from operating-system virtual memory, was to allocate KV in fixed-size blocks (typically 16 tokens) from a free pool, with a per-sequence block table mapping logical positions to physical blocks. Fragmentation collapses, memory utilisation goes from roughly 20–40% to over 90%, and sequences that fork (beam search, parallel sampling) can share blocks copy-on-write.

This one idea is why vLLM took over. Every serious engine now does it.

Prefix and radix caching production

If two requests share a prefix — the same system prompt, the same few-shot examples, the same conversation history — their KV entries for that prefix are identical. Hash the blocks and reuse them.

SGLang's RadixAttention generalises this into a radix tree over token blocks with LRU eviction, so arbitrary shared prefixes hit rather than just exact matches from the start.

The number that reorganised the field

Trace analysis of a coding agent working through SWE-bench Pro: by turn 30, the context is around 80,000 tokens, but each turn adds only hundreds to a few thousand new tokens. Which means over 95% of the prefill work is recomputation of something the system already computed — if you can hit the cache.

But the prefix is gigabytes, it gets evicted under memory pressure, and the next turn may be routed to a different replica that never had it. That single observation is why the 2026 stack converged on cluster-wide RDMA KV pools, cache-aware request routing, and explicit cache pinning. For agentic workloads, cross-request KV reuse is now a bigger lever than anything in the kernel layer.

Tiered KV storage production

Each tier is roughly an order of magnitude slower and an order of magnitude larger. A cache miss at the bottom costs a full re-prefill — which, unusually for a cache, is a cost you can compute exactly in advance.

KV cache in FP8 production

Storing the KV cache in FP8 rather than BF16 halves the bytes streamed during decode, which is precisely the bottleneck. Measured impact: 1–2 points of degradation on reasoning benchmarks, 94–98% quality recovery at contexts up to a million tokens, and inter-token latency scaling improved to 54% of the BF16 slope. This is now a sensible default for long-context deployments.

Sub-8-bit KV cache contested

Below 8 bits it gets ugly, and this is a case where the papers and the production measurements disagree. vLLM's own systematic study found 3-bit KV variants collapsing to roughly 31–33% long-context accuracy and costing 20–34% throughput, because dequantising back to BF16 on every attention call is not free.

The 2026 rule: FP8 KV yes, 4-bit KV only under genuine memory desperation.

KV eviction heuristics research

There is a large literature — H2O, SnapKV, and dozens of descendants — on dropping "unimportant" tokens from the KV cache based on attention scores. Almost none of it is on by default in production, because it breaks exact prefix-cache semantics and has long-tail accuracy failures.

In 2026 it also came under theoretical attack. A result published this year proves that accurate KV compression is impossible in the worst case when attention sharply retrieves individual tokens, and establishes a sharp separation between query-aware and query-agnostic compression. Five years of eviction heuristics are being retroactively bounded, and the bound suggests many of them were exploiting benchmark structure.

Why the standard benchmark misled everyone

"Needle in a haystack" — hide a fact in a long document, ask for it — is the canonical long-context test. But it is a query-known setting: the thing you need to retrieve is exactly what the question asks about. That is precisely the case KV eviction heuristics are tuned to preserve. Passing it tells you almost nothing about whether the heuristic destroys information that matters for a question you did not anticipate.

8. Attention

The FlashAttention lineage production

Standard attention materialises an N × N score matrix in memory. At 128K context that is 16 billion entries per head. FlashAttention's insight was that you never need the full matrix — you can tile the computation, keep tiles in on-chip SRAM, and use an online-softmax formulation to combine tiles correctly without ever writing the matrix to HBM. Memory goes from quadratic to linear and it is faster, because HBM traffic was the bottleneck all along.

FlashAttention-2 fixed work partitioning across warps. FlashAttention-3 exploited Hopper's asynchronous copy engines and FP8. FlashAttention-4 (March 2026) targets a new problem: Blackwell scaled its tensor cores far faster than its exponential units and shared-memory bandwidth, so the softmax became the bottleneck rather than the matmul. The fix involves approximating the exponential with a polynomial evaluated on the FMA units, a ping-pong tile schedule with a dedicated correction warpgroup, and two-CTA matrix instructions. It reaches up to 1,605 TFLOP/s BF16 on a B200 — about 71% utilisation, 1.3× over cuDNN and 2.1–2.7× over Triton.

Notably it is written entirely in CuTe-DSL, a Python kernel DSL, with compile times 20–30× faster than the C++ template equivalent. The tooling for writing world-class GPU kernels got dramatically more approachable in the last year.

Sparse attention actually shipped production

This is genuinely new. For years, sparse attention was a paper-only technique — the theoretical FLOP savings never materialised as wall-clock savings because the memory access patterns were hostile to GPUs. In 2026 that changed, and seven competing designs converged on the same answer: block-level sparsity with a small learned index.

A lightweight indexer scores blocks of tokens; the top-k are attended to; the rest are skipped. Because selection is at block granularity, the memory accesses stay coalesced.

System	Approach	Result
DeepSeek-V3.2 (DSA)	"Lightning indexer" selects top-k	~1.6% of full-attention operations at 128K context
MiniMax M3 (MSA)	Block-sparse with learned routing	~1/20 the per-token compute at 1M context
DeepSeek-V4	128-token sliding window + 4:1 and 128:1 compressed-KV branches	~8.7× less KV than V3.2 at 1M; B200 decode falls only from 199 to 180 tok/s going from 4K to 900K context

The serving work that made this possible is the interesting part

Shipping sparse attention required rebuilding the prefix cache. Sliding-window and compressed-KV branches free tokens at different rates, so a single flat cache no longer works. SGLang built ShadowRadix: a prefix cache that indexes virtual token slots and projects shadows into heterogeneous physical pools, so compressed KV stays shareable even after sliding-window tokens are freed.

They also needed the indexer itself to be nearly free. A Lightning TopK radix-select kernel brought indexer latency from over 100 µs to about 15 µs. At 15 µs it is noise; at 100 µs it eats the savings.

One caveat to carry: there is no unified sparse-attention benchmark, so published speedups across these papers are not directly comparable.

Hybrid linear and state-space attention production

Linear-attention and state-space models (Mamba and relatives) replace the quadratic attention with a recurrent state update — constant memory per token instead of a growing cache. Three separate labs converged on Mamba-Transformer-MoE hybrids in 2026 (Nemotron 3, Qwen3-Next, Kimi K3).

The honest framing: every shipped system is a hybrid, because pure linear models still lose badly on hard retrieval tasks. A few full-attention layers interleaved among many linear ones gets most of the efficiency and keeps the recall.

Serving them is genuinely different work. SSM states update in place, are large, and are all-or-nothing reusable — you cannot slice a recurrent state the way you can slice a KV cache. That means a request-level state pool separate from the token-level KV pool, an elastic allocator to rebalance between the two, and explicit state checkpointing if you want prefix caching to work at all.

Attention sinks production

An odd empirical discovery that turned into infrastructure: the first few tokens of a sequence absorb a disproportionate share of attention mass, apparently acting as a place for the softmax to dump probability it does not want to allocate. Evict them and quality collapses. Keep them plus a sliding window and you can stream indefinitely.

Originally discovered at inference time (StreamingLLM), sink tokens are now trained into models deliberately — gpt-oss ships explicit learned sinks.

9. Speculative decoding, and the 2026 reckoning

The mechanism is elegant. A cheap "draft" model proposes k tokens. The expensive target model verifies all k in a single batched forward pass — which costs barely more than generating one token, since decode is bandwidth-bound and the weights get read once regardless. Rejection sampling then guarantees the output distribution is mathematically identical to what the target model would have produced alone.

You are converting idle FLOPs into tokens. Free speed, provably no quality loss.

The lineage: standalone draft models → Medusa (extra prediction heads plus tree attention) → EAGLE (drafting in the target's feature space rather than token space) → EAGLE-2 (dynamic draft trees) → EAGLE-3 (multi-layer feature fusion) → 2026's parallel drafting, which emits all k drafts in one pass using learnable mask tokens.

And then the reckoning. The most important speculative-decoding paper of 2026 is a critique — "Speculative Decoding: Performance or Illusion?" — which benchmarked systematically at realistic batch sizes rather than batch 1.

Setting	Speedup
EAGLE, batch 1 (Llama-3-70B)	~1.96×
EAGLE, batch 128 (Llama-3.1-8B, GSM8K)	~1.21×
Tree verification, 21 draft tokens, batch 64	below 1.0× — a net slowdown
Share of wall-clock spent in verification	42–95%
n-gram drafting	fails everywhere except code editing

Why the discrepancy

The reason is the roofline again. Speculation works by spending idle FLOPs. At batch 1 you have a 300× surplus. At batch 128 you have already used the surplus on batching — you are near the knee, and the extra verification compute now costs real time. Speculative decoding and batching compete for the same resource.

Production field reports corroborate: real acceptance rates of 0.6–0.8 rather than the 0.95 papers often assume, gains 40–60% below published numbers, benefit largely confined to concurrency below about 4–8, and observed regressions around concurrency 16. There are also specific failure modes where speculation interacts badly with JSON-schema-constrained decoding and with MoE routing.

Against all that, model-native multi-token prediction — speculation heads trained into the model rather than bolted on afterwards, as in DeepSeek and GLM — is decisively economic at scale. DeepSeek-R1 in FP4 goes from $0.251 to $0.057 per million tokens with MTP enabled.

The reconciliation: native MTP with high acceptance is a clear win. A bolt-on drafter at high concurrency is a coin flip you must measure on your own traffic. This is one of the few places where reading the primary literature will actively mislead you about production behaviour.

10. Quantization

The 2026 default recipe, stated plainly: weights and activations in FP8 (or FP4 on Blackwell for models above roughly 30B), KV cache in FP8, and BF16 retained for layer norms, the MoE router, embeddings, and the output head. Those last components are small and disproportionately sensitive.

Scheme	Status	Notes
Weight-only INT4 (GPTQ, AWQ)	Production for memory-constrained, low-batch	Helps decode, not prefill — prefill math stays BF16
W8A8 INT8 (SmoothQuant)	Production	Migrates activation outliers into the weights, where they are easier to quantize
FP8 e4m3	The common production default	Near-lossless, hardware-native since Hopper
MXFP4	Production for models trained in it	32-element blocks, power-of-two (E8M0) scale
NVFP4	Production on Blackwell	16-element blocks, FP8 (E4M3) per-group scale plus an FP32 per-tensor scale
Sub-4-bit / BitNet 1.58	Research	Impressive on CPU; no frontier-quality model, and the ternary silicon never arrived

The NVFP4-versus-MXFP4 distinction is the one worth understanding, because it explains why "4-bit" is not one thing. Both store 4-bit values with a shared scale per block. MXFP4 uses 32-element blocks and a power-of-two scale. NVFP4 halves the block to 16 elements and uses a fractional FP8 scale, giving twice as many opportunities to match the local dynamic range, plus a second global scale on top. Result: about 3.5× memory reduction versus FP16, 1.8× versus FP8, with under 1% degradation.

Quality recovery versus BF16 runs about 99% at 70B–235B, 97–99% around 30B, and 95–98% at 7–14B. FP4 works better the larger the model — which is exactly why frontier deployments use it and small-model deployments do not.

The direction of travel

The interesting shift is toward training in the deployment format. NVIDIA pretrained a 12B hybrid Mamba-Transformer on 10 trillion tokens entirely in NVFP4, matching FP8 quality. DeepSeek-V4 ships FP4 expert weights; gpt-oss is natively MXFP4. When the model is trained in 4 bits, the entire post-training-quantization accuracy debate evaporates. Four-bit crossed over from a compression trick to a training numeric format.

11. Mixture-of-experts and parallelism

A mixture-of-experts model replaces each dense feed-forward block with many "expert" blocks and a router that activates only a few per token. A model with 120B total parameters might activate only 5B per token.

Reread the decode arithmetic and the appeal is obvious: decode bandwidth is set by active parameters, not total parameters. An MoE with 5B active decodes roughly like a 5B model while having the knowledge capacity of a 120B one. MoE is not a compute trick; it is a memory-hierarchy trick, and it is the single most important architectural response to the bandwidth wall.

Expert parallelism production

Rather than sharding every expert's weights across every GPU, shard the experts themselves — GPU 0 holds experts 0–7, GPU 1 holds 8–15, and so on. Each GPU loads far fewer bytes, and the grouped matrix multiplications get higher arithmetic intensity.

Wide expert parallelism is the frontier default: EP32 delivers roughly 1.8× higher per-GPU throughput than EP8 at a matched 100 tokens/sec/user on GB200 NVL72 systems, enabled by that rack's 130 TB/s of coherent NVLink bandwidth.

It needs two things that are genuinely hard. First, high-performance all-to-all communication kernels, because every token must be routed to whichever GPU holds its experts and the results gathered back. Second, an expert-parallel load balancer, because expert selection is data-dependent and popular experts otherwise pile onto one GPU while others idle. The all-to-all collective during decode is latency-bound and load-imbalanced, and it is the wall that essentially every 2026 MoE serving paper is pushing against.

Choosing a parallelism strategy

Mode	When to use it	Communication cost
Tensor (TP)	Default within a node; latency-optimal	All-reduce after every layer — needs NVLink, poor across nodes
Pipeline (PP)	Multi-node scale-out, long prompts	Point-to-point at stage boundaries only — cheapest, but bubbles
Expert (EP)	MoE models — the frontier default	All-to-all; needs a load balancer
Data (DP)	Attention layers in MoE serving, paired with EP for the FFN	Replicated KV cache
Context (CP)	Ultra-long prefill (ring attention)	All-gather per layer; complex

The rule: tensor parallelism inside the coherent high-bandwidth domain, everything else across it. Tensor parallelism requires an all-reduce after every layer — two per transformer block. Over NVLink at 900 GB/s to 1.8 TB/s that is invisible. Over PCIe at 64 GB/s or Ethernet it dominates end-to-end time.

The 2026 result that challenges "tensor-parallel everything"

SGLang's chunked pipeline parallelism splits long prompts into 4–12K micro-batches flowing through pipeline stages with asynchronous point-to-point transfer and dynamically sized chunks. On DeepSeek-V3.1 and Qwen3-235B in FP8, a PP4 × TP8 configuration gave 3.31× the prefill throughput of TP8 and beat pure TP32 by 30.5%, cutting TTFT by 81% — from 55.5 seconds to 10.5 seconds.

The argument is straightforward once stated: pipeline parallelism's communication volume is orders of magnitude below tensor parallelism's, which puts it in the better position for scaling across nodes. The orthodoxy that TP is always the answer came from single-node intuition.

13. Hardware, from first principles

Every accelerator has the same shape: a small amount of extremely fast memory glued to compute, backed by progressively larger and slower tiers. Here is an H100 with real numbers.

Tier	Capacity	Bandwidth	Latency
Register file	~33 MB	100s of TB/s	~1 cycle
Shared memory / L1	~30 MB	10s of TB/s	~30 cycles
L2 cache	50 MB	5–10 TB/s	~200 cycles
HBM3	80 GB	3.35 TB/s	375–500 ns
Host DRAM (PCIe 5)	~2 TB	64 GB/s	~2 µs
Peer GPU (NVLink 4)	80 GB × N	900 GB/s	~2 µs
Network (InfiniBand)	everything else	50 GB/s per NIC	5–10 µs

Two things to notice. First, HBM buys bandwidth, not latency — 375–500 ns is worse than commodity DRAM's 10–15 ns. You are paying for thousands of parallel channels, not for speed. Second, the discontinuity: about 113 MB of on-chip SRAM against 80 GB of HBM. A 700× capacity cliff at a 50× bandwidth cliff. That single fact generates every architectural bet in the silicon section below.

If it helps: L2 is your page cache, HBM is NVMe, NVLink is a rack-local RDMA fabric, and InfiniBand is your WAN. Everyone in this industry is trying to move the boundary where the working set stops fitting in cache.

The divergence that defines the era

≈ 750×Peak compute growth, V100 (2017) → B300 (2025): 20 → 15,000 TFLOP/s

≈ 9×HBM bandwidth growth over the same period: 0.9 → 8 TB/s

Compute grew roughly 750×. Bandwidth grew roughly 9×. Since decode is bandwidth-bound, decode-bound inference gets relatively worse with every hardware generation. This is the memory wall, quantified, and it is why every technique in this article is fundamentally a bytes-reduction technique.

Memory bandwidth across the field

Datacenter Desktop / workstation Apple unified memory

For inference this is the chart that matters. Headline FLOPS mostly tell you about prefill; bandwidth sets the floor on how fast a token can possibly come out.

Execution models

Three different bets on how to organise a chip, and the differences show up directly in what kind of workload each is good at.

SIMT (NVIDIA, AMD) — 32-thread warps, a hardware scheduler, latency hidden by massively oversubscribing threads. Analogous to a thread-pool server: tolerant of unpredictable work, but spending transistors and watts on schedulers, caches and branch predication.

Dataflow (Groq, Cerebras, SambaNova, Tenstorrent) — the compiler statically schedules every instruction and every wire, cycle by cycle. No caches, no dynamic arbitration, no queueing variance. The entire tail-latency distribution collapses to a point. The cost: anything dynamic — variable sequence lengths, MoE routing, speculative decoding — becomes a compiler problem rather than a runtime one.

VLIW (Google TPU, Intel Gaudi) — one very wide instruction per cycle drives the matrix unit, vector unit and memory engines in lockstep, scheduled ahead of time by XLA. High utilisation on static shapes; requires shape bucketing and padding for dynamic ones.

If you want an analogy: SIMT is a work-stealing scheduler, dataflow is a hand-written static pipeline, VLIW is a compiled query plan.

What a tensor core or systolic array actually does

A scalar floating-point unit performs one operation per instruction issue, and issuing the instruction costs roughly 100× more energy than the arithmetic itself. Matrix multiplication is O(N³) compute over O(N²) data — enormous reuse available — but a SIMD machine cannot exploit it, because every operand round-trips through the register file.

A systolic array is a 2D grid of multiply-accumulate units where operands flow between neighbours and never re-enter the register file. Google's TPU v7 uses a 256×256 array: 65,536 MACs per cycle, each input fetched once and reused 256 times. One instruction, 131,072 FLOPs. NVIDIA's tensor core is the same idea at finer granularity — a warp-wide matrix instruction fed from shared memory — which keeps it programmable at the cost of more control overhead.

Bigger tiles amortise control better but waste silicon when your matrices do not tile evenly. That is the entire tradeoff, and it is why TPUs win on large regular workloads and GPUs win on messy ones.

Interconnect decides your parallelism strategy

Fabric	Per-device bandwidth	Coherent domain
PCIe Gen5 ×16	64 GB/s	1 host
NVLink 4 (Hopper)	900 GB/s	8 GPUs
NVLink 5 (Blackwell)	1.8 TB/s	72 GPUs (NVL72)
AMD Infinity Fabric (MI355X)	153 GB/s	8 GPUs
Google TPU ICI v7	1.2 TB/s	9,216 (3D torus)
InfiniBand / RoCEv2	50–100 GB/s	100k+

A GB200 NVL72 rack presents 72 GPUs and 13.4 TB of HBM as one domain at 130 TB/s of aggregate NVLink — enough that a trillion-parameter MoE can do expert-parallel all-to-all without ever touching the network. AMD's MI355X is competitive silicon (288 GB at 8 TB/s, 10.1 PFLOP/s MXFP4) and sits only about 33% behind B200 per-GPU, yet trails GB200 NVL72 by 4.7–5.3× at rack scale, because it has no comparable scale-up domain to disaggregate within.

Scale-up bandwidth, not FLOPS, is NVIDIA's real moat.

A note on AMD, and why "the kernels are fine" is not the same as "it works"

On Kimi K2.5, MI355X was described as "barely usable" with a 150–200 ms latency floor — until a single vLLM pull request fixed kernel dispatch for latent attention on CDNA4. That one change produced 7.7× peak throughput, and up to 15× at matched throughput, in 25 days, landing at 13 ms TPOT.

The lesson generalises. ROCm's problem in 2026 is not kernel quality — the kernels exist and are competitive when they are actually dispatched. The problem is composability: combining FP4 with disaggregation with wide expert parallelism still degrades badly. That composability gap, not peak FLOPS, is what a challenger has to close.

The interesting edge case: unified memory

Apple Silicon and NVIDIA's GB10 both use a single memory pool shared by CPU and GPU. That is excellent for capacity — a Mac Studio can hold 512 GB, more than three B200s — and poor for bandwidth, since LPDDR tops out below 1 TB/s against HBM3e's 8 TB/s.

The two devices land in opposite places despite the shared architecture. Apple's M3 Ultra has 819 GB/s against roughly 26 TFLOP/s — bandwidth-rich, compute-poor, so it decodes well and prefills badly. NVIDIA's DGX Spark has 273 GB/s against roughly 100 TFLOP/s — the reverse, giving it the worst compute-to-bandwidth ratio of any device in this article at ~366 FLOPs/byte.

A result worth knowing about

EXO Labs measured a DGX Spark as 3.8× faster at prefill and a Mac Studio M3 Ultra as 3.4× faster at decode on the same model — then ran them together, Spark doing prefill and Mac doing decode with layer-by-layer KV streaming between them, for a 2.8× end-to-end speedup on Llama-3.1-8B at 8K context.

That is prefill/decode disaggregation — the same technique hyperscalers use across GPU pools — running on two consumer desktop machines. It is the cheapest possible demonstration that the asymmetry at the top of this article is real and exploitable.

The Spark also illustrates the batching argument more vividly than a datacentre GPU can, precisely because it is so bandwidth-starved. Single-stream decode on a 120B MoE runs at 33.5 tok/s; at 256 concurrent streams the same box produces 862 tok/s aggregate. On a Nemotron-49B model the ratio is 5.79 → 695 tok/s, a factor of 120. Nearly every published review of that machine benchmarks it as a single-user chat box and concludes it is slow, which is a bit like reviewing a bus on how fast it carries one passenger.

14. The silicon bet

A dozen companies have raised serious money on variations of one thesis: decode is memory-bandwidth-bound, GPUs run at 30–40% utilisation, and HBM brings packaging complexity and power-scaling problems. Build something else.

What separates them is how much programmability each is willing to surrender.

Every step right buys efficiency and costs generality. The question each company is implicitly answering is how fast model architectures will keep changing.

Company	Architectural bet	Status, Aug 2026
Groq	Statically scheduled deterministic SRAM-only processor, no DRAM at all	Non-exclusive technology licensing agreement with NVIDIA plus a team hire, Dec 2025. Groq remains independent; ~$20B is press-reported and undisclosed by both parties
Cerebras	Wafer-scale — eliminate the reticle cut. 44 GB of on-wafer SRAM at ~21 PB/s	IPO'd May 2026 at $185/share raising $5.55B, the largest US tech IPO since 2019; opened at $350
Etched	The transformer architecture hardwired into silicon	~$800M raised, ~$1B in contracts, racks shipping 2026 — no independent benchmarks published
Taalas	Model weights etched into transistors as mask ROM	First product launched Feb 2026
d-Matrix	Digital in-memory compute on commodity DRAM	$275M raised
SambaNova	Reconfigurable dataflow, three-tier memory for fast model swapping	$1B Series F; sovereign-cloud niche
Tenstorrent	RISC-V plus Ethernet scaling, GDDR6 instead of HBM	Licensing IP rather than only selling chips

The two datapoints that frame the whole category

Groq's economics. 230 MB of SRAM per chip means a 70B model at INT8 needs roughly 300 chips at ~375 W each — about 112 kW just to hold the weights. You buy sub-millisecond, zero-variance latency and you pay for it in silicon area and power. Groq had, by most measures, the best latency in the industry, and the outcome was licensing its technology to NVIDIA and losing its founding team to it.

The counter-argument from software. Together AI's ATLAS pairs a heavyweight statically-trained speculator (guaranteeing a floor) with a lightweight one that learns from live production traffic, arbitrated by a confidence-aware controller that adjusts lookahead depth. It reached 500 tok/s on DeepSeek-V3.1 on stock B200s and explicitly claimed to beat custom silicon on general-purpose GPUs. Hardware advantages in this space have a short half-life, because software redeploys in a week and silicon takes two years.

15. Economics

The formula is trivially simple, and that is precisely the point:

$ per million tokens  =  ( $ per hour ÷ tokens per hour ) × 10⁶

A B200 at $4/hour producing 4,000 tok/s yields 14.4M tokens/hour → $0.28 per million. The same GPU at 1,000 tok/s → $1.11 per million. A 4× software difference is a 4× cost difference — larger than the gap between most hardware choices.

The levers, ranked

Batch size and concurrency. The single biggest lever, and it is a bandwidth-amortisation argument: weight reads are shared across the batch, so tokens/hour rises nearly linearly until you hit the roofline knee or run out of KV memory.
Utilisation. Nothing else matters if the GPU is idle. A $2/hour H100 at 30% duty cycle costs $6.67/hour of useful work. Typical deployments run at 15–30% GPU utilisation, which means a large fraction of the capital currently being financed across the industry is sitting idle.
Precision. FP4 versus FP8 is roughly 2× on both bandwidth and FLOPs.
Model architecture. MoE active-parameter count sets decode bandwidth directly.
Framework and configuration. Routinely 2–3×; the AMD dispatch bug above was 7.7×.
Your interactivity SLO. Every token/sec/user you promise costs throughput. This is a product decision that engineers often inherit without being consulted.
Context length. Long context caps your batch, which caps throughput, which raises cost — this is why million-token pricing is superlinear rather than proportional.

Reference numbers

Configuration	$ / Mtok
B200, gpt-oss-120b, FP4, TensorRT-LLM, 55 tok/s/user	$0.02
B200, typical FP8 production chat	$0.10–0.50
H100, gpt-oss-120b equivalent	~$0.14
B200, Llama-4-70B BF16, batch 8	~$0.57

The one credible production disclosure

DeepSeek published actual numbers from its V3/R1 serving fleet, which as far as I know remains the only credible public unit-economics disclosure in the industry.

Peak nodes (H800)	278
Input tokens per 24h	608 billion
of which cache hits	342 billion (56.3%)
Prefill topology	EP32 across 4 nodes
Decode topology	EP144 across 18 nodes
Cost per day at $2/GPU-hour	$87,072
Theoretical revenue per day	$562,027
Theoretical margin	545%

Three separate load balancers (prefill, decode, expert-parallel), dual-batch overlap hiding all-to-all communication behind compute, and a five-stage decode pipeline. Note especially the 56.3% cache hit rate — over half the input tokens were never actually prefilled.

The conclusion I would draw

Inference is only low-margin if you are bad at it. The engineering gap between a naive deployment and an expert one is roughly an order of magnitude in cost per token. That gap — renewed continuously as models and hardware change underneath you — is the actual product that inference companies sell.

Is inference a commodity?

Mostly yes at the base layer. There is roughly an 8× price spread between the cheapest and the fastest provider serving identical open weights, which is arbitrage rather than differentiation. Token prices have fallen about 1,000× since 2021 — call it 10× per year.

But volume growth has outrun price decline: providers report going from 15 to 40 trillion tokens per day inside a year; aggregators from 5 to 25 trillion tokens per week in six months.

What people credibly claim is durable: supply access (contracted capacity — kernels get copied in a quarter, capacity does not), the train-serve loop (fine-tuning, RL and serving on one substrate creates real switching costs), workload-specific optimisation (speculation tuned to your traffic, prefix caching tuned to your prompt structure), and compliance and sovereignty.

What is clearly not durable: raw tokens per second.

16. Where the remaining headroom is

The field has moved from "make one GPU serve a transformer fast" to "operate a stateful, heterogeneous, power-constrained distributed system whose workload is generated by other programs." Three shifts drove it: reasoning models made output length wildly unpredictable, agentic workloads made statefulness-across-time the central abstraction, and reinforcement-learning post-training turned the inference engine into a component of the training loop.

The single most useful sentence I can offer

Well-tuned production systems run at maybe 30–60% of roofline for decode, and much of the remaining gap is scheduling and fragmentation, not kernels. The agentic results — up to 8× from changing a cache eviction policy — suggest the systems gap is far larger than the kernel gap.

Kernel work is nearly exhausted. Policy work is not.

An opinionated sort

Genuinely promising	Overhyped
KV cache as a first-class distributed system, with theory Agentic serving abstractions FP4 as a training format Hybrid linear and sparse attention RL rollout systems and the determinism thread Power-aware scheduling	Speculative decoding as a general throughput technique BitNet / 1.58-bit absent silicon Semantic caching — approximate-answer reuse with no correctness story LLM-generated kernels beating hand-tuned code (the baseline is usually unfused eager PyTorch — a low bar) Diffusion language models as autoregressive replacements (excellent as drafters, though)

The most underrated: offline and batch inference — enormous in production, almost invisible in the literature — plus fault tolerance for long-running agent inference, and heterogeneous-fleet serving.

Open problems, for anyone looking for something to work on

These are ranked by how accessible they are to someone with a systems background rather than by importance.

1. Scheduling under unknown, heavy-tailed service times

Every good scheduling policy needs to know job size. Reasoning models make output length vary 100× on the same prompt distribution. Measured KV utilisation swings from 3% to 70% for reasoning models versus under 3% for conventional ones, with straggler collapse as batches drain down to a few long-running requests.

Current approaches try to predict length better — entropy-guided predictors, phase-aware heuristics — and they are weak. The framing nobody has done properly: design schedulers with provable competitive ratios under adversarial or heavy-tailed unknown service times, rather than trying harder to predict. This is classical scheduling theory meeting a workload that breaks every assumption it makes.

2. Multi-tenant KV cache as a shared resource

There is no theory of KV admission and eviction across tenants with different SLOs and different reuse probabilities. The structure is unusual and unusually favourable: the "cache miss" cost is recomputation whose cost is known exactly and quadratic in position, and reuse is predictable from program structure rather than statistical. Classical competitive caching analysis should apply cleanly and largely has not been attempted.

3. Serving abstractions for agents — and a public trace corpus

A request is not the unit of work any more; a trajectory is. That invalidates most existing scheduling literature, which assumes independent requests. There is no consensus API, no standard trace corpus, and no accepted metric.

There is no public agentic serving trace corpus at all. That is the single largest benchmarking gap in the field, and building one requires taste and diligence rather than a GPU cluster. Someone will define the equivalent of TPC-C for agentic serving and it will be heavily cited.

4. The verification bottleneck in speculation

42–95% of speculative decoding time goes to target-model verification, not drafting. Nobody has a fundamentally cheaper verifier that preserves the exact output distribution. Approximate verification with controlled, bounded distributional drift is unexplored.

5. Determinism versus throughput

Batch-invariant kernels give reproducibility and unbiased RL gradients but cost performance. Nobody has mapped the Pareto frontier, or established whether statistical equivalence (unbiased log-probabilities) is achievable without bitwise determinism. This began as a reproducibility annoyance and turned out to be a correctness issue for RL training.

6. Certified KV compression, heterogeneous fleets, and power

Certified compression: given the impossibility results, characterise the realistic instance class where compression is safe, and build schemes emitting per-request error certificates a serving system can act on — "this response used compressed KV with bounded divergence; that one needs recompute."

Heterogeneous fleets: real clusters contain several GPU generations at once, and almost all serving research assumes homogeneity. Which shard, on which chip, for which request class, is open.

Power: co-located jobs phase-lock through a shared power cap, like coupled oscillators — a phenomenon only recently identified. Anti-correlating inference load to smooth aggregate draw without violating SLOs is wide open, and inference is the load that keeps growing.

What I would avoid: another KV eviction heuristic, another EAGLE variant, another routing classifier. Those areas are crowded, and the marginal paper in each is being retroactively invalidated by the impossibility results and the benchmark corrections described above.

Where the hard floors are

Floor	What it is	Headroom left
Bandwidth	decode time/token ≥ bytes touched ÷ memory bandwidth	Only by reducing bytes or raising batch size. No software trick escapes it.
Energy	65B model at ~100 ms: dense FP16 ≈ 3.5 J/token → INT4 ≈ 1.2 → latent attention ≈ 1.1 → combined ≈ 0.35 projected	~3× from already-demonstrated techniques, on top of ~3× already captured. Then you hit DRAM access energy of 2–5 pJ/bit, which is physics.
Latency	serialised layers plus collectives — all-reduce is ~5–10 µs, twice per layer	A few ms/token floor regardless of FLOPs. Speculation is the only way around serialisation, and it degrades exactly when you would want it.

A word on benchmarks

Three meta-flaws are worth carrying with you whenever you read a performance claim.

Almost no inference benchmark reports joules per token, despite power being the binding constraint on datacentre buildout. Benchmarks measure steady-state throughput on stationary arrival processes, while real workloads are bursty, multi-turn, and correlated. And quality is evaluated separately from efficiency, so a paper can claim a 4× speedup with "negligible degradation" on tasks chosen after the fact.

MLPerf Inference remains the standardised reference, but its submissions are heavily vendor-optimised, its query distributions are synthetic, and — critically — it has no multi-turn or prefix-reuse scenario at all, which means it does not measure the thing that dominates agentic serving.

17. If you want to learn this properly

A compressed version of a path that works, in dependency order. The most common mistake is starting at step 3.

Step	What	Why
1	Modal's GPU Glossary, cover to cover	You cannot read anything in this field without the words. Two hours.
2	kipply's Transformer Inference Arithmetic, then chapters 1, 4, 7 and 8 of the scaling book	The highest-leverage reading in the field. Then build a spreadsheet that takes a model config and predicts tokens/sec. Do not skip this.
3	GPU Puzzles, then PMPP chapters 1–6, then Modern GPU Programming for MLSys	The execution and memory model, before any kernel writing.
4	Triton tutorials 01 → 02 → 06	Tutorial 06 is FlashAttention-2 in about 200 lines. If you can read and modify it, you are competent.
5	llama2.c in one sitting, then vLLM's `v1/core/sched/scheduler.py`	700 lines of C for a full forward pass and KV cache. Then read a real scheduler — this is where the leverage is.
6	Build a mini engine: paged KV allocator, continuous batching, chunked prefill, prefix cache	~2,500 lines. Benchmark against vLLM and explain every gap. The gaps you cannot explain are your next reading list.

The papers that repay full reading, in roughly this order: Efficiently Scaling Transformer Inference · FlashAttention · Orca · vLLM / PagedAttention · Sarathi-Serve · DistServe · DeepSeek-V2 (MLA) · Mooncake. And DeepSeek's inference system overview, which is not a paper but is worth two readings.

Benchmarking mistakes that invalidate results

If you take away one practical thing, make it this list.

No warm-up. The first requests absorb CUDA graph capture, JIT compilation and weight paging. Discard them.
Not setting ignore_eos when claiming a fixed output length — your "1,024-token" runs quietly become 200-token runs.
Assuming one streaming chunk equals one token. It does not. Concatenate and re-tokenize.
Random-token prompts give an unrealistically low cache hit rate; repeating one prompt gives ~100% and wildly overstates. Use production-shaped traces and always report your hit rate.
Reporting means instead of percentiles. Decode latency is fat-tailed because of preemption and recompute.
Comparing throughput at different latencies. Sweep concurrency and compare curves, never single numbers.
Closed-loop only. Fixed concurrency hides queueing collapse. Use Poisson arrivals to find the real knee.
The client is the bottleneck. A Python asyncio benchmark client saturates at a few thousand tokens/sec.
Not pinning versions. vLLM and SGLang performance moves more than 20% between minor releases.

18. Glossary

Arithmetic intensity: FLOPs performed per byte moved from memory. Compare it to the chip's machine balance to know whether you are compute-bound or memory-bound.
Continuous batching: Scheduling at forward-pass granularity so finished sequences leave and queued ones join immediately, instead of padding a fixed batch to its longest member.
Chunked prefill: Splitting a long prompt's prefill into token-budgeted pieces co-scheduled with ongoing decodes, so a big prompt cannot stall everyone else's token stream.
Decode: The autoregressive generation phase — one token at a time, memory-bandwidth-bound.
Disaggregation: Running prefill and decode on separate pools of workers, shipping the KV cache between them, so each phase can use hardware and parallelism suited to it.
Expert parallelism (EP): Sharding a mixture-of-experts model by placing different experts on different GPUs, rather than splitting every expert across all of them.
Goodput: Requests per second completed within their latency SLO. The metric that matters; throughput alone can be gamed by batching harder and making everyone wait.
GQA / MQA / MLA: Grouped-query, multi-query, and multi-head latent attention — three ways of storing less KV per token by sharing or compressing keys and values across heads.
KV cache: Stored keys and values for every token processed so far, so attention need not recompute them. Frequently larger than the model itself.
Machine balance: A chip's peak FLOP/s divided by its peak memory bandwidth. The arithmetic intensity at which it stops being memory-bound.
MoE: Mixture of experts — many feed-forward blocks with a router activating only a few per token, so decode bandwidth scales with active rather than total parameters.
PagedAttention: Allocating the KV cache in fixed-size blocks from a pool, borrowed from OS virtual memory, eliminating fragmentation and enabling copy-on-write sharing.
Prefill: The phase that processes the input prompt. All tokens in parallel, compute-bound, sets TTFT.
Prefix caching: Reusing KV entries across requests that share a prefix. The dominant optimisation for agentic and multi-turn workloads.
Roofline: A plot of achievable throughput against arithmetic intensity, with a diagonal bandwidth limit and a flat compute ceiling. The standard way to see which resource is binding.
Speculative decoding: A cheap model proposes several tokens; the expensive model verifies them in one batched pass; rejection sampling preserves the exact output distribution.
Tensor parallelism (TP): Splitting individual weight matrices across GPUs. Latency-optimal but requires an all-reduce after every layer, so it needs high-bandwidth interconnect.
TTFT / TPOT: Time to first token (prefill-bound) and time per output token (decode-bound). The two latencies users actually experience.

Notes on sources and confidence

Figures here come from primary sources — vendor specification pages, arXiv papers, and engineering blogs from the teams operating these systems. A few things are worth flagging explicitly rather than burying:

Performance claims from Etched, Taalas and several other silicon startups are vendor-reported with no independent third-party benchmark published.
The ~$20B figure attached to the NVIDIA–Groq arrangement is press-reported; Groq's own release describes a non-exclusive licensing agreement and discloses no value.
Vendor headline FLOPS numbers usually include 2:4 structured sparsity, which almost no production LLM uses. Halve them for dense math.
Sparse-attention speedups across the various papers are not directly comparable — there is no unified benchmark.
The DeepSeek margin figure is theoretical, computed by the authors at list API prices against actual serving cost. It is not a revenue disclosure.

This is a fast-moving field and some of these numbers will be stale within months. The arithmetic in sections 2 and 3, however, will not be — which is the argument for learning that part first.

Thanks for reading. Follow me for more.

← More posts