Why Inference Matters 12x More Than Training | YC Paper Club

Why Inference Matters 12x More Than Training | YC Paper Club — Y Combinator

Y Combinator

AI-assisted summary of Why Inference Matters 12x More Than Training | YC Paper Club — original source. Check the original for full context.

Source published:

Key ideas

  • Serving models accounts for roughly 6 times more annual compute than initial pretraining, meaning that over an estimated two-year model weight lifecycle, inference represents roughly 12 times the economic importance of training. (0:09)
  • The performance variance between the best and worst providers serving the exact same open-weight model is larger than the gap between open and closed models, with low-tier endpoints degrading Llama models down to GPT-3.5 performance. (48:08)
  • Luminal operates a search-driven compiler serving 80 to 90 trillion tokens per month with an eight-engineer team by decomposing PyTorch into 15 atomic primitives and using equality saturation via e-graphs instead of manual heuristics. (64:30)
  • SGLang implements hierarchical KV caching and overlap scheduling to decouple CPU orchestration from GPU kernels, cutting Time to First Token by 56% while delivering up to a 2x throughput speedup on NVIDIA Blackwell architectures. (95:45)

Detailed notes

## Background Published on YouTube by Y Combinator on October 24, 2026, this edition of YC Paper Club focuses on the economics, architecture, and infrastructure of large-scale AI inference. The session explores the complete lifecycle of inference requests, quality disparities across open-weight model providers, search-driven compilation, and advanced memory management for agentic workloads.

## TL;DR While mainstream machine learning research prioritizes training budgets, production models consume roughly six times more compute annually during inference, making serving efficiency the decisive factor in AI unit economics. Experts Dan Fu, Jeff Morgan, Joe Fioti, and Lianmin Zheng dissect the entire inference stack, demonstrating how hardware bottlenecks like the memory-bandwidth wall dictate system performance during prefill and decode stages. They show how provider-level quantization and routing choices can degrade model capabilities more severely than the gap between open and closed weights, how search-based compilers dismantle PyTorch into atomic primitives to out-optimize human-engineered kernels, and how hierarchical Radix tree prefix caching in systems like SGLang drastically reduces latency and expense for multi-turn agentic applications.

## Atmosphere & style The presentation carries the focused, practical energy of an elite systems engineering symposium, devoid of marketing hype and grounded strictly in profiling metrics and hardware realities. The speakers share war stories from production deployments—such as using Flappy Bird test harnesses to catch quantization breakage or tracking down subtle compiler kernel bugs that silently ruin model reasoning. It feels cerebral, pragmatic, and deeply technical, blending systems-level precision with the quiet excitement of builders discovering massive efficiency gains in the lower layers of the AI compute stack.

## Bulletin - [00:09](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=9s) Serving models accounts for roughly 6 times more annual compute than initial pretraining, meaning that over an estimated two-year model weight lifecycle, inference represents roughly 12 times the economic importance of training. - [48:08](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=2888s) The performance variance between the best and worst providers serving the exact same open-weight model is larger than the gap between open and closed models, with low-tier endpoints degrading Llama models down to GPT-3.5 performance. - [01:04:30](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=3870s) Luminal operates a search-driven compiler serving 80 to 90 trillion tokens per month with an eight-engineer team by decomposing PyTorch into 15 atomic primitives and using equality saturation via e-graphs instead of manual heuristics. - [01:35:45](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=5745s) SGLang implements hierarchical KV caching and overlap scheduling to decouple CPU orchestration from GPU kernels, cutting Time to First Token by 56% while delivering up to a 2x throughput speedup on NVIDIA Blackwell architectures.

## Key takeaways - [00:09](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=9s) Inference economics represent the central operational bottleneck in AI startups, having created over $220 billion in enterprise market capitalization over the past decade. - [03:19](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=199s) Prefill compute operates in parallel at roughly 50% model FLOP utilization, while autoregressive decode is strictly memory-bandwidth bound and drops GPU efficiency to roughly 5% MFU. - [06:52](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=412s) Hardware execution is governed by roofline analysis, where attainable performance is strictly bounded by arithmetic intensity and DRAM-to-SRAM memory transfer speeds. - [23:47](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=1427s) Together AI uses autonomous compiler agent loops observing live production traffic to iteratively discover and generate optimal kernel fusions. - [26:12](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=1572s) Serving multi-trillion parameter models across modern 72-GPU NVLink fabrics introduces single-point-of-failure risks if any single device failure crashes the entire 72-device replica. - [29:29](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=1769s) Aggressive quantization recipes risk subtle behavioral failures, which Together AI monitors through automated behavioral canary tests such as generating working Flappy Bird games. - [40:17](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=2417s) Non-transformer architectures such as looped models, diffusion language models, and state space models are expanding inference capabilities without bloating parameter memory footprints. - [48:08](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=2888s) In multi-step agentic workflows, an unvetted endpoint with a 10% per-step drift or deviation compounds into a ~79% to 80% total failure rate across a 15-step task. - [50:50](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=3050s) Inspecting token logprobs against reference instances allows continuous monitoring for unannounced backend software or hardware migrations without storing sensitive user prompt text. - [01:17:38](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=4658s) Superoptimization compilers rely on 'the bitter lesson,' running candidate implementations directly on real silicon to measure runtime rather than attempting to construct theoretical hardware cost models. - [01:27:27](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=5247s) Restricting compiler optimization strictly to dense linear algebra IRs avoids undecidability pitfalls like Rice's theorem while preserving complete compatibility with PyTorch programs. - [01:30:29](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=5429s) RadixAttention structures the KV cache as a Radix tree, preventing redundant prefill computation across parallel sub-agents, multi-step tool calls, and monitor loops. - [01:39:52](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=5992s) Increasing context cache hit rates from 95% to 99% in agentic architectures reduces total inference compute expenses by a factor of five. - [01:48:00](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=6480s) Dynamic memory compaction using low-latency memory copy kernels eliminates physical GPU memory fragmentation caused by heterogeneous context lengths across users and time zones. - [01:49:24](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=6564s) Decoupling rollout sample generation from weight updates via asynchronous execution bridges high-throughput inference stacks directly into reinforcement learning training loops.

## Detailed walkthrough ### [00:09](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=9s) Why Inference Matters: Flops, COGS, and the 12x Ratio

The session opens on an energetic note for "inference night," with the host jokingly acknowledging his slides' Halloween theme—designed to look "spookier," or "pookie" as his three-year-old says—because representing inference visually is surprisingly tricky. He pivots to the foundational question: why care about inference? Y Combinator makes heavy investments in inference, driven by a hard economic reality that the host experienced firsthand when running an AI startup: training models is demanding, but serving them efficiently at scale without suffering devastating Cost of Goods Sold (COGS) is the genuine operational bottleneck. As he quips to the audience, “It's so hard to do inference well at scale and without losing your shirt on cost of goods sold, uh a term called COGS, as some people in software are unaware what that word means.” Over the past decade alone, this infrastructure challenge has created more than $220 billion in enterprise market capitalization.

```mermaid flowchart LR accTitle: Train vs Inference Compute Lifetime Ratio accDescr: Comparison of training compute versus yearly and multi-year inference compute for DeepSeek V3. subgraph Training["Training Phase"] T["DeepSeek V3 Train<br/>2.8M H800 hours"] end subgraph Inference["Inference Phase"] I1["Yearly Serving<br/>43k H800 hrs/day<br/>(~15.7M hrs/yr = ~6x Train)"] I2["2-Year Life Serving<br/>~12x Train FLOPs"] end T -->|"1 year of serving"| I1 I1 -->|"2-year model lifetime"| I2 ```

To illustrate why inference dominates compute budgets, the host points to DeepSeek, the only major player that publicly discloses its complete infrastructure figures: - **Training Budget:** DeepSeek V3 consumed approximately 2.8 million H800 GPU hours during its initial training run. - **Inference Footprint:** DeepSeek burns roughly 43,000 H800 GPU hours every day just running inference—and that number reflects older baseline traffic. - **FLOP Multipliers:** That baseline equals roughly 6x more FLOPs spent on inference per year than the entirety of the model's original training run. - **Weight File Lifetime:** “And if you believe the useful life of a weight file, I don't know, say 2 years, then it should be 12x more important.” - **Paper Metric Blindspot:** Machine learning papers (like the Chinchilla scaling laws) routinely evaluate architectures strictly under "iso-train FLOPs" along the x-axis, whereas researchers should evaluate model efficiency under "iso-inference FLOPs" because inference compute outstrips training compute by an order of magnitude. - **Gross Margins:** DeepSeek recently reported implied gross margins of roughly 80%, a margin profile that caught many observers off guard and stems directly from their deep engineering innovations on the inference stack.

---

[03:19](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=199s) The Realities of LLM Serving: Prefill, Decode, and Key KPIs

While outsiders often assume inference is merely executing `model.forward()`, serving models in production requires managing two completely distinct phases with vastly different hardware efficiencies:

```mermaid flowchart TD accTitle: LLM Request Lifecycle and KV Cache Flow accDescr: Flowchart illustrating how incoming prompts hit the KV cache, prefill phase, and decode loop. subgraph Input["Input Request"] P["User Prompt (Token IDs)"] end subgraph PrefillPhase["Prefill Phase"] C{"Cache Hit?<br/>(DRAM / Redis / SSD)"} Hit["Skip Prefill<br/>Reuse KV Cache"] Miss["Run Prefill Kernel<br/>(Compute-bound, ~50% MFU)"] end subgraph DecodePhase["Decode Loop"] D["Autoregressive Decode<br/>(Memory-bandwidth-bound, ~5% MFU)"] Next["Emit Next Token"] end P --> C C -->|Yes| Hit C -->|No| Miss Hit --> D Miss --> D D --> Next Next -->|"Repeat until EOS"| D ```

1. **Prefill Phase:** The engine processes the prompt token IDs in parallel to compute initial key-value representations. If prompt prefixes are retained in an LRU cache (residing in DRAM, Redis, or local SSDs), recomputing prefill can be skipped entirely. Prefill behaves much like training compute, achieving decent Model FLOPs Utilization (MFU) of around 50%. 2. **Decode Phase:** The engine generates output tokens autoregressively one by one. Decode is heavily memory-bandwidth bound rather than compute bound, collapsing GPU efficiency down to an MFU of roughly 5%.

The host explains the core metrics and benchmarks tracking these phases: - **Time to First Token (TTFT):** Measures prefill latency before generation begins. DeepSeek hits roughly 400 ms median TTFT, with a P99 of approximately 1 second. - **Time Per Output Token (TPOT) / Decode Latency:** The per-step decode duration needed to emit each subsequent token. - **Throughput vs. Goodput:** Throughput measures aggregate tokens generated per GPU hour, while goodput tracks only those tokens generated within agreed Service Level Objectives (SLOs). - **DeepSeek Prefill Throughput:** Hits 60,000 tokens/sec per GPU when excluding cache hits, and skyrockets to 1.5 million tokens/sec per GPU with cache hits. - **DeepSeek Cache Rates & Retention:** DeepSeek operates at a 96% input cache hit rate across petabytes of KV caches, backed by an aggressive 72-hour KV cache retention policy. - **Workload Variance:** High cache hit rates depend on workload patterns. Agentic workflows that repeatedly append context to long system prompts achieve extraordinary cache hits; in contrast, users like the host's mother—who repeatedly opens brand new ChatGPT windows to ask disconnected one-off queries like "What's the weather?"—generate zero reusable KV cache, forcing the engine to re-run full prefill every time. - **Prefill vs. Decode Discrepancy:** On DeepSeek's setup, prefill throughput with cache hits reaches 1,500,000 tokens/sec per GPU, whereas decode throughput caps at only 14,000 tokens/sec per GPU. This stark gap explains why chip architectures focus so heavily on resolving the decode bottleneck.

---

[06:52](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=412s) Roofline Analysis and Three Buckets of Inference Innovation

The host unpacks the physical hardware constraints governing LLM execution: - **Arithmetic Intensity:** Defined as the ratio of compute operations (FLOPs) performed per byte of data transferred from DRAM to fast on-chip SRAM. While diffusion models and recurrent architectures exhibit very high arithmetic intensity (because a single noisy latent or tensor can be repeatedly transformed through the same weights without round-tripping to DRAM), standard autoregressive LLMs suffer from notoriously poor arithmetic intensity during decode. - **MFU and MBU:** Model FLOPs Utilization (MFU) measures the proportion of peak theoretical floating-point compute realized at test time (e.g., getting only a fraction of the 2,000 TFLOPs FP8 rating on an NVIDIA H100). Model Bandwidth Utilization (MBU) measures how close data streaming gets to peak theoretical memory bandwidth (such as saturating a 1 TB/s or higher memory bus). - **Roofline Analysis:** Performance is strictly bounded by the minimum of theoretical peak compute versus the product of peak memory bandwidth and arithmetic intensity: $$\text{Attainable Performance} = \min\left(\text{Peak FLOPs/sec},\; \text{Memory Bandwidth (Bytes/sec)} \times \text{Arithmetic Intensity}\right)$$

To systematically categorize how the AI ecosystem overcomes these physical barriers, the host organizes inference innovations into three distinct buckets:

| Innovation Bucket | Core Techniques | Key Mechanism & Impact | |:--- |:--- |:--- | | **Model Innovations** | Mixture of Experts (MoE), State Space Models (SSMs), Multi-head Latent Attention (MLA), Quantization, Distillation | - **MoE:** Routes tokens across $K$ experts, reducing active FLOPs per token by amortizing compute by roughly $1/K$.<br/>- **SSMs:** Avoid full quadratic self-attention across every layer (e.g., placing attention only every 5th or 10th layer).<br/>- **MLA:** Compresses attention states in the sequence domain to minimize KV cache footprint.<br/>- **Quantization & Distillation:** Drops precision down to 2-bit formats or transfers knowledge into compact, efficient student models. | | **Kernel & Runtime Innovations** | RadixAttention, FlashAttention, PagedAttention | Adapts classical operating systems and database techniques to GPU memory management. Tri Dao introduced FlashAttention via GPU SRAM tiling; PagedAttention brings virtual memory paging to avoid KV memory fragmentation; RadixAttention structures KV caches as prefix/radix trees to maximize reuse. | | **System & Hardware Innovations** | Custom silicon, distributed serving engines | Dedicated chip architectures and serving runtimes tailored to eliminate memory-bandwidth bottlenecks during decode. |

[10:16](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=616s) Harness Innovations and Speaker Introductions

The discussion of inference architecture shifts to higher-level harness innovations that sit above the raw model execution to minimize compute and memory overhead. Routing represents the most accessible entry point: rather than routing every request straight to a massive, computationally expensive model, systems can implement tiered waterfalls or conditional tool-calls that escalate to larger models only when strictly necessary, yielding significant FLOP savings. Speculative decoding speeds up inference by running a fast draft model whose proposed tokens are subsequently verified in parallel by the target model. Beyond these patterns, recursive language models offer a way to bypass storing an entire bloated context inside the KV cache, delegating decisions about what information to fetch and retain directly to the model itself. Memory compaction rounds out these strategies as an essential component of modern inference harnesses, highlighted by research such as the "cartridges" concept developed by Sabri Eyuboglu and Dan Simran.

```mermaid flowchart TD accTitle: Inference Harness Architecture Options accDescr: High-level options for optimizing inference requests before and during execution, including routing, recursive models, and speculative decoding.

subgraph Ingest["Request Ingest"] R["User Request"] end

subgraph Strategies["Harness Strategies"] direction TB Route["Routing / Waterfall<br/>(Escalate to larger models only if needed)"] Spec["Speculative Decoding<br/>(Draft model generates, target validates)"] Recur["Recursive Context Management<br/>(LLM selects KV cache contents)"] Compact["Context Compaction<br/>(e.g., Cartridges)"] end

subgraph Compute["Execution"] GPU["Optimized Model Execution"] end

R --> Route R --> Spec R --> Recur R --> Compact Route --> GPU Spec --> GPU Recur --> GPU Compact --> GPU ```

Before handing over the floor, the host highlights two standout educational resources for mastering inference fundamentals: Percy Liang's Lecture 10 on building LLMs from scratch in Stanford's CS 336, and Dwarkesh Patel's in-depth interview with Reiner Pope. The host then introduces the evening's expert panel: Dan Fu (head of research at Together AI, who graduated from the Hazy Research lab), Jeffrey Morgan (founder and CEO of Ollama), Joe Fioti (founder and CEO of Luminal), and Lianmin Zheng (co-creator and second author of the SGLang paper).

- **Recommended Viewing:** Percy Liang's CS 336 (Lecture 10) on LLMs from scratch, and Reiner Pope on the Dwarkesh Patel podcast. - **Featured Speakers:** Dan Fu (Together AI / Hazy Lab), Jeffrey Morgan (Ollama), Joe Fioti (Luminal), and Lianmin Zheng (SGLang).

---

[12:14](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=734s) The Rise of Open Inference and Together AI

Taking the stage, Dan Fu frames inference optimization as an urgent priority fueled by the explosion of open-source artificial intelligence. Together AI recently hit a run rate of 400 trillion tokens per month. To put that scale into perspective, when Google announced Gemini's traffic volume around that same time, Together AI—operating as an agile startup—was already serving roughly one-eighth of Google's entire Gemini traffic volume. On third-party aggregator OpenRouter, open-weight models rapidly captured more than half of all routed tokens, and on at least one single day, Together AI’s token volume eclipsed OpenAI's.

Beyond sheer volume, economics drive the surge toward open models: serving high-quality open-weight alternatives can achieve parity with—or even outperform—proprietary frontier models while slashing serving costs up to tenfold. To support these demands, Together AI operates as an AI-native cloud. Just as AWS evolved beyond raw CPU rentals to offer databases, object storage, and networking APIs, an AI cloud must provide the full platform stack: distributed pre-training, fine-tuning, sandboxes for agent execution, and high-performance inference engines.

> “So you can get nearly the same quality or in many cases actually the same or better quality but up to 10 times cheaper per year.” — Dan Fu

- **Traffic Milestones:** Serving 400 trillion tokens monthly at Together AI, representing roughly 1/8th of Google Gemini's reported volume at the time. - **Market Dynamics:** Over 50% of tokens flowing through OpenRouter are directed to open models, driven by 10x cost reductions. - **Company Profile:** Together AI fields a research team of approximately 80 researchers and PhDs optimizing GPU systems and workloads.

---

[15:30](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=930s) Lifecycle of an Inference Request: Routing and Tokenization

Analyzing production workloads requires tailoring infrastructure to real-world usage patterns. In interactive coding scenarios—such as Cursor, Claude Code, or open-source coding agents—the operational constraints depend on user turn latency, cache hit ratios, and generation velocity, all bound by strict Service Level Agreements (SLAs). Dan Fu walks through the end-to-end lifecycle that occurs the moment an API call arrives:

```mermaid flowchart LR accTitle: End-to-End Inference Request Lifecycle accDescr: Flowchart tracing an inference request from API arrival through routing, tokenization, batch scheduling, GPU prefill/decode, and response output.

subgraph Ingress["Ingress & Preparation"] A["REST API / DNS"] --> B["Cache-Aware Router"] B --> C["Tokenization & Tool Parsing"] end

subgraph Engine["GPU Engine Core"] C --> D["Batch Scheduler Loop"] D --> E["Prefill Phase"] E --> F["Decode Phase"] end

subgraph Output["Egress"] F --> G["Detokenization"] G --> H["Client Response"] end ```

When an HTTP payload clears DNS and Cloudflare, it lands in a cluster fronting thousands of GPUs. Traditional web services rely on naive round-robin or simple server-load balancing, but LLM serving requires cache-aware routing. Because a coding query frequently sends an entire repository context, dispatching a follow-up request to the exact GPU replica that already retains the prompt's KV cache eliminates the need to re-compute identical tokens.

Once routed to a host, the request undergoes tokenization. While conceptually simple, tokenization carries hidden complexity around function calling and tool use. Neural networks are purely numerical pipelines consisting of matrix multiplications and attention mechanisms; they cannot inherently query APIs, execute bash commands, or grep files. The engine must structure user-provided schemas into rigid formats, ensure generation guarantees so the model can invoke external commands correctly, and manage the system prompts without breaking parsing logic.

> “In this language modeling era, we have something else that we need to account for, which is if you send the same request to the same GPU over and over again, you actually get better performance characteristics because you already processed... that first request you send in your whole code base, it processes it all, figures everything out.” — Dan Fu

- **Cache-Aware Routing:** Load balancers must balance overall GPU saturation against KV cache residency to avoid redundant prefill passes. - **Tool-Call Tokenization:** Models lack autonomous system capabilities; tool definitions must be injected into the token stream, requiring strict schema handling during preprocessing. - **The Core Scheduler:** Incoming parsed requests queue up inside an engine scheduler, balancing concurrency against SLA latency thresholds.

[20:02](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=1202s) Prefill vs. Decode: Two Fundamentally Different Workloads

When an inference system optimizes for snappy responses, the scheduling trade-offs become immediately obvious to end users. A familiar example occurs when using Google Search with Gemini integration: the very first snippet appears almost instantly, but subsequent conversational turns feel noticeably slower. This disparity indicates distinct scheduling requests and potentially entirely different backend models handling different phases of generation.

```mermaid flowchart TD accTitle: Prefill and Decode Processing Stages accDescr: Diagram illustrating the two contrasting stages of LLM inference: compute-bound prefill and memory-bound decode.

subgraph PrefillStage["Prefill Stage (Prompt Processing)"] A["Full Input Context / Codebase"] --> B["Compute-Heavy MatMuls & Attention"] B --> C["High Arithmetic Intensity & High MFU"] end

subgraph DecodeStage["Decode Stage (Token Generation)"] D["Autoregressive Generation (1 Token at a Time)"] --> E["Load All Model Weights from HBM to Compute"] E --> F["Memory-Bound & Low Arithmetic Intensity"] end

C -.->|"Disaggregated Architecture"| D ```

At the core of the engine, the workload splits into two phases: prefill and decode: - **Prefill (Prompt Processing):** The system digests the entire prompt—such as a large codebase—to comprehend the context and determine what is being asked. This phase is heavily compute-bound, resulting in very high Model FLOPs Utilization (MFU) and high arithmetic intensity on the GPU. - **Decode (Token Generation):** The model produces text sequentially, generating one or a few tokens at a time. Each generated token forces the hardware to stream the entire model's weights from GPU High Bandwidth Memory (HBM) into the compute units. Because of this constant reloading, decode is strictly memory-bound.

The final stage converts raw binary token IDs back into human-readable text and streams them to the user. While translating tokens to text is relatively straightforward, substantial operational complexity lies in managing the Key-Value (KV) caches—specifically, determining how long to keep cached tokens resident in memory and precisely when to evict or free them.

---

[21:51](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=1311s) Six Levers for Inference Optimization

Engineers have roughly six primary mechanisms to accelerate inference serving:

1. **Kernels:** GPU-level programs that execute the core operations. Optimizing these routines involves writing specialized kernels for attention, Mixture of Experts (MoE), or compound "mega kernels." 2. **Disaggregation Choices:** Because prefill (compute-bound) and decode (memory-bound) represent opposing workloads, placing them onto separate, specialized GPU pools significantly impacts end-to-end throughput and user experience. 3. **Parallelism:** Frontier open models cannot fit into the memory of a single accelerator. For example, a 3-trillion-parameter model served in FP4 requires 0.5 bytes per parameter, amounting to 1.5 TB of memory. With top-tier hardware like the B200 or B300 offering only around 280 GB of HBM, the model must be partitioned across multiple GPUs. 4. **Routing:** Directing queries once they enter the data center to maximize cache hits and balance load across available clusters. 5. **Speculative Decoding:** A model-layer optimization based on the realization that a massive model is unnecessary for trivial completions (such as evaluating "1 + 1"). A smaller, faster model runs concurrently to predict tokens, allowing the primary model to verify multiple tokens at once for free when the draft model is correct. 6. **Quantization:** Reducing the bit precision required to store each weight and activation without degrading model quality.

---

[23:47](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=1427s) Custom Kernels and the "Together Compiler" Agent Loop

Kernels run the core mathematical operations during both prefill and decode, including matrix multiplications (GEMMs) and attention mechanisms. While major matrix operations can be optimized straightforwardly, models are littered with auxiliary layers—such as LayerNorm, RMSNorm, and various softmax layers—that have very low arithmetic intensity.

If these minor operations are executed in isolation with dedicated kernels, their hardware utilization plunges. Dan Fu explains why this mismatch hurts performance: > "Each of these individual operations has very low arithmetic intensity. So if you write a kernel for it by itself, you get a very low flop utilization on on those kernels. So that's historically one of the reasons why you can only get 50% end-to-end flop utilization during training, 5% during decode."

Traditionally, maximizing performance meant tasking teams of elite kernel engineers with inspecting the model's forward pass line by line and hand-fusing operations. Together has moved toward automating this tedious workflow using autonomous software agents under an internal effort dubbed the "Together Compiler." Dan Fu wryly admitted that the name is temporary: > "I had a shocking realization the other day, which is that Together Compiler is not actually a compiler."

The system pairs a target model, target hardware, and real production traffic, letting agents iteratively optimize the execution inner loop. Incorporating production traffic profiles is critical because real-world workloads—such as large code-completion prompts versus brief conversational chat queries—exercise GPU kernels in fundamentally divergent patterns. By observing what actually executes in production, the system creates a functioning recursive self-improvement (RSI) loop for inference optimization.

---

[26:12](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=1572s) Parallelism Schemes and the 72-GPU Scale-Up Domain

Splitting a model across multiple GPUs unlocks aggregated compute, increased memory capacity, and combined memory bandwidth. System designers rely on several foundational parallelism strategies: - **Tensor Parallelism (TP):** Slices individual matrix multiplications across GPUs and merges intermediate results via map-reduce-style collectives. - **Pipeline Parallelism (PP):** Partitions consecutive layers of the network across different GPUs. - **Data Parallelism (DP):** Replicates the model across devices and routes distinct sequence requests to each copy.

```mermaid flowchart LR accTitle: Evolution of High-Bandwidth GPU Domains accDescr: Flow comparing traditional 8-GPU high-bandwidth scale-up domains with modern 72-GPU NVLink rack-scale domains.

subgraph LegacyDomain["Traditional Architecture"] direction TB G8["8 GPUs via High-Speed Interconnect"] G8 -->|"Split >8 Ways"| SlowPath["Slower Inter-Node Network (Latency Bottleneck)"] end

subgraph ModernDomain["Modern NVLink Rack Architecture"] direction TB G72["72-GPU Scale-Up Domain (GB200/NVL72)"] G72 --> MoE["Assign Single MoE Experts per GPU"] MoE --> LargeGEMM["Tokens Aggregate into Dense Matrix Multiplies"] end ```

The primary engineering constraints have always been GPU memory capacity, interconnect bandwidth, and kernel efficiency. Recent advances—such as the DeepSeek serving benchmarks reaching roughly $38 per hour—are driven by massive scale-up domains. Previously, GPU servers capped high-speed interconnects (such as NVLink) at eight GPUs per node. Distributing a model beyond eight devices forced communication across slower inter-node networking, adding crippling latency.

With newer rack architectures (such as Nvidia's GB200 NVL72), up to 72 GPUs communicate over a unified, high-bandwidth fabric. This allows extreme expert parallelism for Mixture of Experts (MoE) architectures, where individual experts are placed onto dedicated GPUs across the rack. When a single GPU hosts only one expert while handling an enormous aggregate stream of tokens, the computation transforms from a memory-bound lookup into a dense matrix multiplication that GPUs execute with peak efficiency.

However, scaling up an NVLink cluster introduces serious operational fragility. A single 72-GPU model replica will fail if even one accelerator crashes. Dan Fu emphasizes the necessity of automated fault tolerance: > "And if you have a single model replica that depends on 72 GPUs and one of them goes out, if you don't have the right remediation or and reliability pieces behind it, then what that means is a single GPU can going out can take out 72 GPUs worth of of worth of compute, which is not something that you want."

[29:29](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=1769s) Quantization and Rigorous Quality Red-Teaming

Addressing the final piece of the inference lifecycle, Dan Fu highlights quantization as a deceptively simple yet treacherous engineering problem. On paper, it seems straightforward—hundreds or thousands of papers exist outlining different quantization algorithms, allowing engineers to pick a favorite off the shelf. In practice, however, quantization directly impacts the user experience and the downstream behavior of the model. Large language models are typically trained in BF16 (16 bits per parameter, or 2 bytes per parameter), meaning a 3-trillion-parameter model occupies 6 terabytes of memory. Because machine learning computations are inherently fuzzy rather than exact, models do not strictly need that full precision to produce accurate outputs. While compressing parameters down to 8 bits or 4 bits preserves significant quality while boosting hardware throughput, maintaining end-to-end task performance requires deep, obsessive verification.

To prevent performance degradation, Together AI instituted a comprehensive measurement, quality, and red-teaming loop. Before any model variant goes live to external users, it is dog-fooded extensively in-house. This testing catches subtle failures, such as broken tool-calling syntax that disables a Claude Code harness from searching the internet, or benchmark drops indicating a general loss of model intelligence. Among their continuous automated evaluations is an agentic coding loop centered on Flappy Bird. If an aggressive quantization scheme corrupts the model's logic, the model fails to generate functional game code, alerting engineers immediately:

> “So, uh if if something went wrong with the quantization, suddenly the Flappy Bird won't flap anymore and then we have to go and and figure out what went wrong and fix it.”

```mermaid flowchart TD accTitle: Quantization Evaluation and Red-Teaming Loop accDescr: Flowchart illustrating the quality verification pipeline from unquantized BF16 weights down to deployment through continuous red-teaming and agent benchmarks.

subgraph Compression["Weight Compression"] A["BF16 Model Weights<br/>(2 bytes/param)"] --> B["Quantization Engine<br/>(FP8 / NVFP4)"] end

subgraph Verification["In-House Quality Red-Teaming"] B --> C["Layer Activation Checks<br/>(KL Divergence)"] B --> D["Standard Benchmarks<br/>(AIME, GSM8K, MMLU)"] B --> E["Agentic Loops<br/>(Flappy Bird Code Generation)"] B --> F["Kernel Self-Optimization<br/>(Together Compiler Project)"] end

subgraph Production["Serving Layer"] C & D & E & F --> G{"Quality Drop?"} G -->|"Yes"| H["Tune Quantization Recipe"] H --> B G -->|"No"| I["Deploy to High-Throughput Inference"] end ```

[32:35](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=1955s) The MiniMax Success Story and Closing Summary

Dan illustrates Together AI’s quantization expertise through their deep partnership with open-source model builders, specifically MiniMax on their M3 model and the unannounced M3.1. MiniMax struggled to quantize their model directly into NVIDIA's NVFP4 format using standard open recipes. Because NVFP4 maximizes throughput on modern NVIDIA architectures, Together AI developed custom quantization recipes. A slide graphic showed the baseline performance of the original FP8 model in lilac purple, MiniMax’s degraded attempt in hot pink, and Together AI’s orange bars matching the unquantized FP8 baseline performance while running in 4-bit precision. This allowed MiniMax to serve their models with dramatically higher token efficiency, generate sustainable inference revenue, and reinvest in training future architectures.

Summarizing the entire request lifecycle, Dan recaps the path an inference prompt takes through a modern distributed infrastructure: - **Routing:** Directing requests to the optimal data center and replica pool based on current load and hardware topology. - **Tokenization:** Ensuring formatting matches model expectations precisely to preserve structured outputs and tool calling. - **Queuing, Batching, and Scheduling:** Dynamically packing heterogeneous user requests into efficient computational batches. - **Prefill and Decode Execution:** Splitting computational phases across dedicated GPU clusters with distinct kernel tuning and parallelism schemes. - **Quantization and Kernel Tuning:** Squeezing maximum throughput out of modern accelerators while rigorously preserving benchmark scores.

Concluding his presentation, Dan notes that as enterprise CFOs push back on massive proprietary API bills, open models backed by performant inference stacks offer a viable alternative.

> “If your CFO is beating uh, breathing down your neck and saying, 'Hey, uh, we're spending too much money on cloud code, uh, and maybe we should try some of these open models.' Um, one of these is is a sales QR code, um, and then if you if you find this stuff really exciting, you want to come work with us, uh, the other one goes to our careers page.”

---

[35:28](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=2128s) Q&A: Quality Metrics, Self-Optimizing Compilers, and Rack Interconnects

During the audience Q&A, an attendee asks about the specific metrics used to evaluate NVFP4 quantization quality compared to MXFP or FP8, and how interconnect technology allows Grace Blackwell architectures to scale across clusters larger than eight GPUs. Dan explains that standard public benchmarks—like AIME, which is already saturated past 90%—act primarily as tripwires. A movement of just one percentage point signals that the quantized model will fail on complex, real-world tasks. Beyond standard benchmarks, Together AI inspects layer-by-layer activation shifts using KL divergence, runs hundreds of real-world regression cases, and executes self-referential agentic tests where the model attempts to optimize its own kernels inside the Together compiler project:

> “We actually put it put it back into the together to together compiler project so that the model is optimizing its own kernels. And if we see that the model for some reason can't optimize its kernels anymore, then we're like, 'Hmm, okay, we we should take a look at that.'”

Addressing the hardware interconnect question, Dan jokes that his closest hands-on hardware experience was buying a Hugging Face micro-duct. From an infrastructure perspective, modern NVLink racks connect trays of four GPUs using massive, thick copper trunk cables routed through a central NVLink switch chassis rather than simple point-to-point PCB traces. This switch fabric extends low-latency all-reduce operations across 72 GPUs in an NVL72 rack, eliminating the traditional eight-GPU node bottleneck. The session closes as another audience member steps up to ask how Together AI's custom NVFP4 recipe achieved its performance without extensive end-to-end pre-training.

[39:03](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=2343s) Training-Free Quantization and Scale Search

Addressing whether their quantization pipeline simply transforms weight values without fine-tuning or post-training, Dan Fu clarifies that their process is strictly training-free. While many quantization recipes require retraining the model, doing so without the full original training dataset introduces severe catastrophic forgetting. Unless an engineering team is the original model provider, retraining on a subset of data usually degrades model capability.

Instead, their pipeline blends calibration data with techniques like scale search—a method Dan’s group published—to select optimal scales. Because dozens of training-free quantization methods circulate in the academic literature, success depends on meticulously picking the right combination of techniques and being disciplined about empirical validation. As Dan explains, “It’s a mixture of being very careful with your recipes and then also trying to and just being very careful and putting together all the different pieces.”

- Retraining without the full pretraining corpus causes models to forget core capabilities. - The pipeline relies on curated calibration datasets and the published "scale search" framework. - Quality preservation depends on disciplined evaluation across hundreds of candidate training-free methods.

---

[40:17](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=2417s) Next-Generation Architectures: SSMs, Speculative DLLMs, and Looped Models

When asked which emerging architectures might ease the computational burden of token decoding, Dan Fu points out that non-transformer architectures have already quietly permeated production models despite past skepticism. During his PhD work on state space models (SSMs), the common consensus was that alternative architectures were dead ends and that standard transformers would dominate unconditionally. Today, production systems rely heavily on hybrid approaches: Kimi K3 uses KDA (a state-space model/linear attention style layer), production speculative decoders leverage diffusion language models (DLLMs), and sparse attention is widely deployed.

Dan highlights looped models as an especially compelling paradigm recently explored by his students. By routing a token repeatedly through the same computation layers, engineers can scale training FLOPs without expanding parameter counts or memory footprints. Although his team initially worried looped models might be considered too exotic to be adopted, frontier labs had already implemented similar mechanisms:

> “We thought, ‘Oh, this is a weird idea. Nobody’s going to use it.’ And then it turns out they at least that OpenAI and Anthropic were both using it already. Um so yeah, it’s a really exciting time where like if you put out a good idea right now about architectures and go through the leg work of showing that it works, it’ll show up in GPT-6 or GPT-7.”

```mermaid flowchart TD accTitle: Evolution and Adoption of Alternative Model Architectures accDescr: Diagram tracing PhD skepticism of alternative architectures to modern frontier deployments including SSMs, DLLMs, and Looped Models.

subgraph Skepticism["Early Research Consensus"] A["SSMs & Non-Transformers"] -->|"Viewed as dead ends"| B["Transformer Monoculture"] end

subgraph Production["Modern Production Deployments"] C["Kimi K3 (KDA)"] -->|"SSM / Linear Attention"| F["Frontier LLMs"] D["Speculative Decoding"] -->|"DLLMs (Diffusion LMs)"| F E["Looped Models"] -->|"Reuse weights across loops"| F end

B -.->|"Vindicated in modern architectures"| Production ```

---

[42:18](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=2538s) Jeff Morgan on Ollama and the Discrepancy Across Open Model Tokens

Jeff Morgan, co-founder and CEO of Ollama, opens his session by addressing the widening variance across open model providers. While earlier overviews cataloged dozens of major inference hosts, Morgan anticipates an ecosystem containing hundreds or thousands of specialized inference vendors. For builders constructing applications on open weights, understanding and benchmarking output quality across these disparate backends is becoming a critical operational hurdle.

Morgan recaps Ollama’s journey from a 2023 open-source project designed for local execution on developer laptops, standalone H100s, and small local servers to an enterprise infrastructure player. In early 2024, open models experienced a massive leap in parameter counts, with releases like Kimi, GLM, and MiniMax reaching 500 billion to 1 trillion parameters. Because running models of this scale exceeds consumer hardware, Ollama launched its cloud service to replicate the seamless local developer experience in hosted environments, rapidly scaling into one of the largest independent cloud platforms serving open models across the United States and Europe.

- Ollama originated in 2023 to provide local, single-GPU execution with an OpenAI- and Anthropic-level developer experience. - The shift toward 500B–1T parameter models necessitated Ollama Cloud to maintain accessible developer ergonomics. - Ollama is now widely adopted across the Fortune 500, with enterprise teams steadily migrating from closed proprietary APIs to open-weight models for cost and control.

---

[45:11](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=2711s) The Hazards of Inference Routing Across Diverse Backends

Routing traffic across an array of inference vendors promises the best of all worlds in terms of availability, cost, speed, and latency. However, Morgan warns that the underlying infrastructure varies radically between providers. Some providers deploy clusters of high-end enterprise accelerators, while others run unconventional setups such as distributed farms of Mac minis. When aggregating multiple backends under a common routing layer, an application’s user experience often degrades to the lowest common denominator:

> “The challenge with routing across a bunch of providers is obviously that, um, one bad apple can ruin the bunch. And that if you have a single provider that’s underperforming on one of these four, um, criteria, that you, um, you know, you might have a customer experience that’s kind of the least common denominator.”

These quality variances stem from numerous hidden discrepancies across the serving stack. A provider might aggressively over-quantize weights, employ lower-precision KV cache formats, use differing sampling algorithms, or run faulty kernels that corrupt outputs on long-context tasks. Tool calling represents an especially fragile vector: small mismatches in schema enforcement, tokenizer-level parsing errors, or clumsy error recovery can completely break an agentic workflow. In extreme cases, providers have even mislabeled models, serving a completely different architecture than requested while returning standard conversational greetings that conceal the underlying substitution.

| Evaluation Dimension | Common Provider Failure Modes | |:--- |:--- | | **Intelligence & Accuracy** | Aggressive over-quantization, unvetted weight versions, custom kernel precision bugs. | | **Tool Calling Reliability** | Broken JSON schema enforcement, tokenizer parsing faults, poor error recovery. | | **Context Length Fidelity** | Degradation in custom attention kernels and KV cache compression at long context windows. | | **Speed & Throughput** | Highly divergent time-to-first-token (TTFT) and decode throughput across heterogeneous hardware. | | **Metadata Integrity** | Model spoofing, outdated model checkpoints, or serving an entirely different assistant backend. |

[48:08](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=2888s) The Open Provider Quality Gap and the Compounding Agent Problem

The AI industry constantly debates the capability gap separating proprietary frontier models from open-weights models, but Jeff Morgan argues that practitioners are overlooking an even larger variance. When deploying open weights across commercial API hosts, the quality divergence from host to host frequently eclipses the difference between open and closed systems entirely. As Jeff Morgan explains, “It ends up that for open model serving, just open models, that the gap between the worst provider for a model and the best provider is actually bigger than the gap between closed and open models.” Users coming from the ecosystem built by OpenAI and Anthropic take uniform, deterministic endpoint behavior for granted. With open models, however, pointing an application at an unvetted endpoint can degrade performance dramatically. Morgan notes that this divergence is well documented in academic literature: a 2024 study evaluating 31 distinct Llama serving endpoints found that certain sub-par deployments degraded to the point of performing closer to GPT-3.5 than to a reference implementation of that very same Llama model.

```mermaid flowchart TD accTitle: Compounding Error in Agentic Workflows accDescr: Step-by-step propagation showing how a 10 percent per-step deviation compounds into an 80 percent failure rate across a 15-step agent task.

subgraph Horizon ["15-Step Agent Horizon"] S1["Step 1: 10% Deviation Risk"] S2["Step 2: Drift Compounding"] S15["Step 15: Error Cascade"] end

subgraph Impact ["Cumulative Result"] R1["~79% to 80% Final Task Failure Rate"] end

S1 --> S2 S2 -->|"sequential tool calls"| S15 S15 --> R1 ```

This reliability variance poses an existential hurdle for modern agentic and coding workloads. Unlike single-turn completions, agent tasks involve long execution horizons consisting of sequential tool calls and environment feedback loops chained back-to-back. If an endpoint introduces even a slight 10% drift or deviation from the canonical model behavior at each step, the compounding error over a modest 15-step task results in a final failure rate approaching 79% to 80%. When an autonomous agent encounters an 80% error threshold, the entire product experience breaks down.

- A 2024 paper benchmarked 31 distinct commercial endpoints serving the identical Llama model and found severe performance degradation among lower-tier providers. - Compounding per-step inaccuracies turn a manageable 10% single-step token deviation into a catastrophic ~79% failure rate on a standard 15-step task. - Providers like Together AI work closely with downstream teams to preserve reference model fidelity, contrasting sharply with cut-rate or misconfigured host endpoints.

---

[50:50](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=3050s) Methods for Monitoring Inference Endpoints: Evals, Distributions, and Logprobs

Monitoring third-party model endpoints by manually inspecting output streams is physically impossible when ingestion runs at 1.6 million tokens per GPU during preprocessing. Instead, infrastructure teams employ several automated paradigms to detect degradation and quantization defects. The most straightforward approach is brute-force benchmarking. Organizations like Artificial Analysis repeatedly execute standardized evaluation suites against competing providers to publish comparative scoreboards. While benchmark suites successfully flag improper quantization schemes and faulty model configurations, they are prohibitively slow and expensive to run on continuous live traffic. In response to tool-calling failures, model creators have also begun issuing targeted testing harnesses—such as the tool-call verification suites released by Moonshot to force third-party hosts into compliance through public ranking boards.

To monitor traffic more efficiently without running entire evaluation suites, routers utilize distribution testing and token probability tracking. Distribution tests directly compare the output probability distribution of a hosted provider's API against a local reference weight instance. While this costs roughly a dollar per test and exposes divergences cheaply, it requires operating an expensive dedicated reference model and cannot always diagnose the precise architectural root cause of a failure. An even lighter-weight method inspects `logprobs` (log probabilities) over single next-token generations against a fixed historical baseline to measure endpoint drift. As Jeff Morgan points out, “And, you know, they were able to kind of identify these silent changes to the serving stack, which is really interesting if you think about it.” This drift detection exposes unannounced infrastructure modifications—such as an operator updating their backend engine to a new release of vLLM or SGLang, or migrating workloads across cluster hardware—though it remains limited by the fact that many modern inference APIs omit logprob outputs entirely.

- **Brute-Force Benchmarking:** Catches flawed quantization and setup bugs via standard eval suites, but is too computationally expensive for continuous runtime telemetry. - **Verification Harnesses:** Model creators publish targeted test suites (such as tool-calling scoreboards) to publicly benchmark provider compliance. - **Distribution Testing:** Evaluates token probability divergence between a provider API and reference weights for around $1 per run, though it still demands running a baseline instance. - **Logprob Drift Tracking:** Measures single next-token probabilities against prior baselines to catch silent backend updates, though it cannot evaluate baseline correctness if the initial setup was already flawed.

---

[54:38](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=3278s) Request-Shape Introspection and Labs Shipping Full Inference Stacks

When privacy or cost constraints prevent saving token payloads or running parallel reference models, routers resort to non-retentive request-shape introspection. Systems like OpenRouter (using their Exacto routing approaches) continuously evaluate metadata signals: the percentage of tool calls passing without schema violations, sudden drops or spikes in generation length, and API error spikes. In practice, broken model servers—such as faulty DeepSeek deployments on vLLM—often manifest structural bugs where the generation degenerates into an infinite loop. Telemetry flags these anomalies purely by observing abnormal output sequence lengths without inspecting sensitive user text.

```mermaid flowchart LR accTitle: Structural Request Inspection Pipeline accDescr: Telemetry monitoring metadata signals like tool call syntax, generation lengths, and loops to score inference endpoints without storing payload data.

subgraph Signals ["Metadata Telemetry"] A["Tool Call Pass Rate"] B["Output Token Length"] C["Infinite Loop Detection"] end

subgraph Decision ["Routing & Scoring"] D["Running Provider Health Score"] E["Dynamic Router Re-route"] end

A --> D B --> D C --> D D --> E ```

Achieving the high reliability of proprietary systems requires open-source teams to emulate the obsessive quality control seen at closed labs. Morgan recalls an incident where a microscopic bug inside an Anthropic TPU kernel triggered an internal retrospective and an operational outage response purely over subtle degradation in generation quality. To bring that level of consistency to open models, model creators must stop dumping bare weights onto Hugging Face or Ollama accompanied only by toy example scripts authored by external applied teams. Instead, upstream labs are beginning to provide fully verified inference software. When DeepSeek launched DeepInfer, they open-sourced their complete optimized serving stack so operators could run weights on their own hardware at target efficiencies like $38 per GPU hour with guaranteed mathematical parity. Similarly, model creators now collaborate with optimization specialists prior to release to pre-quantize and validate official weights, ensuring the models hosted across external commodity APIs produce outputs identical to the creator's own reference servers.

- Frontier labs treat subtle algorithmic or kernel-level drift with the seriousness of complete site outages; an Anthropic TPU kernel defect prompted an entire service retrospective purely due to a drop in response quality. - Historically, open model weights were dumped online with basic example scripts generated by separate teams, leaving serving stack optimization entirely to provider guesswork. - New initiatives like DeepSeek's DeepInfer package the model weights alongside the exact production inference stack to deliver consistent serving quality and reliable economics. - Labs increasingly pre-quantize models alongside expert compiler engineers before public launch to eliminate fidelity losses between upstream releases and downstream hosted endpoints.

[57:05](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=3425s) Verifying Model Quality and Hardware Attestation

Jeff Morgan concludes the main portion of his presentation by discussing the shifting landscape of model quantization and testing. Historically, quantization was performed post-release by third parties who lacked access to original training data and merely sought to squeeze models onto consumer hardware like MacBooks. Increasingly, AI frontier labs are publishing their own official quantized model releases. Beyond official weights, developers can access models through lower-level APIs, pulling critical logic like tool calling, reasoning extraction, and thinking phases up into the evaluation harness itself to guarantee uniform accuracy across multiple hosting providers rather than delegating it to third-party endpoints.

Looking ahead, Morgan highlights commercial and architectural opportunities to guarantee inference fidelity: - **Startup Opportunities in Quality Monitoring:** A major opening exists for startups to continuously monitor LLM quality, which becomes essential as models are integrated into critical application loops. - **Hardware-Level Attestation:** As detailed in recent research, modern Nvidia GPUs feature a trusted execution environment capable of cryptographically attesting to both the model weights loaded into memory and the software stack serving them, creating a hardware fingerprint that proves serving fidelity.

Jeff Morgan wraps up his presentation by reflecting on the broader ecosystem: “I think there's a big opportunity for startups to start monitoring LLM quality.”

```mermaid flowchart TD accTitle: LLM Serving Quality and Integrity Assurance accDescr: Diagram illustrating client-side extraction, quality monitoring, and hardware attestation.

subgraph Client ["Client Harness"] A["Prompt / Task"] --> B["Harness-Level Logic"] B --> C["Tool Calling & Reasoning Extraction"] end

subgraph Evaluation ["Quality Assurance Loop"] C --> D["Baseline Single-Token Logprob Probes"] D --> E["Bespoke Private Customer Evals"] end

subgraph Hardware ["Inference Execution"] E --> F["Hardware-Level Attestation"] F --> G["Nvidia TEE Weight & Software Fingerprint"] end ```

[58:35](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=3515s) Q&A: Pre-fill Intervals and Preventing Benchmark Gaming

An audience member asks about balancing pre-fill admission interval criteria against decode steps (such as 48 or 64 steps) for multi-turn agent workloads over long sessions. Jeff clarifies that his lightweight quality probing technique relies on issuing a single-token prompt and evaluating the returned logprobs against a known baseline over time, rather than tracking specific token-admission schedules. While context lengths vary, extending the prompt systematically reveals whether logprob distributions diverge as contexts push past 200,000 tokens—a threshold where certain inference providers arbitrarily cap windows or experience next-token degradation relative to reference models.

Another audience member raises the risk of provider manipulation: if inference providers face public open-source quality benchmarks, they might simply hill-climb or train on those specific tests to cut costs without preserving genuine output quality. Jeff explains that the best defense is designing bespoke, private evaluation suites using actual customer traffic data:

> “And that's really hard for a a single provider to game because it's not public. It's specific to that customer's data or their traffic.”

[01:01:45](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=3705s) Q&A: Local Inference and the Hybrid Cloud-Edge Future

The discussion shifts to running models locally versus in the cloud, noting recent developments like Richard Socher's team at Recursive launching the Flint model, alongside local runtimes such as MLX on Apple silicon, Ollama, and Calibri (which streams Qwen MoE architectures directly from disk). Jeff highlights that local execution faces the identical verification challenges as cloud endpoints: whether deploying onto desk-side hardware like an Nvidia DGX, an Apple M5 Ultra, or a private data center, users must navigate up to 20 different quantization variants available on Hugging Face or Ollama.

Addressing where Ollama and client architectures go next, Jeff forecasts an era of unified local-to-cloud request routing:

> “Now, with these ultra-capable local models that can run on, you know, this MacBook, we're going to see, you know, this routing problem go from just being from a bunch of providers on the cloud to like local and cloud routing.”

[01:04:30](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=3870s) Joe Fioti on Compilers as the Hardware-to-Token Engine

Joe Fioti, co-founder and CEO of Luminal, opens the next session by framing the economics of model compilation. A compiler serves as the direct conversion function transforming raw GPU FLOPs or GPU hours into revenue-generating tokens, meaning compiler quality directly dictates whether an inference provider operates at a loss or achieves massive profitability.

To solve the complex mapping of models onto diverse, heterogeneous silicon, Luminal rejects standard heuristics in favor of equality saturation—a search-driven rewrite methodology: - **Search vs. Heuristics:** Traditional compilers rely on rigid optimization heuristics; Luminal leverages equality saturation to explore vast rewrite spaces for hardware optimization. - **Industrial Scale:** Luminal stands as the largest industrial implementation of equality saturation applied to machine learning inference. - **The Ecosystem Moat:** Joe emphasizes that Nvidia’s multi-trillion-dollar market capitalization largely stems from the CUDA ecosystem, underscoring that mapping modern deep learning graph representations down to hardware remains one of the hardest software challenges in modern computing.

Joe Fioti captures the financial stakes of low-level optimization: “And how good of a compiler you have greatly determines whether you lose money, whether you make money, or whether you make a lot of money.”

[01:05:37](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=3937s) Leveraging The Bitter Lesson: Compile-Time Search Over Heuristics

Joe Fioti explains that traditional compiler development has reached an impasse when it comes to supporting modern machine learning workloads across evolving hardware architectures. Instead of relying on human engineers to maintain brittle, hand-crafted heuristics, complex rewrite rules, or sprawling kernel libraries that must be manually updated every time a new neural network architecture drops, their compiler embraces Rich Sutton’s "Bitter Lesson" by substituting human cleverness with massive compile-time search. By relying on search, an engineering team of just eight people is able to maintain a compact, shared compiler core capable of targeting multiple data center inference ASICs—spanning Nvidia and AMD GPUs, Google TPUs, and AWS Trainium chips—while serving approximately 80 to 90 trillion tokens per month.

The overall approach departs significantly from conventional MLIR-based compiler stacks. Rather than trying to out-engineer hardware diversity with thousands of specialized lowering passes, the system pairs an equality saturation engine called Egglog with a clean intermediate representation (IR) structure designed to scale across heterogeneous backends and fully saturate target compute engines.

- The compiler replaces hand-tuned heuristics with exhaustive compile-time search to discover optimal execution schedules. - Hardware backends supported include Nvidia GPUs, AMD GPUs, Google TPUs, AWS Trainium, and specialized inference ASICs. - A lean team of eight engineers operates a stack running between 80 and 90 trillion tokens monthly.

> “And our solution is to leverage the bitter lesson. So, we do a lot of compile time search.” — Joe Fioti

---

[01:07:15](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=4035s) Reducing PyTorch Complexity to 15 Primitive Operations

When examining modern deep learning workloads—whether large language models, diffusion models, or biological models—Fioti notes that their underlying computations are fundamentally straightforward: they are almost entirely dense linear algebra consisting of a surprisingly small palette of distinct operations. While frameworks like PyTorch expose around 2,000 distinct API operations, Fioti’s compiler strips away this superficial complexity by decomposing every PyTorch operation into a tiny RISC-style primitive operation set containing only 14 or 15 atomic ops, such as addition, multiplication, and sum reductions.

Crucially, the intermediate representation does not even feature a native matrix multiplication (`matmul`) operator. Instead, a matrix multiply is lowered into its fundamental mathematical components: an outer broadcasted element-wise multiply followed by a sum reduction. By creating this narrow intermediate bottleneck, the compiler eliminates the need for an intractable all-to-all mapping between thousands of PyTorch functions and hundreds of vendor-specific hardware instruction sets. If an incoming model can be expressed in PyTorch, it can be mapped into this minimal primitive set without any loss of mathematical expressiveness.

```mermaid flowchart TD accTitle: PyTorch to Hardware Lowering Pipeline accDescr: Diagram illustrating the decomposition of PyTorch operations into primitive ops, expansion in an e-graph, and execution on hardware. subgraph Frontend ["Frontend Ingestion"] A["PyTorch Graph (~2,000 Ops)"] -->|"Decompose"| B["Primitive RISC Op Set (14-15 Ops)"] end

subgraph Optimization ["Equality Saturation"] B -->|"Apply Rewrite Rules"| C["Egglog E-Graph (Saturated Space)"] C -->|"Extract Candidates"| D["Concrete Program Executables"] end

subgraph Hardware ["Execution & Profiling"] D --> E["Run on Hardware (GPU / TPU / Trainium)"] E -->|"Profile Runtime"| F["Select Minimum Execution Time"] end ```

- PyTorch's library of approximately 2,000 functions maps completely down into 14 to 15 primitive ops (e.g., `add`, `multiply`, `sum`). - High-level operations like `matmul` are synthesized from a broadcasted multiply and a sum reduction. - This RISC-style bottleneck insulates the compiler core from changes in high-level model definitions and downstream hardware targets.

> “And so, you know, we're not Turing complete, but I like to say we're PyTorch complete, meaning if you can represent your model in PyTorch, we can ingest it and compile it uh to all of our supported backends.” — Joe Fioti

---

[01:08:50](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=4130s) Superoptimization and Equality Saturation via E-Graphs

Fioti connects this design philosophy to the academic concept of "superoptimization" introduced in the 1990s. Classical superoptimization systematically searches through all valid permutations of a program to locate the optimal sequence according to a cost metric. While historically confined to niche research because of state-space explosions, it solves the primary scaling dilemma that burdens industrial compiler teams: as the combinations of model architectures, data types (FP32, FP16, BF16, FP8, INT4), and target hardware backends multiply, the engineering headcount required to manually author lowering passes explodes into 500 to 1,000 engineers without ever achieving optimal code generation.

Instead of writing monolithic, destructive lowering passes, developers write small, atomic equality rewrite rules specifying mathematical equivalencies. In a traditional compiler pass, matching a subgraph pattern replaces it in-place, permanently destroying the original representation; if that heuristic guess proves suboptimal downstream, the compiler cannot backtrack. Fioti’s stack utilizes an E-graph (equivalence graph) engine powered by Egglog. E-graphs execute non-destructive rewrites: applying a rule attaches the equivalent right-hand side representation directly into the data structure alongside the original left-hand side, retaining all historical and newly generated representations simultaneously.

- Conventional compiler team headcounts scale exponentially with hardware and model targets, often reaching 500 to 1,000 engineers. - Destructive graph rewrites overwrite patterns permanently, forcing compilers to rely on rigid heuristics with no recourse if choices degrade performance. - E-graphs store both the input pattern and transformed variants simultaneously in equivalence classes without discarding information.

> “Because there is an exponentially, you know, large number of possible ways you can do that lowering.” — Joe Fioti

---

[01:11:47](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=4307s) How E-Graphs Work: From Expression Simplification to Hardware Extraction

Fioti illustrates the mechanics of an E-graph using a simple arithmetic expression: `(A * 2) / 2`. In an E-graph representation, the expression tree begins with a division root pointing to two sub-branches: the numerator `(A * 2)` and the denominator `2`. The compiler iteratively fires rewrite rules in every applicable permutation. One rule may match `A * 2` and add an equivalent bit-shift representation `A << 1`. Another algebraic rule recognizes the division of a multiplied factor by itself.

Rather than picking one path over another, the E-graph accumulates these equivalent forms until no new mathematical information can be deduced—a state known as "saturation." At saturation, the equivalence class boundary enclosing the initial root division expression encompasses the raw variable `A` itself. The compiler has proven equivalence through pure search without an engineer writing a targeted heuristic for that exact pattern. Once an E-graph is fully saturated, the system extracts multiple concrete executable candidate programs, compiles them into machine executables, benchmarks them directly on the target hardware, and selects the variant that yields the lowest measured execution time.

- Expressions expand non-destructively through rewrite rules until reaching fixed-point equality saturation. - Arithmetic simplifications (such as `(A * 2) / 2 = A`) emerge naturally from local rewrites rather than top-down pattern matchers. - Instead of using synthetic cost heuristics alone, the system compiles candidate graphs and times their real execution on the target hardware to select the fastest implementation.

> “And so, we basically apply all of our rewrite rules in every possible order until they can add no more information, and that effectively has grown the e-graph as much as it can go. We call that saturated.” — Joe Fioti

---

[01:13:07](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=4387s) Writing Rewrite Rules in Lisp and Targeting Hardware Backends

Fioti shows what these rewrite rules look like in practice: compact, declarative declarations written in S-expression / Lisp syntax, a convention inherited because the authors of the underlying Egglog engine prefer Lisp. Rules declare fundamental algebraic identities, such as basic commutativity (`x + y = y + x`), alongside transformations like constant folding—evaluating known constants at compile time rather than leaving redundant arithmetic operations in the runtime compute stream. Fioti notes that while commutativity seems trivial on paper, handling it rigorously in numerical computing requires special care due to non-associative and non-commutative edge cases inherent to IEEE 754 floating-point standards.

This shared, primitive intermediate representation forms the universal substrate of the compiler. However, each target processor—whether an Nvidia GPU, an AMD ROCm-compatible accelerator, a Google TPU, or an AWS Trainium—brings its own specific instruction set architecture (ISA) and native matrix/vector compute primitives. Fioti transitions into how the compiler bridges the gap between this hardware-agnostic 15-op IR and the specialized, low-level instructions native to individual inference ASICs.

- Rewrite rules are authored declaratively using Lisp S-expression syntax dictated by the Egglog engine. - Elementary transformations include commutativity axioms and compile-time constant folding. - Floating-point representations present subtle hazards for algebraic transformations because IEEE floating-point arithmetic does not strictly obey real-number associativity. - The platform uses this unified IR to map directly down to the native instructions provided by distinct hardware vendors.

> “And so, this is kind of what our rewrite rules look like. They're very, very straightforward and simple. They're written in a Lisp syntax. The reason for that is because the people who made the engine that we use inside our compiler really like Lisp.” — Joe Fioti

[01:14:39](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=4479s) Hierarchical Intermediate Representation and Mega Kernels

To compile efficiently across target chips, Lumina's compiler treats low-level hardware primitives as operations with specific memory semantics. For example, a Tensor Core matrix multiply-accumulate (MMA) operation reads its inputs from SRAM and writes its outputs back to SRAM. Adding a new hardware target does not require rewriting the whole compiler pipeline; the team simply characterizes the hardware primitives, writes rewrite rules equating them to existing intermediate representation (IR) operations, and lets those equivalences propagate throughout the system.

```mermaid flowchart TD accTitle: Multi-Level Intermediate Representation Hierarchy accDescr: Hierarchy of hardware execution levels captured in Lumina's intermediate representation.

subgraph Cluster_IR ["Lumina Multi-Level IR"] L1["Node Level"] L2["Device / GPU Level"] L3["Thread Block / Streaming Multiprocessor (SM) Level"] L4["Warp Level"] L5["Single Thread Level"] L1 --> L2 L2 --> L3 L3 --> L4 L4 --> L5 end

subgraph Cluster_Libs ["Vendor & Community Libraries"] cuBLAS["cuBLAS (Host / Device Matmul)"] BlockLibs["CUTLASS / CuBLASDx / ThunderKittens (Block Level)"] MMA["Tensor Core MMA (Instruction Level)"] end

cuBLAS -.-> L2 BlockLibs -.-> L3 MMA -.-> L5 ```

Because modern execution environments are strictly hierarchical, Lumina's IR is multi-level by design. It can explicitly express instructions at the single-thread, warp, thread block (a full Streaming Multiprocessor or GPU core), entire GPU, or full node level. By relying on equality saturation, the compiler easily integrates third-party and vendor primitives across these tiers: - Host-invoked GPU-wide operations (e.g., cuBLAS for high-performance matrix multiplications). - Thread-block-level operations that control a single GPU core (e.g., CUTLASS, CuBLASDx, or Hazy Research's ThunderKittens). - Instruction-level primitives executed by warps or individual threads.

Because the compiler reasons across hierarchical levels, it can group subgraphs of low-level instructions together into massive "mega kernels." However, as Joe Fioti emphasizes, mega kernels are not a universal panacea: “And so, we can have at any given point a forward pass that is implemented in both kernels and mega kernels. And so, that's great because as much as you know, mega kernels are neat and cool and everyone likes to talk about them, they don't always they aren't always the optimal solution, especially when you're dealing with very large models.”

---

[01:17:38](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=4658s) Cost Models and "The Bitter Lesson": Real Hardware Benchmarking

Defining a combinatorial search space creates the challenge of evaluating candidate programs without spending decades building theoretical cost models. Accurate analytical models for complex modern GPUs are notoriously difficult to create. Lumina sidesteps this challenge by applying Rich Sutton’s "bitter lesson" directly to compiler development: rather than approximating performance with flawed models, they measure programs directly on real hardware.

Joe Fioti makes the case plainly: “And so, you know, in the spirit of the bitter lesson here, the best cost model is no cost model and we should actually just go and run these things on the real hardware that we're going to run on and see how long they take.” Program transformations that human engineers might assume are counterproductive frequently prove to be fastest when measured directly on the silicon.

To prevent benchmark evaluation times from exploding, the search engine utilizes early pruning: - Uses roofline models on program subgraphs to establish an upper performance bound. - Discards candidate paths immediately if their theoretical maximum performance cannot beat an already measured, functionally equivalent program. - Concentrates profiling passes strictly on candidates capable of yielding end-to-end performance gains.

---

[01:19:18](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=4758s) Commercializing a Compiler: Lumina's Cloud and Model Reach

After briefly playing a screen recording of the compiler generating kernel code, Fioti pivots to Lumina’s commercial traction. Software compiler companies historically have a bleak business survival rate, but Lumina solved monetization by bundling its technology directly into an inference cloud.

Key company metrics and operational details include: - **Team Size:** 8 engineers, actively and aggressively hiring. - **Throughput:** Serving approximately 80 trillion tokens per month alongside millions of generated images. - **Monetization:** Profitable inference cloud infrastructure running across heterogeneous hardware backends. - **Model Diversity:** General-purpose IR allows them to support unconventional workloads beyond standard LLMs, including protein folding, biomedical models, and diverse computer vision architectures.

As Fioti summarizes their core mission: “We use wacky compiler tech to commoditize wacky hardware. And yeah, we think the future of compute is commoditized. So come visit us at lumina.com.”

---

[01:20:37](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=4837s) Q&A: Numerical Precision, Search Space Scaling, and Profiling Speed

Following the presentation, audience members probe how the compiler handles mathematical properties like associativity and how the search strategy remains computationally feasible.

#### Floating-Point Associativity and Numerical Fidelity An audience member asks whether Lumina exploits floating-point associativity in IEEE 754 arithmetic, noting that machine learning practitioners often treat intermediate operations as swappable. Fioti counters that order of operations has real downstream consequences: “Meaning if you reduce your you know if you're doing a sum reduction and you're doing a tree sum reduction that is going to give you a different answer than a purely sequential sum reduction. And you know to the point about provider model quality if you screw up numerical accuracy that really does impact your output token quality.” Maintaining accurate models of numerical behavior is essential, though algebraic rewrites still offer significant performance advantages when mapping inputs to specialized memory spaces (such as routing one tensor through SRAM and another through register or tensor memory).

#### Benchmarking Overhead and Search Algorithms Another attendee asks how real-device evaluation scales without bogging down, noting that engines like TensorRT are slow even when relying on heuristics, and that evaluating full models makes hardware contention hard to account for.

Fioti explains their runtime realities and search engine design: - **Fast Profile Iterations:** A full model forward pass profiling run takes under 10 milliseconds, allowing roughly 1,000 profile evaluations per second. - **Genetic Search:** The compiler traverses the search graph using standard genetic algorithms rather than exhaustive exploration. - **Future Search Approaches:** The team sees potential for learned search guided by algorithms like MuZero. - **Targeting Local Optima:** The compiler does not attempt to locate the absolute global optimum; searching merely "1% of 1%" of the total state space reliably uncovers fast local optima that outpace hand-tuned kernels.

[01:23:33](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=5013s) Trading Compile Time for Runtime and Profile-Guided Search

In inference workloads, the economics of compilation diverge sharply from traditional development cycles. Joe Fioti explains that finding an optimization point that is 99% as good as a global optimum is vastly easier to discover, making it an effective search target. Because a compiled model is deployed once and served across millions of GPUs, customers are completely indifferent to compile latency if it yields faster execution. Joe notes that users can allocate an hour or ten hours of compilation budget; while returns eventually diminish, prolonged search scales remarkably well.

When asked whether actual production profiles guide this search, Joe clarifies that running an inference cloud gives them precise visibility into customer workloads:

> “Unfortunately, we run a cloud so we know the exact kind of shapes that customers care about. So we you know in terms of tensor shapes, we know roughly exactly what they're going to feed us at run time.”

Since rewrite rules guarantee semantic correctness, verification against ground truth is largely unnecessary during search. Providing the compiler with industry-relevant tensor shapes and realistic data distributions is sufficient to generate highly accurate profiling feedback.

- Compiling once to serve across millions of GPUs justifies trade-offs where hours or an entire day are spent compiling for a 1%, 5%, or 10% runtime speedup. - Users configure an explicit time budget for compilation, letting the search run until diminishing returns set in. - Real-world customer tensor shapes and data distributions are used during search to maintain accurate kernel timing.

---

[01:25:09](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=5109s) Debugging Compilers with Equality Saturation and Proofs

Debugging lowerings across disparate model architectures and hardware backends is notoriously difficult. Joe details how Luminex manages this complexity by building on top of PyTorch’s built-in instrumentation. The compiler tags high-level operators and threads that metadata through each stage of lowerings, utilizing standard PyTorch profilers to track down performance bottlenecks.

When correctness bugs do arise, the equality saturation framework provides an explicit trace mechanism through `egglog`:

> “And then what equality saturation does, what egglog specifically does for us quite nicely is it lets us say, 'Okay, we got an output program out and this program is wrong. Give me a proof that connects this output program back to our input program which basically says like here are the chain of rewrites that were applied to go from the input to the output.'”

By extracting the concrete chain of rewrite rules connecting the input specification to the malfunctioning output graph, engineers can isolate the invalid rewrite rule without manually stepping through intermediate representation steps.

---

[01:26:22](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=5182s) Operator Expansion vs. Kernel Fusion

An audience member asks about the latency trade-off incurred by expanding high-level unsupported operators down into smaller primitives rather than falling back to CPU execution or maintaining fused operations. Joe emphasizes that operator expansion is merely an intermediate representation phase, not the final lowering state:

> “So, we absolutely do fusion. We expand them to get to this like core set of primitive IR, but then when we actually go and do lowerings, the compiler searches through various forms of lowerings, and what we end up with is, yeah, quite fused ops as well.”

The intermediate representation expands complex operations to reveal common underlying primitives. The search process then analyzes the space of legal lowerings to re-fuse those primitives into optimized kernels. A simple roofline heuristic guides this phase: if a candidate kernel moves an order of magnitude more bytes to and from DRAM than an alternative, the compiler rejects it to avoid memory-bandwidth bottlenecks.

---

[01:27:27](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=5247s) Exploiting Non-Turing Complete IRs and Search Scope

Addressing an audience observation about the compiler language not being Turing complete, Joe explains that avoiding general-purpose computing semantics was a deliberate architectural choice. A key reason superoptimization struggled during the 1990s was its attempt to optimize languages like C, where Turing completeness and Rice's theorem impose severe limits on static reasoning. In contrast, deep learning workloads deal strictly with dense linear algebra:

> “We compile dense linear algebra programs, and that's a very very narrowed space versus like all possible programs. So, you know, again, PyTorch complete but not Turing complete.”

Because the compiler targets a bounded domain, it can analytically derive memory movement and roofline bounds without executing the code. This analytical modeling prunes unpromising branches early, dramatically reducing the number of costly profiling iterations needed on physical hardware.

When asked whether this search encompasses higher-level inference optimizations such as quantization, speculative decoding, and tensor parallelisms, Joe states that their compilation engine intentionally scopes itself to the forward pass graph. Engines like vLLM and SGLang already excel at orchestrating speculative decoding and serving infrastructure; the compiler plugs directly into those engines, dedicating its search optimization exclusively to maximizing execution performance within the forward pass.

```mermaid flowchart TD accTitle: Compiler Search Pipeline accDescr: High-level PyTorch graphs undergo operator expansion, roofline analysis, and profile-guided fusion within vLLM and SGLang.

subgraph IR ["Representation & Proofs"] In["Input Forward Graph (Dense Linear Algebra)"] --> Exp["Mid-Level IR: Expand to Core Primitives"] Exp --> Sat["egglog Equality Saturation"] Sat -.->|"Failure"| Proof["Extract Proof Trace to Pinpoint Invalid Rule"] end

subgraph Search ["Search & Lowering Engine"] Sat --> Anal["Analytical Roofline Pruning (Byte Movement)"] Anal --> PerfSearch["Search Fused Lowerings (Real Tensor Shapes)"] end

subgraph Serving ["Serving Runtime"] PerfSearch --> Engine["Deploy to Engine Forward Pass (vLLM / SGLang)"] end ```

---

[01:30:29](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=5429s) Lianmin Zheng on SGLang and Memory for the Agent Era

Following the applause for Joe Fioti, Lianmin Zheng takes the stage to introduce the motivation behind SGLang. Reflecting on the landscape in late 2023, open-source AI was dominated by basic Llama deployments, and developers were only beginning to anticipate what future agentic workloads would demand. Lianmin and his team foresaw that emerging agents—akin to systems like Claude Code—would rely heavily on parallel sub-agents, multi-step tool calls, and supervisory monitor loops.

```mermaid flowchart LR accTitle: Radix Tree Prefix Cache in SGLang accDescr: Shared prefix tokens branch into multiple agent forks without duplicating KV cache memory.

Root["Shared Context (System Prompt / Tool Defs)"] --> Fork1["Sub-Agent A Suffix"] Root --> Fork2["Sub-Agent B Suffix"] Root --> Mon["Monitor Query Suffix"] ```

In these workflows, when a supervisor forks multiple sub-agents, each branch operates against a shared base context inside the same KV cache pool. Furthermore, monitor loops continuously submit short queries, re-entering and exiting the serving engine repeatedly. At the time, existing serving backends had no mechanism to reuse these shared prefixes without redundant prefill computation:

> “So, at that case we notify like in that at that time there is there is no way engine can serve the prefix reusing efficiently. So, at that time we brought the Radix attention which can automatically reuse the share prefix and compute with the new with the new suffix.”

Under RadixAttention, the memory management system structures all active and past context tokens as a Radix tree. By matching incoming query prefixes against cached tree paths, the engine reuses pre-computed KV tensors directly from memory and computes attention only for the distinct suffix tokens.

[01:33:11](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=5591s) Multi-Agent Prefix Sharing and Radix Tree Memory Management

The multi-agent paradigm introduces unique challenges for memory management during model serving. In a representative setup, multiple agents share a massive system prompt while executing distinct roles—for example, one agent handling user authentication while a second conducts reviews. Both agents operate against the same foundational context, yet generate divergent outputs. Because GPU memory is fundamentally constrained, serving engines must evict suffixes from memory when capacity is exhausted. SGLang solves this challenge by structuring its memory subsystem as a RadixAttention radix tree, using this data structure to coordinate all core designs, prefix protection, and request scheduling.

```mermaid flowchart TD accTitle: Radix Tree Cache Management and Request Scheduling accDescr: Diagram showing how shared agent prompts are organized in a Radix tree and batched in cache-aware order to minimize evictions.

subgraph Memory["Radix Tree Memory Subsystem"] SP["Shared System Prompt / Prefix"] A1["Agent 1: Authentication Suffix"] A2["Agent 2: Review Suffix"] SP --> A1 SP --> A2 end

subgraph Scheduling["Cache-Aware Batch Scheduling"] R1["Interleaved Requests: A, B, A, B"] --> RO["Reorder Engine"] RO -->|"Group by Prefix"| R2["Cache-Aware Batches: A, A then B, B"] end

A1 -.->|"Evict Unused Leaves When Memory Is Full"| Evict["Memory Eviction"] A2 -.->|"Evict Unused Leaves When Memory Is Full"| Evict ```

To optimize cache retention, SGLang implements leaf eviction, dropping unused leaf nodes when fresh memory allocations are required. Request scheduling plays an equally critical role: when disparate agent queries arrive in an interleaved sequence, standard FIFO execution risks constant cache thrashing and premature evictions. By reordering pending prompts into a cache-aware sequence, requests sharing identical prefixes (such as context A) are batched together rather than toggling repeatedly between context A and context B.

- RadixAttention structures KV cache management around a unified radix tree to track token sequences across requests. - Eviction policies prioritize removing inactive leaf nodes while locking shared prefix branches. - SGLang transitioned from a research prototype into a general-purpose serving engine with its version 0.2 release in mid-2024, introducing production support for models ranging from Llama 3 8B to Llama 3 405B across both online and offline workloads. - Cache-aware scheduling reorganizes request queues dynamically to maximize context reuse and mitigate eviction risk.

---

[01:35:45](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=5745s) Overlap Scheduling: Eliminating CPU Bottlenecks

While GPU execution speeds have accelerated dramatically with each hardware generation, host CPU efficiency has remained largely stagnant. Under vanilla scheduling paradigms, serving engines follow a strict sequential loop: prepare the batch on the CPU, execute on the GPU, synchronize data back to the host, and only then begin preparing the subsequent batch. While the resulting GPU idle bubbles hovered under an acceptable 10% on older, slower accelerators, modern GPU performance causes CPU orchestration overhead to become a glaring bottleneck.

To eliminate these idle cycles, SGLang implements overlap scheduling, completely removing device-to-host synchronizations from the critical execution path by delaying sync operations to subsequent cycles. Because the host pipeline avoids blocking on immediate return values, the scheduler prepares batch $N+1$ concurrently while batch $N$ executes on the GPU. Similar zero-overhead scheduling techniques have gained industry-wide adoption, including in DeepSeek's DeepEP / DeepGEMM-based DeepSeek-V3 inference infrastructure.

```mermaid flowchart LR accTitle: Comparison Between Vanilla and Overlap Scheduling accDescr: Sequence diagram illustrating GPU bubbles in vanilla execution versus continuous utilization via overlap scheduling.

subgraph Vanilla["Vanilla Scheduling (GPU Bubbles)"] direction LR P1["Prep Batch 1"] --> G1["Run Batch 1"] G1 --> S1["Sync Host"] S1 --> P2["Prep Batch 2"] P2 --> G2["Run Batch 2"] end

subgraph Overlap["Overlap Scheduling (Zero Bubbles)"] direction TB subgraph CPU["CPU Track"] OP1["Prep Batch 1"] --> OP2["Prep Batch 2 (Concurrently)"] end subgraph Dev["GPU Track"] OG1["Run Batch 1"] --> OG2["Run Batch 2"] end OP1 --> OG1 OG1 -.->|"Delayed Sync"| OP2 OP2 --> OG2 end ```

The throughput gains from overlapping scheduling scale directly with GPU raw compute power: - On NVIDIA Hopper architectures (H100), overlap scheduling delivers approximately a $1.1\times$ throughput speedup over vanilla synchronous scheduling. - On NVIDIA Blackwell architectures (B200 / GB200 systems), hiding CPU dispatch latency yields up to a $2\times$ throughput improvement.

---

[01:38:25](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=5905s) Hierarchical KV Cache (L1, L2, L3)

Real-world agentic applications frequently feature long multi-turn interactions where minutes, hours, or even days elapse between user turns. Storing the active key-value (KV) cache purely in scarce high-bandwidth GPU memory (HBM) is economically non-viable. SGLang adopts a hierarchical caching architecture (HiCache) that links the internal radix tree to a tiered memory hierarchy:

1. **L1 Cache:** On-device GPU High-Bandwidth Memory (HBM). 2. **L2 Cache:** Host System Memory (CPU RAM). 3. **L3 Cache:** Local high-speed persistent storage (NVMe SSDs).

Tokens within the prefix radix tree can reside across any of these tiers. Intelligent background paging policies orchestrate prefetching and migration between L3, L2, and L1 based on access frequency and eviction rules. In benchmarked evaluations, hierarchical cache offloading and rapid retrieval reduced Time to First Token (TTFT) by 56% by avoiding expensive full-prompt prefill computations when warm contexts are re-invoked.

---

[01:39:52](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=5992s) The Unified Radix Tree for Hybrid Transformer Architectures

Traditional standard self-attention architectures fail to scale efficiently for extreme context lengths, driving the adoption of architectural innovations such as sliding window attention (used in Mistral and various proprietary models) and recurrent linear-attention state spaces (such as Mamba states and recurrent linear layers). However, supporting hybrid memory layers creates deep algorithmic complexity for prefix caching.

```mermaid flowchart TD accTitle: Cache Hit Requirements for Hybrid Architectures accDescr: Demonstrates how full attention, sliding window, and recurrent layers have differing requirements for a cache hit at position 10.

subgraph Target["Target Cache Hit at Position 10"] T["Token Index: 10"] end

subgraph Requirements["Layer Cache Constraints"] FA["Full Attention Pool"] -->|"Requires contiguous KV cache"| R1["Tokens 1 through 10"] SWA["Sliding Window Pool (Size = 4)"] -->|"Requires minimum window history"| R2["Tokens 7 through 10"] REC["Recurrent State Pool (e.g., Mamba)"] -->|"Requires compressed hidden state"| R3["Exact state at Token 10"] end

T --> FA T --> SWA T --> REC ```

For instance, hitting a cache prefix at token position 10 requires different data invariants across layer types: - **Full Attention:** Requires matching all key-value tokens from index 1 through 10. - **Sliding Window Attention:** If the window size is 4, the memory pool must contain at least the preceding four contiguous tokens (positions 7 through 10). If the retained window only extends back to position 8, the cache misses completely and requires a full attention recomputation for that window. - **Recurrent State:** Requires the exact recurrent state vector snapshot captured precisely at position 10.

Managing separate, standalone radix trees for distinct memory pools inevitably triggers compatibility breakage and eviction conflicts. Furthermore, the economic impact of optimizing cache hit rates is exponential rather than linear:

> “Imagine when you do the cache sharing from 95% to 99%, you reduce the expense by by five times, not by only 5% because it's you use four times less computation for all the inference workload.”

To eliminate the brittle handshakes previously required to coordinate prefix locking across diverse layer policies, SGLang implements a **Unified Radix Tree**: - The tree is anchored to canonical, full token sequences, reflecting the invariant that regardless of internal layer mechanics, external prompt representations remain token-based. - The unified radix tree acts as the master coordinator and controller across all low-level memory pools (full KV, sliding window buffers, and recurrent state pools). - Each component implements its own localized eviction and matching rules within the unified framework, resolving edge cases and preventing memory leaks or allocation waste.

[01:45:09](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=6309s) Unified Memory Pools for Heterogeneous Context Lengths

The speaker explains a critical architectural challenge in serving diverse edge and analytic inference workloads: managing different types of context state across wildly varying user behaviors. When handling distinct memory requirements—such as full token pools, sliding-window attention tokens, and recurrent states—static allocation breaks down. Different users and workloads display vastly different context footprints: one user may only require around 4K tokens of context, a second might need 128 tokens, while another consistently pushes 1 million tokens. The 1M-token workload constantly saturates the full token pool, whereas a sliding-window workload might only ever need the most recent 128 tokens.

```mermaid flowchart TD accTitle: Unified Memory Pool Dynamic Allocation accDescr: Diagram showing dynamic bidirectional memory allocation and defragmentation via compaction. subgraph UnifiedPool["Single Physical Memory Buffer"] FullTokens["Full Token Pool<br/>(Grows from Left)"] Holes["Idle Holes / Fragments"] WindowTokens["Window / Recurrent State<br/>(Grows from Right)"] end FullTokens -->|"Long Context"| UnifiedPool WindowTokens -->|"Short Context"| UnifiedPool Holes -->|"Compaction Kernel"| Compact["Repacked Contiguous Memory & Remapped Virtual Page IDs"] ```

Because context length varies dramatically across users, workloads, and even time zones, reserving token pools using fixed, pre-computed ratios is inherently inefficient:

> “So, that is a ratio difference because because you cannot reserve the token pool by the fixed ratio. If you have to design an engine to fix all the different ratios of all the different workload context, then it may you have to use a lot of different types. But, if you only have one engine or limited resources or you want to serve a lot of different workloads, then you must to have some unified way to manage all the low-level memory pools.”

To solve this, the engine moves away from pre-allocating fixed memory slices at static ratios. Instead, it provisions a single contiguous block of physical memory without pre-set capacity limits for individual pools. Memory is treated as a uniform collection of basic byte units: full token buffers grow inward from one side of physical memory, while windowed key-value caches and recurrent states grow from the opposite side, dynamically adjusting boundaries based on instantaneous demand. When requests feature long contexts, more physical memory expands toward the full token pool; when contexts are short, space dynamically reallocates toward sliding-window and recurrent state buffers.

[01:48:00](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=6480s) Physical Fragmentation and Memory Compaction

A fundamental constraint in GPU inference is that high-performance compute kernels and low-level backends require memory buffers to be physically contiguous. When requests finish and memory eviction occurs dynamically across mixed workloads, empty "holes" and idle gaps inevitably form throughout physical GPU memory. Because these disjoint holes cross the boundaries between separate memory pools, they cannot be readily reused by contiguous kernels.

To reclaim this unusable fragmented space, the engine introduces an explicit compaction routine:

> “So at that case we introduce one of our operation called compaction. We just reorganize all the memory fragments in the physical memory space. Then we pack them with low latency memory copy kernels. Then we remapping all the virtual—not token ID—it’s actually the memory page IDs. Then we just do the remapping and all the virtual page IDs they do not to be changed. We just change the low level mapping itself. That bring much more flexibility for the whole serving workload.”

By running specialized, low-latency memory copy kernels, the system repacks physical memory fragments into contiguous layouts. Crucially, the virtual page IDs exposed to the upper layers of the runtime remain completely stable—only the underlying physical page tables are updated, insulating the rest of the execution engine from the compaction overhead.

[01:49:24](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=6564s) Bridging Serving and RL Training Loops

The presentation broadens from standalone inference serving to unifying inference directly within model training pipelines, specifically reinforcement learning (RL) loops. In RL training, efficient inference is responsible for fast rollout generation, acting as the critical engine feeding generation steps back into training. By supporting fully asynchronous execution, the runtime decouples sample generation from model weight updates, ensuring that generation and training run continuously in parallel without stalling hardware.

Reflecting on the broader open-source trajectory, the speaker highlights the journey from introducing RadixAttention to building full-scale inference engines and expanding into RL training frameworks and diffusion pipelines. The session wraps up by inviting attendees to the upcoming summit in the following month via an on-screen QR code before opening the floor to audience questions.

## Notable quotes > [01:04:30](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=3870s) How good of a compiler you have greatly determines whether you lose money, whether you make money, or whether you make a lot of money. — Joe Fioti

> [48:08](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=2888s) It ends up that for open model serving, just open models, that the gap between the worst provider for a model and the best provider is actually bigger than the gap between closed and open models. — Jeff Morgan

> [01:17:38](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=4658s) The best cost model is no cost model and we should actually just go and run these things on the real hardware that we're going to run on and see how long they take. — Joe Fioti

> [01:20:37](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=4837s) If you screw up numerical accuracy that really does impact your output token quality. — Joe Fioti

> [57:05](https://www.youtube.com/watch?v=mHNKuH-9OxU&t=3425s) I think there's a big opportunity for startups to start monitoring LLM quality. — Jeff Morgan

## Mentioned AMD, Anthropic, Artificial Analysis, AWS, AWS Trainium, Calibri, Claude Code, CuBLASDx, cuBLAS, Cursor, CUTLASS, Dan Fu, Dan Simran, DeepEP, DeepGEMM, DeepInfer, DeepSeek, DeepSeek V3, Dwarkesh Patel, Egglog, FlashAttention, Google, Google Gemini, Google TPU, Hazy Research, Hugging Face, Jeffrey Morgan, Joe Fioti, Lianmin Zheng, Llama 3, Luminal, Mamba, MiniMax, MiniMax M3, MiniMax M3.1, Mistral, MLX, Moonshot, MuZero, Nvidia, NVIDIA B200, NVIDIA B300, NVIDIA GB200 NVL72, NVIDIA H100, NVIDIA H800, NVLink, Ollama, Ollama Cloud, OpenAI, OpenRouter, PagedAttention, Percy Liang, PyTorch, RadixAttention, Recursive, Reiner Pope, Rich Sutton, Richard Socher, Sabri Eyuboglu, SGLang, Stanford CS 336, TensorRT, ThunderKittens, Together AI, Tri Dao, vLLM, Y Combinator

## Video description This week’s Paper Club is all about inference. For every FLOP spent training a model, roughly six go to serving it each year — and getting that part right has produced over $220 billion in market cap in the last decade. We cover the full lifecycle of an inference request, from the moment it hits the API down to the kernel running on the GPU; why the gap between the best and worst provider serving the same open model is bigger than the gap between open and closed models; how a compiler built on search rather than heuristics turns GPU hours into tokens; and how SGLang manages memory for the agent era.

Chapters: 00:00 - Why inference matters more than training 08:24 - Three kinds of inference innovation 12:14 - Dan Fu: the lifecycle of an inference request 42:18 - Jeff Morgan: not all open model tokens are the same 1:04:31 - Joe Fioti: compiling models with search, not heuristics 1:30:30 - Lianmin Zheng: SGLang and memory for the agent era

Apply to Y Combinator: https://www.ycombinator.com/apply Work at a startup: https://www.ycombinator.com/jobs

[Source](https://www.youtube.com/watch?v=mHNKuH-9OxU)

中文简报

为何推理的重要性是训练的12倍 | YC论文俱乐部 尽管主流机器学习研究往往优先考虑训练预算,但生产模型在推理阶段每年消耗的算力大约是训练的六倍,这使得服务效率成为AI单位经济效益的决定性因素。专家Dan Fu、Jeff Morgan、Joe Fioti和Lianmin Zheng全面剖析了整个推理技术栈,展示了内存带宽墙等硬件瓶颈如何在预填充和解码阶段决定系统性能。他们指出,服务商层面的量化和路由选择对模型能力的削弱程度,可能超过开源与闭源权重之间的差距;基于搜索的编译器如何将PyTorch拆解为原子基元以超越人工编写的内核优化;以及SGLang等系统中的分层基数树前缀缓存如何大幅降低多轮智能体应用的延迟与成本。 服务模型占用的年算力大约是初始预训练的6倍,这意味着在模型权重预计两年的生命周期内,推理在经济重要性上大约是训练的12倍。 (0:09) 提供完全相同开源权重模型的最佳服务商与最差服务商之间的性能差异,大于开源与闭源模型之间的差距,低质端点甚至会将Llama模型的性能拉低至GPT-3.5的水平。 (48:08) Luminal运营着一个基于搜索的编译器,由八名工程师组成的团队每月支撑80万亿至90万亿个token的服务,其做法是将PyTorch分解为15个原子基元,并通过e-graph利用等价饱和技术取代人工启发式规则。 (64:30) SGLang实现了分层KV缓存和重叠调度,以将CPU编排与GPU内核解耦,在NVIDIA Blackwell架构上实现高达2倍吞吐量加速的同时,将首字生成时间缩短了56%。 (95:45)

Reading summary

At Y Combinator's Paper Club, infrastructure practitioners argued that operational and economic viability in AI relies far more on inference optimization than model training. Overcoming memory-bandwidth limits, provider output variance, and orchestration bottlenecks requires specialized architectures across compilers, runtime memory management, and serving stacks. The Economics and Physical Realities of Inference (0:09) Training models requires immense capital, but serving them at production scale dictates ongoing Cost of Goods Sold (COGS). LLM serving splits into two contrasting execution phases governed by hardware roofline constraints. The prefill phase processes the prompt in parallel, behaving much like training compute with high arithmetic intensity and roughly 50% Model FLOPs Utilization (MFU). In contrast, the decode phase generates tokens autoregressively one by one, resulting in memory-bandwidth-bound execution that often drops MFU to roughly 5%. Because serving outlasts training over an architecture's deployment life, system designers must aggressively target decode bottlenecks. With models scaling beyond single-GPU high-bandwidth memory capacities—such as a 3-trillion-parameter model requiring 1.5 TB in FP4 precision—operators rely on combinations of tensor, pipeline, and data parallelism, as well as disaggregating prefill and decode phases across dedicated hardware pools. Managing Open-Model Divergence and Quantization Dan Fu of Together AI and Jeff Morgan of Ollama emphasized that while open-weight models offer substantial cost savings over proprietary APIs, serving quality fluctuates widely across third-party providers. Morgan cited a 2024 study of 31 Llama serving endpoints where improper quantization and poor configurations degraded certain deployments down to GPT-3.5 performance levels. In multi-step agentic workflows, even minor per-step degradation rapidly compounds into high task failure rates. To preserve fidelity without fine-tuning, Together AI employs training-free calibration and scale search pipelines, avoiding catastrophic forgetting while matching unquantized FP8 baselines at 4-bit precision (NVFP4). Because continuously running full benchmark suites on live traffic is cost-prohibitive, serving routers monitor metadata signals, compare output token logprob distributions against baseline references, and look toward hardware-level attestation inside trusted execution environments to guarantee endpoint integrity. Compile-Time Search Over Heuristics (65:37) Joe Fioti of Lumina outlined an alternative compiler paradigm that replaces brittle, hand-authored heuristics with compile-time search, aligning with Rich Sutton's 'Bitter Lesson.' Instead of mapping thousands of high-level PyTorch operators directly to hardware instruction sets, Lumina decomposes workloads into a minimal intermediate representation consisting of roughly 15 atomic primitive operations, excluding even native matrix multiplication. Using the Egglog equality saturation engine, the compiler applies declarative rewrite rules non-destructively within an E-graph until reaching saturation. The compiler explores candidate execution programs using genetic algorithms, taking advantage of short profiling runs under 10 milliseconds to evaluate candidate kernels on physical hardware. Because inference schedules deploy across millions of serving queries, spending hours at compile time to discover a local optimum 99% as good as theoretical peak is economically advantageous. Memory Management and Scheduling for Agent Workflows (90:29) Lianmin Zheng detailed how agentic workloads—characterized by branching sub-agents, multi-step tool interactions, and recurring monitor loops—force a redesign of serving engines. SGLang uses a RadixAttention radix tree to retain and share KV cache prefixes across requests, eliminating redundant prefill computation. To keep GPUs saturated despite lagging CPU performance, overlap scheduling decouples batch preparation from device-to-host synchronization, preparing future batches while the GPU executes current operations. Because maintaining long-context KV caches in high-bandwidth GPU memory is cost-prohibitive, SGLang employs a three-tier hierarchical cache across GPU HBM (L1), host CPU RAM (L2), and NVMe SSDs (L3), reducing Time to First Token by 56% on re-invoked contexts. To handle heterogeneous context lengths without wasting memory, unified memory buffers dynamically allocate token pools and use low-latency memory copy kernels for physical memory compaction, preserving virtual page IDs while assembling contiguous blocks required by high-performance kernels. As inference workloads expand into reinforcement learning rollout loops and agent-driven interactions, hardware bottlenecks are shifting the software ecosystem toward non-destructive compilers, hierarchical memory systems, and automated quality monitoring to bridge the divide between theoretical silicon performance and real-world serving efficiency.

精炼阅读摘要

Y Combinator研讨会指出,模型生命周期中推理算力消耗已达训练的数倍,企业正面临严重的销货成本(COGS)瓶颈。来自Together AI、Ollama、Lumina和SGLang的从业者强调,必须通过全栈编译优化、端点质量校验以及分层显存管理来应对推理扩展难题。 推理的经济现实与两阶段物理瓶颈 (0:09) 在模型部署生命周期中,推理阶段的硬件算力消耗与成本远超训练阶段,过去十年围绕推理基础设施的挑战已催生超过2200亿美元的企业市值。生产环境中的大语言模型服务被严格划分为两个阶段:Prompt处理的Prefill阶段是算力受限型(Compute-bound),具有较高的算术强度,Model FLOPs Utilization(MFU)通常可达50%左右;而逐字生成的Decode阶段是显存带宽受限型(Memory-bandwidth-bound),算术强度极低,硬件MFU往往跌至5%左右。 根据Roofline分析,推理吞吐受到芯片理论峰值算力与显存带宽乘算术强度的双重物理限制。为了打破这一边界,工程优化主要依靠六项核心杠杆:底层Kernel融合、Prefill与Decode解耦部署(Disaggregation Choices)、多卡并行切分、感知缓存的动态路由(Routing)、投机解码(Speculative Decoding)以及权重与激活值的量化压缩。 开源生态扩容与端点质量离散挑战 (12:14) 开源权重模型正在迅速普及,Together AI的月度标记吞吐量已达400万亿Token,在OpenRouter聚合平台上开源模型路由量已超半数。然而,Ollama联合创始人Jeff Morgan指出,当前开源模型在实际落地中面临被忽视的质量鸿沟:不同第三方API提供商服务同一开源模型时产生的表现差异,往往超过了开源与闭源前沿模型本身的差距。2024年针对31个Llama端点的学术评测显示,部分劣质端点性能甚至退化至接近GPT-3.5水平。 在由15步构成的多步智能体(Agent)工作流中,每一步微小的输出偏差都会发生级联放大,进而导致整体任务失败率剧增。由于线上高吞吐场景无法依靠人工检验输出,平台开始采用Moonshot工具调用验证套件、输出概率分布测试以及非保留请求特征内省(如监测输出死循环和Schema违规)进行实时质检。此外,硬件级可信执行环境(TEE)对显存中模型权重与运行栈的密码学证明,正成为验证服务真实性的新兴方向。 等式饱和与编译期搜索取代人工启发式 (65:37) 针对跨多芯片架构部署推理时的工程冗杂问题,Lumina团队践行Rich Sutton的“苦涩的教训”(The Bitter Lesson),使用由Egglog驱动的等式饱和(Equality Saturation)和E-graph替代传统编译器中脆弱的手写启发式规则与下沉Pass。该系统仅由8名工程师维护,每月服务80万亿至90万亿Token。该方案将PyTorch中约2000个高阶算子拆解为仅包含14至15个基础原语(如加法、乘法和求和规约)的极简中间表示(IR),甚至原生不包含矩阵乘法算子。 在非破坏性的E-graph中,声明式的Lisp重写规则不断推导数学等价结构,直至达到信息饱和。系统结合遗传算法在硬件上进行快速性能基准剖析,由于单次前向剖析耗时低于10毫秒,系统可实现每秒约1000次实机测试。在一次编译、服务于数百万GPU的云场景下,团队通过投入数小时乃至整天的编译期搜索,换取运行期1%至10%的性能提升;同时利用中间展开与最终降级融合的机制,避免了内存带宽瓶颈。 面向智能体时代的显存层级与调度创新 (90:29) 智能体系统普遍具备多分支子代理共享系统提示词、长会话交互以及监督循环频繁进出的特征。SGLang提出了RadixAttention基数树显存管理架构,对分支间共享的KV Cache前缀进行重用与保护,避免重复的Prefill计算。为了解决GPU速度飞跃而主机CPU控制逻辑停滞带来的调度空泡问题,系统引入了重叠调度(Overlap Scheduling),在GPU执行第N批次时并发准备第N+1批次,将CPU与GPU之间的同步移出关键执行路径。 针对长期多轮交互带来的显存容量瓶颈,SGLang建立了贯穿GPU显存(L1)、主机内存(L2)和本地NVMe SSD(L3)的三级分层缓存体系(HiCache),在评测中将首字延迟(TTFT)降低了56%。此外,为了应对不同长短上下文混用导致的物理显存空洞,系统开发了低延迟物理内存碎片整理内核(Compaction),在保持上层虚拟页ID不变的前提下重排底层物理布局,并以全异步执行将推理框架与强化学习(RL)的Rollout循环深度融合。 推理系统的优化不仅涉及底层算子融合与编译期搜索,还深度依赖分布式动态路由与跨层级的显存管理。随着模型规模突破万亿参数且智能体链路不断拉长,如何防止端点量化引起的质量漂移并平抑长期会话带来的显存碎片,仍是产业界实现可持续商业化变现必须攻克的关键课题。

Source quotes

[64:30] Joe Fioti: How good of a compiler you have greatly determines whether you lose money, whether you make money, or whether you make a lot of money. [48:08] Jeff Morgan: It ends up that for open model serving, just open models, that the gap between the worst provider for a model and the best provider is actually bigger than the gap between closed and open models. [77:38] Joe Fioti: The best cost model is no cost model and we should actually just go and run these things on the real hardware that we're going to run on and see how long they take. [80:37] Joe Fioti: If you screw up numerical accuracy that really does impact your output token quality. [57:05] Jeff Morgan: I think there's a big opportunity for startups to start monitoring LLM quality.

Original transcript

[0:01] [music] [0:09] >> How's everyone doing today? [0:11] Woo! All right, it's inference night. [0:14] Come on, you got to get excited. Okay, I [0:15] didn't really know how to uh [0:17] to represent an inference in a [0:21] pictorial, so I did Halloween. So, I [0:23] made it uh spookier, as my 3-year-old [0:25] would say, [0:27] pookie. Um this is my wife over here. [0:30] Okay, so why inference? Who cares about [0:32] inference? [0:33] Um well, I think we do a lot of [0:36] inference investments at YC. We look at [0:38] this a ton. [0:40] Um I used to run an AI company where I [0:42] would train and we would do a lot of [0:44] inference, and inference was always the [0:45] bottleneck. It's so hard to do inference [0:48] well at scale and without losing your [0:51] shirt on cost of goods sold, uh a term [0:53] called COGS, as some people in software [0:56] are unaware what that word means. [0:58] Um and so far, in just the last 10 [1:01] years, it's produced 220 billion dollars [1:03] of market cap. Um and we have two really [1:06] cool companies Sorry, we have about [1:08] three four really cool companies talking [1:10] tonight. [1:12] Um [1:13] And number eight here, actually, [1:13] together, we have the head of research, [1:15] Dan Foo. [1:17] Um and so, why is inference important? [1:19] Why does this really matter? Uh the only [1:22] the this compa- I like to think about it [1:24] from a comparison of train flops versus [1:26] inference flops, and the only uh company [1:29] that actually publishes all their uh [1:31] data is DeepSeek. And so, just to think [1:34] about the the ratio of how much we [1:36] should care about train flops versus [1:37] inference flops, um uh [1:41] uh DeepSeek reports that it took about [1:43] 2.8 uh million H1 H800 hours to train uh [1:47] DeepSeek V3, [1:49] and they do about 43,000 H800 hours per [1:53] day, and that's like old data, right? [1:55] This is not recent data. So, it's [1:56] probably much more than that. [1:58] But even those numbers is about 6x [2:01] more flops on inference than train for [2:04] that weight file [2:06] per year. And if you believe the useful [2:07] life of a weight file, I don't know, say [2:10] 2 years, then it should be 12x more [2:11] important. And as more users of the that [2:15] weight file come online as AI is [2:17] proliferating, then it actually gets [2:20] more and more important. And so, I think [2:22] about this a lot when [2:24] I'm writing papers and things like that, [2:25] but the the train flops [2:28] that when you're doing a selection [2:31] across different models in your Nerf [2:32] paper, you should be looking to minimize [2:35] train flops, yes, but 12 times more you [2:37] should care about inference flops. [2:40] But typically, this is like the [2:41] Chinchilla paper, for example, in the [2:43] x-axis for most Nerf papers, you're [2:45] going to see just train flops to be ISO [2:48] train flops. We should really start [2:49] caring about ISO inference flops. [2:53] Um [2:54] and then, yeah, for the people that have [2:57] done SAS, they don't know what COGS [2:58] means, but it really matters if you are [3:01] a model server. [3:02] And if you just these are some numbers [3:04] that came out yesterday, some people [3:07] kind of were shocked that that their [3:09] margins could be this good, but it [3:11] implies that Deep Seek's margins are 80% [3:13] and it's because they have innovated so [3:14] much on the inference side. [3:19] Um well, first of all, why is it so [3:20] complicated? It's just like [3:21] model.forward. This is [3:23] it's not that hard. Um [3:25] it's actually really, really hard. And [3:26] so, there's some for the non-LLM people [3:29] in the in in the crowd, basically what [3:31] happens is you get served some [3:34] prompt, some set of of tokens, token [3:37] IDs, and then you have to call prefill, [3:40] and you'll run that on the [3:43] the the initial prompt. If you already [3:45] have some of that pre-cached, you [3:47] should, you shouldn't run prefill again [3:48] if you've already done it. And and and [3:51] even better, if it's if it's hot in uh, [3:54] in um, [3:55] DRAM or if it's in some Redis cache, [3:58] some some uh, [3:59] um, [4:00] SSD somewhere. [4:02] And then uh, uh, so make sure make make [4:05] so that you do some LRU cache kind of [4:06] thing like that. So this is already a [4:08] big opportunity for savings. And then [4:11] after that you'll have this decode loop. [4:13] And if you think about uh, we'll define [4:15] these terms in a second, but like model [4:16] flop utilization, like how much of the [4:18] GPU you're really using for a train and [4:20] prefill. And in train time you're almost [4:23] always doing uh, prefill is essentially [4:25] the same thing. Your um, MFU might be [4:27] 50%. So if Jensen Huang is guaranteeing [4:30] you 2 trillion FP8 flops, you're not [4:33] getting that. It's just not not even [4:35] close. Maybe you get half of that. [4:37] That'd be great if you can get half of [4:38] that. And in decode it's even worse. It [4:40] might be like 5% or something like that. [4:43] Um, and so it's really really important [4:45] that you're you're careful with these [4:46] things. And so the terms of art, the [4:49] things that that KPIs that um, [4:51] Dan Fu and others are are constantly [4:53] trying to optimize is time to first [4:55] token. So this is how long it takes [4:57] takes to get through prefill. And then [4:59] uh, time per output token or time per [5:01] decode step um, to produce the next [5:03] token. And then throughput, you want [5:05] really good throughput for a GPU hour, [5:08] how many tokens are you actually [5:09] producing? [5:10] Uh, and then good put is like how many [5:12] you're producing at some service level. [5:17] And just to put example numbers, these [5:19] are the the the the DC uh, numbers that [5:21] just came out I think yesterday. Um, the [5:23] throughput per GPU excluding cache hits [5:26] is 60,000 tokens tokens per second on [5:29] prefill. Including cache hits, they're [5:31] at 96% uh, uh, input cache hit rate, [5:34] which is really really good. Um, that [5:36] will change based on your traffic. If [5:38] you're doing agents, you're going to [5:39] have really good uh, input cache, but if [5:41] you have like my mom is constantly [5:43] starting brand new chats with ChatGPT [5:45] that have nothing to do with the past [5:47] and like they're asking, "What's the [5:47] weather? What's this? What's that?" [5:49] There is no no KV cache. You [5:52] constantly have to run the prefill again [5:53] and again and again. [5:54] Um and so that will be much worse. But [5:56] throughput per GPU, when there is cache [5:58] hits, is 1.5 million tokens a second, [6:00] which is really really good. Um and then [6:02] time uh to first token about 400 [6:04] milliseconds, time to to first token P99 [6:07] about 1 second. And then uh um [6:11] throughput per GPU [6:13] uh [6:14] on decode is 14,000 tokens per second. [6:17] And you can see the throughput for [6:19] decode versus prefill 1.5 million to [6:21] 14,000. So you're going to see decode is [6:23] really really the bottleneck and that's [6:25] where lots of chip innovations uh are [6:27] coming out to try to speed that up. [6:30] Um and then uh [6:32] this is actually the the most [6:34] interesting one that uh people clued in [6:36] on, which is the KV cache retention [6:38] policy. And so they have petabytes of of [6:42] users that are using this data for the [6:44] their KV caches and they retain it for [6:46] 72 hours, which is how they're achieving [6:49] such high uh cache hit rates. [6:52] Um some other terms of art here you'll [6:54] you'll see, which is arithmetic [6:56] intensity. So as I go from DRAM to SRAM, [6:59] if I if I transfer, you know, uh 4 [7:01] bytes, how much how many how much flops [7:04] or how much compute can I am I doing on [7:06] those 4 bytes? If I'm doing an RMN or if [7:09] I'm doing diffusion, these are really [7:11] high cuz I transfer an image and then I [7:14] have to or even noise and then I have to [7:16] keep calling the same uh um weight file [7:19] on on a same network again and again and [7:22] again. I don't have to go back to DRAM [7:23] unless it's a really big um image and I [7:25] have to tile or something like that. And [7:27] so uh but for LLMs, the the arithmetic [7:29] intensity is quite poor. [7:31] And this is the argument that Volo gave [7:33] when he came up here for diffusion LLMs. [7:35] Um and [7:37] and there's a next concept model that um [7:39] Herman out of my lab just published as [7:41] well, which is also has much higher [7:43] arithmetic intensity. [7:45] And the two other ones [7:47] model bandwidth utilization, model flop [7:49] utilization. So, when you buy an H100, [7:51] they say that you get 2 trillion FP8 [7:55] flops, [7:57] and what are you actually going to get [7:59] at test time? And it's it's usually [8:01] going to be much less than that, and [8:02] that's the MFU is how many you actually [8:05] get. And then the and model bandwidth is [8:07] how much what is peak memory bandwidth? [8:09] You have 1 terabyte of bandwidth or [8:11] whatever it is, how much you're actually [8:12] getting at test time on average. [8:15] And then [8:16] there's this idea of like of roofline [8:20] analysis, which is like how [8:22] good could you possibly get? So, it's [8:24] going to be the min of your chips peak [8:26] flops per second versus the memory bytes [8:29] per second times flops. [8:31] And so, there's lots of innovations [8:33] people have come up with. I bucket it [8:34] into these three buckets. Some people [8:37] might, you know, bicker about which one [8:39] would where buckets should live, but the [8:41] way I think about it is there's model [8:43] innovations where [8:45] you do this thing called MOE where you [8:48] if you have K experts, [8:51] then you can amortize reduce the amount [8:52] of flops at test time [8:54] over all of the experts by 1 over K [8:56] essentially. [8:58] There's SSMs where you don't need full [8:59] attention on every single layer. Maybe [9:01] you only need it every fifth layer or [9:03] every 10th layer, and so you can really [9:05] reduce the amount of [9:07] inference needed at test time. [9:09] And then [9:10] multi-headed latent attention where I'm [9:12] actually compressing in the sequence [9:14] domain [9:15] to [9:16] reduce the amount of flops. [9:18] And then there's I have this [9:20] quantization. I think we've all seen [9:21] these [9:22] like even as low as 2 bytes [9:25] models, and then distillation into [9:27] smaller models. So, once we get some big [9:29] beefy model, [9:30] and we we [9:32] we don't want to necessarily serve that, [9:33] we can distill it into a smaller model [9:35] if that's all we need. [9:37] So, those are uh model innovations. Then [9:39] there's kernel and runtime innovations [9:41] like radix attention, flash attention, [9:43] page attention. These are They're kind [9:45] of a lot of these are are just [9:47] uh taking from classic OS and database [9:50] theory. If you have taken uh intro [9:52] databases, where how do you make a [9:54] production database really fast? And so, [9:56] there's, you know, are you cache? You [9:58] can do prefix trees. You can do radix [9:59] trees. That's kind of the idea of radix [10:01] attention. Uh TriDao did this GPU tiling [10:04] idea called flash attention, which is uh [10:06] huge innovation. And then page [10:08] attention, which uh is um [10:11] mutually exclusive from radix attention, [10:13] but is basically this idea of a virtual [10:16] uh uh um [10:17] uh memory paging and buffer pool that [10:20] you can that you can page into um [10:22] into the DRAM. [10:24] And so, and the last one is the harness [10:25] innovations. So, of course, the easiest [10:27] one is routing. We did this at Focal. [10:29] And so, not everything should go to the [10:31] big gory model. We can do much more uh [10:33] clever things like that. Maybe you can [10:35] waterfall to the or tool call to the [10:36] bigger model to the bigger model to the [10:37] bigger model. And that will save a lot [10:39] of flops. Speculative decoding, we have [10:41] a draft model that is validated. And [10:44] then, my favorite one is is the [10:45] recursive uh language model, where never [10:48] put the full context into uh um into the [10:52] the KB cache. Just let the LLM decide [10:56] what it needs and when. And so, you can [10:57] really reduce the amount of KB cache [11:00] with with those ideas. And then last [11:02] one, uh compaction, which um [11:05] is a mandatory part of every every [11:06] harness. And there's a really cool [11:08] innovations that Sabri um and Dan Simran [11:11] in my lab came up with called [11:12] cartridges. [11:13] >> [clears throat and cough] [11:15] >> Um if you want to learn more, I think [11:16] the the two best uh lectures I've seen [11:19] on this, um Percy gives an amazing [11:21] lecture on on his lecture 10 on CS 336 [11:24] LLMs from scratch. And then uh the Ryner [11:26] Pope and Dwarkesh I thought were just [11:28] super super cool and super good. [11:30] Um enough for me. You have the absolute [11:33] experts in the world on this. Uh, first [11:36] Dan Fu, who runs research at uh, [11:38] Together AI. [11:40] Uh, and came out of my lab at Hazy. Uh, [11:43] Jeffrey Morgan, who is found founder and [11:45] CEO of Alama. Uh, [11:47] Joe Fiaty, who's founder and CEO of [11:49] Luminol. And then Lang Shang, I think [11:52] he's coming. He says he's late. He's [11:54] hopefully shows up. Uh, is one of the [11:55] core developers and second author of the [11:57] SG Lang paper. [11:59] Well, uh, let me or please uh, help me [12:03] uh, um, uh, welcome Dan Fu. [12:07] >> [applause] [12:14] >> Great. [12:15] All right. Uh, [12:17] thanks everyone. Thanks so much for [12:18] having me. Thanks Francois for that that [12:20] great uh, [12:21] uh, [12:22] you know, overview of inference. Um, so [12:24] today I'm going to talk about that uh, [12:27] whole life cycle of a token, life cycle [12:29] of an of an inference request in a [12:31] little bit more detail. Kind of walk you [12:33] through when you send the request in and [12:35] it gets to the API, makes it to a data [12:37] center, all the way down to the to the [12:39] root of the kernels and how to actually [12:41] executes on a GPU and then how you get [12:43] uh, tokens back out. Um, and uh, this [12:46] was a talk that I cut down from before. [12:48] So it says, you know, there's six things [12:50] that we can do to make it faster. I uh, [12:52] crossed out half of them and we'll talk [12:53] about three of them. Um, some of which [12:55] uh, Francois mentioned a little bit. [12:58] So, um, okay. Uh, [13:01] so why do we care about inference right [13:02] now? So we are in a really exciting um, [13:06] moment in in inference. So open source [13:09] AI has never been bigger. Um, [13:12] this this figure is a bit old, but a few [13:13] months ago at some point we're doing 400 [13:16] uh, trillion tokens a month at Together. [13:18] Um, what was interesting about that is [13:20] around the same time that we put out [13:21] that number, Google said, "Hey, this is [13:23] how much traffic Gemini is doing." And [13:25] we scrappy little startup we're doing [13:27] 1/8 of total traffic that Google is [13:29] doing which is pretty cool. [13:32] By [13:34] by by earlier this year more than half [13:36] the tokens through open router were [13:38] through open models [13:40] and the marketing team wants me to [13:41] remind you that at some point a couple [13:43] weeks ago on for one day we had more [13:45] tokens than open AI as well together. [13:49] And these open models are are also [13:52] one of the drivers behind them is that [13:54] they're a lot cheaper. So you can get [13:56] nearly the same quality or in in many [13:58] cases actually the same or better [13:59] quality but up to 10 times cheaper per [14:02] year. [14:04] So there we're in this really [14:06] interesting area where the open source [14:08] models are really improving and there's [14:10] a lot of [14:12] a lot more happening every day. [14:14] So in today's talk I'll talk a little [14:16] bit about how they're served and then [14:18] how do you get more tokens? So at [14:19] together we spend a lot of time thinking [14:20] about how can we create more tokens in [14:23] the world [14:24] and and serve serve more tokens for [14:26] everyone who wants them. [14:28] A bit about together. We are an AI [14:31] native cloud so the same way that AWS [14:33] you know you can rent CPUs and and get [14:35] some compute but you can also set up SQL [14:38] databases and storage and rest APIs and [14:41] all the other stuff that you need to run [14:43] a web service. At together we're [14:45] interested in saying hey what are all [14:47] the things that you need to do to build [14:49] a [14:51] AI native company and so that's of [14:54] course there's pre-training there's [14:55] fine-tuning there's inference which [14:57] we'll talk about. And then you know [14:59] sandboxes for your for your agents to to [15:01] try out all sorts of things and at [15:03] together we're building the the cloud to [15:05] do it all [15:07] and [15:08] and especially focused on inference. [15:12] At together so we we have a pretty [15:14] correct team we're about 80 researchers [15:17] bunch of PhDs [15:19] we did a bunch of stuff in the past, [15:22] and now we're we're focused on inference [15:25] and making GPUs more performant. [15:29] Okay, [15:30] Francois talked about this a little bit, [15:32] but one of the things that we think [15:33] about is what do the workloads look [15:35] like? What do you different people's [15:37] workloads look like? [15:38] And so we'll spend a lot of time [15:39] thinking about if you're running a [15:41] coding workload, what is your cache hit [15:43] rate? What is the time between your [15:44] turn? So when you go ask your Claude [15:47] code or your open code or your cursor to [15:49] say, you know, build me this thing, how [15:51] long do you wait before you go and talk [15:53] back to it? [15:55] How responsive does that need to be? If [15:58] you're generating a lot of code, how [15:59] fast does that code need to happen? [16:01] And subject to certain SLAs, so like how [16:04] fast you need your response, how quickly [16:06] you need the whole thing to be finished, [16:08] we'll go and and try to optimize these [16:10] pieces. [16:12] So what happens when a request actually [16:14] comes in? So imagine you send in a [16:16] request to some REST API, it goes [16:19] through the normal parts of the internet [16:21] through Cloudflare and DNS and all that [16:23] stuff. Eventually it hits a series of [16:26] data centers. So let's say you have [16:29] thousands of GPUs, you have this thing [16:30] called the KV cache that that Francois [16:32] was talking about. [16:34] And then at the first thing that you do [16:35] is this basic [16:38] request level routing piece. So you're [16:40] going to I'll go over all these [16:42] momentarily, so you're going to you need [16:44] to figure out which GPU replica to send [16:46] your request to, then you have to take [16:47] this high-level human text, turns it [16:49] turn it to ones and zeros for your model [16:51] to actually process, [16:53] then you're going to start the big batch [16:55] processing loop where there's this [16:57] prefill thing and this decode thing, [16:59] and then at the end you get tokens out, [17:01] which you then translate back into into [17:03] human text that you send back to the [17:04] user. [17:06] So that first bit, request and routing. [17:08] So [17:09] the way to think about this is so you've [17:11] got thousands of GPUs, you have all [17:13] these requests coming in from from all [17:15] over the world, [17:17] And really it's a mix of both load [17:20] balancing, but also trying to think [17:21] about things like the KV cache. So, if [17:24] you're familiar with the with the old [17:26] world of [17:27] of databases or production web services, [17:31] you would just say, "Okay, which of my [17:33] computers is serving too much traffic? [17:35] I'll just round robin and send it [17:37] somewhere else." [17:38] In [17:39] in this in this language modeling era, [17:42] we have something else that we need to [17:43] account for, which is if you send the [17:45] same request to the same GPU over and [17:47] over again, you actually get better [17:49] performance characteristics because you [17:51] already processed [17:53] let's say in the code in the code piece, [17:55] that first request you send in your [17:56] whole code base, it processes it all, [17:58] figures everything out. But then in the [18:00] subsequent requests, you've already done [18:02] all that work, so you don't need to redo [18:03] that. So, this routing actually has a [18:05] little bit of has a little bit of an art [18:08] to it, which is you have to figure out [18:09] not only how do you balance the load [18:11] across your data centers, across your [18:12] GPUs, but also where can you send the [18:14] request so that the KV cache will [18:16] already be resident. [18:19] Next, there's the stage of tokenization. [18:21] So, this part you probably imagine [18:22] should be fairly straightforward, but [18:24] there's actually a lot of stuff that [18:25] happens here, too. And that's because [18:28] your requests, when you use these [18:29] models, come with a lot of things that [18:31] we call tool calls. So, these tool calls [18:33] are basically like, "Hey, I'm a I want [18:36] to ask my model to do this. And by the [18:37] way, you can [18:39] search, you can use grep to search [18:41] through my file system, or you can run [18:43] this command to build the app and test [18:44] everything, or you can search the [18:46] internet." And the way that that gets [18:48] passed in, so the model is just some [18:50] abstract, you know, combination of like [18:53] mathematical things. It's like a bunch [18:55] of matmuls and some attention modules. [18:57] So, it actually it has no ability to [18:59] search the internet or do anything. So, [19:01] the way that that happens is when you [19:02] send in your request, you say, "Hey, [19:04] here's a list of 10 things that you can [19:05] do." And then the model needs to go and [19:07] say, "Hey, I want to call the internet [19:10] right now, or I want to search through [19:12] this. Um and there's actually a lot of [19:15] correctness things that you need to [19:16] ensure when you get there. So, I think [19:19] Jeff is going to talk about this in a [19:20] little bit how you manage that from the [19:22] user side, but there's quite a bit of [19:24] complexity just at this tokenization [19:26] setting [19:27] of how to [19:29] of how to get this stuff to be correct [19:31] and and work well. [19:33] Um next you have what you can think [19:35] about as the core scheduler loop of your [19:38] model. So, you have all these requests [19:40] streaming in. They're going to wait to [19:42] batch together [19:44] and either because either wait for a [19:46] free GP or some free cycles or wait for [19:49] the previous request to finish. [19:51] The choices that you make here on the [19:52] scheduler will decide will determine [19:55] what's your total throughput. So, are [19:57] you trying to optimize for more users on [19:59] on the same GPUs or are you trying to [20:02] optimize for very snappy response. So, [20:04] is this like [20:06] if you ever use the Google search and [20:08] see the Google Gemini, the very first [20:09] response is always really really fast. [20:11] Then if you try to chat with it, it's [20:12] kind of slow afterwards. So, that's [20:14] probably very different scheduling [20:16] requests, possibly even different models [20:17] on the back end. [20:20] Now we get to these two stages called [20:21] prefill and decode. So, you can think of [20:24] these these things as prefill is sort of [20:27] like prompt processing. So, when you [20:30] pass in your whole code base, you need [20:32] to figure out exactly what's in your [20:33] code base, exactly what's being asked [20:35] and and understand all those pieces. [20:38] That's a very compute bound operation. [20:40] So, it's going to be very heavy on your [20:43] GPU and it's going to have some of these [20:47] these words that Francois mentioned MFU. [20:49] Your MFU is going to be very high. The [20:51] arithmetic intensity is going to be very [20:52] high. [20:53] On the other hand, you have the second [20:55] stage where the model is just generating [20:57] one or a few tokens at a time to get [20:59] back to you. So, one way to think about [21:01] this is every time you generate one of [21:03] those words, you have to load the entire [21:05] model weights from uh GPU memory to the [21:08] compute unit. So, that that's this [21:09] decode stage. It's a very memory-bound [21:11] operation and it's a very different type [21:13] of operation. So, we spend a lot of time [21:15] thinking about how can we serve these [21:18] two different workloads in different [21:19] ways? How do we separate them? [21:21] Disaggregate them? Are there different [21:23] parallelisms or kernels or or anything [21:26] that that we can use to to make this [21:28] part more efficient? [21:30] Finally, so now you get a bunch of ones [21:32] and zeros, these token IDs that then get [21:34] translated back into text and get [21:35] streamed back to you. This part is uh is [21:38] is going to be relatively simple. [21:40] There's of course complexities around uh [21:42] when do you free up your KV caches? When [21:44] are you How long do you keep those [21:45] around? Uh how how long do you try to [21:48] keep um those tokens resident? [21:51] All right. So, um and together we think [21:54] about there's roughly six or so things [21:56] that you can use to to to make things [21:58] things faster. So, um in in in some uh [22:02] in no particular order, you have [22:03] kernels. Kernels is just a fancy word [22:05] for a GPU program. So, ultimately at the [22:08] end of the day, when you're running the [22:09] model, all these things have to actually [22:11] run on the GPU, so you can spend a lot [22:13] of time optimizing them, um writing [22:15] different attention kernels, MOE [22:17] kernels, um mega kernels to make this [22:19] all faster. Uh there's disaggregation [22:21] choices, so the prefill and the decode [22:23] are very different workloads, so you [22:25] want to put them on different GPUs. How [22:27] you make those choices can affect a lot [22:29] of things about user experience, total [22:31] throughput of your GPUs. [22:33] Um there's going to be parallelism, so [22:35] uh Chimney K3, for example, is a 3 [22:36] trillion parameter model in FP4 at half [22:39] a byte per parameter, that's going to be [22:41] 1 and 1/2 terabytes. Uh your best B200 B [22:44] your best B300 only holds 280 GB of [22:47] memory, so that means, okay, you have [22:49] this big model, you have to split it [22:50] across GPUs. How do you split it across [22:53] those GPUs? That's a parallelism [22:54] question. Um there's routing in the [22:57] sense of once your request enters the [22:59] data center, which uh which pieces do [23:01] you send it to? A lot of choices to make [23:03] there. [23:04] And then the last two almost more model [23:06] layer optimization. So speculative [23:08] decoding. [23:09] The basic idea is you have this really [23:11] big model. [23:13] It doesn't you don't really need it to [23:14] answer you know, what is one plus one. [23:16] So you always have a little model [23:18] running at the same time that's trying [23:20] to preempt the big model. [23:22] And if you get it right, you can just [23:23] get those tokens for free. [23:25] And then lastly quantization which I'll [23:27] talk about is so you have these big [23:29] models. How many per token do you [23:32] actually need to store? How many bits do [23:34] you need to store these [23:36] these weights? [23:38] For sake of time, I'm going to talk [23:39] about kernels, parallelism, and [23:40] quantization today. And go into these in [23:44] a little bit more detail. [23:47] So first kernel. So [23:49] the place where kernels matter is at [23:50] this core at these core prefill and [23:52] decode parts of the computation. So this [23:55] is where the the math is actually [23:57] running the matmuls, the attention, [24:01] and all those other pieces. [24:04] So what are some things that you think [24:06] about when optimizing kernels? So when [24:08] you think about the model, there's some [24:10] pieces that are very relatively easy to [24:12] optimize. And then there's a lot of [24:14] other pieces [24:15] that are harder to optimize. So like all [24:17] the random layer norms and RMS norms [24:19] that are that are in your model, random [24:21] soft maxes here and there. [24:23] Each of these individual operations has [24:25] very low arithmetic intensity. So if you [24:28] write a kernel for it by itself, you get [24:31] a very low flop utilization on on those [24:33] kernels. So that's historically one of [24:35] the reasons why you can only get 50% [24:38] end-to-end flop utilization during [24:40] training, 5% during decode. [24:43] And the way that that you would do this [24:45] that you do this [24:46] that we have been doing for years is you [24:49] go look at your model forward pass, try [24:50] to figure out every place where you can [24:52] squeeze, have a team of really talented [24:55] kernel engineers go and optimize this. [24:58] Um and for a long time this was um this [25:00] was the modus operandi for everyone in [25:02] the industry, which is you just have a [25:03] bunch of really smart people um work on [25:05] this really hard. [25:07] Um at Together recently we've been [25:09] thinking about how to use agents to [25:11] actually optimize automatically optimize [25:13] a lot of this piece. Um so this is an [25:15] internal project called Together [25:17] Compiler. We're going to be talking [25:18] about it more um in the next coming [25:20] weeks. I think probably by the time we [25:22] release we'll name it something slightly [25:23] different because um I had a I had a [25:26] shocking realization the other day, [25:27] which is that Together Compiler is not [25:29] actually a compiler. Um but uh yeah, so [25:32] we'll we'll have to rename it. But the [25:33] basic idea is you have a target model, [25:36] some target hardware, a particular [25:38] amount of traffic, and you're going to [25:39] use agents to go optimize this inner [25:41] loop. Um so a couple key ideas here is [25:44] that um you really want this loop to be [25:47] driven by real traffic because uh [25:49] different traffic or coding workload [25:50] versus chat workload is actually going [25:52] to activate kernels in a very different [25:54] different way. Um and we found it's been [25:56] really important to do the full [25:58] end-to-end research loop. So actually [26:00] look at what is running in production to [26:02] go um iterate on that. So we what we [26:05] started to build here is almost the [26:06] beginnings of an RSI loop, uh recursive [26:09] self-improvement loop for inference. [26:12] Okay, so so that was kernels. Um the the [26:15] next the next thing that I'll talk about [26:17] is how to optimize your parallelism [26:19] scheme. So the basic idea here is you [26:22] have all these weights. They're split uh [26:24] they're too big to fit onto a single [26:26] GPU, and if you can split up your model [26:29] across multiple GPUs you get more [26:31] compute, more bandwidth, more of all the [26:32] things that we love to um to to do these [26:35] things. Um a couple basic parallelism [26:38] schemes that have been around. So [26:39] there's this thing called tensor [26:40] parallelism, which is you just take the [26:42] matrix multiply and you split it up and [26:44] then you do a basic map reduce style [26:46] reduction at the end. Um pipeline [26:48] parallelism is this idea where you take [26:50] your many layers and you cut it um layer [26:53] by layer. So, you put different layers, [26:55] full layers onto GPUs. Data parallelism [26:58] is another idea where you just send [27:00] different sequences to different GPUs, [27:02] but you keep you keep whatever parts of [27:04] your model there are the same. [27:07] Um the the game here has always been [27:09] how much memory do you have on your [27:11] GPUs? How fast can you write your [27:13] kernels to support a particular [27:15] parallelism scheme? [27:16] And and what things are available? [27:19] One thing that's become quite popular [27:21] recently that we've been investing a lot [27:23] in and one of the core drivers behind [27:25] that DeepSeek $38 per hour number that [27:29] Francois threw out earlier is leveraging [27:31] much larger scale up domains. So, what [27:34] this means is that if you're familiar, [27:36] if you try to go out and buy GPUs [27:37] recently, you've heard about these [27:39] Blackwell chips and these GB chips. [27:42] And the biggest difference between them [27:43] is that these Blackwell chips and also [27:46] all the Nvidia chips before them, you [27:48] could only have eight of these these [27:50] GPUs really talk to each other in a high [27:52] bandwidth way. So, that basically meant [27:54] that you take your model and split it up [27:55] eight ways. And if you try to split it [27:57] up more, if you try to split it up 16 [27:59] ways, you'd go over a much slower [28:01] bandwidth path and then that would make [28:03] your serving just unbearably slow. [28:06] Um recently Nvidia has started selling [28:09] building these these racks that can [28:10] connect 72 GPUs in much higher in with [28:14] much higher bandwidth pieces. So, what [28:17] one thing that you can start to do is [28:18] you can take all the experts in your [28:20] model and put the experts individually [28:22] onto [28:24] onto different onto different GPUs. So, [28:27] what this means is that [28:29] you can end up starting to serve it [28:31] means that you can serve a lot higher [28:33] volumes on the same GPU. And the basic [28:36] reason also comes back to these ideas [28:37] about arithmetic intensity. Basically, [28:39] if you only have one expert on your GPU [28:42] serving a ton of tokens, now that looks [28:44] like a big matrix multiply operation [28:46] that GPUs are really good at serving [28:48] again. [28:49] Um a couple pieces here. So, um I think, [28:52] you know, people have known about these [28:53] ideas for a couple years now. You can go [28:55] back 2 3 years and and find these demos [28:58] of expert parallelism. [29:00] Um Recently, uh we've figured out how to [29:02] actually make them more reliable. Um [29:04] these these GB racks are really awesome, [29:06] but they do tend to uh have a node or [29:09] GPU go out at different times. And if [29:11] you have a single model replica that [29:13] depends on 72 GPUs and one of them goes [29:16] out, um if you don't have the right [29:18] remediation or and reliability pieces um [29:21] behind it, then what that means is a [29:23] single GPU can going out can take out 72 [29:26] GPUs worth of of worth of compute, which [29:28] is not something that you want. So, [29:29] we've been spending a lot of time [29:30] figuring out how to make this reliable, [29:32] resilient to failures, um gracefully uh [29:35] gracefully keep going even through these [29:37] through these hardware failures. [29:41] And the last bit that I'll talk about [29:42] again um is is just quantization. And [29:45] and the this one is going to be [29:46] interesting because in some sense it's [29:48] it's quite simple, you know, there are [29:49] quantization algorithms that you go, uh [29:52] there's probably hundreds or thousands [29:54] of papers on it. So, you just go pick [29:55] your favorite one and use it. Um in [29:57] another sense, it's it has actually ends [29:59] up being quite complicated because it it [30:02] really touches the quality and the [30:03] quality of experience um of of your of [30:07] your end user at the end of the day. [30:09] So, the the basic idea here is that your [30:12] models will be trained in something like [30:14] BF16. That's 16 bits per parameter, 2 [30:17] bytes per parameter. 3 trillion [30:19] parameter model's going to take 6 [30:20] terabytes. Um [30:22] and uh but the nice thing about machine [30:25] learning is that the entire computation [30:27] is kind of fuzzy, right? It's not exact. [30:29] It's not um [30:31] uh you don't necessarily need all those [30:33] bits of precision. You can still get the [30:35] right answer at the end of the day. Um [30:38] so, for a while we've known that there [30:39] are these algorithms that I can use to [30:41] take those 16 bits, that be of 16, [30:43] shrink it down to eight bits or four [30:44] bits and still have still have decent [30:47] quality. And and there's a lot of [30:48] technical detail there that I'll gloss [30:50] over, but the the main thing that we [30:52] care about is [30:54] how do we ensure the right quality and [30:56] the right uh user experience for for for [30:59] the user. So, we're we're actually we're [31:01] pretty obsessed with this at together. [31:03] Um we So, we we know both from the [31:06] system side what is the right [31:07] quantization for the hardware, what is [31:09] the thing that's going to get you the [31:11] most throughput, the most tokens per [31:13] GPU. But at the same time we know that [31:15] people are using these increasingly to [31:17] build companies and and run high um you [31:20] know, very very important things. So, [31:22] what we've built at together is a sort [31:24] of a um a measurement quality red [31:27] teaming loop. So, before we send out any [31:29] model to any user, um we are internally [31:32] dog fooding that model for quite a bit [31:34] internally. So, this does a number of [31:36] things. So, it's going to catch all [31:37] those tool call things that I talked [31:38] about earlier cuz if suddenly your your [31:40] Claude code harness can't uh [31:43] can't search the internet, you're going [31:44] to notice that as a user. Um but you're [31:46] also going to um notice if the model is [31:50] not quite as smart as it should be. If [31:52] it if it doesn't quite get the right [31:54] scores on the benchmarks or on [31:56] end-to-end um evaluation. So, one of the [31:58] things that we have running all the time [32:00] is actually a little Flappy Bird loop um [32:02] where we have all our models trying to [32:04] um trying to code up Flappy Bird um in [32:07] in a rapid loop. So, uh if if something [32:09] went wrong with the quantization, [32:10] suddenly the Flappy Bird won't flap [32:12] anymore and then we have to go and and [32:13] figure out what went wrong and fix it. [32:16] Um but basically the name of the game [32:17] here is uh of course there there's a lot [32:20] of algorithmic uh [32:21] optimization you can do, a lot of [32:23] algorithmic innovation, but really what [32:25] it comes down to is really rigorous um [32:27] evaluation loop. So, benchmarking, [32:29] agentic evaluations, uh a lot of [32:31] in-house dog fooding, etc. [32:35] Um one success story here. So, uh we [32:38] were pretty lucky to work with a lot of [32:39] the open source model providers pretty [32:41] deeply. Um, one of uh, and we worked [32:44] with MiniMax to serve their M3 model. [32:47] Um, and uh, also their upcoming M3.1 [32:50] model. Um, not sure if I was supposed to [32:53] say that out loud, but now now you all [32:55] know. [32:55] Um, and one of the one of the [32:57] interesting bits here was that um, [32:59] MiniMax was actually having some trouble [33:01] quantizing their model to NVFP4. [33:04] Um, but we were quite interested in it [33:05] because we know from the kernel [33:06] perspective, the serving perspective, [33:08] um, and also general uh, experience like [33:10] this is kind of the right quantization [33:12] you want to use on Nvidia GPUs. Um, so [33:15] on the left there that uh, what color is [33:17] that? That that lilac purple, um, that's [33:20] the original um, model in in FP8. In the [33:23] middle in this [33:25] hot pink. Yeah, in the in the hot pink [33:26] in the middle of that was the uh, the [33:29] the MiniMax attempt at uh, at quantizing [33:31] with their own open with their own [33:32] recipe. Um, and but with our recipes in [33:35] the orange on the right here, we were [33:36] actually able to match the original [33:38] model performance and then we were able [33:40] to serve that, um, get MiniMax a lot [33:42] more tokens, um, and help them kind of [33:44] close the loop from training the model [33:45] to serving it profitably and and uh, and [33:48] help them, you know, recognize the [33:50] revenue, get the money so that they go [33:51] and go train their next model that [33:54] uh, that we'll all um, know and love. [33:56] Um, [33:56] so yeah, that this this was a a fun [33:58] little success story where we were able [34:00] to um, use our quantization expertise to [34:03] to also help help the people who who are [34:05] doing the the core model development. [34:08] All right, so that that's the that's [34:11] most of it for for the presentation. So [34:12] just to recap, I talked a little bit [34:14] about how the request flows through the [34:16] system. Um, we didn't go into too many [34:18] detail details about most of it, but [34:20] there's routing, you know, which replica [34:22] do you send it to, which data center. [34:24] Uh, there's the tokenization, how do you [34:25] make sure that the uh, that the tool [34:27] calling and all the pieces that you [34:28] expect to be correct are are [34:30] Um, the core queuing and batching, uh, [34:33] scheduling work, um, prefill and decode [34:35] often on different GPUs with different [34:37] kernels, uh, different parallelism [34:39] schemes, um, and then finally, once you [34:41] go through all that, uh, and make the [34:43] GPU really, you know, prove its worth, [34:45] you get the tokens back, um, to do your [34:47] coding and your chat, um, and and all [34:50] the other things that that we know and [34:51] love. [34:53] Um, so with that, if, uh, [34:55] you know, just to conclude, if you found [34:57] this stuff interesting, if you want to [34:59] go use open models, if your CFO is [35:01] beating uh, [35:02] breathing down your neck and saying, [35:04] "Hey, uh, we're spending too much money [35:06] on cloud code, uh, and maybe we should [35:08] try some of these open models." Um, one [35:10] of these is is a sales QR code, um, and [35:12] then if you if you find this stuff [35:14] really exciting, you want to come work [35:15] with us, uh, the other one goes to our [35:17] careers page. Um, so yeah, thanks again [35:19] so much for listening in for having me. [35:21] >> [applause] [35:28] >> And I think we have a few minutes for [35:29] questions. Is that right? [35:31] All right. [35:34] >> Hey Dan, thank you for the talk. Uh, two [35:36] questions. So in the NVFP 4, I think [35:38] quantization that you did, uh, can you [35:40] talk about how did you kind of like [35:42] measure the quality? Was it GSM 8K or it [35:44] was it like perplexity, MMLU, uh, with [35:47] respect to the MXFP FP8? [35:50] And the other question is with respect [35:51] to the, I think, the interconnect [35:53] technology that enabled basically, uh, [35:56] Grace Blackwell class of chips to [35:58] connect beyond like like a in a mesh [36:01] form uh, beyond eight uh, [36:04] eight [36:05] GPUs. [36:06] >> Yeah. Um, yeah, so, uh, let's see. These [36:09] three are kind of three, uh, open [36:11] benchmarks that are out there. Um, I [36:13] think, uh, some of them are quite [36:16] saturated. Like Amy is at 90 plus [36:18] percent. Um, really, I think the thing [36:20] there is if you see one point or more of [36:22] movement, then you know that something's [36:24] gone wrong. Like the model's actually [36:25] going to struggle with some, um, with [36:27] some real world activity. So, we have [36:29] these kind of stock standard benchmarks [36:31] that we release and we talk about that [36:34] you know, we also [36:36] we will look at like KL divergence of [36:39] the individual layers. So, [36:41] how how much did the activations [36:43] actually change? [36:45] And then of course we'll look at all the [36:47] end-to-end agentic pipeline. So, we've [36:48] built up like I think at this point [36:50] hundreds of different test cases where [36:52] So, every time anybody said, you know, [36:55] anything went wrong with any model ever, [36:56] we put that into the test case, make [36:57] sure make sure that it's okay. And then [36:59] we also have these agentic evaluation [37:01] loops like Flappy Bird. [37:04] We actually put it put it back into the [37:06] together to together compiler project so [37:08] that the model is optimizing its own [37:10] kernels. And if we see that the model [37:12] for some reason can't optimize its [37:13] kernels anymore, then we're like, "Hmm, [37:15] okay, we we should take a look at that." [37:17] The second question was [37:20] a little bit about Oh, yeah, the the the [37:22] NVLink. So, [37:23] um [37:24] I'm personally not a hardware person, so [37:26] I've never built or designed hardware. [37:27] The closest I've come is I bought one of [37:29] those Hugging Face micro ducts. So, that [37:32] that's my closest you know, you know, [37:34] play play with hardware recently. [37:37] Basically, if you look at one of these [37:39] racks, it's really crazy. So, each tray [37:42] of four GPUs has these really thick [37:43] cables that go to a really big switch. [37:46] And then all the trays are connected [37:47] that way. The you know, the bits go [37:49] through that really thick cable, go [37:51] through the switch, and then you can do [37:52] your all reduces and and all the pieces [37:54] there. I I I'm not super familiar with [37:57] you know, what was the jump from the 8 [37:59] GPU technology to the 72. [38:02] But I'm you know, I love it and I using [38:05] it. [38:06] >> So, if I would like you know, [38:08] batteries they be versus one. So, you [38:10] have your phone and you don't know what [38:11] kind of interview that kind of like you [38:13] know, what they test against and kind of [38:15] really like with the this possible that [38:17] I'm on. [38:18] >> Yeah, yeah. So, yeah, we we do we we do [38:20] the public benchmarks, we do our own [38:22] agent loops and and everything in [38:23] between. Yeah. [38:25] Yeah. [38:27] Sure if there's a microphone for Yep. [38:31] >> Okay. [38:33] Thank you. [38:34] >> Um my question uh is about this graph [38:37] where um [38:39] you found a recipe to do better [38:41] quantization. [38:42] So, I'm curious uh what was the [38:44] difference? [38:46] What did you [38:47] were you [38:49] training from scratch or were you [38:50] running [38:51] um multiple pre-training epochs after [38:54] introducing quantization? Or was it that [38:58] um there was no training done on the hot [39:00] pink one and just like straight up uh [39:03] just, you know, transforming the numbers [39:06] uh without doing any [39:07] training uh fine-tuning on it? [39:09] >> Yeah, so uh so yeah, so for context [39:11] there there's a lot of quantization [39:13] recipes that out there that require you [39:15] to train the model after it's done. Um [39:18] we don't like using them because we find [39:20] that unless you have [39:22] uh the entire [39:24] uh original training data set. So, [39:26] unless you're the model provider um [39:28] yourself. [39:29] >> You don't have the [39:30] >> Yeah, we Yeah, we we don't do that [39:31] because we find that when you do that [39:33] you can actually like the model will [39:34] forget things. So, this is entirely a [39:36] training-free approach. Um it's uh it's [39:39] a combination of stuff that is, I think, [39:41] publicly out there. So, uh you know, [39:44] there's some calibration data involved. [39:46] Um there's this paper called scale [39:47] search that that we put out um uh a few [39:50] months ago that that's used pretty um [39:52] pretty deeply, but it's a it's a mixture [39:54] of being very careful with your um with [39:58] your recipes and then also um trying to [40:01] uh and and just being very careful and [40:03] and putting together all the different [40:05] pieces. So, like there's a hundred [40:06] training-free quantization methods out [40:08] there. It's being very careful about [40:09] which ones you use and and being uh very [40:12] sedulous about how you evaluate them. [40:14] >> Thank you. [40:14] >> Yeah. [40:15] >> All right, I think we got to move on [40:16] here. [40:17] >> Um last question for you actually would [40:19] be [40:19] >> Yes. [40:20] >> Are there any model types that you're [40:22] excited about? I know there's lots to [40:23] talk about SSMs, lots to talk about [40:26] DLLMs, there's all these other types of [40:28] architectures that might dramatically uh [40:31] make your life a lot easier specifically [40:32] on the decode side. Is there any ones [40:34] that you're still holding out hope for? [40:37] >> Um yeah, I think uh the funny thing [40:39] about all the all the new architectures [40:42] is that so during my PhD I worked on [40:45] some of these like SSMs and things and [40:46] and the thing that we always got was [40:48] like nobody's ever going to use this [40:49] thing. It's it's a dead end. [40:51] Transformers are going to take over the [40:52] world. It transform AGI all that stuff. [40:55] Um a few years later [40:57] and all those technologies are used in [40:59] in all of the big models. So Kimmy K3 [41:01] has this thing called KDA which is [41:03] basically an SSM layer. Um it might be [41:06] slightly controversial for me to say [41:07] that, but it's an SSM linear attention [41:09] style layer. Um so it's being used. It's [41:12] it's in the one of the best open models. [41:14] Uh there's sparse attention being used [41:16] which um for the longest time nobody was [41:18] using. Uh we use the diffusion language [41:20] models the DLLMs in the in all the [41:22] speculators. That's the best thing to do [41:24] now um to to do it. So um what's [41:27] interesting is that all these things are [41:28] being used. Um my students have recently [41:31] been very interested in looped models. [41:33] So the idea here is um you know, in this [41:35] in this basic computational graph you [41:38] send the token back through the same [41:39] layers a bunch of times and then the [41:41] idea is you can ramp up the training [41:42] flops without increasing your your [41:44] weights. Um and we thought, "Oh, this is [41:46] a weird idea. Nobody's going to use it." [41:48] Um and then it turns out they at least [41:49] that OpenAI and Anthropic were both [41:51] using it already. Um so yeah, it's the [41:54] it's it's a really exciting time where [41:55] like if you if you put out a good idea [41:57] right now about architectures and and go [42:00] through the leg work of showing that it [42:01] works, it'll show up in GPT-6 or GPT-7. [42:05] Um and it the the the time to feedback [42:08] and and time to impact is is so so quick [42:10] now. [42:11] >> Thank you, Dan. [42:12] >> Cool. [42:16] >> [applause] [42:18] >> Hi everybody. I'm Jeff. I'm the [42:19] co-founder and CEO of Ollama and um [42:23] uh zooming out a little bit, um I uh [42:25] wanted to talk about how all open model [42:28] tokens, not all open model tokens are [42:31] the same. And how to think about in a [42:33] world where there can be, you know, I [42:36] think we saw a list on François's deck [42:37] of just how many inference companies [42:40] just the top set are exist today. Um in [42:43] the foreseeable future, there'll be [42:45] hundreds if not thousands of potential [42:47] inference providers. And how to think [42:50] about measuring quality across these [42:51] inference providers if you're a builder [42:53] and you're trying to like build an [42:54] application on top of open models. [42:57] Um a little bit of background. Uh I have [42:59] Ollama here just finished dinner uh with [43:01] all of us. Um [43:03] uh Ollama's an open source project that [43:05] originally released in 2023 for running [43:07] uh open models locally. Back in 2023, um [43:11] the most popular open models like Llama, [43:13] for example, were [43:15] generally smaller models that you could [43:17] run on either like a single H100 or on a [43:21] laptop or like a small server. Um and [43:24] the core goal of the project has always [43:25] been to give an OpenAI and Anthropic [43:26] level developer experience but for open [43:28] models. Um since then it's grown a lot [43:31] um and because the models got bigger and [43:33] bigger and bigger, really culminating [43:35] with the start of this year when Kimmy [43:37] and Jill M and MiniMax uh all launched [43:39] these massive 500 billion to 1 trillion [43:42] parameter models, um we launched [43:44] Ollama's Cloud um to [43:46] really provide the same experience that [43:49] that developers got locally but in the [43:50] cloud uh and that's really exploded uh [43:52] over the course of this year um and [43:54] it's, you know, pretty much becoming uh [43:56] one of the largest if not the largest [43:58] independent uh cloud service for [44:00] accessing open models across the US and [44:02] Europe. Um [44:03] and so all of this is important because [44:05] we spend a lot of time thinking about [44:07] how to bring um a great developer [44:10] experience across many different models [44:12] and over many different providers. [44:15] And so, if you think about the [44:16] architecture of our cloud service, [44:19] um [44:20] unlike kind of a traditional inference [44:22] provider, [44:23] um we really focus on a few things, like [44:25] getting the latest models, making sure [44:26] they're really fast, making sure you get [44:28] a reliable experience. A lot of these [44:30] inner workings A lot of this is in that [44:32] router box that Dan was just talking [44:34] about. Um and we do this for a bunch of [44:36] these reasons. Um I think, you know, [44:38] reliability and having, you know, early [44:41] access to models is probably the most [44:42] important because different providers [44:44] excel at serving different models. But [44:46] in the end, we have this architecture [44:47] which is very similar to a lot of [44:50] customers that use Ollama. Um Ollama's [44:52] deployed by a huge chunk of the Fortune [44:54] 500. And our cloud service and what From [44:57] talking to [44:58] teams that are moving from closed models [45:01] to open models due to cost or control or [45:03] other reasons, is that they've all kind [45:05] of have the similar problem. And so, I [45:06] wanted to spend some time uh talking [45:08] about about this tonight. Um but the [45:11] main reasons to route are really that [45:12] you kind of get in the In the ideal [45:14] case, the best of all the worlds by [45:15] combining a bunch of these providers [45:17] together. And the definition of of an [45:20] inference provider is kind of loose, [45:21] like this is even a problem that you [45:23] might have if you are an inference [45:24] provider and you have different [45:25] deployments on different hardware in [45:26] different regions from different teams [45:28] that are using different uh inference [45:30] engines, for example. [45:32] Um and so, when you adopt an [45:34] architecture like this, when you're [45:35] thinking about, you know, pooling a [45:36] bunch of inference providers together, [45:39] um what are some of the challenges? [45:40] Well, the biggest one in summary is that [45:43] not everyone serves models the same way. [45:45] In fact, that's a good thing because [45:47] different teams have excel at different [45:49] uh performance optimizations. [45:51] Um [45:52] Some folks run on specific kinds of [45:53] hardware. I think there's like a new [45:55] inference provider that just runs on [45:57] like Mac minis all over the world. Um [46:00] and so, [46:01] you know, if you think about it, um, [46:04] ultimately customers care about a few [46:06] different dimensions, right? Cost, [46:07] intelligence, speed, and then ultimately [46:10] like region and privacy. Um, and the [46:12] challenge with routing across a bunch of [46:14] providers is obviously that, um, one bad [46:17] apple can ruin the bunch. And that if [46:19] you have a single provider that's [46:20] underperforming on one of these four, [46:23] um, criteria, that you, um, [46:27] you know, you might have a customer [46:28] experience that's kind of the least [46:30] common denominator. Um, and this can be [46:32] for a variety of reasons. Uh, one of [46:34] them could be the model precision. You [46:35] know, you may have a provider that has [46:37] over quantized a model or is running it [46:40] on faulty hardware or they're using a [46:43] different version of the weights. Um, [46:45] tool calling is a big area of concern [46:47] like Dan was talking about. Doing tool [46:50] calling effectively is a really hard. [46:52] Um, while it seems simple because it's [46:54] like parsing JSON and strings, um, [46:57] there's actually a lot that can go wrong [46:58] from recovering from tool calling [46:59] failures, um, and you know, enforcing [47:02] specific schemas for different tool, uh, [47:05] tool, uh, parameters and names, um, and [47:07] also just making sure that the parsing's [47:10] actually happening correctly at the [47:11] tokenizer level. Um, different providers [47:14] may have different sampling techniques, [47:16] um, and then inference engines have such [47:18] depth to them and to what you can edit. [47:20] So, whether it's custom kernels or KB [47:21] cast precision, um, [47:23] a lot of these, [47:25] you know, changes can really be only [47:27] seen in like really bespoke situations [47:28] as well, like when you read a really [47:30] long context. Um, and there's even silly [47:33] things that can happen like, you know, [47:34] models being served with the wrong name. [47:36] Um, you think you're accessing one [47:37] model, it's actually a different model, [47:39] but it's hard to know. Um, they're all a [47:40] friendly assistant trying to help you. [47:42] And so, a lot of the things can go wrong [47:44] in the serving stack and and it can [47:45] really affect, uh, you know, if you're [47:47] trying to serve a large customer base [47:48] across different regions and providers, [47:50] it can really all it can be tricky to [47:52] track that down. And so, I spent some [47:54] time ahead of this talk trying to think [47:55] through what we do and also what have [47:57] what has the industry done for this? [47:59] Um, and [48:01] it ends up it's not a new problem, and [48:03] there's been a lot of research done [48:05] about well, how big of a problem is [48:06] this? And the crazy thing is that, you [48:08] know, we all talk about the open model [48:10] to closed model gap, you know, how far [48:12] behind are the open models? It ends up [48:14] that for [48:16] open model serving, just open models, [48:18] that the gap between the worst provider [48:20] for a model and the best provider is [48:22] actually bigger than the gap between [48:25] closed and open models. And so it's a [48:27] really big problem, and you know, I the [48:29] experience you can have with a great [48:31] provider, you know, like [48:33] Together, for example, who are good [48:35] friends and partners with, [48:36] um, [48:38] can be a lot different than [48:40] uh an endpoint served by, you know, a a [48:42] different provider. And it's really hard [48:44] to know this right up front because, you [48:47] know, we come from a world of where Open [48:49] AI and Anthropic really kind of [48:50] kickstarted um a lot of the kind of [48:54] you know, ultimately the demand for [48:56] tokens. And in those worlds you you kind [48:57] of expect a consistent experience no [48:59] matter how you're accessing uh those [49:01] models. Um, but for open models that's [49:03] definitely not the case. In fact, it can [49:04] even be a bigger problem than the [49:07] difference in performance you'd get [49:08] between the closed and open models. [49:10] Um, and the good news is it's not a new [49:12] problem. There's actually been a lot of [49:13] great work done around this. Um, there's [49:15] a paper in 2024 that tested, you know, [49:17] 31 different Llama endpoints, and [49:21] very similarly they found that [49:23] um you know, some models behaved were [49:26] even closer to like the performance of [49:27] GPT 3.5 than another Llama model of the [49:30] same that were supposed to be exactly [49:31] the same way. It's [49:33] Um, and so I had a good time digging [49:34] through this and thinking through well, [49:35] if you're going to use open models, how [49:37] do you really make sure you're getting a [49:38] good quality experience? [49:40] Um, the last thing I'll say before [49:41] jumping into like some of the techniques [49:43] and ideas is that this problem gets [49:45] significantly worse with coding agent [49:47] and agent workloads because they're such [49:49] long context or long horizon workloads. [49:52] they generally involve multiple calls [49:54] back-to-back-to-back-to-back. [49:56] So, if you have a model that's serving [49:57] an inaccurate call an inaccurate [49:59] response, and then you're compounding [50:01] that um [50:03] a pretty like you know, modest coding [50:07] task, for example, that might take 15 [50:08] steps could be up to, you know, 79% [50:11] wrong if like 10% if you have a 10% [50:14] chance [50:15] deviation from kind of what the model [50:17] was output was designed to be. Um which [50:20] can obviously ruin your experience if [50:21] you're ultimately getting it a [50:23] an error rate of 80%. So, it's a big [50:25] problem that gets worse as our agent [50:27] workloads get harder and harder. [50:30] Um so, how do you then go about [50:32] monitoring and introspecting and looking [50:34] at inference endpoints? Um [50:38] you know, obviously, you can't read [50:39] through all the tokens like Alamas [50:41] attempting to do here. That'd be really [50:43] difficult, especially when you're doing, [50:44] you know, what is it? 1.6 million tokens [50:46] per GPU and and preprocessing. [50:49] Um and well, the first thing you can [50:51] obviously do is run evals, um which is [50:53] kind of like the brute force way and um [50:55] probably the most expensive way to do [50:57] it. I mean, artificial analysis actually [50:59] does this on a regular basis where [51:01] they'll just run benchmarks a few times [51:03] um over the different providers and kind [51:05] of rank them. I think here this is a [51:07] small picture, but you can see here that [51:09] by and large for the Kimik 3, um most [51:12] providers get it pretty much right. [51:13] There's some on the right that aren't [51:15] quite there. Um there's one that's um [51:17] uh not doing too well, but generally, [51:19] they're all in in good shape. The [51:21] challenge with this is that um [51:24] ultimately, you know, it'll catch like [51:25] bad quantization, um the challenge and [51:28] like a variety of issues cuz you're [51:29] ultimately testing how the model was [51:31] benchmarked in the first place. The hard [51:33] part is that it's really expensive and [51:35] it takes a long time to run these, so [51:36] you can't do it on an ongoing basis. [51:39] Um the other thing we saw rise at the [51:41] end of last year was, you know, the [51:43] folks at Moonshot at one point were kind [51:45] of ultimately fed up about how different [51:48] providers were doing tool calling with [51:49] the Gemini models, which are incredible [51:51] family models. And so they released the [51:53] Gemini K2 tool verifier, which was and a [51:56] scoreboard for which providers did best. [51:59] And this was awesome because it caused [52:00] everyone to really kind of enter a race [52:02] to like make sure that the tool calls [52:04] were accurate. Gemini models has since [52:05] done something similar. So if you're [52:08] building a model, this is a great way to [52:09] make sure that people actually run it [52:11] correctly. Give them the tooling to test [52:13] it. Another more interesting way that [52:16] these are some of the papers that I've [52:17] read about and we're thinking about [52:19] integrating into how we do [52:21] Olama's routing and model aggregation is [52:25] a distribution test where you basically [52:27] compare the outputs of a reference [52:29] weight and [52:32] a provider API and then really just look [52:34] at the distribution [52:36] between the two and understand you know [52:37] how much of a gap there is. This is [52:39] great because it's generally pretty [52:40] cheap unlike the the benchmarks which [52:42] could be very expensive. It's maybe like [52:44] a dollar per test. [52:47] The hard part is that you know while [52:50] it it might only be may not catch some [52:54] of the more intricate problems such as [52:56] like poor quantization. [52:58] It doesn't really tell you why [52:59] something's going wrong whereas if you [53:01] have a more comprehensive eval suite you [53:03] can like kind of pick out a specific [53:04] test [53:06] and then you really do need the [53:07] reference model running. So which can be [53:09] an extra expense you have to you have to [53:11] go under. And take under. [53:13] Another one is to use the same called [53:15] log probs [53:17] which is [53:19] a way to see kind of ultimately the [53:21] probability distribution of the [53:22] different token IDs that for future [53:25] tokens. [53:26] And this is great because it's generally [53:28] a pretty cheap call. You're just getting [53:29] like a single you know the next token [53:32] log probs and comparing it with a [53:33] previous baseline. So you're kind of [53:34] measuring the drift from that provider [53:37] which then can tell you if things have [53:39] changed. For example, did they deploy it [53:41] somewhere else? Did they upgrade their [53:42] inference engine? [53:44] Um [53:45] this is really cheap. Again, it's [53:47] potentially even much cheaper than the [53:48] previous test. Um it's only like a [53:50] single token request basically. Um and [53:53] so, you know, the example in the paper [53:54] found a ton of changes across um [53:58] different endpoints. And, you know, they [54:01] were able to kind of identify these [54:02] silent changes to the serving stack, [54:04] which is really interesting if you think [54:05] about it. Everyone's serving stack is [54:06] evolving, it's getting patched, they're [54:08] deploying new versions of popular [54:09] inference engines like VLM or SG Lang. [54:12] Um so, that can be a good way to monitor [54:13] it over time to make sure your customers [54:14] are getting [54:15] a consistent experience. Um Log Probs [54:18] isn't on a lot of the APIs these days, [54:21] so that's one downside. And you're [54:23] really only detecting drift from the [54:24] past versus, you know, and you could [54:26] potentially compare it to uh another uh [54:29] reference model as well. But in this [54:31] approach, you're really only getting [54:32] drift from the past. So, if you're [54:33] starting from a place that wasn't [54:34] serving it correctly, this isn't super [54:36] helpful. [54:38] Um another thing a lot of popular [54:40] routers do, and this is what we're we're [54:42] kind of at this stage right now, is [54:44] looking through just, you know, without [54:46] any kind of data retention, can you just [54:48] introspect how many tool calls are [54:50] passing, how many of them are incorrect, [54:52] are there tool calling errors that are [54:53] occurring, how's the shape of the [54:56] request changing over time? Um different [54:58] providers may serve different input and [55:00] output sorry, different output lengths, [55:02] and that could really give you an [55:03] indication of whether there's, you know, [55:05] a common issue. I've seen from like the [55:08] DLM or DeepSeek serving have run on [55:10] problems where the model gets into an [55:12] infinite loop. Um and so, that's where [55:14] you can really catch these problems [55:15] ahead of time just by looking at the [55:16] shape of the request without even having [55:18] to look at the data. [55:19] Um and then you can keep a running [55:20] score, and then ultimately, you know, [55:22] keep an eye on that and monitor it [55:23] request to request. Um Open Router does [55:26] something similar to this, uh like for [55:28] example with like the exacto model [55:30] approach that they're taking. [55:32] Um and so, zooming out, you know, one [55:35] thing you know, we're always working [55:36] backwards from how do you get to an [55:38] experience with open models like you get [55:39] with open AI and Anthropic? You think [55:41] about those, you know, they run on [55:43] different hardware like TPUs. I think [55:45] there was a a bug last year at one point [55:47] with Anthropic where like a tiny bug in [55:49] a TPU kernel caused like potentially [55:51] like a whole retrospective of an outage [55:53] just in terms of quality. And so they [55:56] have such a strong care for the serving [55:58] quality. How do we get there but in the [56:00] open model ecosystem? And so I think one [56:03] way is that the labs can share a much [56:05] bigger stack for how to run their model. [56:07] Kind of most commonly people will put [56:09] their weights on the internet on Hugging [56:11] Face for example or upload them to a [56:12] llama and then they'll say here's some [56:14] example code. But it's actually not what [56:16] they're serving in their own inference [56:18] stack. It's generally like another [56:19] applied AI team that generates that so [56:21] that the community can use it at large. [56:23] DeepSeek yesterday launched DeepInfer [56:25] which is their own inference stack. So [56:27] now you can take the DeepSeek weights [56:29] and run it hopefully at $38 a GPU hour [56:33] on your own GPUs. And what's great about [56:35] that is you can use a consistent basic [56:38] basically a consistent stack. So if [56:40] you're just a new upstart inference [56:41] provider you can set that up and have [56:43] some confidence that you're serving it [56:44] with high quality. [56:45] A lot of labs are quantizing the model. [56:48] Um [56:49] and they're working with you know pros [56:51] like you know Dan his team [56:53] to quantize the model before launch so [56:55] that the model that they're serving on [56:57] their API is the same as what an open [56:59] model [57:00] that you would take and serve are. [57:03] This is very different than what [57:04] happened originally which is [57:05] quantization tend to be done by post you [57:07] know post model release by folks who [57:09] don't have the training data or we're [57:11] ultimately doing it to you know fit it [57:13] onto like a MacBook or something like [57:15] that. [57:15] This is much more increasingly common [57:17] that the labs will actually publish [57:18] their own quantized versions of the [57:19] model. Um you can access models through [57:22] lower level APIs. You know one thing we [57:24] realized is that some of the tool [57:26] calling you can actually pull up into [57:28] the harness and get a lot of [57:30] accuracy by doing that consistently [57:32] across providers versus, you know, [57:34] delegating that to the providers. And [57:36] same with reasoning, uh, extraction and [57:38] thinking. Um, [57:40] I think there's a big opportunity for [57:42] startups to start monitoring LLM [57:44] quality. This is something that's super [57:45] important if you think about LLMs being [57:47] part of a critical loop of a [57:48] application. I don't know if this exists [57:50] yet, but it's a really interesting idea. [57:53] And then one of my favorites was this [57:54] paper that talked about how, you know, [57:55] every modern Nvidia GPU actually comes [57:58] with a trusted execution environment [58:00] that can attest to the weights that are [58:02] running and the software that's serving [58:03] the model. And this can be a way to get [58:05] a fingerprint telling you that it's [58:07] being served in a correct way. I think [58:08] this is pretty advanced, but I thought [58:10] it was a really interesting idea. [58:13] Um, yeah, that's a little bit about how [58:15] to think about, you know, once there's [58:17] inference everywhere and it's open [58:19] models are being served everywhere, how [58:20] to really measure quality and like make [58:22] sure that, um, you're getting the best, [58:25] uh, tokens out of these awesome open [58:27] models that are coming out. Um, thanks [58:28] for having me. [58:31] >> [applause] [58:35] >> Hey Jeff, thanks for the talk. My [58:36] question is, uh, regarding the pre-fill [58:38] admission, [58:40] uh, you know, interval criteria [58:42] for quality. How do you like decide [58:44] between tool calls and like large [58:46] context, running context, [58:48] uh, special especially specifically for [58:51] multi-turn agent like workloads that [58:53] span like [58:54] long sessions. Uh, like which would be [58:57] better? I I am assuming it's the longer [58:59] context pre-fill admission, [59:02] uh, interval. [59:02] >> Yeah, I'm not sure exactly, [59:06] uh, if I have the best background to [59:08] answer that. You're asking about the [59:09] admission. [59:10] >> Yeah, yeah, yeah. The the interleave [59:12] admission between like, let's say [59:14] certain decode after like after say say [59:16] 48 or 64 decode steps, you kind of take [59:19] every API endpoint tool call [59:22] like admission versus, [59:25] uh, from the running context you take [59:27] the the pre-fill admission. [59:29] So, there was there was a there was a [59:30] slide in your in your deck which kind of [59:32] like, you know, went over that. [59:34] >> Oh, yeah. This is just saying, you know, [59:36] you can basically send a single token [59:38] prompt [59:40] and then measure the log probs that are [59:41] coming out the other side and then [59:43] really measure that with like a a [59:44] baseline. So, you can do that as the [59:47] over time. [59:49] And so, it's not really specific to like [59:50] a specific context length. I think one [59:52] thing that you could do though is do [59:54] this with a consistently long like a [59:56] longer and longer prompt and then [59:59] start to measure what is the deviation [60:01] in log probs as the prompt gets longer. [60:03] Cuz what we've seen with a lot of [60:05] providers, I mean, first off, some [60:06] providers will just cap the context [60:08] window for a model. So, that's like one [60:09] silly functional thing to look at. But [60:11] as you get especially past the 200, 000 [60:14] token range in context, be able to test [60:16] that the next token accuracy is [60:18] consistent with what [60:20] the other with, you know, the reference [60:22] model providers or what it was [60:24] previously returning. [60:26] I think it's the general approach here. [60:27] I'm not [60:28] super familiar with each of these [60:29] papers, but that was I thought this was [60:31] a pretty elegant approach to do it that [60:33] doesn't require running an entire eval, [60:35] for example. [60:36] >> Okay, thanks. [60:36] >> What do you think about quantizers then [60:39] just hill climbing the evals? So, if [60:41] there's an increasing number of like [60:43] open source tests to ensure model [60:45] quality that [60:47] they just train on that and we're not [60:48] really preserving the model quality. [60:50] Like, how do we get around that? [60:52] >> Yeah, it's a great question. I've heard [60:53] a lot about this how, you know, [60:54] increasingly in the in in, you know, [60:56] when there's a large marketplace [60:58] especially of inference providers [60:59] fighting over who can be the cheapest [61:00] token, you're incentivized to match some [61:03] traffic pattern, some use case. [61:06] I think what we've done a lot with [61:07] customers is that they, you know, they [61:09] come to us and they say, "Hey, we've got [61:10] this this specific use case and we can [61:12] design with them a [61:15] you know, a set of tests that you can [61:16] run consistently or a set of evals." And [61:18] that's really hard for a a single [61:19] provider to game because it's not [61:21] public. It's specific to that customer's [61:23] data or their traffic. And so, that can [61:26] be a way where just because they don't [61:27] have access to it or it's more bespoke, [61:30] um that you can kind of catch a provider [61:31] for specific customer not delivering, [61:34] uh you know, at at the performance that [61:35] it should. Um and that's good because [61:37] generally that's where we hear issues on [61:39] quality is because some a customer would [61:40] say something. So, it kind of completes [61:42] the loop. [61:43] >> Thank you. [61:45] >> I have a quick question. Over here, [61:47] Jeff. [61:48] >> Oh, hey. [61:49] >> Um you know, like the when I first tried [61:51] a llama, it was a local model. [61:53] And then we have like gotten way off of [61:56] local models. And then as of today with [61:58] Flint and Richard Socher's uh recursive [62:00] launch, uh we're back to local models. [62:03] What do you kind of see for the future [62:05] of a llama? Do you think it's going to [62:06] be Yeah, how how do you think this [62:07] impacts uh a llama? [62:10] >> Yeah, you know, with the launch of a lot [62:13] of these local models, there's a lot of [62:14] ways to run models locally as well. [62:16] Different inference engines, um you [62:18] know, there's the MLX framework that [62:19] powers a llama now on Apple silicon. [62:22] >> Calibri, too. I don't know if you have [62:23] seen Calibri. It does Qwen 3.5 MoE uh [62:27] locally. And it's like serves it from [62:29] disk. It's pretty cool. [62:30] >> Yeah, there's so many of these [62:31] approaches. And I think all the problems [62:32] we talked about here are also problems [62:34] for local inference model serving. Um [62:37] and that can be desktop, it can be like [62:39] desk side, like, you know, the DJX [62:41] Spark. Um they're you know, if you have [62:43] the new M5, we just got ours today, M5 [62:45] Ultra. Um [62:47] or, you know, data center uh serving for [62:50] a lot of businesses who do have GPUs. I [62:52] think it's a problem, too. It's not just [62:53] a problem with the provider, but it's [62:54] also what's the technology stack you're [62:56] using to serve the model. And you need [62:57] to be able to [62:59] test that reliably. Um and so, I think [63:02] two things. One, with local I mean, [63:03] local this new model is incredible. [63:04] We're super excited to be launch [63:05] partners um with recursive. Um the Flint [63:08] model. [63:09] Um [63:10] but I do think, you know, it is a [63:12] challenge whether it's a cloud provider [63:14] serving a serving you tokens or it's an [63:16] inference engine you're running yourself [63:17] serving tokens that you're updating [63:19] every week. This is a common problem. [63:21] The reality with open models is it's [63:23] really the wild west, which is an [63:24] incredible thing cuz these models can be [63:25] customized and changed and quantized. I [63:28] think there's like 20 different [63:29] quantizations online for almost every [63:32] model that you can just pick up, you [63:33] know, off of a llama or hugging face and [63:35] download. The challenge really is how do [63:37] you benchmark these? But then, if you're [63:39] serving it to customers live, you need [63:40] to wait to do you know, you need a [63:42] lighter test than benchmarks. So, I [63:43] think it's a very equal equally [63:45] important problem locally. [63:47] >> Awesome. Thank you, Jeff. [63:49] Future of Ollama, I think um is probably [63:52] people running a mix of local and cloud [63:53] models. We've seen this new model is [63:55] really good at coding locally. And so, I [63:58] think, you know, the number of tasks you [64:00] can now run [64:01] uh you know, frankly, a lot of the best [64:03] coding problems were only solvable or [64:05] hardest coding tasks were only solvable [64:06] by cloud models that really grew popular [64:08] this year. Now, with these ultra-capable [64:12] uh local models that can run on, you [64:14] know, this MacBook, um we're going to [64:16] see, you know, this routing problem go [64:18] from just being from a bunch of [64:19] providers on the cloud to like local and [64:21] cloud uh routing. So, anyhow, awesome. [64:24] Well, thank you guys so much. [64:25] >> Thank you. [64:27] >> [applause] [64:30] >> All right. Yeah, my name is Joe. Uh I'm [64:32] one of the co-founders and CEO of a [64:34] company called Luminal here. And we are [64:36] a compiler company, which is kind of [64:39] unusual, but throughout this talk, I [64:41] think you'll see why. [64:42] Uh we really see compilers as being a [64:44] very critical part of the stack and kind [64:46] of defining your conversion function [64:48] between GPU flops or GPU hours and [64:51] tokens. And how good of a compiler you [64:53] have [64:54] greatly determines whether you lose [64:56] money, whether you make money, or [64:57] whether you make a lot of money. And so, [65:00] we make a compiler that's based on a [65:02] very wacky technology called equality [65:04] saturation. It's a fairly new technique. [65:07] We are the largest, I think, [65:09] industrial application of equality [65:10] saturation, and we think it's the [65:12] correct approach for efficient [65:15] compilation of models to a bunch of [65:17] different types of hardware. And so, [65:20] we really think, you know, in a large [65:21] part of why Nvidia is worth $4 trillion [65:25] comes from the fact that CUDA is a very [65:26] well-supported software ecosystem. Uh [65:29] it's very it's a very difficult problem [65:31] to map models down to complex hardware. [65:35] It's a very, you know, I think fairly [65:37] unsolved problem at this point. And our [65:40] solution is to leverage the better [65:43] lesson. So, we do a lot of compile time [65:46] search. Uh this lets us basically not [65:49] maintain a lot of different complex [65:50] rewrite rules or complex heuristics that [65:53] ordinary compilers maintain. Uh and it [65:56] lets us not maintain a big kernel [65:57] library that we need to constantly [65:59] modify every time a new model comes out. [66:01] And so, we're able to take a very, very [66:04] small shared compiler core, build [66:07] backends for different types of [66:09] inference ASICs in the data center, and [66:11] then run a whole lot of inference on [66:13] this across both GPUs from Nvidia and [66:15] AMD, but also TPUs and Trainiums, [66:18] and a whole host of different inference [66:20] ASICs as well. And so, you know, at any [66:22] given point, we're running, I believe, [66:24] something like 80 or 90 trillion tokens [66:26] a month. Uh we're not together, but [66:29] we're also eight people. So, uh we'll [66:31] get there someday. [66:34] So, this talk, yeah, it's going to be [66:35] largely about how we think about [66:37] compilers, um kind of a unique approach [66:40] versus your typical MLIR-based compiler [66:43] stack, and why this tends to work quite [66:46] well for ML applications for large [66:49] language models and diffusion models and [66:51] bio models and all the ones that at [66:53] least make money today. [66:55] Um specifically, how we use the egg log [66:58] engine, which powers our equality [66:59] saturation, what what equality [67:01] saturation really even means to begin [67:03] with. Uh and then how we structure our [67:05] IRs in order to target a bunch of [67:07] different hardware and fully utilize all [67:10] of that hardware. [67:11] Uh and then how we scale up search. [67:15] And so, you know, really what we do is [67:17] we we kind of take a look at all of the [67:19] neural networks out there, all the deep [67:20] learning applications, and we realized [67:22] they're not that complex. Uh they are [67:25] effectively dense linear algebra [67:27] and not even that many different [67:28] possible operations of dense linear [67:30] algebra. And so, we take models and we [67:32] actually boil them down to something [67:34] like 14 or 15 primitive ops. So, those [67:38] operations are like add, multiply, sum, [67:42] um [67:43] things like that. And it turns out you [67:44] can decompose the whole wide world of [67:47] PyTorch ops into this very simple op set [67:51] and not lose any sort of information. [67:54] And so, you know, we're not Turing [67:56] complete, but I like to say we're [67:57] PyTorch complete, meaning if you can [67:58] represent your model in PyTorch, we can [68:00] ingest it and compile it uh to all of [68:03] our supported backends. [68:05] So, we take a matmul, for instance. We [68:07] don't even have a matmul operation. Uh [68:09] we take a matmul and we decompose that [68:11] into a broadcasted multiply and then a [68:13] sum reduction. [68:15] Um so, yeah, this this really helps us [68:17] to bottleneck the complexity of new [68:19] models in PyTorch, of which, you know, [68:21] PyTorch has something like 2,000 [68:23] operations, um and and, you know, new [68:26] models are very complex, and hardware [68:29] implements tons of different operations [68:31] as well. And so, doing this complex [68:32] mapping, you don't really want to go [68:34] all-to-all. You want to have some sort [68:36] of a bottleneck in between with which [68:38] you can map all your PyTorch operations [68:40] down to and then map back down into [68:43] hardware operations. And so, this [68:45] RISC-style primitive op set is that [68:47] bottleneck for us. [68:50] So, [68:51] there's a kind of term in the compiler [68:53] world called super optimization. Uh it's [68:55] from, you know, the '90s. Uh it it [68:58] didn't really ever take off except for [69:00] in compiler research. Uh but basically [69:02] the concept is that you should [69:04] try to search through every possible [69:07] form of a given program and use some [69:09] cost function to determine which is [69:11] best. And that way you don't need to, as [69:13] the compiler engineer, know how to lower [69:16] some given program onto a hard on a [69:18] piece of hardware uh ahead of time. [69:20] Because there is an exponentially, you [69:23] know, large number of possible ways you [69:25] can do that lowering. And so, what you [69:27] often see from these companies is they [69:29] hire compiler teams and their compiler [69:31] team head count uh scales exponentially [69:34] with how many different applications [69:36] they want to target. So, you end up with [69:37] a compiler team of 500 or 1,000 people. [69:40] Uh and even then, a very imperfect [69:41] compiler. And so, [69:44] once you try to scale up the number of [69:45] models or types of model architectures [69:47] you want to support and the types of [69:48] back ends and the types of data types [69:50] you want to support, you get this big [69:52] multiplicative explosion. [69:54] So, what you really want to do is you [69:56] want to write these very small atomic [69:58] rewrite rules that kind of specialize or [70:00] that kind of specify uh different [70:02] equivalencies. And then you want to let [70:04] the search go and figure out what to [70:06] apply when and and not not tell it to [70:09] apply these these different rewrite [70:11] rules. So, [70:12] what equality saturation lets us do is [70:14] do exactly that. Write these very small [70:16] equality rewrite rules, build up the [70:19] search space, and then go and walk [70:20] through that search space. [70:22] So, we have a compute graph as our IR [70:26] structure. So, it's a DAG. You can [70:28] picture just like, you know, your your [70:29] add operations, your multiplies, and you [70:31] get your dependencies between those ops. [70:34] And we use this this concept called an [70:36] E-graph. So, it takes a graph and it [70:39] applies these non-destructive rewrite [70:41] rules, which means that it an ordinary [70:43] rewrite rule does a pattern match onto a [70:46] graph and then replaces what it matched [70:49] with with something else on the [70:51] right-hand side of the rule. And so, [70:52] that effectively destroys the [70:54] information that it matched on. And so, [70:55] when you apply that rule, you better be [70:57] sure that that was the right rule to [70:59] apply because there's really no going [71:00] back. And what we're allowed to do is [71:02] we're allowed to say, "Well, we can [71:03] apply a rule and we're just going to add [71:05] the right-hand side information to this [71:07] e-graph, but this e-graph kind of [71:09] maintains the both sides of the [71:11] equation, what came before and what [71:12] comes after." [71:13] And so, we basically apply all of our [71:15] rewrite rules in every possible order [71:17] until they can add no more information, [71:20] and that effectively has grown the [71:21] e-graph as much as it can go. We call [71:23] that saturated. And once that saturated, [71:26] then we actually go and we extract a [71:28] whole bunch of different concrete [71:30] graphs, which represent programs. We [71:32] compile those all down to executables on [71:34] the hardware, and then we actually go [71:36] and run those on the hardware, and we [71:38] basically time how long it takes to run [71:41] on the hardware, and we go and try to [71:42] minimize that time. It's very [71:44] straight-forward and simple uh process. [71:47] But this basically is a kind of visual [71:49] representation of what an e-graph [71:50] actually looks like. So, [71:53] you kind of have uh if you look on the [71:55] very left-hand side there, uh you have a [71:58] graph that represents the equation A * 2 [72:01] / 2. It's kind of uh difficult to to [72:04] imagine, but basically the top node [72:06] there is the root of your expression, [72:08] and then you kind of go down, and one of [72:10] the arrows coming out of the root is the [72:11] one side of the division, and the other [72:13] is the other side. And so, you have A * [72:15] 2 / 2. And what we're doing here is [72:17] we're actually applying different rules, [72:19] and these applications grow the graph. [72:22] So, they specify that, "Okay, well, [72:24] actually A * 2 is equivalent to A bit [72:27] shifted over by 1." [72:29] Um and then you can keep applying these [72:31] different rules, [72:32] and eventually what you get is the fully [72:34] saturated graph on the right-hand side [72:36] of the picture there. And what that does [72:39] actually is it has expanded the dotted [72:42] line that contains both that division, [72:45] which we started out as our root. That [72:47] that box there contains A as well. And [72:50] so, what we've actually worked out is [72:51] that a * 2 / 2 is equivalent to just a. [72:56] So, it's kind of a way that we're able [72:58] to simplify down this expression without [73:00] having known, you know, do a do a [73:02] heuristic-based rewrite rule that has [73:04] looked for that pattern ahead of time. [73:07] And so, this is kind of what our rewrite [73:09] rules look like. They're very, very [73:10] straightforward and simple. They're [73:12] written in a Lisp syntax. The reason for [73:15] that is because the people who made the [73:18] engine that we use inside our compiler [73:20] really like Lisp. And so, we have to [73:22] write all our rules in Lisp. [73:24] But, they they kind of look like that. [73:26] So, you kind of see the first rule there [73:27] is is basic commutativity. So, you know, [73:30] x + y = y + x. Um [73:34] Notably, that is not true with [73:35] floating-points. But, that, you know, [73:37] the associativity and commutativity are [73:40] kind of tricky when you get to IEEE [73:42] floating-points, but that's another [73:44] topic. Um and then, you know, what you [73:46] actually have in the bottom there is a [73:47] technique called constant folding where [73:50] if you have like a + b and at compile [73:52] time you know what a + b are, you can [73:54] actually just go and add them and then [73:56] remove that entire addition from your [73:57] program. You don't need to run that at [73:59] runtime because you can do that at [74:00] compile time, so to speak. [74:03] So, [74:05] you know, we've kind of talked about the [74:07] our core IR set. And that core IR, that [74:09] that very simplified primitive IR set is [74:12] shared across all of our back-ends. But, [74:14] we want to actually go and compile to a [74:17] bunch of different back-ends. So, we [74:18] want to be able to compile to TPUs and [74:20] training ins and GPUs. And so, how do we [74:23] actually go and specify each chip's [74:26] implemented operations? So, what a [74:28] hardware does is it kind of brings a set [74:30] of implemented operations. And these are [74:32] generally quite low-level. So, they [74:34] could be in a kind of varying levels, [74:36] but they could be high-level, they could [74:37] be low-level. But, on a GPU, you can [74:39] imagine that, you know, a tensor core [74:41] MMA is one of the operations that you [74:43] can specify. And that kind of reads [74:45] inputs from SRAM and writes outputs to [74:48] SRAM. [74:49] Um [74:51] and so you we go in and we can well [74:53] characterize a given new hardware [74:55] backend that we want to target [74:57] uh and then we go and specify, "Okay, [74:59] what do each of these operations [75:00] actually mean in terms of equating them [75:03] to operations that our compiler already [75:05] understands." So, we can write those [75:07] little rewrite rules to do that equality [75:09] and that kind of propagates out to the [75:12] rest of the equalities that we've [75:13] already specified in the compiler. So, [75:16] bringing in a new hardware backend is a [75:17] fairly low-effort thing even if it looks [75:20] very different from the hardware [75:21] backends we already support. [75:24] So, yeah, this all gets around the the [75:27] typical need for a very massive compiler [75:29] team and a lot of complex rewrites. [75:32] So, and you know, hardware fundamentally [75:34] is hierarchical, meaning you have a GPU [75:37] and that GPU, you know, you're actually [75:39] possibly programming a single thread, [75:41] you may be programming a warp, you may [75:43] be programming a thread block, which is [75:45] essentially programming a full CUDA core [75:48] or GPU core and SM. Uh or you may be [75:50] programming an entire GPU or you may [75:52] even be programming an entire node with [75:54] an operation. And so, we want to be able [75:56] to represent that explicitly. So, our IR [75:59] is fundamentally multi-level. Uh we can [76:01] model all of those different levels [76:03] explicitly and that lets us bring in [76:05] operations that may come from the [76:07] community, from vendor libraries that [76:09] are implemented at these different [76:11] levels. And so, equality saturation lets [76:13] us kind of bridge all of these different [76:15] semantics in this like sane way and just [76:18] specify what is equal to what and then [76:20] have the engine swap out and do all this [76:22] do all this search. [76:25] And so, this is kind of what it looks [76:26] like uh when you swap out a bunch of [76:28] different implemented operations. And [76:30] so, uh you know, you can have [76:33] a a cuBLAS [76:36] uh host operation that basically is [76:37] called from the CPU and it does a full [76:39] matrix multiply. Uh and cuBLAS is a [76:42] matmul library uh authored by Nvidia and [76:45] it implements very good matmuls. Uh you [76:47] can have a thread lo- thread block level [76:50] operations like uh as implemented in [76:53] CUTLASS, which is another Nvidia [76:54] library, or CuBLAS DX, which is also an [76:57] Nvidia library, or ThunderKit, which is [76:59] from Hazy Research. [77:01] Uh and those are thread block level [77:02] operations, meaning they actually [77:04] program one GPU core at a given time. [77:07] And our compiler, because it can reason [77:08] about these different levels of [77:10] operations, we can actually go and take [77:12] subgraphs of lower-level operations and [77:15] kind of combine them all into a big mega [77:16] kernel. And so, [77:18] we can have at any given point a forward [77:20] pass that is implemented in both kernels [77:23] and mega kernels. And so, that's great [77:24] because as much as you know, mega [77:26] kernels are neat and cool and everyone [77:28] likes to talk about them, they don't [77:29] always uh they aren't always the optimal [77:31] solution, especially when you're dealing [77:33] with very large models. [77:37] And so, [77:38] you know, the final kind of problem here [77:40] is once you've defined your search [77:42] space, how do you actually search it? [77:44] Meaning, what do you know? How do you [77:46] say program A is better than program B, [77:48] so to speak? And this is kind of a [77:49] problem uh where you you need a cost [77:51] model, essentially. You need some [77:53] fitness function to evaluate a program [77:55] and say this is more expensive than [77:57] that. [77:58] And having a good cost model, especially [78:00] of complex hardware, you know, take a a [78:02] modern GPU or something, [78:04] is extraordinarily hard. This is like [78:06] many, many, many PhDs worth of work to [78:09] in order to, you know, take a GPU [78:11] program and say like how long do you [78:12] actually take to run um without actually [78:15] running it. And so, we as a little [78:18] little company here can't possibly uh do [78:21] accurate cost models for every type of [78:23] hardware we want to target. And so, you [78:25] know, in the spirit of the bitter lesson [78:26] here, the best cost model is no cost [78:28] model and we should actually just go and [78:30] run these things on the real hardware [78:32] that we're going to run on and see how [78:34] long they take. And it's, you know, it's [78:36] kind of a dumb solution, but it's a very [78:38] very effective solution. Uh Where you [78:40] know, we can have different different [78:42] moves in our search space that we may [78:44] not as programmers think are beneficial, [78:46] but when we actually go and run them on [78:47] the GPU, they end up being very [78:49] beneficial. [78:51] Um [78:52] So yeah, we can do we don't actually [78:54] have to run everything in our search [78:55] space. There are techniques where you [78:56] can prune the search space because you [78:58] know some roofline evaluation of a given [79:01] subprogram and then you can kind of [79:03] early prune it because you know that it [79:04] will be in the best case slower than [79:08] something that you've already measured [79:09] that is equivalent to that. So we do we [79:11] do some early pruning here and there, [79:12] but you know, largely we try to measure [79:15] everything on real hardware. [79:18] And oh yeah, this is a little video. [79:20] This is just [79:21] what it looks like when it's code [79:22] generating a kernel. [79:24] There's not a whole lot of information [79:25] you can get from this video, but it's [79:26] just like kind of neat to see. [79:29] But yeah, and obviously obligatory plug. [79:32] We have done the very rare feat of [79:34] solving compiler company monetization. [79:36] There's a very long list of dead [79:38] compiler companies in the graveyard [79:40] because they haven't figured out how to [79:41] make money from a compiler. [79:43] We've finally closed the loop on [79:44] profitability. But yeah, like I said, we [79:47] run a cloud powered by a bunch of [79:49] different heterogeneous hardware. We're [79:51] at yeah, something like 80 trillion [79:52] tokens a month, [79:54] many many millions of images a month, [79:56] and one of the benefits of being a [79:58] compiler like general purpose compilers [80:00] that we can not just run the most [80:01] popular models, but we can run all these [80:03] sorts of wacky models. So we run protein [80:06] folding models, bio models, image [80:08] models, all this all this all this [80:09] stuff. [80:10] We're growing really really rapidly. So [80:12] we're eight people. We're not staying [80:14] that way. So we are we are aggressively [80:16] hiring. And then [80:17] we use wacky compiler tech to [80:20] commoditize wacky hardware. And [80:23] yeah, we think the future of compute [80:24] compute is commoditized. So come visit [80:27] us at lumina.com. [80:30] >> [applause] [80:37] >> Hey, [80:38] I really like your risk style approach. [80:41] I was wondering how much of you adapted [80:43] to [80:44] machine learning workloads because [80:47] you don't need for instance you [80:48] mentioned that you also network about [80:51] associativity in IEEE 754. [80:54] But ML doesn't care. You can even do [80:57] different sequences hardware style. How [80:59] much of that do you bake into your does [81:02] that enter into the cost model or how [81:04] you model the language itself? [81:06] >> So, you know, associativity [81:07] commutativity [81:09] they do matter so much as they do affect [81:11] numerical accuracy. [81:13] Meaning if you reduce your you know if [81:15] you're doing a sum reduction and you're [81:17] doing a tree sum reduction that is going [81:18] to give you a different answer than a [81:19] purely sequential sum reduction. And you [81:22] know to the point about provider model [81:25] model quality if you screw up numerical [81:27] accuracy that really does impact [81:30] your your output token quality. So, we [81:32] do we want to be very very explicit in [81:34] how we model numerical accuracy. [81:37] Um [81:37] and then in in the reasons why some of [81:40] those rewrite rules still are [81:41] advantageous though is that [81:43] often times you have hardware that [81:45] implements different functionality for [81:46] input A versus input B. So, you know, [81:49] for instance Nvidia tensor cores may [81:51] only actually consume input A from SRAM [81:53] and input B from you know team M or [81:55] something like that some other memory [81:57] space. [82:00] >> Hey, uh this is uh [82:02] Yeah. [82:03] So, [82:04] I didn't quite understand when you say [82:06] why can't we can just run it on the [82:09] actual hardware because that seems [82:10] impossible if you're talking about the [82:12] entire model. [82:14] And if you are only doing that for a [82:16] particular kernel, [82:17] I don't know how you model contention [82:19] and stuff like that. [82:20] Um so, that's part one and part two [82:23] you [82:25] meaningfully expanded search space uh a [82:27] much larger search space, right? [82:28] >> Mhm. [82:29] >> Yep. [82:31] >> Even like a stupid compiler like a TRT, [82:33] for example, is extremely slow, even [82:35] with this heuristic. [82:36] So, how does this search scale at all [82:38] given like you have basically expanded [82:39] it to be an order of magnitude bigger [82:41] from what I can tell? [82:42] >> Yeah, so two different questions, but [82:44] the first answer would be that the is it [82:46] still fairly quick to run a full model [82:48] pass like on the order of I guess it [82:50] depends what you consider to be slow, [82:52] but we can run a forward pass, you know, [82:54] and do a profiling iteration on the [82:55] order of like under 10 milliseconds. So, [82:58] you can still do, you know, a thousand [83:00] odd profile runs in 1 second. And so, [83:04] that's still like reasonably quick. [83:06] Um for the second question, yes, the [83:08] search space does get quite large and we [83:09] use pretty standard search algorithms to [83:11] walk through that. So, we use like a lot [83:13] of like genetic algorithms. I think [83:15] there's a lot of room for innovation [83:16] here where we can use something like [83:18] like MuZero uh to do more like a learned [83:21] guided search, but so far we haven't had [83:22] to do that. We end up only searching [83:25] like 1% of 1% of the space, but uh and [83:27] people are I think surprised to hear [83:29] that, but we never are targeting a [83:31] global optima. We are targeting a local [83:33] optima that's like 99% as good as a [83:36] global optima. And so, that that ends up [83:37] being much easier to find. And then in [83:39] terms of compile time, um [83:42] that is an explicit trade-off we've made [83:44] that customers in inference don't really [83:46] care about compile time if it gets them [83:49] a faster run time. And so, because you [83:51] compile once and then you serve across, [83:53] you know, millions of GPUs, uh if it [83:56] gets you 1% or 5% or 10% better run time [83:59] performance, you won't you don't care if [84:01] you compile for a full day [84:02] um sometimes. And so, the way we [84:04] generally model this is we let the user [84:06] set a time budget or we set a time [84:07] budget. So, you can compile for an hour, [84:09] you can compile for 10 hours. [84:11] Eventually, you'll hit some point of [84:12] diminishing returns, but that's kind of [84:14] how we we think about this. And scaling [84:16] it oftentimes it scales pretty well. [84:18] Again, you hit some diminishing returns [84:19] after after quite a while, but it ends [84:22] it ends up working fairly well. [84:24] >> Thanks. I think the the kicking this [84:25] profiling makes makes a lot of sense. Do [84:27] you guys use actual profiles from [84:30] inference as well to sort of guide the [84:32] search or this is just pure [84:33] um on in form search in some sense? [84:36] >> So we try to use as close to real inputs [84:38] as possible. Unfortunately, we run a [84:40] cloud so we know the exact kind of [84:42] shapes that customers care about. So we [84:43] you know in terms of tensor shapes, we [84:45] know roughly exactly what they're going [84:47] to feed us at run time. [84:49] Um [84:50] and uh you know data distributions and [84:52] such like that. And so we don't need to [84:54] check for correctness because we already [84:56] know that you know unless our rewrite [84:58] rules are wrong that we're going to get [84:59] the correct answer out but uh as long as [85:01] we try to represent the shapes to be [85:03] this like industry relevant tensor [85:04] shapes, we end up with pretty accurate [85:06] timing. [85:09] >> Uh hi, thank you for sharing and uh [85:11] another compiler need to compare the [85:13] different archi different model [85:14] architecture to different hardware and I [85:16] feel like it's pretty hard to do the [85:18] debugging especially we need to identify [85:21] the which layer that we're trying to [85:23] lowering the ops health issue. So I'm [85:26] curious about how uh Luminos solve this [85:28] issue. Thank you. [85:30] >> Yeah, I mean general debugging is always [85:31] very difficult. Um we use a lot of tools [85:33] that PyTorch provides where they're [85:35] they're you know really nice about [85:36] tagging uh the high level ops and then [85:38] we kind of track that metadata as it [85:40] flows through our compiler. Uh we use a [85:42] lot of like the standard PyTorch uh [85:44] profiler uh to go and and work out you [85:47] know bottlenecks and issues. And then [85:49] what a quality saturation does, what [85:51] egglog specifically does for us quite [85:53] nicely is it lets us say, "Okay, we got [85:55] an output program out and this program [85:57] is wrong. [85:58] Give me a proof that connects this [86:01] output program back to our input program [86:04] which basically says like here are the [86:06] chain of rewrites that were applied to [86:08] go from the input to the output." And [86:09] then we can narrow down pretty quickly [86:11] like which which rule is invalid [86:13] effectively. [86:22] >> Yeah, um, I actually like this solution [86:23] a lot because usually unsupported [86:26] operations we just kick them back to the [86:27] CPU very elegantly. Um, but because [86:30] there is no free lunch, what kind of [86:32] latency, because now we're expanding an [86:34] operation when we usually want to fuse [86:36] operations at the compiler level, what [86:38] kind of latency impact do you usually [86:40] see? [86:42] >> So, we absolutely do fusion. Um, we [86:44] expand them to get to this like core set [86:46] of primitive IR, but then when we [86:48] actually go and do lowerings, uh, the [86:50] the compiler searches through various [86:52] forms of lowerings, uh, and what we end [86:54] up with is, yeah, quite fused ops as [86:56] well. And so, we generally try as a, you [86:59] know, roofline heuristic, like you can [87:01] see, okay, if this kernel is is moving [87:04] uh, in order of magnitude more data back [87:06] and forth from DRAM than this other [87:07] kernel, then like, yeah, probably that's [87:08] bad. Um, so we end up with quite fused [87:11] operations. It's just that in the [87:13] mid-level it kind of expands those out, [87:15] and then our, you know, kind of search [87:17] has to go and find those patterns and [87:19] refuse them back down again. [87:20] >> Okay. [87:21] >> All right, thank you. [87:27] >> So, one of the comments I picked up on, [87:29] uh, during, uh, [87:30] the the slides was that the language you [87:34] work with is actually not Turing [87:36] complete. Um, and that's kind of [87:38] interesting because that's that says [87:41] it's not going to be subject to like [87:42] Rice's theorem and things like this. And [87:44] so, you can reason about, uh, things [87:47] like what will the program actually do [87:49] without without even running it, at [87:50] least for for some properties. And so, [87:53] I'm really curious, cuz I don't think [87:55] you mentioned this by accident, so I'm [87:56] very curious where in the stack do you [87:58] exploit this, and how? [88:00] >> So, [88:02] very good point. It is not subject to [88:04] Rice's theorem, and this was kind of an [88:06] early realization we had because this is [88:08] one of the reasons super optimization [88:09] failed in the '90s is that they tried to [88:11] apply this to general-purpose [88:13] programming languages, meaning like C. [88:16] And C is Turing complete, and there's a [88:18] near infinite number of possible [88:20] compilations. And and so, you know, you [88:22] face a lot of issues when you try to [88:24] start reasoning about Turing complete [88:25] programs. We compile dense linear [88:28] algebra programs, and that's a very very [88:30] narrowed space versus like all possible [88:32] programs. So, you know, again, PyTorch [88:34] complete but not Turing complete. Uh one [88:36] of the exploitations we can do is we can [88:38] know exactly, you know, analytically how [88:41] much uh byte movement we do between [88:43] different memory spaces. And that's how [88:45] we're able to analytically derive this [88:47] like roofline uh of performance without [88:49] actually running uh a model. And so, we [88:51] actually do a lot of early pruning, and [88:53] we end up not needing to run very many [88:56] uh profiling iterations relative to the [88:58] full size of the search space. [89:02] >> Okay. Thank you. [89:03] Uh [89:04] >> Uh [89:06] >> Yeah. [89:07] So, um [89:11] the solution space that you're searching [89:13] for, does it also include optimizations [89:15] like quantization, speculative decoding, [89:17] and [89:18] parallelisms? [89:20] >> Uh I'm sorry. I didn't I didn't hear [89:21] you. [89:21] >> Sorry. Uh [89:23] so, I guess that the like the solution [89:25] space that you're searching inside is at [89:27] a kernel level. Does that include uh [89:30] other techniques like speculative [89:32] speculative decoding uh the the [89:34] >> So, generally no. We try to basically [89:37] just handle the forward pass. So, the [89:39] you know, you you specify like a forward [89:40] pass graph, and we compile that. And so, [89:43] you know, it's end-to-end within the [89:45] span of a forward pass. And then we live [89:47] inside of, you know, the great uh [89:49] inference engines of vLLM and SGLang. [89:52] And so, the great solutions they already [89:54] provide for speculative decoding, we're [89:56] able to, you know, just go and directly [89:57] piggyback on them. And our, you know, [90:00] thesis is that the large part of the [90:02] alpha lives in being able to handle very [90:04] complex, you know, parallelism [90:06] topologies and being able to do this, [90:09] you know, handle the complexity of the [90:10] forward pass and then leave the rest of [90:12] the innovation out to, you know, VL and [90:14] SG length. So, we don't really search [90:15] that level. [90:20] >> Cool. Thank you so much, Joe. [90:22] >> Thank you. [90:25] >> [applause] [90:29] >> Uh hi everyone. Uh I'm Leon Shaw and I [90:32] am president of SG length and ready [90:34] start. And I'm very honored to uh [90:37] be invited here. Uh I thought this is a [90:40] like [90:41] paper reading club. Yeah. Uh it's like a [90:44] previous uh talking a lot of uh [90:46] technologies rather than a single paper. [90:49] I actually uh prepared maybe [90:52] for a 50% of the contents talking about [90:55] the SG length's original paper and other [90:56] 15% talking about our open source and [90:59] our [91:00] uh current uh optimizations, but I I [91:03] just uh added a [91:05] uh few more pages, but uh yeah, I can uh [91:09] it's not in the uh slides. Yeah, I can [91:11] start. So, [91:13] uh [91:15] So, uh that's uh the story begins when [91:20] uh the agent workload. Actually, the [91:22] story begins in the [91:23] uh in the late 2023. At that time, uh we [91:26] only have the Llama. Uh we all have like [91:29] for all the open source models. So, uh [91:32] at that case, uh people just a guess [91:34] what the uh agent workloads looks like. [91:37] So, at that time, we imagine that there [91:39] will be some tool calling and we imagine [91:41] there are some parallel sub-agents uh [91:43] for uh for your like something like uh [91:47] the future cloud code to do that. Uh [91:50] So, imagine when you fork a agent uh [91:53] fork a sub-agent in the cloud code, [91:54] there will be a lot of a context. [91:56] They're sharing the a big uh context uh [92:00] in in the same KV pool. And when you are [92:03] the monitor to query each the query [92:05] result and each the prompt you entered [92:08] is a very short prompt that it will be [92:11] uh get preview into the age into the [92:14] engine multiple times and get out of [92:17] of the engine multiple times. So, at [92:19] that case we notify uh like [92:22] in that at that time there is [92:25] there is no way engine can serve the [92:28] prefix reusing efficiently. So, at that [92:31] time we brought the Radix attention uh [92:34] which can automatically reuse the share [92:37] prefix and compute with the new [92:40] with the new suffix. [92:42] So, uh the idea is we propose the Radix [92:46] attention. Uh all the agents' contacts [92:48] they were uh managed inside the engine [92:51] system [92:52] in the memory system uh just like a [92:54] tree. So, we call it like a Radix tree [92:57] in the Astra inner implementation. Uh [93:00] so, here is the illustration for how we [93:04] how we do the memory management. It's [93:06] just like the figure says. Uh uh we we [93:11] share the setup like there is a long [93:13] system prompt and uh uh one do the [93:17] authentication, one do the review and [93:19] there two agents they they share the [93:21] same prompt to do the same [93:23] with the with the same within the same [93:25] context and do the different things. [93:27] Then uh one agent they may want to give [93:30] some results and because the memory is [93:32] limited, so why uh the suffix of one [93:36] agent that got evicted. So, that is [93:37] basically how we do the uh memory [93:41] management. It's it's it's really a [93:43] Radix tree in our management system and [93:45] all our uh uh all our core designs [93:49] including our scheduling are around this [93:52] topic. [93:54] So, you [clears throat] can see [93:56] uh [93:57] when when the uh [94:00] uh [94:01] Uh so, the left side is how we uh do the [94:05] uh do the prefix protection uh for the [94:07] agents. Uh [94:09] there is a uh [94:10] when [94:12] uh [94:13] So, so we do the eviction. Uh the [94:15] eviction policy is we uh we evict the [94:17] unused leaves when the memory is needed. [94:20] And uh there will be uh a lot of uh [94:22] algorithms uh the scheduling eviction [94:24] algorithm uh in our [94:27] in our scheduling uh base to to fit your [94:31] uh workload and fit your needs. And the [94:33] right describes why we need this. Uh for [94:36] example, uh we have uh uh the prompts or [94:39] the agents. They are sending to the [94:41] engine uh interve- in interleavingly. [94:45] So, if you can uh reorder uh the prompt, [94:48] then they can be uh batched in uh [94:51] cache-aware order uh then to reuse the [94:54] context uh A rather than uh interleave [94:57] the A and B uh [94:59] then there may be the risk of eviction. [95:04] So, uh that is the story how Islam [95:07] becomes the a serving engine. Uh [95:11] Uh [95:12] The original paper is an idea to build a [95:16] red tree inside the memory management [95:18] system. And uh the second page is we uh [95:21] make this uh universal and general uh LM [95:24] inference engine uh in the middle 2024. [95:27] Uh so, uh we at the mid-2024, we [95:31] published the 0.2 version. Uh and uh [95:34] with the with supporting to uh at that [95:37] time, uh the Llama 8B2 Llama uh 4.405B [95:42] uh for both online and offline serving. [95:45] >> [snorts] [95:45] >> So, you may uh you may notice that for [95:48] uh for all the for all the uh batching [95:50] operations, for all the [95:52] uh memory uh management eviction [95:54] matching. Uh you have to order all the [95:57] waiting request, uh get the priorities [95:59] for them. [96:00] Uh so, and we know that uh the GPU their [96:04] inference time is getting faster and [96:05] faster. Uh but but at the meanwhile, the [96:08] the CPU they didn't um [96:10] their efficiency didn't change. So, [96:13] uh for this time we can say uh we can [96:16] we can notice that uh if we use the [96:20] vanilla [96:21] uh scheduling uh strategy, then there [96:23] will be a preparation for the batch one, [96:26] then run the batch one, then get all [96:27] data, then preparation for the batch [96:29] two. The so so, when the GPU is quite [96:33] slow, it's okay. Uh the GPU bubbles is [96:36] uh acceptable, normally with within 10%, [96:40] but uh when the GPU they are getting [96:43] faster and faster, [96:44] uh the CPU the the bubbles the GPU [96:47] bubbles by the GPUs bottleneck has been [96:50] uh [96:51] has become noticeable. So, for this [96:53] case, we [96:55] we designed our uh uh scheduling system [96:59] uh with the overlap scheduling, uh which [97:02] means uh we we delay all the uh device [97:06] to host the synchronization, [97:08] uh like we last it for another round, uh [97:11] to make sure that there's no blocking [97:13] operations uh in the critical path. Uh [97:17] at the meantime, uh we do the [97:19] preparation in advance. Uh because the [97:21] GPU they are running, uh there's nothing [97:23] blocked by the I Once we do the delay uh [97:26] operations, then there is no [97:28] uh [97:29] there is nothing uh [97:30] there is nothing blocking uh the next [97:33] batch launch. So, we just uh launch the [97:36] batch two uh when the batch one is [97:38] running. So, uh we still using this, and [97:41] a lot of uh uh [97:43] uh new engines in uh including uh the [97:46] yesterday uh DeepSeek [97:49] DeepSeek they they published their [97:51] DeepInfra. They also used a [97:53] uh [97:54] the uh the zero over the zero over life [97:57] scheduling uh to support their internal [97:59] inference engine. Uh [98:02] Uh [98:03] Like uh in the Hopper uh [98:06] in the Hopper days [98:07] uh it's only in uh 1.5 uh 1.1 uh [98:12] uh times of throughput comparing to the [98:14] vanilla scheduling, but uh in the [98:16] Blackwell days uh it's it's maybe [98:18] sometimes it can achieve two times in [98:20] the GB settings. [98:24] Uh [98:25] So, [98:26] uh the agent and the cache story [98:29] continues [98:31] uh because uh our GPU memory is always [98:34] limited. Um [98:36] when we are doing serving uh with long [98:39] contexts and the people there [98:41] uh like uh sending their prompts, [98:43] sending their next round of response uh [98:46] within few minutes or few hours, even [98:48] few days, [98:49] uh you have to [98:51] uh find some place to restore the [98:54] uh their previous uh prefix uh [98:57] somewhere. [98:58] Uh so, that is why we bring the high [99:00] cache uh which is also um pretty [99:03] standard uh technology used in many uh [99:06] engines. Uh we connect the register I [99:10] mentioned before [99:11] uh with the GPU memory, host memory, and [99:13] the internal storage we call it L1, L2, [99:16] and L3 cache. Uh every token they can be [99:21] uh exist uh existing on the L1 cache or [99:25] L2 cache or L3 cache. Uh all these they [99:28] can managed by [99:30] uh [99:30] managed by some uh specific policy to [99:33] increase uh the cache hit rate and also [99:37] uh the the loading time from the L2 [99:39] cache to the L1 cache or from the L3 [99:42] cache to the L2 cache. [99:43] >> [snorts] [99:44] >> Uh in our experiment, this can lower the [99:48] TRFT by 56%. [99:50] >> [snorts] [99:52] >> So, [99:53] uh as as [99:55] as nowadays, [99:57] uh [99:58] like [99:59] the the traditional transformer [100:02] doesn't make sense for for long context [100:04] and the for efficient deployment. So, uh [100:07] there will be a more and more uh things [100:10] things about the [100:12] some some some innovative designs about [100:15] the transformer architectures. For [100:17] example, there is a [100:19] sliding window. [100:21] As we all know that a lot of open source [100:23] and maybe closed source model they use [100:26] the open window they use the sliding [100:28] window to to to serve their models. And [100:30] also there is a current recurrent states [100:33] like a member state, KDH states. They [100:35] were they were just putting [100:38] the information into the fixed shape of [100:41] state like a token by token. So, at that [100:45] time you may think that uh [100:47] like it's very easy we just need to add [100:49] different layer at different kind of [100:51] memory pool to do all the hybrid memory [100:54] pool, hybrid storage, hybrid [100:56] computation. Uh but that is not so [100:59] so so trivial because in that case you [101:03] have to consider [101:04] how to build the the cache share uh [101:08] strategies. [101:09] Uh [101:10] Uh for example, if you want to [101:14] if you want to cache hit, let's see this [101:17] figure. If you want to cache hit at the [101:20] position 10, uh let's starting from one. [101:23] Like you have to match the full KV [101:26] tokens for the full attention pool. And [101:28] also if the in in the example, [101:32] if the if the sliding window size is [101:34] four, you have to contain at least the [101:37] four tokens uh starting reversibly from [101:40] 10 uh to seven. Uh if the if the if the [101:45] retained token is only reaches eight, [101:48] then there there will be a no cache hit [101:50] because it's new it's need a full [101:52] attention calculation for the sliding [101:54] window. [101:55] And also, if you want to start share the [101:57] prefix from the position 10, you have to [102:00] have the uh member state KDS state [102:02] exactly at the position 10 to share the [102:05] prefix. So, uh [102:08] so if we [102:10] uh manage all the memory pools [102:12] separately, then there will be a huge [102:14] challenge to how to uh do the efficient [102:17] cache sharing. Imagine when you do the [102:19] cache sharing from uh 95% to 99%, you [102:23] reduce the uh the expense by by five [102:27] times, not by only 5% because it's you [102:31] uh use uh four times less computation [102:35] for all the inference uh for all the [102:36] inference workload. Uh so, so [102:40] previously, we used different um radix [102:44] tree like a which was implemented uh [102:47] separately stand-alone [102:49] uh without consideration of these [102:51] different uh low-level uh memory pool. [102:54] And uh for that case, this will [102:58] uh [102:59] uh bring a lot of compatibility and [103:02] flexibility issues. So, [103:05] uh also, uh when you try to uh [103:09] try to protect your history, try to lock [103:12] the uh prefix node, you also have to do [103:14] a lot of coordination uh across the full [103:17] layers, across the uh [103:19] uh sliding window layers, and across the [103:20] recurrent layers. Uh each uh it seems [103:24] like uh for each component, they have to [103:27] reach the [103:28] uh some some specific protocols to to to [103:32] coordinate to to satisfy an external [103:35] policy. So, [snorts] [103:37] that is a No. Oh, that is a why we why [103:40] we brought the unified radix tree. [103:43] Uh in the newly designed unified radix [103:46] tree, we didn't [103:48] we didn't treat all the [103:50] components like [103:51] as a as a standalone components. We used [103:55] We used a semantic [103:57] token ideas tree to [103:59] which was based on the full token. [104:01] Because [104:03] based on full token, that is an [104:04] interesting because [104:06] what matter [104:07] how the architecture changes, there will [104:09] always be the full tokens. Yeah, at [104:11] least in nowadays. So, we built the [104:13] semantic token ideas [104:15] uh [104:16] radix tree for the full tokens and all [104:18] the other tokens, they were based on the [104:22] based on the internal maintained [104:24] component. Uh and the the [104:29] the [104:30] the unified radix tree will be a [104:32] controller and a coordinator [104:35] for each component to make sure we can [104:37] achieve the best [104:39] cache matching efficiency. For this, [104:42] and also for each component, they were [104:45] designed to fit the unified radix tree. [104:48] They have their own rules. They have [104:49] their own matching eviction and the [104:52] memory rules. Uh yeah. [104:55] So, this will bring all the all the [104:59] corner cases together to avoid any [105:03] memory leakage or memory waste [105:06] in the [105:07] uh [105:09] in the virus workloads of edge edge [105:11] analytic inference. [105:12] Also, [105:14] there's something I didn't post on the [105:16] slides, but you may notice that uh [105:19] if we have the the three types of heavy [105:22] posts like the full token, the window [105:23] token, and the recurrent states like a [105:26] states. But, [105:27] uh you can you can imagine like when the [105:30] context they varies, uh they they varies [105:33] across different users or different [105:35] workloads or even different time zones. [105:38] Uh for example, a user A, they only use [105:41] context around 4K. Uh user B, they use [105:44] context maybe 128. And another user, [105:47] they use context for 1 million tokens [105:49] always. So, at that case, they always [105:52] require the full token pool to feed [105:54] their context. But, the sliding window, [105:56] as for some purpose, it only feed the [105:58] last 128 tokens for their working load. [106:02] So, that is a ratio difference because [106:05] because you cannot reserve the [106:07] uh token pool uh by the fixed ratio. If [106:10] you have to design an engine to fix all [106:13] the different ratios of diff- all the [106:14] different to workload context, then it [106:17] may uh may you have to use a lot of uh [106:20] different types. But, uh if you only [106:22] have one engine or limited resources or [106:25] you want to see you want to serve a lot [106:26] of different workloads, then you must [106:29] have uh you must to have some uh [106:32] some unified way to to to to manage all [106:35] the all the low-level memory pools. So, [106:39] so besides the unified Redix tree, we [106:41] also proposed our uh unified memory [106:45] pools to serve all the different [106:47] workloads. Uh before, uh we just uh [106:51] pre-allocated all the workloads at the [106:53] at the precomputed ratio and uh and the [106:57] fixed the size is. And uh but for now, [106:59] we only allocate one uh one [107:02] uh one uh [107:04] uh like a [107:05] like a [107:06] like a one pile of a uh physical memory, [107:09] then we we do not uh set the limit for [107:14] each kind of uh uh [107:16] memory pools. Like we didn't we do not [107:19] set the limit for the 4K. We we did we [107:22] do not set the limit for the window KW. [107:24] We just treat them as all of the [107:26] combination of the smallest units like [107:28] bytes. [107:29] They were just [107:31] keep growing from the left side of the [107:34] physical memory and the from the right [107:35] side of the physical memory they were [107:37] like growing [107:39] like [107:40] according to our usage. [107:42] At that time [107:44] uh [107:44] when the when the contest getting like a [107:47] pretty long we may allocate more memory [107:50] for the for the full token pool. When [107:52] the contest is pretty short we just [107:54] allocate more memory for the recurrence [107:55] state and the window KW. So uh [107:59] but for most of the kernels and back-end [108:01] implementations all the all the all the [108:04] operation they need to be continuously [108:06] in physical memory. [108:08] If we do the memory allocation the [108:11] memory eviction before then there will [108:14] be some [108:15] there will be some holes in the in the [108:17] in the GPU in the GPU's physical memory. [108:20] So the holes like the empty the the idle [108:24] available ones they were they cannot be [108:28] they cannot be [108:30] reused because because once you have to [108:34] like like [108:36] you may hit the the boundary with two [108:38] different memory pools. So at that case [108:40] we introduce one of our operation called [108:43] compaction. We just reorganize all the [108:46] memory fragments [108:49] in the physical memory space. Then we [108:51] pack them with low latency memory copy [108:54] kernels. Then we remapping all the [108:58] virtual token I am virtual [109:01] token IDs. It's not token ID it's [109:04] actually the uh the memory [109:08] page IDs. Then we just do the remapping [109:10] and all the virtual page IDs they do not [109:12] to be changed. We just change the the [109:15] low level mapping itself. [109:17] >> [snorts] [109:18] >> That bring much more uh flexibility for [109:21] the whole serving workload. [109:24] Yeah, and also uh start from serving, we [109:27] also introduce the Miles [109:29] uh inference [109:30] uh inside the training loop. XLA and [109:32] Miles we uh [109:34] we are both uh [109:36] uh contributed mainly contributed by Red [109:38] Hat Arc. And uh XLA is responsible for [109:42] the rollouts and uh is the perfect back [109:44] end for the uh for for for Miles RL [109:47] training. [109:48] And uh Miles supported the fully [109:50] asynchronous execution, uh keep the [109:52] generation and the and and the training [109:55] moving. [109:56] Yeah, as actually I took more things uh [109:58] about the paper. Yeah, we start from the [110:01] Red Hat casual attention and we grow up, [110:04] we do the open source inference engine, [110:06] and we do the Miles, even XLA diffusion. [110:09] We do a lot of other things uh to boost [110:11] the open source uh AI infra uh at that [110:15] time. Uh [110:17] So [110:19] So at the last page, like we will be [110:21] holding uh the XLA summit uh in the next [110:24] months. And there will be a lot of uh [110:26] feature speakers. Uh we'll be there. You [110:29] can scan this QR code to register for [110:31] this. Yeah, and uh that is [110:34] uh [110:35] my talk and uh any Q&As. [110:39] >> [applause] [110:44] >> All right, thank you so much. All right.