Why Inference Matters 12× More Than Training

YC Paper Club walks the inference stack from an API request down to GPU kernels. Four talks from Together AI, Ollama, Luminal, and SGLang argue that serving quality, not just training FLOPs, is where most of the compute and the money now sit.

DFDan FuPonente · Together AIJMJeff MorganPonente · Ollama
Canal
Y Combinator
Publicado
24 oct 2026
Vídeo
1:50:48
Lectura
30 min
Diapositivas
32
Original
inglés
Abrir en YouTube
Al desplazarte, el vídeo se reduce a la esquina inferior derecha y sigue reproduciéndose. Los subtítulos siguen tu idioma de lectura.
Idioma de lecturaDisponible en 2 idiomas
Este artículo aún no está en español, así que se muestra el original en inglés.
RESUMEN

En 30 segundos

François Chaubard opens YC Paper Club with DeepSeek’s published numbers: about 2.8 million H800-hours to train DeepSeek V3 versus 43,000 H800-hours of inference per day, or roughly 6× more inference FLOPs per year—12× over a two-year weight-file lifetime. He frames serving as a stack of prefill, decode, KV cache, and three buckets of innovation (model, kernel/runtime, harness), then hands the night to practitioners. Dan Fu of Together AI traces a request from DNS to tokens, covering cache-aware routing, tool-call tokenization, prefill/decode disaggregation, kernels, expert parallelism on 72-GPU scale-up domains, and training-free quantization. Jeff Morgan of Ollama shows that the quality gap between the best and worst open-model providers can exceed the closed-versus-open gap. Joe Fioti of Luminal describes a search-based compiler built on equality saturation. Liangsheng Yin of RadixArk closes with SGLang’s radix-tree KV cache, overlap scheduling, and unified memory for hybrid architectures.

  1. 1Inference FLOPs dwarf trainingDeepSeek reports 2.8 million H800-hours to train V3 and about 43,000 H800-hours of inference per day. That is roughly 6× more inference FLOPs per year, or 12× if a weight file lives two years—and the ratio grows as more users share those weights.
  2. 2Serving is not model.forwardA request hits prefill, a decode loop, and a KV cache that may live in GPU DRAM, host memory, or slower storage. Training-like prefill might reach ~50% MFU; decode can sit around 5%. Time-to-first-token, time-per-output-token, throughput, and goodput are the KPIs.
  3. 3Expert parallelism needs a reliable 72-GPU domainFu argues that putting one MoE expert per GPU on NVIDIA’s 72-GPU scale-up racks turns decode back into a fat matmul GPUs are good at. The catch is reliability: one failed GPU can take out 72 GPUs of a replica unless the serving stack remediates gracefully.
  4. 4Provider gap can beat the open–closed gapMorgan cites work showing that the worst Llama endpoint can look closer to GPT-3.5 than to the best Llama endpoint. For multi-step coding agents, a 10% per-step deviation can compound to ~79% task failure over 15 steps.
  5. 5Compilers turn GPU hours into tokensFioti treats the compiler as the conversion function between GPU-hours and tokens. Luminal decomposes PyTorch into ~14 primitive ops, saturates an e-graph of equality rewrites, then profiles candidate programs on real hardware instead of maintaining huge heuristic kernel libraries.
  6. 6RadixAttention was built for forked agentsYin describes SGLang’s original bet: forked sub-agents share a long prefix and send short follow-ups. A radix tree over the KV cache reuses that prefix automatically, and cache-aware scheduling reorders requests so shared prefixes are not evicted between turns.

Conclusiones

  1. Over a typical two-year life of a weight file, inference can consume about 12× the FLOPs of training, so model comparisons should optimize iso-inference-FLOP, not just iso-training-FLOP.
  2. Prefill is compute-bound and decode is memory-bound; KV-cache hits, disaggregation, and decode kernels decide whether a GPU hour actually produces tokens.
  3. The same open weights can differ more across serving stacks than open models differ from closed ones, and agent loops compound a small per-step error into a failed task.
  4. Cheap ongoing checks—logprob drift, tool-call error rates, request-shape monitors—catch silent serving changes that expensive one-off benchmarks miss.
  5. Search-based compilers (equality saturation plus real-hardware profiling) can map a tiny primitive IR onto many chips without a thousand-person heuristic compiler team.
  6. Agent serving lives on prefix reuse: radix trees, overlap scheduling, hierarchical KV cache, and a unified memory pool across full attention, sliding windows, and recurrent states.
TRANSCRIPCIÓN

Lee la transcripción completa por capítulos

01

Why inference matters more than training

François Chaubard, a YC visiting partner, uses DeepSeek’s published H800-hour numbers to argue that inference is about 6× training compute per year and 12× over a two-year model lifetime. He walks through prefill versus decode, KV cache, MFU, and why COGS decides whether a model server keeps 80% margins.

François Chaubard opens YC Paper Club as inference night. He jokes that he didn’t know how to picture inference, so the slides are Halloween-themed, then turns to the question he actually cares about: why should anyone fuss over serving at all?

YC looks at inference constantly, he says, because he used to run an AI company where training and serving both mattered and serving was always the bottleneck. Doing it well at scale without losing your shirt on cost of goods sold—COGS, a word some software people never learned—is hard.

In the last ten years, getting inference right has produced about $220 billion of market cap. Four companies are on stage tonight; Together AI is on the board as well, represented by head of research Dan Fu.

Chaubard likes to compare training FLOPs with inference FLOPs, and DeepSeek is the rare lab that publishes both. Training DeepSeek V3 took about 2.8 million H800-hours. Daily inference was about 43,000 H800-hours—and that figure is already old, so the true load is higher.

Even those stale numbers are about 6× more inference FLOPs than training FLOPs per year for that weight file. If you give the file a two-year useful life, inference is roughly 12× more important, and the ratio keeps climbing as more users share the same weights.

That is why, when he is writing papers or picking models, he wants people to minimize training FLOPs—and to care twelve times more about inference FLOPs. Most NeurIPS scaling plots, Chinchilla included, still put iso-training-FLOP on the x-axis. The field should start plotting iso-inference-FLOP.

COGS is what a model server actually lives on. Numbers that had just come out implied DeepSeek was running around 80% margins, which shocked some people; Chaubard’s reading is that the margin comes from how much they have invented on the serving side.

Why is any of this complicated if serving is just `model.forward`? For people who don’t live in LLM stacks, a request arrives as token IDs. You run prefill on the prompt—unless that prefix is already in cache, ideally hot in DRAM, or sitting in something like Redis or SSD behind an LRU.

Then comes the decode loop. Model FLOP utilization (MFU) during training and prefill might reach 50%: if Jensen Huang promises 2×10¹² FP8 FLOPs, you are not getting that, and half would already be a win. Decode can sit around 5%. Being careless here is expensive.

The KPIs Fu and others chase are time to first token (how long prefill takes), time per output token on each decode step, throughput (tokens per GPU-hour), and goodput (tokens you actually deliver at a service level).

DeepSeek’s just-published datacenter numbers: about 60,000 tokens/s per GPU on prefill excluding cache hits, a 96% input-cache hit rate, and 1.5 million tokens/s per GPU when cache hits. Time to first token is about 400 ms, P99 about a second. Decode throughput is 14,000 tokens/s per GPU—versus 1.5 million on cached prefill—so decode is the bottleneck chip designers are chasing.

Cache hit rate is traffic-shaped. Agents reuse prefixes and look great in these numbers. A parent starting brand-new ChatGPT threads—“what’s the weather?”—has no KV cache at all and pays full prefill every time, so the 1.5 million figure is not what every product will see.

The number people in the room latched onto is KV-cache retention. DeepSeek keeps petabytes of user KV state for 72 hours after inactivity, which is how a 96% input-cache hit rate is even possible. That is a storage and privacy policy as much as a kernel trick.

Arithmetic intensity is the other term of art: if you move 4 bytes from DRAM to SRAM, how many FLOPs do you perform on them? RNNs and diffusion are intense—you transfer an image or noise and hammer the same weights. LLMs are poor, which is the argument a previous Paper Club speaker, Volo, made for diffusion language models. A next-concept model from Herman in Chaubard’s lab is in the same spirit.

Model FLOP utilization and model bandwidth utilization ask what share of an H100’s advertised 2×10¹² FP8 FLOPs and terabyte-class memory bandwidth you actually get at test time. Roofline analysis is the min of peak FLOPs versus bytes per second times FLOPs per byte—an upper bound on how good you can possibly be.

02

Three kinds of inference innovation

Chaubard buckets the field into model changes (MoE, SSMs, MLA, quantization, distillation), kernel and runtime work (FlashAttention, PagedAttention, RadixAttention), and harness tricks (routing, speculative decoding, recursive language models, compaction). He then introduces the four speakers.

Chaubard buckets inference work into three piles, with the usual fights about which idea belongs where. First, model innovations: mixture-of-experts, where K experts let you amortize test-time FLOPs by about 1/K; SSMs and sparse attention so you do not pay full attention every layer; and multi-head latent attention, which compresses in the sequence dimension.

Quantization—down to 2-bit in some demos—and distillation sit in the same bucket. Once you have a large model you would rather not serve, you distill into a smaller one if that is all the job needs.

Second, kernel and runtime innovations: RadixAttention, FlashAttention, PagedAttention. A lot of this is classic OS and database theory. Prefix trees and radix trees are how you make a production database fast; Tri Dao’s GPU tiling in FlashAttention was a huge jump; PagedAttention is virtual memory and a buffer pool for GPU DRAM, mutually exclusive with radix attention in the way Chaubard tells it.

Third, harness innovations. The easy one is routing—Focal did this—so not every request hits the gory big model; you can waterfall or tool-call up only when needed. Speculative decoding runs a draft model that the large model validates. Chaubard’s favorite is recursive language models: never dump the full context into the KV cache; let the model fetch what it needs, which shrinks cache a lot. Compaction is mandatory in every harness; a lab project with Sabri and Dan Simran produced “cartridges.”

For people who want a deeper lecture, he points to Percy Liang’s CS 336 lecture 10 and a Dwarkesh interview with Ryner Pope. Then he cedes the night to the people who do this for a living: Dan Fu, research lead at Together AI and an alumnus of Hazy; Jeffrey Morgan, founder and CEO of Ollama; Joe Fioti, founder and CEO of Luminal; and Liangsheng Yin, a core SGLang developer and second author of the paper, who had texted that he was running late.

03

Dan Fu: the lifecycle of an inference request

Together AI’s head of research follows a request from the public internet through cache-aware routing, tool-call tokenization, the scheduler, and the split between compute-bound prefill and memory-bound decode. Open models, he says, now carry a large share of token volume at much lower cost.

Dan Fu thanks Chaubard for the overview and says he will walk the whole life cycle of an inference request: from the API down through a datacenter to the kernels on a GPU, then tokens back out. The talk is a cut-down version of a longer one; six speedups are on the slide, and he has crossed out half of them.

Open-source AI has never been bigger. A slightly stale Together figure: 400 trillion tokens a month. Around the same time Google published Gemini traffic, and the “scrappy little startup” was doing about one-eighth of that volume. Earlier this year more than half of OpenRouter tokens were open models; marketing wants him to mention that Together once out-tokened OpenAI for a single day.

Those open models are also cheap. You can get nearly the same quality—sometimes the same or better—at up to 10× lower cost per year. Together spends a lot of time asking how to create more tokens in the world for everyone who wants them.

Together is trying to be an AI-native cloud: the AWS analog, but for the pieces you need to build an AI company—pre-training, fine-tuning, inference, and sandboxes for agents. The company is about 80 researchers, many of them PhDs, now focused on making GPUs more performant.

Workloads differ. For coding, they care about cache hit rate, the idle time between turns when you ask Claude Code or Cursor to build something, how fast generated code has to stream, and the SLAs on first token versus finishing the whole reply. Those constraints drive the optimizations.

When a request hits a REST API it traverses the ordinary internet—Cloudflare, DNS—and lands in a fleet of thousands of GPUs that hold a KV cache. First comes request-level routing: which replica should see this job. Then human text becomes ones and zeros. Then a batch loop of prefill and decode. Then token IDs become text again.

Routing is load balancing plus cache affinity. In the old web-service world you round-robin away from hot machines. In language models, sending the same user back to the GPU that already ate their codebase is a win, because subsequent turns skip that work. The art is spreading load across datacenters without missing a resident KV cache.

Tokenization looks simple and is not, because requests arrive with tool calls—search, grep, a build command, the internet. The model is matmuls and attention; it cannot browse. You hand it a list of tools, it emits a call, and you have to get schemas, recovery, and tokenizer-level parsing right. Jeff will pick this up from the user side.

The core scheduler batches incoming streams: wait for a free GPU, or wait for the previous request to finish. Those choices set the throughput-versus-latency tradeoff. More users on the same GPUs, or a snappy first token. Fu’s example is Gemini in Google Search: the first answer is very fast, and chatting afterward feels slower—probably different scheduling, maybe even different models.

Prefill is prompt processing. Dump in a whole codebase, and the GPU has to understand it; that is compute-bound, with high MFU and high arithmetic intensity. Decode emits one or a few tokens at a time. Each word reloads the entire weights from GPU memory onto the compute units: memory-bound, a different animal. Together spends a lot of time disaggregating the two onto different GPUs, parallelisms, and kernels.

Once the GPU has done its work, token IDs are detokenized and streamed back as text. Fu calls this the simple-looking end of the pipe. The remaining policy questions are when to free KV cache and how long to keep those tokens resident so the next turn can reuse them.

04

Kernels, parallelism, and quantization

Fu focuses on three of Together’s levers: agent-written kernels driven by production traffic, expert parallelism on 72-GPU Blackwell scale-up domains, and training-free quantization with a dogfooding loop that includes agent evals. A MiniMax NVFP4 recipe is his quality-matching example.

Together thinks about roughly six levers. Kernels are GPU programs: attention, MoE, mega-kernels. Disaggregation puts prefill and decode on different GPUs, which changes user experience and tokens per GPU. Parallelism is forced by size: Kimi K3 at 3 trillion parameters in FP4 is about 1.5 terabytes, and a B300 holds 280 GB, so the model has to be split—how you split it is the parallelism question.

Routing, after the request is inside the datacenter, is which replica gets it. The last two levers sit closer to the model: speculative decoding, where a small draft model tries to guess the big model’s next tokens and you keep them if they match; and quantization—how many bits per weight you actually store. For time, Fu will cover kernels, parallelism, and quantization.

Kernels matter in the prefill and decode math: matmuls, attention, the rest. Some ops are easy to tune; layer norms, RMSNorms, and stray softmaxes have tiny arithmetic intensity, so a kernel written for one of them alone wastes the GPU. That is a historical reason end-to-end training sits around 50% FLOP utilization and decode around 5%.

The old method was to stare at the forward pass, find squeezable spots, and send in a team of kernel engineers. Together has been using agents to do that automatically—an internal project called Together Compiler, which Fu admits is not actually a compiler and will get a better name.

The loop is driven by real traffic, because a coding workload and a chat workload light up kernels differently, and they insist on the full production research loop. It is the beginning of a recursive self-improvement loop for inference.

Parallelism exists because weights do not fit on one GPU; splitting buys compute and bandwidth. Tensor parallelism shards a matmul and reduces at the end. Pipeline parallelism cuts the network layer by layer. Data parallelism sends different sequences to copies of the same weights. The game is memory per GPU, which kernels you can write, and what the interconnect allows.

A core driver of DeepSeek’s $38-per-GPU-hour number is large scale-up domains. Until recently you could only have eight NVIDIA GPUs talking at high bandwidth, so you split eight ways; sixteen ways dropped onto a slow path and serving died. New racks connect 72 GPUs at high bandwidth, so you can put individual MoE experts on different GPUs.

One expert on a GPU serving a ton of tokens looks like a fat matrix multiply again—the thing GPUs are good at. Expert-parallel demos are years old; the new work is making them reliable. GB racks drop nodes. If a 72-GPU replica depends on every chip and one dies, you have lost 72 GPUs of compute unless you remediate. Together has been building that resilience.

Quantization looks simple: hundreds of papers, pick a recipe. It is not, because it hits end-user quality. Models train in BF16, two bytes per parameter; a 3-trillion-parameter model is 6 terabytes. The computation is fuzzy, so you can often shrink to 8-bit or 4-bit and still get the right answer. The question Together obsesses over is which recipe preserves the experience.

They know from the systems side which quantization the hardware wants for tokens per GPU. They also know people are building companies on these endpoints. Before a model ships, Together dogfoods it: broken tool calls show up when Claude Code cannot search; a dumber model shows up on benches and end-to-end evals.

One loop has every model coding Flappy Bird on repeat. If the bird stops flapping, quantization broke something. The real game is a rigorous evaluation loop: benchmarks, agentic evals, in-house dogfooding—not just a clever algorithm.

A success story: they worked with MiniMax to serve M3, and he lets slip M3.1. MiniMax was struggling to quantize to NVFP4, which Together wanted because it is the right format on NVIDIA GPUs. On the slide, the original FP8 model is lilac; MiniMax’s own recipe is hot pink; Together’s orange recipe matches the original, which meant more tokens and a path from training to profitable serving so MiniMax can fund the next model.

Recap: a request is routed to a replica, tokenized with tool calls kept correct, queued and batched, prefilled and decoded—often on different GPUs with different kernels and parallelism—and only then do tokens come back for coding and chat.

If your CFO is breathing down your neck about Claude Code spend, he has a sales QR code for open models; if you want to work on this, the other code is careers. Applause, then questions.

05

Q&A with Dan Fu

Audience questions cover how Together measures NVFP4 quality (saturated public benches, KL on activations, agent loops), NVLink on GB200 racks, training-free versus fine-tuned quantization, and which “dead-end” architectures (SSMs, sparse attention, dLLMs, looped models) already show up in frontier systems.

An audience member has two questions. First, how did Together measure NVFP4 quality against the MXFP8 original—GSM8K, perplexity, MMLU? Second, what interconnect jump let Grace Blackwell chips talk in a mesh beyond the old eight-GPU domain?

Fu says they publish a few open benchmarks, some of them saturated: AIME in the 90s. A one-point move is a real-world warning. They also look at KL divergence of layer activations, then hundreds of end-to-end agent cases scraped from every past failure, plus loops like Flappy Bird and feeding the model back into Together Compiler. If the model can no longer optimize its own kernels, something is wrong.

On NVLink he is not a hardware person—the closest he has come is a Hugging Face mini box. A GB rack is wild: each tray of four GPUs runs thick cables into a big switch, and all-reduces ride that path. He is not an expert on the jump from 8 to 72 GPUs; he just loves using it.

Pressed again on what they actually test against—because a buyer never sees the private suite—Fu repeats the mix: public benches, their own agent loops, and everything in between. There is no single score that tells you the quantization held.

Another question is about the MiniMax quantization graph: did they train from scratch, run extra pre-training epochs, or just transform the numbers with no fine-tune? The hot-pink bar looked like a naive cast.

Many recipes want a training pass after quantization. Together avoids those unless you own the original dataset—the model will forget things. Their MiniMax work is training-free: calibration data, a paper they released called Scale Search, and being sedulous about which of the hundred training-free methods you actually combine.

Chaubard squeezes in one last question before they have to move on: which new architectures is Fu still holding out for on the decode side—SSMs, diffusion language models, and the other ideas that might make decode less memory-bound?

Fu’s PhD was on SSMs, and the line then was that nobody would use them, transformers would eat the world. A few years later they are in the big models. Kimi K3’s KDA is, slightly controversially, an SSM-style linear-attention layer. Sparse attention, long unused, is in production. Diffusion language models are now the default speculators.

His students have been playing with looped models: send the token through the same layers several times so you can ramp training FLOPs without growing the weights. They thought it was a weird idea. Then it turned out OpenAI and Anthropic were already doing it. Put out a working architecture idea now, he says, and it can show up in GPT-6 or GPT-7; time-to-impact is that short.

06

Jeff Morgan: not all open-model tokens are the same

Ollama’s co-founder and CEO argues that routing across many providers creates a least-common-denominator problem: quantization, tool calling, sampling, and even wrong weights can silently ruin quality. The spread across Llama endpoints can exceed the closed-versus-open gap, and agents make it worse.

Jeff Morgan, co-founder and CEO of Ollama, zooms out: not all open-model tokens are the same. Chaubard’s deck already listed a crowd of inference companies; soon there will be hundreds or thousands of providers. If you are building on open models, you need a way to measure quality across them.

Ollama started in 2023 as an open-source project for running open models locally. The popular models then—Llama and friends—fit on a single H100, a laptop, or a small server. The goal has always been an OpenAI- and Anthropic-level developer experience for open weights.

Models then got huge. Early this year Kimi, GLM, and MiniMax shipped 500-billion- to trillion-parameter models, and Ollama launched Cloud so developers could keep the local UX in the datacenter. It has become one of the largest independent clouds for open models in the US and Europe.

That means they spend a lot of time bringing a consistent developer experience across many models and many providers: latest weights, speed, reliability. A lot of that lives in the router box Fu just described. Different providers excel at different models; reliability and early access matter most.

The architecture looks like what Ollama’s Fortune 500 customers build when they move off closed models for cost or control. Routing, in the ideal case, is best-of-all-worlds by pooling providers. The same problem shows up inside one provider that has different hardware, regions, teams, and engines.

The headline challenge: not everyone serves a model the same way, which is also a feature, because teams optimize differently and some even run on Mac minis around the world. Customers care about cost, intelligence, speed, and region/privacy. One bad apple on any of those four becomes the least common denominator.

Failure modes include over-quantization, faulty hardware, the wrong weight version, and tool calling. Parsing JSON looks simple; recovering from failed calls, enforcing schemas, and getting tokenizer-level parsing right is not. Sampling recipes differ. Custom kernels and KV-cache precision only show up in long context. Sometimes a model is even served under the wrong name, and every chatbot still claims to be a friendly assistant.

This is not new, and the scale of it is worse than the folklore. Everyone talks about the open-versus-closed gap. For open models alone, the gap between the worst provider and the best can be bigger than that closed-open gap. Together can feel like a different product from some other endpoint, and you cannot see it up front, because OpenAI and Anthropic trained people to expect a consistent token.

A 2024 paper tested 31 Llama endpoints and found some behaving closer to GPT-3.5 than to another copy of the same Llama. The problem gets much worse for coding agents: long context, long horizon, call after call. A modest 15-step coding task with a 10% per-step deviation from the intended output can be ~79% wrong. An 80% error rate ruins the product, and it gets worse as agent workloads get harder.

07

How to measure serving quality

Morgan surveys evals, lab-published tool-call scoreboards, distribution tests against reference weights, logprob drift monitors, and request-shape checks. He wants labs to ship the stack they actually serve—quantized weights, DeepSeek-style inference code, harness-side tool calling—and floats TEEs as a way to attest running weights.

You cannot read 1.6 million tokens per GPU by hand. The brute-force move is evals. Artificial Analysis reruns benches across providers and ranks them. For Kimi K2, most providers land in a cluster; a few on the right are off. Evals catch bad quantization because they test how the model was originally scored. They are expensive and slow, so you cannot run them continuously.

Late last year Moonshot got fed up with how providers did tool calling on Kimi and shipped a Kimi K2 tool-call verifier and scoreboard. It started a race to get calls right. Other labs have done something similar. If you train a model, giving people a test harness is how you make sure they run it correctly.

A cheaper method Ollama is looking at for routing is a distribution test: compare a provider API to reference weights and measure the gap. About a dollar a test, versus huge bench bills. It may miss subtle quantization damage, it does not tell you why, and you still need a reference model running.

Logprobs—the next-token probability distribution—are even cheaper: one token, compared to a baseline, to see whether the endpoint drifted. Did they move regions, upgrade vLLM or SGLang, patch the stack? A paper found a pile of silent serving changes this way. The downsides: many APIs no longer expose logprobs, and you only see drift from the past, so a bad starting point stays invisible.

Where Ollama is today is request-shape introspection without retaining user data: tool-call pass rates, errors, output-length distributions. DeepSeek-style engines have gone into infinite loops; you can catch that from the shape of the request. Keep a running score, request to request. OpenRouter’s Exacto-style routing is in the same family.

The north star is an OpenAI/Anthropic-like experience on open weights. Those labs run on TPUs and care enough that a tiny TPU-kernel bug became a quality retrospective. How do you get there in the open ecosystem? Labs can ship a much bigger serving stack, not just Hugging Face weights plus example code written by an applied team that is not what they run internally.

DeepSeek had just open-sourced DeepInfer, its custom inference engine, so you can take the weights and aim for that $38/GPU-hour number with some confidence. Labs are also quantizing before launch with people like Fu’s team, so the API and the open file match—unlike the old world where a third party quantized after the fact to fit a MacBook. Publishing official quantized weights is becoming normal.

You can also pull tool calling and reasoning extraction up into the harness so every provider is parsed the same way. Morgan thinks there is a startup in monitoring LLM quality for apps where the model is in the critical loop. And a favorite paper: modern NVIDIA GPUs have a trusted execution environment that can attest to the weights and the serving software—a fingerprint that you are running the real thing. Advanced, but interesting.

08

Q&A with Jeff Morgan

Questions hit logprob tests at long context, providers hill-climbing public evals, and Ollama’s local-plus-cloud future as capable local models return. Morgan’s answer is private customer evals, lighter tests than full benchmarks, and routing that will soon include the laptop as well as the datacenter.

An audience question about prefill-admission intervals for multi-turn agents—tool calls versus long-context prefill—does not quite land. Morgan points at the logprob slide: send a one-token prompt, measure logprobs against a baseline over time. You can grow the prompt and watch next-token accuracy as context passes 200,000 tokens, which is where some providers silently cap the window. It is not a full eval, and that is the point.

Someone asks about quantizers hill-climbing public evals: if the scoreboard is the objective, you are not preserving the model. Morgan has heard the worry, especially in a marketplace racing to the cheapest token. Ollama’s answer with customers is private, traffic-specific tests the provider cannot see, which is also how quality complaints usually arrive—closing the loop.

Chaubard notes that he first used Ollama as a local runtime, the industry ran to the cloud, and with Recursive’s Flint launch (Richard Socher) local models are back. What does that do to Ollama?

Morgan says there are many local stacks now—MLX on Apple silicon powers Ollama there. An audience member mentions Calibri serving Qwen 3.5 MoE from disk. The same quality problem exists on the desktop, on a DGX Spark, on an M5 Ultra they had just unboxed, and in a company’s own GPU closet. It is not only the provider; it is the engine you update every week.

They are excited to be launch partners on Recursive’s Flint. Open models are the wild west: twenty quantizations of everything on Ollama and Hugging Face, which is the point, and also why you need lighter tests than full benches if you serve customers live. Locally, that problem is just as important.

The future of Ollama, he thinks, is people running a mix. The hardest coding tasks this year needed cloud models; ultra-capable local models on a MacBook pull some of that work back. The routing problem stops being cloud-provider versus cloud-provider and becomes local versus cloud. Thanks, applause.

09

Joe Fioti: compiling with search, not heuristics

Luminal’s CEO presents equality saturation: a RISC-style primitive IR, non-destructive rewrite rules in an e-graph, multi-level ops from cuBLAS to ThunderKittens, and hardware profiling instead of a hand-written cost model. An eight-person team compiles onto NVIDIA, AMD, TPUs, Trainium, and inference ASICs.

Joe Fioti, co-founder and CEO of Luminal, says they are a compiler company, which is unusual until you notice that the compiler is the conversion function between GPU-hours and tokens. How good that function is decides whether you lose money, make money, or make a lot of money.

Their compiler is built on equality saturation, a fairly new technique; they believe they are the largest industrial user of it, and the right way to compile models onto many kinds of hardware. A large part of NVIDIA’s multi-trillion-dollar value, he argues, is that CUDA is a well-supported software ecosystem. Mapping models onto complex chips is still unsolved.

The better lesson, for them, is compile-time search. That lets them skip a zoo of rewrite heuristics and a kernel library that has to be rewritten every time a new model ships. A small shared compiler core, backends for inference ASICs, and they run on NVIDIA, AMD, TPUs, Trainium, and other accelerators.

At any given point they are doing something like 80 or 90 trillion tokens a month—not Together’s volume, but they are eight people. The talk is why this approach beats a typical MLIR stack for LLMs, diffusion, and the bio models that actually make money: egglog, what equality saturation means, how they structure IRs, and how they scale search.

Neural nets are not that complex: dense linear algebra, and not even that many ops. Luminal boils models down to about 14 or 15 primitives—add, multiply, sum, and so on. The whole of PyTorch, some 2,000 ops, decomposes into that set without losing information. They are not Turing-complete; Fioti likes to say they are PyTorch-complete.

They do not even have a matmul primitive: a matmul is a broadcasted multiply plus a sum reduction. That RISC-style bottleneck is how you avoid an all-to-all map from every PyTorch op onto every hardware op.

Superoptimization, a 1990s compiler-research idea that never took off in production, says you should search every form of a program with a cost function instead of knowing ahead of time how to lower it. The lowering space is exponential. Ordinary vendors hire compiler teams that grow with every new model, backend, and data type—500 or 1,000 people, still an imperfect compiler.

The alternative is tiny atomic rewrite rules that state equivalences, then let search decide what to apply when. Equality saturation does exactly that: write small equality rules, grow the search space, walk it.

The IR is a compute-graph DAG. An e-graph applies those rules non-destructively. A normal rewrite pattern-matches and replaces, destroying what it matched—you had better have been right, because you cannot go back. In an e-graph you add the right-hand side and keep both sides. Apply every rule in every order until nothing new appears: saturated. Then extract concrete graphs, compile them, run them on the hardware, and time them.

A visual: start with A × 2 / 2. Rules grow the graph—A × 2 is a left shift by one, and so on—until the saturated e-class contains both the original division and A itself. You have shown A × 2 / 2 = A without a human writing that heuristic.

The rewrite rules are Lisp, because the engine authors like Lisp: commutativity, x + y = y + x (not quite true in IEEE floats, a long side topic), constant folding so a known a + b disappears from the runtime program.

That primitive IR is shared across backends. Each chip brings its own implemented ops, at various levels: a GPU tensor-core MMA that reads and writes SRAM is one. You characterize a new backend by writing rewrite rules that equate those ops to things the compiler already understands; the equalities propagate. Adding hardware that looks nothing like the last one is comparatively little work, which is how you avoid a massive compiler team.

Hardware is hierarchical: a thread, a warp, a thread block (an SM), a GPU, a node. The IR is multi-level so vendor libraries and community kernels can plug in at the level they actually program. Equality saturation is the sane way to say what equals what across those semantics.

Examples: a host-level cuBLAS matmul from the CPU; thread-block ops from CUTLASS, cuBLASDx, or Hazy Research’s ThunderKittens that program one GPU core. Because the compiler can reason across levels, it can fuse lower-level subgraphs into a mega-kernel—or not. Mega-kernels are fashionable and not always optimal on very large models, so a forward pass can mix both.

The last problem is the cost model: how do you say program A is better than B? An accurate GPU cost model is many PhDs of work. An eight-person company cannot write one for every chip. In the spirit of the bitter lesson, the best cost model is none: run candidates on the real hardware. Moves that look dumb on paper sometimes win on the GPU.

They do not profile every program in the e-graph. Roofline bounds on a subgraph let them prune anything that could not beat a measured equivalent even in the best case. Mostly, though, they still measure on the chip they will actually serve.

A short video of kernel codegen does not carry much information; it just looks neat. Then the obligatory plug: they have done the rare thing of making a compiler company profitable. They run a heterogeneous cloud—80 trillion tokens a month, millions of images—and because the compiler is general they will serve protein-folding and other “wacky” models, not only the popular LLMs. Eight people, hiring, commoditizing wacky hardware. lumina.com, he says, meaning Luminal.

10

Q&A with Joe Fioti

Fioti addresses floating-point associativity versus token quality, why profiling full forward passes is fast enough, genetic search toward a local optimum, compile-time budgets, e-graph proofs for debugging, fusion after expanding to primitives, and why the search stays inside the forward pass rather than speculative decoding.

Someone who likes the RISC-style IR asks how much of machine learning’s relaxed numerics they bake in. Fioti had mentioned associativity in IEEE 754; ML often does not care about reduction order. Does that enter the language itself, or only the cost model?

Associativity and commutativity matter because they change numerical accuracy. A tree sum is not a sequential sum. If you screw up numerics you screw up output-token quality, the same provider-quality problem Morgan described. They model accuracy explicitly. Some rewrites still help because hardware is asymmetric: NVIDIA tensor cores may take input A from SRAM and input B from a different memory space.

Another questioner does not see how you can run the whole model on hardware for every candidate, and if you only profile a kernel you miss contention. Also, the search space looks an order of magnitude bigger than TensorRT, which is already slow.

A full-model profiling pass, Fioti says, can be under 10 milliseconds, so you can do about a thousand runs a second. Search uses standard tools, mostly genetic algorithms; learned search like MuZero is future work. They explore maybe 1% of 1% of the space and are not hunting a global optimum—just a local one that is ~99% as good, which is much easier.

Compile time is an explicit trade: inference customers do not care if you compile for a day when the result runs across millions of GPUs 1–10% faster. Users set a time budget—an hour, ten hours—until diminishing returns. It scales reasonably until it doesn’t.

Do they feed production profiles back into search, or is it pure compile-time? They use shapes as close to real traffic as a cloud operator can: they know the tensor sizes customers will send. Rewrite rules are assumed correct, so they are not checking answers so much as timing industry-relevant shapes.

Debugging a compiler that lowers many architectures onto many chips is painful. Luminal keeps PyTorch’s high-level op tags as metadata, uses the PyTorch profiler for bottlenecks, and leans on egglog’s gift: if an output program is wrong, ask for a proof—the chain of rewrites from input to output—and find the invalid rule.

A commenter likes kicking unsupported ops back to CPU, but expanding ops in the IR usually hurts fusion. What is the latency hit? Fioti says they do fuse. Expansion is to reach the primitive IR; lowering searches through fused forms again. A roofline check that one kernel moves an order of magnitude more DRAM traffic is enough to reject it. The mid-level looks expanded; search has to find the patterns and fuse them back.

Someone picked up that the language is not Turing-complete, so it is not subject to Rice’s theorem: you can reason about what the program will do. Where do they exploit that?

That was on purpose. Superoptimization died in the 1990s on C, which is Turing-complete, with a near-infinite compilation space. They compile dense linear algebra, a tiny slice of all programs—PyTorch-complete, not Turing-complete. They can analytically count bytes moved between memory spaces and derive a roofline without running the model, which is why they prune so aggressively and profile so little of the space.

Does the search include quantization, speculative decoding, and parallelism, or only kernels? Generally no. They take a forward-pass graph and compile it end-to-end inside that pass, then live inside vLLM and SGLang, piggybacking on those engines’ speculative decoding. The thesis is that most of the alpha is complex parallelism topologies and a hairy forward pass; they do not search the engine layer.

11

Liangsheng Yin: SGLang and memory for the agent era

The SGLang co-author tells the 2023 origin of RadixAttention, then overlap scheduling for GPU/CPU bubbles, HiCache across GPU/host/storage, and a unified radix tree plus compacting memory pool so full attention, sliding windows, and recurrent states can share cache without fixed ratios. He closes with Miles RL and an upcoming SGLang summit.

Liangsheng Yin, a core SGLang developer at RadixArk, jokes that he thought this was a paper-reading club and had prepared half on the original SGLang paper and a slice on current open-source work, then added extra slides on the fly.

The story starts in late 2023, when Llama was the open model people had and agent workloads were still a guess: tool calls, parallel sub-agents, something like a future Claude Code. If you fork a sub-agent, a lot of context is shared in the same KV pool, and the follow-up prompts are short—in and out of the engine many times. No serving engine then reused prefixes efficiently.

That is why they built RadixAttention: automatically reuse a shared prefix and compute only the new suffix. All agent contexts live in the engine’s memory system as a tree—a radix tree in the implementation—rather than as isolated per-request KV blobs.

The picture is a long system prompt, then one agent doing authentication and another doing review, sharing that prompt, doing different work. Memory is finite, so a suffix gets evicted. Memory management really is that radix tree, and scheduling is built around it.

On eviction they drop unused leaves when memory is needed, with policies tuned to the workload. Cache-aware scheduling matters because prompts arrive interleaved. If you reorder them, you can batch in an order that reuses context A instead of interleaving A and B and risking an eviction. That idea is how SGLang became a serving engine: a radix tree inside the memory system.

By mid-2024 they had made it a general LLM engine—version 0.2—covering Llama 8B through Llama 3.1 405B for online and offline serving. Batching, eviction, and prefix matching all mean ordering waiting requests and assigning priorities.

GPUs got faster; CPUs did not. Vanilla scheduling prepares batch one, runs it, copies results back, then prepares batch two. When the GPU was slow, a <10% bubble was fine. On faster GPUs the CPU-side bubble became obvious.

Overlap scheduling delays device-to-host synchronization by a round so nothing blocking sits on the critical path, and prepares the next batch while the GPU is still running the current one. Launch batch two while batch one runs. DeepSeek’s just-released inference stack uses the same zero-overhead idea. On Hopper it was maybe 1.1× versus vanilla; on Blackwell in GB settings it can hit 2×.

The agent-and-cache story continues because GPU memory is always too small. Users come back minutes, hours, or days later. You have to put the prefix somewhere. That is HiCache, now a standard trick: GPU memory, host memory, and internal storage as L1, L2, and L3. Every token can live on one of those tiers, with policy for hit rate and load time. In their experiments it cut TTFT by 56%.

Plain transformers are a poor fit for long context and cheap deployment, so architectures keep changing: sliding-window attention in open and closed models, recurrent or Mamba-style states, Kimi’s KDA, fixed-shape memory updated token by token. The naive fix is a different memory pool per layer type. It is not trivial, because cache sharing has to be consistent across them.

To hit cache at position 10 you must match the full-attention KV up to there. If the sliding window is four, you also need tokens 7–10; if you only retained through 8, there is no window hit and you recompute. To share a prefix at 10 you need the recurrent state exactly at 10. Manage those pools separately and efficient sharing becomes a mess.

Moving cache hit rate from 95% to 99% is not a 4% win: it can cut compute about 5×. Separate radix trees per pool caused compatibility headaches, and locking a prefix meant coordinating full-attention layers, window layers, and recurrent layers under one external policy.

The unified radix tree stops treating those components as standalone. It is a semantic tree over full token IDs—whatever the architecture, there are still tokens—and the other memories hang off that controller. Each component keeps its own match, eviction, and memory rules, but the tree coordinates them so agentic workloads do not leak or waste memory.

Workloads disagree on mix. User A lives at 4K context, user B at 128K, someone else at a million tokens. Full-token cache has to cover the long ones; a sliding window may only need the last 128. You cannot pre-allocate pools at a fixed ratio and still serve everyone on one engine.

So besides the unified tree they built unified memory pools. Previously they pre-allocated at a computed ratio. Now they allocate one pile of physical memory with no per-pool cap: full-token cache grows from one side, window KV and recurrent state from the other, as usage demands. Long context gets more full-token pages; short context gives more to windows and states.

Most kernels want contiguous physical memory. Allocate and evict freely and you punch holes the two pools cannot reuse when they meet in the middle. Compaction packs fragments with low-latency copy kernels and remaps virtual page IDs—the virtual IDs stay put; only the low-level map changes. That is what makes mixed serving flexible.

They have also brought SGLang-style inference into the training loop. Miles, an RL framework from RadixArk, runs fully asynchronous so generation and training keep moving; SGLang does the rollouts. The arc he wants you to see: RadixAttention, then a general open engine, then Miles and even diffusion, boosting open AI infra.

He closes with an SGLang summit the following month—featured speakers, a QR code to register—and opens the floor. The clock is gone, so there is no Q&A. Chaubard thanks Yin and the rest of the night, and Paper Club is over.