Why inference is the piece

Scale drove the capability explosion (hundreds-of-millions of params in 2018 → multi-trillion today), and the transition is faster than expected — like Manhattan going from 130,000 working horses (1902) to cars outnumbering them by 1912. For LLMs that "1912 moment" was roughly last year. If GPUs are the new oil, inference is the engine that turns it into useful motion. Training is a one-time cost; inference is the repeated cost that actually serves users — and ML models are just DAGs of operations that inference engines/kernels map onto hardware.

The pipeline

Production traffic looks like neither training data nor a made-up workload. Each application has its own token-shape distribution — a coding agent sends tens of thousands of input tokens (the whole codebase) then a short or thinking output; summarizing a book looks different again; quick chat different still. Workloads are turn-based and agentic (tool calls, multiple turns, variable inter-turn gaps), each with its own SLA: interactive chat wants first token in <1s; batch jobs just want throughput. The session shape (tokens/turn, generation length, stickiness, gaps) defines the serving problem.

Prefill vs decode — two different beasts

After tokenizing and a cache lookup, the ML computation splits into two phases with opposite characteristics:

Because they're so different, you disaggregate them onto different workers — even different chips: prefill on flop-heavy GPUs, decode on memory-bandwidth chips (NVIDIA buying Groq's LPUs; OpenAI's Cerebras partnership; SambaNova). The engine runs a loop: schedule → execute → sample, repeat. Continuous batching keeps it full — requests of varying length come in, finish, and are evicted dynamically, contending for both compute and KV-cache memory.

The KV cache is the heart of serving

Many users start with "hi ChatGPT," or paste the same book, or continue a conversation — so you cache activations and reuse them via prefix sharing (a radix tree finds which tokens you've seen). You want this cache as large as possible, so it spills down a memory hierarchy:

Scale bugs & a simple win

At trillions of tokens/day, things that work at small scale break at 0.001% frequency: a subtly-wrong kernel turns logits to NaN → the model loops "hi hi hi"; a broken tool-call handler → doom loops of tens of thousands of tokens; an off-by-one read of uninitialized GPU memory → the model randomly emits a Chinese character then "decides" to think in Chinese (misblamed on quantization).

Megakernels — beat the per-op overhead

Decode runs the whole model to produce one token, turning a massively parallel GPU into a "glorified memory loader." Worse, we write one kernel per operation (norm, matmul, attention) — easy to program, but it injects downtime: kernel launch/teardown gaps, tail effects (waiting on the longest item), and gaps between kernels that add up across the model.

Parcae — scaling by recurrence, stabilized

Beyond scaling params and data, can you scale a third axis? Parcae (UCSD) revisits looped transformers: run some blocks in a loop, keeping parameters constant but adding a FLOPs dial — more quality-per-parameter and more expressivity. The catch: naive loop transformers blow up if you nudge the learning rate (loss spikes, NaNs) — a sign something deeper is wrong.

# Look at the residual across loop steps (it changes little). Isolate the
# nonlinear block R; what's left is a linear system in matrices A, B:
h[t+1] ≈ A^t · (stuff) + …            # dominated by powers of A
# spectral radius of A > 1  →  explodes;  prior choices were (marginally) unstable
# Parcae: A = negative-diagonal (decays), B = norm-constrained  →  spectral radius < 1 → stable

With this SSM-style reparameterization, training is stable even at learning rates that wrecked prior loop models, and Parcae beats both a prior loop transformer and a strong transformer baseline on perplexity and downstream quality. The scaling-law punchline: holding params fixed, as you scale data you should also scale recurrence (the curves go "down and to the right") — yet today's models have zero recurrence (the far-left of the curve), hinting we might be leaving quality on the table. And fewer parameters is an inference win: more room for KV cache, less cross-GPU communication, maybe even a recurrent-block megakernel.

Co-design notes (from Q&A)

• Design for the serving hardware, memory first. Size the model + KV cache to fit the target chip (Cerebras wafer, Huawei chip). Quantization is hardware-specific: NVFP4 for NVIDIA (Nemotron), MXFP4 for AMD.
• Use case dictates architecture. Agentic loops want a hot KV cache → aggressive compression (DeepSeek MLA, FP8/FP4 KV cache). Batch processing that sees each doc once may not need the cache at all — and may prefer bidirectional (BERT-style) over causal attention.
• Megakernels + multi-GPU: you can fuse NCCL communication into the kernel (DeepSeek shipped a fused MoE-inference megakernel), but you typically megakernel a part of the model, not the whole thing, unless you pay the full cost.

Golden rules & takeaways

• Inference is the engine turning GPUs into intelligence — and it's a repeated cost, so efficiency compounds.
• Prefill (compute-bound) and decode (memory-bound) are different beasts — disaggregate them, even onto different chips.
• The KV cache is central: maximize it via prefix sharing and an OS-style memory hierarchy (GPU → DRAM → SSD) with LRU/prefetch.
• Scale exposes rare, nasty bugs (NaN loops, doom loops, off-by-one Chinese characters) — and simple systems wins (cache-aware routing → 40%).
• Megakernels fuse many ops to kill per-op downtime, reaching near speed-of-light decode — at a heavy engineering cost.
• Looped models (Parcae) add a recurrence axis; SSM-style stabilization (spectral radius < 1) fixes their instability, and scaling laws suggest recurrence should grow with data.
• Understanding inference + kernels enables full-stack innovation — routing, kernels, and architectures all become design levers.

A guest perspective rounding out the course: how everything you've built gets served — and where the open research problems are.

References

Guest lecture by Dan Fu (UCSD / Together). Related works & tools:

1. Yu et al. (2022) — Orca (continuous batching)
2. Kwon et al. (2023) — vLLM / PagedAttention
3. Zhong et al. (2024) — DistServe (prefill/decode disaggregation)
4. ThunderKittens (Hazy Research) · Megakernels blog
5. Dehghani et al. (2018) — Universal Transformers (looped)
6. Geiping et al. (2025) — Scaling test-time compute via recurrent depth
7. Together AI · UCSD Hao AI Lab (inference research)