# InferenceX Articles — Full Content > By SemiAnalysis This file contains the full text of all articles from InferenceX (https://inferencex.semianalysis.com/blog). It is intended for consumption by large language models and AI assistants. --- # Ultra-High Interactivity on NVIDIA GPUs? TileRT on InferenceX > Can TileRT software on NVIDIA GPUs compete with Cerebras, Groq LPU, and SambaNova? Batch size 1, disaggregated engine, high-throughput prefill engine, high-interactivity decode engine - **Author**: SemiAnalysis - **Date**: 2026-08-10 - **URL**: https://inferencex.semianalysis.com/blog/ultra-high-interactivity-on-nvidia - **Tags**: benchmark, gpu, inference, nvidia, b200, gb300, tilert, vllm, glm5 - **Reading time**: 18 min _Originally published on the [SemiAnalysis newsletter](https://newsletter.semianalysis.com/p/ultra-high-interactivity-on-nvidia) on August 10, 2026._ Premium-priced "fast modes" are proving that users will pay more for lower latency and faster tokens, potentially yielding higher gross margins. Frontier AI labs such as OpenAI are therefore evaluating purpose-built inference systems, including Cerebras and NVIDIA Groq LPUs, that prioritize ultra-high interactivity over maximum batched throughput. Ultra-low latency matters most in interactive workloads, including real-time assistants and full-duplex voice. [OpenAI GPT-Live](https://x.com/OpenAI/status/2080378182469857576), for example, can listen and speak simultaneously, making response delay immediately perceptible to the user. GPUs perform exceptionally well at high throughput and low-to-medium interactivity, but their architecture is less suited for ultra-low-latency inference. An 8-GPU HGX B200 server provides a theoretical HBM memory bandwidth of 64 TB/s in aggregate. At batch size 1, GLM-5 at NVFP4 requires only approximately 21 GB of active-parameter traffic per generated token. The B200 HBM bandwidth roofline would therefore suggest up to 3,047 tokens/s/user without speculative decoding. In practice, GPUs come nowhere close to this limit. The gap comes from latency rather than bandwidth. The traditional GPU programming model launches and synchronizes many individual kernels, whose setup and teardown overhead becomes significant at ultra-high levels of interactivity. While these latency costs are less visible at conventional serving speeds, even with CUDA graphs, they dominate as token latency approaches the sub-millisecond Time Per Output Token (TPOT) range. Furthermore, although [GPU memory bandwidth increases by roughly 2–3× each generation, memory latency has not improved at all](https://newsletter.semianalysis.com/p/vera-rubin-nvl72-vs-gb200-nvl72-inference). While using alternative hardware is popular, there are ways to use GPUs to do this too. This is where [TileRT](https://github.com/tile-ai/TileRT)'s persistent engine comes in. TileRT statically compiles the entire decode graph into a single persistent kernel on NVIDIA GPUs, maximizing overlap across computation, memory loads and stores, and communication. On the InferenceX GLM-5 FP8 744B benchmark on a single B200 decode server, TileRT has been verified to reach up to 500 tokens/s/user, approximately 3× faster than GB300 NVL72 running traditional inference engines. Iso-cost per output token, TileRT can achieve up to 2× faster interactivity than traditional engines. We thank the TileRT maintainers for collaborating on TileRT InferenceX benchmarks, and we are also grateful to the [vLLM community](https://vllm-project.github.io/) for their design on the V1 connector. TileRT comes from the same community maintainer organization that built the widely popular [TileLang DSL](https://github.com/tile-ai/tilelang). With PD disaggregation, the hyperspecialized TileRT engine handles latency-sensitive decode while throughput-optimized engines such as vLLM and SGLang continue to serve prefill. The TileRT decode engine is already being deployed in production at Xiaomi for [MiMo V2.5 Pro UltraSpeed](https://mimo.mi.com/docs/en-US/news/latest/1000tps) and ZAI with [GLM 5.1 HighSpeed](https://www.tilert.ai/). In this article we deep dive into the TileRT InferenceX results, what TileRT is, and how it composes with the existing inference ecosystem, along with the tradeoffs and challenges of TileRT. We will also elaborate on the tradeoffs of using TileRT on standard GPUs versus ultra-low-latency specialized chips like NVIDIA Groq LPU, Cerebras, and SambaNova, weighing in on whether there is potential for TileRT software running on GPUs to disrupt these specialist chips' TAM. The [SemiAnalysis Accelerator Model](https://semianalysis.com/accelerator-hbm-model/) provides quarter-by-quarter estimates of NVIDIA LPU30, LPU40, Cerebras WSE-3 and WSE-4 shipments and much more. Click to see the full InferenceX dashboard → ## InferenceX InferenceX is our open-source, vendor-neutral, continuously updated AI inference benchmarking and research platform. We measure leading models, inference frameworks, and hardware across the latency-throughput Pareto frontier, tracking how real-world inference performance and economics improve over time. Our benchmark has been [widely reproduced, validated and/or supported](https://inferencemax.semianalysis.com/quotes) by almost every major buyer of compute from [Google Cloud](https://cloud.google.com/blog/products/compute/scaling-moe-inference-with-nvidia-dynamo-on-google-cloud-a4x) to [Microsoft Azure](https://blog.aks.azure.com/2025/10/24/dynamo-on-aks#enterprise-scale-inference-experiments--dynamo-with-gb200-running-on-aks) to Oracle, to Meta and many more. Furthermore, it has the support of the ML community including from vLLM, LMCache, SGLang, PyTorch, Hugging Face, and the support of major labs like OpenAI, MiniMax, ZAI, Qwen, Moonshot Kimi, etc.
Star the [InferenceX GitHub repository](https://github.com/SemiAnalysisAI/InferenceX) if you find the open-source benchmark and data useful. As previously mentioned, NVIDIA has committed to submitting verifiable Vera Rubin numbers to InferenceX. We will have Google TPUv7 results soon, and AMD has committed to MI455X UALoE72 this year too.
## Throughput vs Interactivity Curve Every inference system must balance two competing goals. - **Interactivity (tok/s/user)** measures how quickly a single user receives tokens, the inverse of time per output token (TPOT). It determines whether a response feels snappy or sluggish. - **Throughput (tok/s/GPU)** measures how many tokens the system produces in total across all users. It largely determines the cost per token. Batching increases aggregate throughput by processing more requests together, but each user typically waits longer for each token. Small batches do the opposite: they improve per-user speed while reducing the amount of useful work each GPU completes in aggregate. A bus amortizes its cost across many passengers but makes each passenger wait for shared stops. A race car carries only one or two people and reaches the destination faster, but at much higher cost per passenger. Inference has the same trade-off: batching improves aggregate throughput and cost per token, while small batches improve per-user responsiveness. There is no one-size-fits-all operating point. In the configuration shown below, increasing interactivity from roughly 25 to 260 tokens/s/user reduces per-GPU throughput from about 5,900 to 200 tokens/s/GPU. That is roughly a 30× reduction in aggregate throughput for a 10× increase in per-user speed.
## TileRT Results As we describe in the next section, GPUs already perform well in high-throughput scenarios but struggle in high-interactivity ones. This weakness has created an entire market segment for dataflow chips. TileRT targets the same weakness and therefore focuses exclusively on high-interactivity operating points. TileRT on B200 is in a class of its own. For the 8k/1k input/output token scenario, TileRT reached 340 tokens/s/user on an eight-GPU B200 node. The fastest result in the current dataset was previously 181.4 tokens/s/user on GB300 NVL72 with NVFP4 and MTP, making TileRT 1.9× faster on this metric. Of course, this is on batch size 1, where all that extra trouble to set up the complicated copper backplane in the case of the GB300 NVL72 does not come into play at all in boosting interactivity. Meanwhile, the fastest FP8 result was 113.6 tokens/s/user on B300 with MTP, making TileRT 3.0× faster at the same precision.
At 1k/1k input/output, TileRT FP8 reached 494.2 tokens/s/user. That was 1.9× the best conventional result, at 256.3 tokens/s/user using FP4, and 3.6× the best conventional FP8 result, at 136.3 tokens/s/user. TileRT doesn't yet have FP4 support, but it is already beating non-TileRT FP4 implementations. The result is also notable because it comes from an eight-GPU B200 node rather than the 72-GPU NVLink scale-up domain of GB200 or GB300 NVL72. This comparison concerns per-user interactivity, not aggregate throughput or cost.
However, there are always tradeoffs when it comes to inference. TileRT's interactivity advantage comes with lower aggregate throughput. Conventional engines can amortize weight loads and fixed kernel costs across more users as concurrency rises. At 8k/1k input/output, the GB300 FP4+MTP point at concurrency 12 delivers approximately 240 total tokens/s/GPU while maintaining 154 tokens/s/user. TileRT delivers 160.4 total tokens/s/GPU while reaching 340 tokens/s/user. The trade-off is therefore: TileRT provides much higher per-user speed, but the conventional GB300 point completes more aggregate work per GPU. TileRT as of publication also serves only one in-flight request per decode node, making this a deliberately specialized operating point rather than a general throughput configuration. Thus, with support only for a batch size of one user, TileRT is not just a race car, it is more like a private rocket ship with room for just one passenger. Engineering TileRT to support more passengers might be possible, but it is an ambitious goal. For end-to-end latency, TileRT at FP8 outperforms the best previously recorded GLM-5.1 result by 4.5× at 1k/1k and 3.0× at 8k/1k. As expected, TileRT's time to first token (TTFT) is good but not exceptional. The decisive advantage comes from the decode tail: 3.01 seconds, compared with 6.54 seconds for the best NVFP4 + MTP competitor and 18.18 seconds for MI355X.
## But What Exactly is TileRT? We briefly introduced what TileRT does and showed some benchmark results, but let's pause and explain more deeply what TileRT is and how it works. Traditional serving engines run as thousands of separate GPU kernels launched one after another. All that setup and teardown means the GPU spends a surprising amount of time waiting, and while this setup/teardown time might not matter for low-to-medium interactivity inference, it definitely does for ultra-high interactivity inference (aka low-latency inference). Worse, each kernel writes its half-finished work out to HBM. At small batch sizes this is a bigger problem, as kernels aren't large enough to amortize launch latency, synchronization, and scheduling overhead. As mentioned earlier, when running TileRT at batch size 1 for just a single HGX H200 server (38.4 TB/s of aggregate HBM memory bandwidth), the active parameter memory bandwidth stands at 42 GB per token at MXFP8. In theory, if we were only bound by memory bandwidth, then even without speculative decoding, inference should be able to reach up to 1,000 tok/s/user interactivity. This is obviously not the case in the real world. The roadblock is that GPUs' programming and architecture model is traditionally not built for low latency. Even though memory bandwidth per GPU increases 2-3× each generation, memory latency has not improved at all, even as HBM prices continually increase.
Instead of continuously launching kernels, TileRT has the GPU continuously execute a persistent pipeline, statically compiling the whole model ahead of time into a persistent Engine Kernel: the host launches once, execution stays resident on the GPU for the whole decode lifecycle, and most runtime orchestration moves into compile time. This is different from CUDA graphs, which capture the DAGs (directed acyclic graphs) of kernel launches and memcpys once, then replay them with a single `cudaGraphLaunch`. But the kernels themselves are still separate kernels; this boundary between kernels carries device-side costs and the on-chip state is wiped at every boundary. A CUDA graph optimizes the launching of kernels, while TileRT abolishes the kernel as the unit of execution.
Also, through decomposing work into tile-level tasks with warp and block specialization, the runtime dynamically reschedules computation, I/O, and communication in a highly overlapped way. Inside the Engine Kernel, different warp groups take on different jobs: asynchronous data movement, tensor computation, and communication overlap. Where stages used to run serially as load → barrier → compute → barrier, they now overlap at tile granularity, and intermediate results flow forward through registers, shared memory, and L2 instead of repeatedly spilling to global memory. Effectively, each CTA (Cooperative Thread Array) becomes a small heterogeneous factory rather than a uniform SIMT (Single Instruction Multiple Threads) worker.
The next optimization TileRT introduces is specialization extended to whole GPUs. Most TP frameworks assume all ranks execute identical logic synchronously, but sparse routing, Top-K selection, dynamic indexing, long-context attention, and MTP don't fit homogeneous scale-out well; they're not compute-heavy but depend on global information, so forcing every rank through them adds redundant work and synchronization amplification. So, if warps can specialize, so can GPUs. In GLM-5.1's attention layer, GPU 0 becomes a Sparse Indexer worker handling Top-K selection, sparse index construction, and routing, while GPUs 1 through 7 run the MLA workers doing RMSNorm, GEMM, flash sparse attention, and AllReduce.
Finally, instead of treating communication as an external stage, broadcasts, reductions, and synchronization execute directly inside the tile-level flow; with TileRT, an entire attention layer corresponds to a single kernel launch at the host, and execution shifts from compute → sync → compute toward a continuously overlapping compute ↔ communication ↔ compute pipeline. ## PD Disaggregated Engine with vLLM and TileRT LLM inference consists of two distinct phases: prefill and decode. Prefill processes the input prompt in parallel and is primarily compute-intensive, making aggregate throughput the key performance metric. Decode generates tokens sequentially and repeatedly accesses the growing KV cache, making it memory-intensive and highly sensitive to per-token latency.
TileRT does not replace vLLM. vLLM remains the high-throughput prefill engine and the surrounding serving layer, including its scheduler, chunked prefill, prefix caching, OpenAI-compatible API, and operational tooling. Only latency-critical decode traffic moves to TileRT. TileRT is engineered to be a single-passenger rocket ship, and vLLM remains the plane, car, bus, and train. The prefill and decode phases can be [disaggregated](https://arxiv.org/abs/2401.09670) into separate nodes. With disagg, one shared vLLM prefill pool can feed two entirely different decode pools. - **Pool A: Ultra-high interactivity decode with TileRT.** Latency-critical requests pass through the TileRT PD Router, which instructs vLLM to generate the first token and marks the request with the destination TileRT node in `kv_transfer_params`. - **Pool B: General low-to-medium interactivity decode with vLLM decode.** General traffic continues through vLLM's native disaggregation proxy to a conventional vLLM decode pool.
This is done via [vLLM's MultiConnector API](https://vllm-project.github.io/2026/07/14/vllm-tilert-pd.html) that composes the TileRTConnector with its native connector. The TileRT connector claims only marked high-interactivity traffic-class requests and becomes a no-op for everything else, meaning both traffic classes can share the same prefill server. Between the prefill and decode, TileRT uses Mooncake Transfer Engine and NIXL Transfer Engine to move KV cache. In TileRT v0.1.5, each decode node serves one in-flight request at a time. The router gates dispatch and applies back-pressure when the node is occupied. ## How does TileRT compare to Cerebras/Groq/SambaNova? Purpose-built inference vendors identified the same execution bottleneck years ago, but encoded more of the solution in hardware. The [SemiAnalysis Accelerator Model](https://semianalysis.com/accelerator-hbm-model/) has our quarter-by-quarter estimates of NVIDIA LPU30, LPU40, Cerebras WSE-3 and WSE-4 shipments. Groq uses deterministic, compiler-orchestrated execution and a large on-chip SRAM hierarchy. Cerebras maps computation spatially across a wafer-scale processor; the CS-3 provides approximately 900,000 cores, 44 GB of on-chip SRAM, and 21 PB/s of memory bandwidth. SambaNova maps model graphs onto reconfigurable dataflow units backed by a tiered SRAM, HBM, and DDR memory system. The silicon differs, but the systems share the same idea: latency-sensitive inference benefits from reducing runtime scheduling, operator boundaries, synchronization, and unnecessary movement through external memory. At large batch sizes, those costs are easier to amortize. At batch size 1, they occupy a much larger share of each token's latency.
TileRT imports software analogues of several dataflow ideas: AoT scheduling, persistent execution, specialized workers, and tighter overlap between communication and computation. The resemblance is architectural rather than literal. TileRT still runs on a SIMT GPU with dynamic hardware scheduling, HBM, and a model-specific compiled schedule. However, TileRT is still software only: dataflow is imposed on a machine that was never specially designed for it. A GPU carries dynamic warp schedulers, a SIMT model, and an HBM hierarchy, and TileRT gets its numbers by spending enormous compiler effort convincing that machinery to impersonate a spatial pipeline through statically expanded persistent kernels, hand-carved warp specialization, and per-model compilation against pinned driver stacks. Native dataflow silicon never fights its own substrate. Purpose-built accelerators encode more of the execution model in hardware and can avoid some of the overhead TileRT must hide in software. Their advantage still depends on the model, precision, memory hierarchy, compiler quality, system scale, and serving configuration. That is why Cerebras serves a dense 70B at speeds no eight-GPU node can reach regardless of scheduling: software can approach the HBM roofline, but it cannot raise it.
The market's early answer is that purity is negotiable. TileRT's decode engine is already in production behind Xiaomi's MiMo V2.5 Pro UltraSpeed and Z.ai's GLM-5.1 HighSpeed, and the deployment pattern is the tell. Neither company procured a new dataflow chip. They carved a speed tier out of the accelerator cluster they already ran, with vLLM keeping prefill, scheduling, and the API while TileRT takes over decode behind the same endpoint. Good enough on hardware you already own tends to beat architecturally pure on hardware you have to buy. That points at the deeper structural problem: fungibility and flexibility in prefill-decode (PD) ratio. A GPU pool is one liquid resource, excellent at prefill, excellent at high-to-medium-batch decode, and now somewhat credibly strong at ultra-interactive decode, with capacity moving between those roles as a software scheduler decision that can follow demand hour by hour. An ASIC fleet is the opposite: the ratio of speed-tier capacity to everything else is fixed in hardware the day the purchase order is signed. Changing the ratio of the physical fleet will take months to physically re-rack and re-cable. That would be fine if the workload mix were stable and known. Unfortunately, the split between users who need ordinary conversational latency and users, increasingly agents, who will pay for extreme-interactivity SLOs has a lot of different variables at play when estimating. Guess wrong with GPUs and you rebalance in software. Guess wrong with dedicated silicon and you either strand capital in idle speed machines or turn away the exact premium traffic you bought them for. On top of that, requirements may shift over time, so a correct guess will only be right for a limited period of time. Going back to the shared prefill pool mentioned earlier, providers do not need to pay the TileRT premium for all traffic. General requests can stay on throughput-optimized vLLM or SGLang decode pools, while only latency-critical requests are routed to the TileRT decode pool.
None of this kills the top of the speed market. The SRAM roofline is still better, certain sizes of models still favor it, and some workloads will always want maximum tokens per second at any price. But TileRT reframes what most buyers need: not a speed machine, but a speed tier, provisioned dynamically out of the fleet they were going to own anyway. Cerebras, Groq, and SambaNova are no longer competing against a clumsy kernel-launcher. They are competing against their own execution model, running on fungible hardware, reallocated by a config file. TileRT may be a single-passenger rocket ship, but it allows providers to strap solid rocket boosters to your Metro Bus instead of having to design an entirely new launch vehicle. ## Why Is TileRT Development Slow? GLM-5.1 is a generation behind, and has already been deprecated on mainline InferenceX. TileRT's model catalog is very limited, currently supporting GLM-5/5.1 and DeepSeek-V3.2. MiMo-V2.5-Pro-UltraSpeed is the result of a co-design partnership and has yet to be open-sourced. TileRT inherits ASIC vendors' biggest weakness. Static ahead-of-time compilation means a tiny model catalog, hard-pinned dependencies, and real engineering effort per new architecture. There is no fully generic path. A persistent engine kernel means the model is statically expanded ahead of time into one resident program, so decisions have to be made on tile shapes, pipeline depth, buffer residency across registers/shared memory/L2, how warp groups split between loading, compute, and communication, where collectives get fused into the tile flow, and which GPUs take specialized roles like GLM-5.1's dedicated sparse indexer rank. Change the attention mechanism or the routing scheme and much of that schedule is invalidated. Dataflow chips also face this same issue; good compilers can be notoriously difficult to create. Work is being done to simplify this, especially as software development can be accelerated with AI. [TileOPs](https://github.com/tile-ai/TileOPs) is intended to reduce this burden. Each operator is declared in a machine-readable manifest specifying its signature, workloads, and roofline model. The manifest drives code generation, testing, and benchmarking against hardware bounds rather than only against earlier implementations. AI coding agents accelerate tuning within known templates, but novel transformations still require expert judgment. A monolithic persistent kernel also reduces the usefulness of conventional per-kernel profiler timelines, making automated feedback loops more difficult. ## Next steps with TileRT and InferenceX We are actively working on moving TileRT benchmarking from InferenceX's single-turn 8k/1k to our new agentic coding benchmark, which we call AgentX. This scenario replays real Claude Code and Codex traces with long-context, multi-turn requests, realistic subagent activity, and dynamic tool-use delays. Its median input length is 140k tokens, while the theoretical median cache-hit rate roofline reaches 99.2%.
This workload will test the entire TileRT and vLLM system, not just decode speed, including incremental KV transfer, prefix-cache reuse, cache retention and offloading, routing, and scheduling. The critical question is whether TileRT can transfer only the newly introduced context between turns while preserving its ultra-high interactivity advantage.
The second step is to move beyond just batch size one. We will also benchmark TileRT at batch sizes 2, 4, and 8. The goal is to map its throughput-interactivity Pareto frontier and identify the point at which the persistent Engine Kernel's latency advantage begins to flatten. The full perf-per-TCO analysis — TileRT's cost per million output tokens at ultra-high interactivity compared against decode at normal lower-interactivity operating points, using the [SemiAnalysis AI Cloud TCO Model](https://semianalysis.com/ai-cloud-tco-model/) as the capex and opex baseline — is available to subscribers in the [full article on the SemiAnalysis newsletter](https://newsletter.semianalysis.com/p/ultra-high-interactivity-on-nvidia). Click to see the full InferenceX dashboard → {`{ "@context": "https://schema.org", "@type": "FAQPage", "mainEntity": [ { "@type": "Question", "name": "How fast is TileRT on NVIDIA B200 compared to conventional inference engines?", "acceptedAnswer": { "@type": "Answer", "text": "On the InferenceX 8k/1k input/output scenario, TileRT reached 340 tokens/s/user on an eight-GPU B200 node, against a previous best of 181.4 tokens/s/user on GB300 NVL72 with NVFP4 and MTP, making TileRT 1.9x faster. Against the fastest FP8 result of 113.6 tokens/s/user on B300 with MTP, TileRT is 3.0x faster at the same precision. At 1k/1k, TileRT FP8 reached 494.2 tokens/s/user, 1.9x the best conventional result of 256.3 tokens/s/user using FP4 and 3.6x the best conventional FP8 result of 136.3 tokens/s/user." } }, { "@type": "Question", "name": "What is TileRT and how does it differ from CUDA graphs?", "acceptedAnswer": { "@type": "Answer", "text": "TileRT statically compiles the entire decode graph ahead of time into a single persistent Engine Kernel on NVIDIA GPUs. The host launches once and execution stays resident on the GPU for the whole decode lifecycle, moving most runtime orchestration into compile time. CUDA graphs capture the directed acyclic graph of kernel launches and memcpys once and replay it with a single cudaGraphLaunch, but the kernels remain separate, so each kernel boundary still carries device-side cost and wipes on-chip state. A CUDA graph optimizes the launching of kernels, while TileRT abolishes the kernel as the unit of execution." } }, { "@type": "Question", "name": "What is the trade-off for TileRT's higher interactivity?", "acceptedAnswer": { "@type": "Answer", "text": "TileRT's interactivity advantage comes with lower aggregate throughput. At 8k/1k, the GB300 FP4 plus MTP point at concurrency 12 delivers approximately 240 total tokens/s/GPU while maintaining 154 tokens/s/user, whereas TileRT delivers 160.4 total tokens/s/GPU while reaching 340 tokens/s/user. TileRT as of publication serves only one in-flight request per decode node, making this a deliberately specialized operating point rather than a general throughput configuration." } }, { "@type": "Question", "name": "Does TileRT replace vLLM?", "acceptedAnswer": { "@type": "Answer", "text": "No. vLLM remains the high-throughput prefill engine and the surrounding serving layer, including its scheduler, chunked prefill, prefix caching, OpenAI-compatible API, and operational tooling. Only latency-critical decode traffic moves to TileRT, via vLLM's MultiConnector API composing the TileRTConnector with the native connector. One shared vLLM prefill pool can feed both a TileRT ultra-high interactivity decode pool and a conventional vLLM decode pool. TileRT uses Mooncake Transfer Engine and NIXL Transfer Engine to move KV cache between prefill and decode." } }, { "@type": "Question", "name": "Can TileRT on GPUs displace Cerebras, Groq, and SambaNova?", "acceptedAnswer": { "@type": "Answer", "text": "TileRT imports software analogues of dataflow ideas including ahead-of-time scheduling, persistent execution, specialized workers, and tighter compute-communication overlap, but it still runs on a SIMT GPU with dynamic hardware scheduling and an HBM hierarchy. Software can approach the HBM roofline but cannot raise it, which is why Cerebras serves a dense 70B at speeds no eight-GPU node can reach. The structural argument favoring GPUs is fungibility: a GPU pool can shift capacity between prefill, general decode, and ultra-interactive decode as a scheduler decision, while an ASIC fleet fixes that ratio in hardware on the day the purchase order is signed." } }, { "@type": "Question", "name": "Why is TileRT development slow?", "acceptedAnswer": { "@type": "Answer", "text": "Static ahead-of-time compilation means a tiny model catalog, currently GLM-5/5.1 and DeepSeek-V3.2, plus hard-pinned dependencies and real engineering effort per new architecture. A persistent engine kernel requires up-front decisions on tile shapes, pipeline depth, buffer residency across registers, shared memory and L2, how warp groups split between loading, compute and communication, where collectives fuse into the tile flow, and which GPUs take specialized roles. Changing the attention mechanism or routing scheme invalidates much of that schedule. TileOPs is intended to reduce this burden by declaring each operator in a machine-readable manifest that drives code generation, testing, and benchmarking against hardware bounds." } } ] }`} --- # Kimi K3: The Manos, The Mythos, The Legendos > Kimi K3's architecture: compressed memory, attention across depth, latent expert routing, and serving performance - **Author**: SemiAnalysis - **Date**: 2026-08-03 - **URL**: https://inferencex.semianalysis.com/blog/kimi-k3-the-manos-the-mythos-the - **Tags**: inference, benchmark, gpu, kimi, vllm, nvidia, b200, b300, dynamo - **Reading time**: 25 min _Originally published on the [SemiAnalysis newsletter](https://newsletter.semianalysis.com/p/kimi-k3-the-manos-the-mythos-the) on August 3, 2026._ Kimi K3 took the world by storm at its announcement, sweeping leaderboards and establishing itself as the open frontier model. While the community is eager to understand how Kimi K3 works, many have been surprised by the unconventional techniques driving its performance. This article serves as a primer to understanding the core techniques of the Kimi K3 model architecture. ## Kimi Delta Attention Kimi Delta Attention (KDA) is the linear attention layer in Kimi K3's hybrid attention mechanism. We trace the origins of KDA, starting from linear attention, DeltaNet, Gated DeltaNet (GDN), then to KDA. ### Linear Attention The derivation of linear attention stems from **removing the softmax operation in the standard softmax attention**. Below we compare the iterative inference formulas, which show the computation of the output vector at token position t:
By removing the softmax operation, we can reorder the operations and reduce the computation complexity of attention from quadratic to linear:
The new equations are as follows:
Vectors q, k, v, have dimensions _L_ by _d_. The computational complexity of both equations are O(*Ld*²), thereby making the computation linear. Comparing the new equations with softmax attention's equation, we see that **softmax attention requires accessing all past key and value vectors**, whereas **linear attention compresses all past key and value vectors** **into one hidden state S**. We reinterpret the new equations as an **online learning objective**. We view matrix S as an associative memory that stores the associations between key vector k and value vector v, and we retrieve v by multiplying S with k. We can then interpret the first equation as **continuously updating the matrix S at every position to perfect the retrieval**. Finally, we can interpret the vt @ kt.T term as the gradient of loss function -(S @ kt.T) @ vt with respect to S. $$ \begin{aligned} \textbf{Objective:}\quad \mathcal{L}_t(\mathbf{S}) &= -\left\langle \mathbf{S}\mathbf{k}_t,\mathbf{v}_t \right\rangle \\[8pt] \textbf{SGD update:}\quad \mathbf{S}_t &= \mathbf{S}_{t-1} -\beta_t\nabla\mathcal{L}_t(\mathbf{S}_{t-1}) \\ &= \mathbf{S}_{t-1} +\beta_t\mathbf{v}_t\mathbf{k}_t^\top \end{aligned} $$ ### DeltaNet Under the online learning objective view, we see the values of matrix S will grow unboundedly: old and new information gets blurred together in S as the sequence grows, which destabilizes learning. Without softmax giving well-scaled and bounded outputs, linear attention typically lags behind softmax attention on long-range recall tasks. DeltaNet improves upon linear attention by **changing the loss function to minimizing the L2 norm of the value retrieval**. Unlike linear attention's loss function, DeltaNet's loss function regularizes the growth of S. This creates a new matrix S update rule, the Delta Rule, as below: $$ \begin{aligned} \textbf{Objective:}\quad \mathcal{L}_t(\mathbf{S}) &= \frac{1}{2} \left\| \mathbf{S}\mathbf{k}_t-\mathbf{v}_t \right\|^2 \\[6pt] \textbf{SGD update:}\quad \mathbf{S}_t &= \mathbf{S}_{t-1} -\beta_t\nabla\mathcal{L}_t\!\left(\mathbf{S}_{t-1}\right) \\ &= \mathbf{S}_{t-1} -\beta_t \left( \mathbf{S}_{t-1}\mathbf{k}_t-\mathbf{v}_t \right) \mathbf{k}_t^\top \end{aligned} $$ Source: [Linear Attention and Beyond (Interactive Tutorial with Songlin Yang)](https://www.youtube.com/watch?v=d0HJvGSWw8A) The Delta Rule becomes the basis of DeltaNet's attention equation: $$ \mathbf{S}_t = \mathbf{S}_{t-1} - \beta_t \left( \mathbf{S}_{t-1}\mathbf{k}_t-\mathbf{v}_t \right) \mathbf{k}_t^\top $$ Conceptually, Sₜ-1 @ kₜ - vₜ represents the associations irrelevant to the current key and value, and DeltaNet performs targeted removal of those associations. ### Gated DeltaNet GDN and KDA are adaptations of DeltaNet. Gated DeltaNet applies the LSTM forget gate _alpha_ on the matrix S, allowing the model to control memory lifespan with weight decay. KDA further expands _alpha_ into a diagonal matrix that enables fine-grained per-channel memory decay and positional awareness.
### FlashKDA Algorithm Moonshot developed FlashKDA, their custom kernels for KDA, and [open-sourced it](https://github.com/MoonshotAI/FlashKDA/tree/master). Here we explain the algorithm and derive the arithmetic intensity. #### Algorithm First, let's start from an alternative formulation of the recurrence formula: ```text u_t = beta_t * (v_t - (D_t @ S_t-1).T @ k_t) S_t = D_t @ S_t-1 + k_t @ u_t.T o_t.T = q_t.T @ S_t ``` Here, D_t is the diagonal matrix of the alpha forget gate, and u_t is the delta in the delta rule. For decode, the kernel roughly follows the formula. For prefill, we parallelize the operation by unrolling the recurrence formula in chunks of tokens, in order to efficiently execute the operations on GPUs. Assume we unroll token i to j, and the starting state is S_i-1, we get: ```text S_j = D_j:i @ S_i-1 + sum(D_j:t+1 @ k_t @ u_t.T, t=i:j) o_j.T = q_j.T @ S_j = q_j.T @ D_j:i @ S_i-1 + sum(q_j.T @ D_j:t+1 @ k_t @ u_t.T, t=i:j) ``` D_j:i refers to the cumulative decay from token i to j: D_j @ D_j-1 @ D_j-2 @ … @ D_i. In FlashKDA's matrix form, the formula becomes: ```text S_out = D_j:i @ S_in + K_restore.T @ U M_qk = tril(Q_decay @ K_inv.T) O = Q_decay @ S_in + M_qk @ U ``` The vector to matrix mapping is as follows: - S_in refers to the state at the starting position of a chunk - S_out refers to the state at the end position of a chunk - K_restore is the matrix form of D_j:t+1 @ k_t - Q_decay is the matrix form of q_j.T @ D_j:i - Q_decay @ K_inv.T the matrix form of q_j.T @ D_j:t+1 @ k_t, derived from (q_j.T @ D_j:i) @ (D_t:i^-1 @ k_t) - M_qk is the causal mask, so it's a lower triangular matrix U is the matrix form of unrolled u_t. To compute this, we apply UT transform and compute the following: ```text B = Diag(beta) @ (V - K_decay @ S_in) L = StrictTril(Diag(beta) @ K_decay @ K_inv.T) U = (I + L)^-1 @ B ``` Please consult [Songlin Yang's blog post](https://sustcsonglin.github.io/blog/2024/deltanet-2/) and [Kimi Linear paper section 3.1](https://arxiv.org/abs/2510.26692) for the full derivation. Note that here U corresponds to the pseudo-value term in the Kimi Linear paper. Implementation-wise, FlashKDA launches 2 kernels: K1 and K2. K1 prepares chunk-level tensors in parallel, including: ```text a = exp2(cumsum(g)) K_decay = Diag(a) @ K Q_decay = Diag(a) @ Q K_inv = Diag(a)^-1 @ K K_restore = a[-1] * K_inv L = StrictTril(Diag(beta) @ K_decay @ K_inv.T); INV = (I + L)^-1 M_qk = tril(Q_decay @ K_inv.T) ``` Here `a` is the cumulative decay, where each element is the cumulative decay at a token position. K2 performs chunk-level recurrent computation: ```text U = INV @ Diag(beta) @ (V - K_decay @ S) O = Q_decay @ S + M_qk @ U S = Diag(a[-1]) @ S + K_restore.T @ U ``` #### Complexity Analysis Here we analyze the complexity of an attention head. For decode, the critical path computations are: - D_t @ S_t-1: Element-wise multiplication, D × D - S_t-1.T @ k_t: D × D × 1 - k_t @ u_t.T: D × 1 × D - q_t.T @ S_t: 1 D × D The decode kernel roughly performs `7*D²` FLOPs. Reading and writing the FP32 recurrent state dominates the memory traffic, so the memory traffic is roughly `8*D²` bytes. For prefill, the critical path of K1 is at computing L, INV, and M_qk. - L: C × D × C - INV: [Neumann factorization](https://github.com/MoonshotAI/FlashKDA/blob/1ce47ea3bb22c84eb9cc665028399cf35e8ffb0b/csrc/smxx/utils.cuh#L190), performs 6 C × C × C matrix multiplications - M_qk: C × D × C For K2, - K_decay @ S: C × D × D - Q_decay @ S: C × D × D - M_qk @ U: C × C × D - INV @ B: C × C × D - K_restore.T @ U: D × C × D Combining K1 and K2, FlashKDA performs `12*C^3 + 8*C²*D + 6*C*D²` FLOPs. Since we analyzed at the chunk level (chunk size C), assuming sequence length T `>>` C, the overall FLOPs is `T/C * O(C*D²) = O(T*D²)`. For memory traffic: - K1 read Q, K, g: C × D - K1 write and K2 read Q_decay, K_decay, K_restore: C × D - K1 write and K2 read INV, M_qk: C × C - K2 read V and write O: C × D - K2 read and write S once per kernel: D × D In total, FlashKDA accesses `3 * 2*C*D + 2 * (3 * 2*C*D + 2 * 2*C*C) + 2 * 2*C*D = 8*C² + 22*C*D` bytes. At the kernel level, it accesses `T/C * (8*C² + 22*C*D) + 8*D² ~ O(TC + TD + D²)`. This concretely shows that the computational complexity of KDA: - Prefill: Linear to sequence length for both computation and memory - Decode: Constant to sequence length for both computation and memory ## Kimi Linear Moonshot trained Kimi Linear models as proof of concept for their KDA design, so we can infer Kimi K3's architecture design from Kimi Linear. Comparing the K3 release tech blog with Kimi Linear, we see Kimi K3 shares the shared expert count, the hybrid linear attention ratio, and the general attention module design.
The diagram above shows the operations performed on the inputs of KDA. For the query, key, and value, we apply linear transformation and short convolution. Applying short convolution effectively capturing local token dependencies, and doing a left padding convolution avoids breaking causality. We additionally apply L2 norm to the query and key to stabilize the eigenvector of the transition and the output matrices. For the decay memory gates, alpha is a low rank projection, and beta is a down projection. The KDA output is normalized per head and controlled by an output forget gate, implemented as a linear transformation in K3, instead of a low rank projection in Kimi Linear. Finally, we apply a linear layer to mix per-head information. Kimi Linear interleaves KDA with full attention Multi-head Latent Attention (MLA). Kimi Linear showed that 3:1 is the ideal KDA to MLA ratio that balances performance and efficiency. KDA also serves as a strong position-aware operator, replacing the RoPE in MLA. Keeping MLA as full attention is an interesting choice, as all other open weight models move to Grouped Query Attention (GQA). MLA uses an absorption trick to reduce the computation of a decode step at the cost of extra computation during the prefill step. This is a sensible trade-off for decode-dominant reasoning workloads, but for prefill-dominant agentic workloads, extra computation becomes a high cost with little benefits. As a result, all frontier open weight models use GQA-based attention mechanisms: GLM 5.2 DeepSeek Sparse Attention, DeepSeek V4 Compressed Sparse Attention, MiniMax M3 MiniMax Sparse Attention, and MiMo V3 HySparse are all based on GQA. We suspect Moonshot's future models such as Kimi K4 will feature attention mechanisms that replace MLA. ## KV Cache Efficiency We argue that **one should not infer KV cache efficiency solely based on KV cache space complexity**. KV cache size is not a standalone factor but a property of the model design: no open weight models are released with static KV cache compression techniques, and model architecture inference efficiency affects KV cache efficiency. The effects of KV cache size also vary, depending on the total memory capacity of a deployed model instance. For example, deploying a model with wide expert parallelism has very different memory profiles than doing so with tensor parallelism, which affects the memory capacity left for KV cache. Thus, we propose considering both the model architecture system efficiency and the KV cache size to understand the KV cache efficiency, and we quantify that with **KV throughput**. ### KV Throughput KV throughput is defined as KV cache size divided by the prefill time (Time to first token), given a specific sequence length. KV throughput represents the minimum bandwidth required to reliably serve a model with PD disaggregation, but it is also a good proxy for understanding KV cache efficiency. Prefill time encapsulates the efficiency of the model architecture, and as the sequence length increases, we will see the memory-bounded and the compute-bounded situations. As shown in the table below, we can see the benefits of hybrid linear attention become more pronounced as sequence length increases.
This is also a good way to understand the bandwidth requirements of KV cache offloading to different memory tiers in a cluster. ### KV Cache Residency The location where KV cache is stored follows the memory hierarchy. First, KV cache resides in HBM, the fastest memory in a GPU cluster, consuming whatever capacity is left by model weights and activations. As KV cache size exceeds the HBM capacity, it spills into server DRAM, a higher capacity but lower bandwidth memory pool. Finally, when KV cache exceeds DRAM capacity, it spills to disk storage such as SSD. This is analogous to the computer architecture cache hierarchy: register, cache memory, main memory, disk storage. The analogy continues for memory coherency. Popular distributed KV cache framework Mooncake Store supports write-through and write-back policies for KV cache loading. Mooncake Store features a distributed KV cache pool that makes all KV cache visible to all workers. Implementing write-through policy between DRAM and the lower-level distributed KV cache pool offers multiple benefits in multi-node scenarios, including sharing prefix cache across nodes, avoid KV cache duplication for tensor parallelized MLA, and KV cache redundancy when a node goes down.
### KDA Prefix Cache Management At each token position in a request, Kimi K3 KDA's recurrent state is fixed in size, whereas standard attention KV cache grows with sequence length. This KV cache space reduction comes at the cost of complicating prefix caching, especially when Kimi K3 is a hybrid attention of KDA and MLA. Roughly speaking, modern inference engines identify prefix cache hits by matching the longest token prefix in the existing cache.
Identifying the longest prefix becomes a problem for linear attentions like KDA. Without prior knowledge of where the boundary of a prefix is, we will have to cache KDA's recurrent state at every token position. This means every token has a cache, and the KV cache memory usage regresses to growing with sequence length, defeating the purpose of using linear attention. To tackle this problem, Moonshot saves recurrent states at a coarse granularity, e.g. vLLM caches every 32K tokens. vLLM additionally caches at prompt boundaries, since for agentic workloads, a new turn typically starts at the end of a prompt.
This shows that even though linear attentions like KDA greatly reduce KV cache memory consumption, **realistically during serving, they do not consume a constant amount of KV cache memory**. ## Attention Residuals ### Residual Connections Residual connections are one of the key innovations that allowed us to build bigger deep neural networks through scaling model depth. The deeper the neural network, the more expressive they become but training them naively is hard. Signals from the earlier layer need to be preserved till the last layer and gradient need to survive from output to first without vanishing. Instead of modeling whole networks as a single function, passing information only through nonlinear transformations, residual networks connect smaller blocks with identity paths. Each block fᵢ ​ learns a change to its input xᵢ, given by the recurrence: $$ x_{l+1} = x_l + f_l(x_l) $$ The identity mapping allows features to carry from shallower units to any deeper unit and gives the gradient a path highway so they do not vanish. $$ \frac{\partial x_{l+1}}{\partial x_l} = I + f_l'(x_l) $$ While residual connections allow us to build deeper networks, they come with challenges. Early layers heavily influence residual stream to have effect on final output. Because of which residual stream has irreversible information loss with increasing depth. Later layers increase output gain to have effect on this modified residual stream which can destabilize training. Another variant like highway networks allow gating mechanisms for information flow but they suffer from the same crucial problem. Layers don't have selective access to information from earlier layers. ### Recurrence In Time and Depth Sequence modeling dominated by recurrent neural networks had the same recurrence formulation. $$ h_{t+1} = h_t + f(h_t, x_{t+1}) $$ Where each step has identity mapping with previous state for direct information flow and the sequence model faced the same challenge: depth in the time axis dilutes signal.
Attention machines transformer removed this constrained by retrieving any token in past with powerful and expensive attention mechanism ### Attention on residual stream Motivated by attention mechanism in sequence modeling, kimi developed attention residual, where they take attention over depth blocks,
Standard causal self-attention computes the output of token _t_ as a weighted sum of previous token representations: $$ \mathbf{o}_t = \sum_{i=1}^{t} \alpha_{i\rightarrow t}\,\mathbf{v}_i, \qquad \alpha_{i\rightarrow t} = \frac{\phi(\mathbf{q}_t,\mathbf{k}_i)} {\sum_{j=1}^{t}\phi(\mathbf{q}_t,\mathbf{k}_j)} $$ Attention Residuals use the same attention mechanism, but replace the sequence dimension with the depth dimension. Instead of attending over previous tokens, each layer attends over representations produced by previous layers.
Unlike standard attention, the query is a learned parameter for each layer rather than being generated from the current token. $$ \alpha_{i \to l} = \frac{\phi(\mathbf{q}_l,\mathbf{k}_i)} {\displaystyle\sum_{j=0}^{l-1}\phi(\mathbf{q}_l,\mathbf{k}_j)} $$ For each layer ℓ, we define: $$ \mathbf{q}_l=\mathbf{w}_l, \qquad \mathbf{k}_i=\mathbf{v}_i= \begin{cases} \mathbf{h}_1, & i=0,\\ f_i(\mathbf{h}_i), & 1\le i Attention residual allows the model to get fine grained control over what inputs to pick from past layers making the model more expressive. ### Block Attention Residuals Attention residual need to all past layer outputs for attention. For large models distributed over many GPUs this creates O(Ld) communication overhead. To overcome this block attention residual dividends layers L into N blocks of S layers. Block AttnRes applies attention over completed block outputs and for current block its evolving partial sum.
Block Attention has minimal tread over full attention residual but they cut down communication from O*(Ld)* to O*(Nd).* Let bₙⁱ denote the partial sum over the first _i_ layers in block _n_, such that $$ \mathbf{b}_n=\mathbf{b}_n^S, \qquad \mathbf{b}_0=\mathbf{h}_1. $$ For the _i_-th layer in block _n_, the available block representations are $$ \mathbf{V}_l= \begin{cases} [\mathbf{b}_0,\mathbf{b}_1,\ldots,\mathbf{b}_{n-1}]^\top, & i=1,\\[4pt] [\mathbf{b}_0,\mathbf{b}_1,\ldots,\mathbf{b}_{n-1}, \mathbf{b}_n^{i-1}]^\top, & i>1. \end{cases} $$ Unlike standard attention, the query is not input-dependent. Each layer learns a query vector: $$ \mathbf{q}_l=\mathbf{w}_l $$ Attention weights over the available block representations are computed as $$ \boldsymbol{\alpha}_l = \operatorname{softmax} \left( \mathbf{K}_l\mathbf{w}_l \right). $$ The output is the weighted sum of previous layer representations $$ \mathbf{h}_l = \boldsymbol{\alpha}_l^\top\mathbf{V}_l. $$ Rather than depending only on the residual stream to preserve information, Attention Residuals give every layer direct, selective access to earlier representations. This block based variant of attention residuals greatly reduces communication overhead while having competitive performance. Block residuals show better scaling compared standard residual connection achieving 1.25× compute efficiency. Consistently lower validation loss compared to baseline and gap widening with decay phase. Unlike standard residual networks where output magnitude increases as depth increases. selective aggregation of block attention has bounded output. And consistent gradient magnitude. ### Training Unlike standard residual networks, attention residuals need all N-1 block input for computation of the Nth layer. This becomes a problem for pipeline parallelism as all N layer blocks output need to be transferred across stages. With clever cross stage caching and activation checkpointing, Kimi reduced overhead to only 4% compared to standard architecture for pipeline parallelism. #### Cross-stage caching For _P_ physical stages and _V_ virtual stages. Each block _N_ needs _C=PV_ communication for each chunk. Naively this needs transferring all accumulated blocks for each stage. This is quadratic cost growth for each physical and virtual stage $$ \mathrm{Comm}_{\mathrm{naive}} = \sum_{j=1}^{C-1} jN_p \cdot d = \frac{C(C-1)}{2}N_p d $$ This high communication can be reduced by caching input across virtual stages. Blocks computed in earlier layer can be stored in local memory,
For the first virtual stage all block embedding needs to be transferred in the physical stage, each completed block is stored on respective rank. For all subsequent virtual stages all cached blocks can be reused for computation. Only the block not present on rank need to be transferred for attention to residual computation.
These split communication costs for first and subsequent virtual stages. For the first virtual stage its need incur the same quadratic cost for all physical layers. In subsequent virtual stages we need cached inputs from local devices and Transfer of only PNp chunks needed cutting down total communication from O*(C)* to O*(P)* $$ \mathrm{Comm}_{\mathrm{cached}} = \underbrace{\frac{P(P-1)}{2}N_p d}_{\text{first virtual stage}} + \underbrace{(V-1)P^2N_p d}_{\text{subsequent virtual stages}} $$ The cutdown of communication is directly proportional to virtual stages V. Because of this for full stage of one forward and backward pass all computation and communication can be overlapped #### Memory overhead Due to cross stage caching all blocks are stored once across all V virtual stages. With Activation checkpointing all inter-block chunks for attention are eliminated. Each stage activation checkpoint _Pl_ matches memory size of H*l* of standard architecture and has no extra memory cost. ### Inference Because Attention Residuals need the output of all previous blocks to compute attention, a naive implementation has excessive memory accesses. To reduce overhead, inference is split into two phases which mirror prefill and decode stages of autoregressive attention. This computation is divided into inter-block attention for completed blocks and intra-block attention for evolving attention in the running block. #### Phase 1: Parallel Inter-Block Attention
During decoding, we have to output the completed block and the query vector learned per layer. All inter block layers simultaneously with a single batched query against the completed block representations, returning both outputs and softmax statistics which can be reused for further computation. This phase is similar to prefill phase decoding #### Phase 2: Sequential Intra-Block Attention
This phase is analogous to the decode phase, Similar to flash attention, evolving sum can be computed with online softmax for intra blocks combined with precomputed inter block results. Which reduces redundant memory access. With this two phase design, the IO footprint is similar to standard residual architecture, with only the addition of phase ones inter block computation, amortized by batching all queries in the block. ## LatentMoE LatentMoE compresses the routed tokens before the dispatch operation, and then decompresses them after the aggregation operation. In Kimi K3's Stable LatentMoE, they apply an RMSNorm before the up-projection (decompressing) operation to reduce sensitivity to scale variations and improve model performance.
Here we explain the design principles behind LatentMoE regarding MoE communication. As shown in the LatentMoE paper, the communication volume is proportional to total routed tokens _t_, number of active experts _K_, and expert input dimension _d_, while being inversely proportional to the expert parallel size _E_. This is potentially the reason behind Kimi K3's latent MoE dimension size and active expert count configuration. Kimi K2 series feature 8 active experts with input dimension size 7168, so Kimi K3's latent input dimension size being 3584 (half of 7168) would allow the active expert count to double to 16 without increasing the communication volume. However, the **ratio of communication to computation** **time** is arguably more important for estimating system efficiency (Discussions [here](https://x.com/chhillee/status/2077966168304787769) and [here](https://x.com/chhillee/status/2078130513546723531)). The ratio indicates the roofline of how well MoE kernels can overlap communication with computation at a throughput-bound regime, and **expert intermediate dimension size** is the only model configuration that affects the ratio. Specifically, increasing the expert intermediate dimension size would decrease the ratio, meaning that the theoretical maximum fraction of communication that can be hidden is higher. Here we derive the formula: - _t_: total input tokens across the expert parallel (EP) domain - _K_: number of active experts per token - _N_: number of total experts - _E_: Ranks in the EP domain - _d_: Expert input dimension - _m_: Expert intermediate dimension - _P_: Aggregate bytes communicated per activation element (dispatch + combine) - _F_: Effective FFN expert (modeled as SwiGLU) computation throughput per GPU, FLOP/s - _B_: Effective uni-directional network bandwidth per GPU, B/s 1. Assuming uniform expert routing, each GPU is assigned _t \* K / E_ tokens 2. Assuming uniform expert routing, an average _1 / E_ tokens are local to the source GPUs, so each GPU dispatches _(t\*K/E) \* (1-1/E)_ tokens 3. Each token is a d dimensional vector, so the communication volume per token is _d \* P_ 4. The communication volume per GPU is _(t\*K/E) \* (1-1/E) \* d \* P_ 5. The communication time T*comm = *(t \* K \* d \* P) / (E \* B) \* (1-1/E)\_ 6. The SwiGLU computation involves 3 matrix multiplications: 1. Up (First) projection: _d_ to _m_ 2. Gate projection: _d_ to _m_ 3. Down (Second) projection: _m_ to _d_ So the computation is `2*d*m + 2*d*m + 2*m*d = 6*d*m` FLOPs per token 7. The computation time per GPU is `T_comp = (6*d*m) * (t*K/E) / F` 8. The communication to computation time ratio is `T_comm / T_comp` `= ((t * K * d * P) / (E * B) * (1-1/E)) / ((6*d*m) * (t*K/E) / F)` `= (P*F) / (6*m*B) * (1-1/E)` We believe this formula also motivates an increase in expert intermediate dimension to 3072 in not just Kimi K2 to K3, but all recent open weight models, including DeepSeek V4 Pro, MiniMax M3, MiMo V2.5 Pro, and Inkling. As hardware improves and expert weight precision reduces to save memory capacity, the compute throughput increases, so one way of reducing the ratio is by increasing the expert intermediate dimension. ### Quantile load balancing (QB)
Many previous load balancing methods require careful hyperparameter tuning. Quantile balancing is hyperparameter free aux-loss free load balancing technique developed by Jianlin Su in [Feb 2026 blog post](https://kexue.fm/archives/11619) Base principle QB is the same as auxfree load balancing where router biases are updated dynamically based on the system's load. But instead of updated bias by some small coefficient like aux-free lb, QB directly computes the next bias from the distribution of router scores relative to routing cutoff threshold. Bias updates become small naturally when the router balances load evenly.
QB tries to find the bias that would have approximately balanced under the current cutoffs and routing on the current batch, solving constraint optimization problems and applying these updates for the next batch. The first constraint is that each token is routed to exactly k experts. The second constraint is a batch of m tokens each picks k experts, gives (mk) assignments in total, to spread load evenly across n experts each expert should process _q=mk/n_ tokens. Each token finds the cutoff threshold as the (k+1)-th highest biased router score and uses it to calculate the bias update needed to balance load for each expert. For each expert, QB sorts the margins between its router score and every token's cutoff. It sets negative bias to q+1 the largest margin, leaving exactly q margin above the threshold. Since q/m=k/n, this is (1-k/n) quantile of the margin, which is why it's called Quantile Balancing. ## Inference performance We are actively tracking Kimi K3's inference performance on [InferenceX](https://inferencex.semianalysis.com/). As of 30th July, all providers on OpenRouter have a floor of $3 per million tokens input and $15 per million tokens output. Both Nvidia and AMD had Day 0 recipes on vLLM, boasting DRAM offload and DSpark speculative decoding.
On InferenceX, we benchmark Kimi K3 serving performance directly on recorded internal claude code traces. We replay an hour of these traces as they reach a steady state. There is a median of 142k input tokens and a median of 444 output tokens per turn with a median of 65 turns per session. The short output tokens per turn is typical for workloads on agentic harnesses, where the agent calls tools frequently, even edits are tool uses. This benchmark is a big step up from our previous 8k1k/1k1k benchmark, as it truly reflects real-world agentic use cases. From a systems perspective, it is also realistic and closest to production systems. It can reflect KV cache behavior, including prefix cache and KV offloading to DRAM.
For Kimi K3, Day 0 bringup was easier than DSv4 due to better documentation and preparation ahead of weights release. Appropriate images and a speculative decoder model were released at the same time as the weights. _Related: [DeepSeekV4 1.6T Day 0 to Day 43 Performance Over Time — Huawei, GB300 NVL72, MI355X, B200](/blog/deepseekv4-16t-day-0-to-day-43-performance)_ For Nvidia, bringup was simple. But due to the models' sheer size, it doesn't fit on a single B200 node. We had to use PP to get it working. DSpark also didn't work with PP.
For B300, the model fits on 1 node and serves well. After accounting for the weights, GPU HBM can only hold 3.25M tok. In the graph below, throughput goes up as batch sizes increase until concurrency increases above 8. This roughly correlates to the 3.25M tok KV cache budget, and cache starts to thrash, resulting in hit rates falling to `<` 10% when theoretical hit rate is 95%. Click to see the full InferenceX dashboard → _The article continues with the B300 concurrency sweep, the AMD MI355X and MI455X bringup, and the full cross-hardware serving comparison for Kimi K3, in the [subscriber edition on the SemiAnalysis newsletter](https://newsletter.semianalysis.com/p/kimi-k3-the-manos-the-mythos-the)._ {`{ "@context": "https://schema.org", "@type": "FAQPage", "mainEntity": [ { "@type": "Question", "name": "What is Kimi Delta Attention (KDA)?", "acceptedAnswer": { "@type": "Answer", "text": "KDA is the linear attention layer in Kimi K3's hybrid attention mechanism. It descends from linear attention and DeltaNet: linear attention removes the softmax so all past keys and values compress into a single hidden state S, DeltaNet regularizes the growth of S by minimizing the L2 norm of the value retrieval, Gated DeltaNet adds an LSTM-style scalar forget gate alpha, and KDA expands alpha into a diagonal matrix for fine-grained per-channel memory decay and positional awareness. KDA also replaces RoPE as the position-aware operator in the full attention layers." } }, { "@type": "Question", "name": "What is the KDA to MLA ratio in Kimi K3?", "acceptedAnswer": { "@type": "Answer", "text": "3:1, the same ratio Kimi Linear showed balances performance and efficiency. Kimi K3 interleaves three KDA layers with one full-attention Gated MLA layer. Keeping MLA as the full attention mechanism is unusual: GLM 5.2, DeepSeek V4, MiniMax M3, and MiMo V3 all use GQA-based sparse attention instead, because MLA's absorption trick trades extra prefill computation for cheaper decode, which suits decode-dominant reasoning but not prefill-dominant agentic workloads." } }, { "@type": "Question", "name": "What is the computational complexity of FlashKDA?", "acceptedAnswer": { "@type": "Answer", "text": "For decode, both computation and memory traffic are constant with respect to sequence length: roughly 7*D^2 FLOPs and 8*D^2 bytes per head, dominated by reading and writing the FP32 recurrent state. For prefill, FlashKDA runs two kernels (K1 prepares chunk-level tensors, K2 does the chunk-level recurrence) performing 12*C^3 + 8*C^2*D + 6*C*D^2 FLOPs per chunk of size C, so overall FLOPs is O(T*D^2) and memory traffic is O(TC + TD + D^2) — linear in sequence length T." } }, { "@type": "Question", "name": "Does linear attention actually give Kimi K3 a constant-size KV cache during serving?", "acceptedAnswer": { "@type": "Answer", "text": "No. KDA's recurrent state is fixed size at any single token position, but prefix caching requires storing state at positions a future request might resume from. Caching every position would regress memory usage to growing with sequence length, so Moonshot saves recurrent states at coarse granularity — vLLM caches every 32K tokens, plus at prompt boundaries because agentic turns typically start at the end of a prompt. Realistically, KDA greatly reduces KV cache memory but does not consume a constant amount of it." } }, { "@type": "Question", "name": "What are Attention Residuals and Block Attention Residuals?", "acceptedAnswer": { "@type": "Answer", "text": "Attention Residuals apply softmax attention across the depth dimension instead of the sequence dimension: each layer attends over the representations produced by previous layers, using a learned per-layer query vector rather than one generated from the current token. This gives every layer direct, selective access to earlier representations instead of relying on the residual stream alone. Block Attention Residuals group the L layers into N blocks and attend over completed block outputs plus the current block's evolving partial sum, cutting communication from O(Ld) to O(Nd) while staying competitive on quality — about 1.25x compute efficiency over standard residual connections." } }, { "@type": "Question", "name": "What is LatentMoE and why does expert intermediate dimension matter?", "acceptedAnswer": { "@type": "Answer", "text": "LatentMoE compresses routed tokens before dispatch and decompresses them after aggregation; Kimi K3's Stable LatentMoE adds an RMSNorm before the up-projection. Halving the expert input dimension from 7168 to 3584 lets the active expert count double from 8 to 16 at the same communication volume. But the ratio of communication to computation time, which sets how much communication a MoE kernel can hide, works out to (P*F) / (6*m*B) * (1-1/E) — expert intermediate dimension m is the only model configuration in it. That is why recent open weight models, Kimi K3 included, raised the expert intermediate dimension to 3072." } }, { "@type": "Question", "name": "How does InferenceX benchmark Kimi K3?", "acceptedAnswer": { "@type": "Answer", "text": "On recorded internal Claude Code traces rather than fixed 8k1k/1k1k synthetic shapes. An hour of traces is replayed at steady state, with a median of 142k input tokens and 444 output tokens per turn across a median of 65 turns per session. Short outputs per turn are typical of agentic harnesses where the agent calls tools frequently. Replaying real traces exercises production KV cache behavior including prefix cache hits and KV offloading to DRAM." } }, { "@type": "Question", "name": "How does Kimi K3 serve on B200 and B300?", "acceptedAnswer": { "@type": "Answer", "text": "Both Nvidia and AMD shipped Day 0 vLLM recipes with DRAM offload and DSpark speculative decoding, and Day 0 bringup was easier than DeepSeek V4 because images and a speculative decoder model landed alongside the weights. Kimi K3 does not fit on a single B200 node, so it needs pipeline parallelism, and DSpark does not work with PP. On B300 the model fits on one node; after weights, HBM holds about 3.25M tokens of KV cache, and throughput improves with concurrency only up to 8, past which the cache thrashes and hit rates fall below 10% against a theoretical 95%." } } ] }`} --- # Vera Rubin NVL72 vs GB200 NVL72? Inference TCO & Architecture Analysis > Rubin LUT Based Tensor Core, Feynman, Rack Scale, Perf Per MegaWatt, Perf Per Dollar, Software Improvements, Public Rubin Software, PyTorch, vLLM, OpenAI Triton - **Author**: SemiAnalysis - **Date**: 2026-07-23 - **URL**: https://inferencex.semianalysis.com/blog/vera-rubin-nvl72-vs-gb200-nvl72-inference - **Tags**: benchmark, gpu, inference, nvidia, rubin, gb200, gb300, deepseek, trtllm, dynamo - **Reading time**: 20 min _Originally published on the [SemiAnalysis newsletter](https://newsletter.semianalysis.com/p/vera-rubin-nvl72-vs-gb200-nvl72-inference) on July 23, 2026._ [Vera Rubin NVL72 is the second generation of Nvidia’s rack-scale Oberon architecture, and its gains on inference come from extreme co-design](https://newsletter.semianalysis.com/p/vera-rubin-extreme-co-design-an-evolution). Early results from engineering samples are encouraging. Vera Rubin NVL72 running DeepSeek R1 delivers 5.4x performance per MW and 5x performance per dollar over GB200 NVL72 today, and the gap is even wider against GB200 NVL72 during its early bringup in 2025. Vera Rubin is still in the early bringup stage now, so we expect the gap to continue to widen. Rubin’s inference performance will keep improving as software matures, the same pattern we demonstrated for Blackwell in our [InferenceX benchmarks](https://github.com/SemiAnalysisAI/InferenceX), and Rubin still has a long runway ahead. [Nvidia has also recently made available their first public release of the Rubin (SM_107) software stack with CUDA 13.4](https://docs.nvidia.com/cuda/developer-preview/13.4/cuda-toolkit-release-notes/index.html) and has upstreamed Rubin PRs to PyTorch, vLLM and OpenAI Triton Compiler. Blackwell was not able to reuse Hopper WGMMA kernels, but Rubin is able to reuse Blackwell’s kernels, which makes the software bring up process much smoother. For speed of light (SOL) performance, engineers will still need to tune and rewrite kernels but for those that are focused on time to market, Blackwell kernels can be reused. We will also explain Rubin’s new 3-bit programmable LUT tensor core. NVIDIA has also released on GitHub that Feynman is SM_140. [Unlike Blackwell to Rubin, Rubin to Feynman will be a much more complex transition on the kernel front.](https://semianalysis.com/accelerator-hbm-model/) _Related: [Vera Rubin – Extreme Co-Design: An Evolution from Grace Blackwell Oberon](https://newsletter.semianalysis.com/p/vera-rubin-extreme-co-design-an-evolution)_ The early metrics gathered on VR NVL72 come from CoreWeave. We have not independently verified them. Nvidia has committed to submitting verifiable numbers to InferenceX by Q3 CY2026. Google should submit TPUv7 results in the next couple of months, and AMD has committed to MI455X UALoE72. Once those land, the ecosystem gets an objective comparison across systems. In this article we break down Nvidia’s Rubin claims against several baselines, showing where Rubin clearly leads Blackwell and where the lead is thinner. We will also analyze Rubin’s performance per total cost of ownership using our already existing estimates for Rubin’s total cost of ownership (TCO). The TCO for Rubin and many other systems is sourced from our [AI TCO model, which tracks the total cost of ownership of different AI chips, factoring in capex, opex and different other expenses.](https://semianalysis.com/ai-cloud-tco-model/) We also consider performance per watt using our [All-in Utility Provisioned Power Estimates from our Datacenter Model](https://semianalysis.com/datacenter-industry-model/). Finally, we will present [a component by component build up of the Bill of Materials (BoM) for the VR NVL72. This is available in our upcoming SemiAnalysis Bill of Materials (BoM) Model.](https://semianalysis.com/vr-nvl72-model/) Another area where Rubin Oberon NVL72 will fare better than Blackwell Oberon NVL72 is in a much faster production ramp period. This is thanks to Rubin’s simpler cableless compute tray design and learnings from Nvidia’s experience with deploying a rack-scale copper backplane, having invested much effort into ironing out issues with Blackwell’s copper backplane. [Our Accelerator Model tracks quarter by quarter shipments of Rubin at both the package level and the rack level.](https://semianalysis.com/accelerator-hbm-model/)
## A Brief Breakdown of Rubin Chip-Level Microarchitecture Features Going through a complete breakdown of Rubin microarchitecture will have to wait until we obtain ssh access to Rubin systems, allowing us to run [benchmarks similar to those we carried out when first analyzing Blackwell.](https://newsletter.semianalysis.com/p/dissecting-nvidia-blackwell-tensor) However, there are still a few interesting points we can still make. We expect that Rubin bringup will be much more seamless compared to the transition from Hopper to Blackwell, where engineers expended much effort just to port kernels to Blackwell. This simplicity comes from the fact that Rubin is able to run Blackwell SM100-family kernels across all the important kernel libraries in DeepGEMM, FlashMLA, CUTLASS, among others. Moving from Hopper to Blackwell meant rewriting kernels from scratch. Hopper’s kernels don’t run on Blackwell at all. Reusing Blackwell SM100 kernels means a clear time to market advantage, but for speed of light (SOL) performance, engineers will still need to tune and rewrite kernels specifically for the Rubin architecture, though kernel reuse buys them time to focus much more on this kernel tuning. Turning to architectural details, Rubin’s SMEM increased to 328 KiB compared to Blackwell’s 228 KiB. While the default SMEM capacity is 228 KiB, [Rubin comes with an oversized shared memory mode](https://github.com/triton-lang/triton/blob/24fcd59d53e42c7fe7b696c235d12ce039af1015/third_party/nvidia/backend/driver.c#L984-L993) that allows an increase to 328 KiB. Furthermore, TMEM has been increased to 288 KiB up from 256 KiB in Blackwell as the number of columns increased from 512 to 576. The additional columns will allow stashing block scale factors, while keeping the TMEM region for accumulators disjointed from it. [This greatly simplifies block-scaled kernel logic](https://x.com/ReubenConducts/status/2078514481261400109): it saves the kernel writers from carefully pipelining and overlapping MMA matrix and block scaling factor loads.
[Source](https://github.com/triton-lang/triton/pull/10936) Rubin’s TMA now supports inline descriptor updates. There are tons of use cases for this. For example, in an MoE layer, each expert is a separate weight matrix at its own address in HBM, so the TMA descriptor has to point somewhere new on every expert switch. On Blackwell, that meant rewriting the descriptor in memory and synchronizing before the next load. Now, with Rubin, the per-expert offset is passed inline to the TMA instruction so that one descriptor covers all experts, with no in-memory rewrite between them. This removes overhead during token dispatch and improves decode speed at low batch sizes.
[Source](https://developer.nvidia.com/blog/inside-nvidia-rubin-gpu-architecture-powering-the-era-of-agentic-ai/) Inline TMA descriptor update corresponds to the [ISA feature `.override` qualifiers](https://docs.nvidia.com/cuda/developer-preview/13.4/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-overriding-tensor-property-value). TMA instructions require a `tensorMap` object that specifies layout and format metadata. By using the `.override` qualifier, bulk asynchronous copy instructions can reuse the `tensorMap` object as a template, but replacing certain metadata fields, such as strides. In the case of MoE, expert weights have identical shapes, data types, and properties. By overriding the global address, kernel writers can avoid duplicating or replacing `tensorMap` objects when loading different experts. Rubin doubles BF16/FP16 exponential throughput per clock per SM again, which helps overlap Tensor Core work with softmax during attention. FP32 throughput is unchanged from Blackwell Ultra.
Compared to Blackwell NVFP4/MXFP4 Tensor Cores which could only accept UE4M3/UE8M0 block scale factor format, Rubin Tensor Cores will now be able to accept the UE5M3 8-bit block scale factor too. This additional block scale format will allow for more flexibility and less quantization error in certain cases due to its wider range.
It is also important to note that latency has improved for SM-driven NVLink communications through the use of counted writes, which reduces the number of back and forth messages required to send data between GPUs over the copper backplane. This is a huge deal because Blackwell NVLink latency is multiple times higher than that of TPU and Trainium.
Rubin’s Tensor core delivers twice the throughput for FP8 and FP4 as compared to Blackwell. A key change driving this is doubling the k dimension as this means in theory, a GEMM takes half the number of clock cycles to execute. Additionally, the awkward K=96/3xFP4 instructions from Blackwell Ultra are still present, but alongside a new K=128 variant.
In Blackwell, PDL allowed for overlap at the grid level, where the dependent grid needs to wait for all threadblocks in the previous kernel to complete before starting. This allowed for some overlap and hiding of the ramp-down and ramp-up time between kernels, but didn’t come close to the extremely fine-grained overlap that trendy megakernel authors look to achieve. In Rubin, finer-grained overlap is enabled, where the dependent kernel can synchronize with the previous kernel at the threadblock level instead.
Rubin has 2.8x higher global memory bandwidth than Blackwell Ultra through the use of 3D-stacked HBM4 memory. It is unlikely Rubin delivers any improvement versus Blackwell in memory system latency. [Our Accelerator & HBM Model has a complete breakdown of the memory volumes estimates & vendor used in Rubin.](https://semianalysis.com/accelerator-hbm-model/)
Rubin adds 2:4 sparsity support for activations. In every group of four values, two are kept and two are zeroed. The pattern is regular, so the Tensor Core knows where the survivors are, skips the rest, and runs the MMA at twice the rate. A small metadata field tracks which slots were kept. Nvidia shipped 2:4 on weights back in Ampere and nobody used it, because it meant pruning the model and retraining. Rubin applies it to activations at runtime, so no retraining is needed. In attention, QK^T runs dense, then the scores get compressed on the way out of Tensor Memory. Softmax processes only the survivors, and the following GEMM against V runs sparse. The output stays dense, so nothing else in the model changes. It works on MLP activations too. Nvidia has published no accuracy data, and throwing out half the attention scores before softmax is not obviously free. CoreWeave’s DeepSeek R1 results do not appear to use it either, which makes it another Rubin feature with silicon today and no tuned kernels behind it yet.
### Lookup Table Weight Decompression in Rubin SM107 Tensor Core Rubin adds LUT B, a Tensor Core MMA mode that decompresses the weight operand from a lookup table. In this mode, the B operand is a compressed matrix of indices. In a standard inference GEMM, the B operand holds the weights. The weight values live in a lookup table in Tensor Memory. The Tensor Core reads each index and reconstructs the weight value inside the MMA. There is no separate dequantization pass. After lookup, the multiply runs at FP8. In LUT B, every weight position stores a 3-bit index rather than a complete numerical value. The index selects one of eight E4M3 values in the lookup table shared by that weight’s 8×64 block. For example, if the stored index is 5, the Tensor Core uses entry 5 from that block’s lookup table as the weight. The lookup happens inside the MMA, so the kernel never has to construct a separate decompressed weight matrix.
[Source](https://docs.nvidia.com/cuda/developer-preview/13.4/pdf/ptx_isa_9.4.pdf) The lookup table (LUT) does not have to follow the spacing of a conventional uniform or floating-point grid. A quantization algorithm can therefore place more values around dense clusters of weights, use uneven spacing for long tails, or choose an asymmetric codebook when positive and negative weights have different distributions. This flexibility creates the possibility of better accuracy per stored bit. Rubin LUT B has fewer individual codes than FP4, but it can place those codes where a particular weight group needs them rather than accepting the fixed ratios of E2M1. It is not automatically more accurate than MXFP4 or NVFP4, however. One codebook is shared across 512 weights, whereas NVFP4 adapts its scale over much smaller groups of 16. The result will depend on the codebook-fitting algorithm, calibration data, quantization-aware training and whether sensitive layers remain at higher precision. Each index is 3 bits while the lookup table has 8 entries. Each entry is one byte, an E4M3 8-bit float in the reference kernel. This results in 3.125 bits per weight: 3 bits for the index, plus 64 bits of codebook spread across 512 weights. The codebook sits in HBM with the indices, so 3.125 bits per weight is the full stored footprint. The instruction loads the compressed weights into the collector buffer. The Tensor Core can hold them there and reuse them across a run of activation tiles. This is a weight stationary pattern. The mode also has limits as it does not support transpose of the B matrix. The block-scaled formats, NVFP4 and MXFP, also decompress inside the MMA. But they apply one uniform scale per block, not a codebook. Software methods like AWQ reach low bit counts by running a separate dequantization step before the matmul, Rubin is the first NVIDIA Tensor Core input format that reconstructs a non-uniform codebook inside the MMA.
A lower bit rate cuts the HBM capacity that the weights need. It also cuts the bytes that the GPU reads for each weight. At low batch size, weight bandwidth limits the decode step. Fewer bytes per weight then raise decode throughput. A non-uniform codebook also holds accuracy better than uniform rounding at the same bit count. This feature should also have an impact on power efficiency, as fewer bits will need to move through the memory system for each flop. Using Kimi K3 2.8T as an example, at about 4.25 bits per weight, MXFP4 stores 2.8e12 x 4.25 / 8 = about 1,487.5 GB, where GB = 1e9 bytes. At 3.125 bits per weight, the Rubin lookup-table format stores 2.8e12 x 3.125 / 8 = about 1,094 GB (about 1.09 TB). The difference is about 393.5 GB. These figures cover the raw weight payload only, and exclude the KV cache, activations, and any parallelism replication. At 288 GB of HBM4 per Rubin package, the weights alone need about 6 packages in NVFP4 and about 4 packages in the new Rubin format. ## Feynman Architecture Sneak Peek From Blackwell (SM100)/Blackwell Ultra to Rubin (SM107), the jump is relatively small in terms of the microarchitecture, so Rubin can be thought of as a Blackwell kicker architecture. In comparison, Feynman (sm_140) is a completely new architecture family. This will require rewriting lots of kernels from Rubin to Feynman, which is similar to what happened from Hopper WGMMA to Blackwell tcgen05. [Our Accelerator & HBM Model provides a full breakdown of Feynman quarter by quarter volume estimates.](https://semianalysis.com/accelerator-hbm-model/) Feynman’s 3D stacking will be similar to what AMD has been doing with 3D stacking since their MI300X with CDNA3. One of the new features of the Feynman architecture is that it will contain sparsity aware data movement ops. These can be used in sparse GEMMs to improve performance by avoiding pointless loads, stores and FMAs.
## Nuances of CoreWeave VR NVL72 Results Yesterday, [CoreWeave published their benchmarked Vera Rubin NVL72 Inference results](https://www.coreweave.com/blog/nvidia-vera-rubin-nvl72-on-coreweave-10x-more-tokens-per-megawatt-than-blackwell) expressed in units of performance (tokens/sec) per power used (MW). We will break down the nuances of their data and compare their results against Blackwell’s performance using our own InferenceX July 2026 results as a baseline. Click to see the full InferenceX dashboard →
The first notable claim on the CoreWeave-Nvidia chart is that VR NVL72 achieves 10x better token throughput per megawatt than GB200 NVL72 at the iso-interactivity of ~150 tok/s/user. This is about 50% faster than today’s “fast mode” on frontier models. Three things about their chart. The benchmark is single-turn, 8k in and 1k out. The y-axis is output token throughput per megawatt, not total throughput. And their power number covers **both prefill and decode GPUs**, even though only output tokens are counted. InferenceX measures output throughput against **decode GPU watts only**, so we renormalized our data to match theirs for this comparison. It is important to point out that CoreWeave claims to have enabled all of the following inference optimizations on both their baseline GB200 NVL72 and Rubin NVL72 performance results, including but not limited to: - NVFP4 Precision - Speculative Decoding (Using MTP) - Disaggregated Serving (Using Dynamo) - Wide Expert Parallelism - via TensorRT-LLM
The above results seem to suggest that Rubin comes to market with a strong performance gain vs Blackwell out of the gate. However, there are a few nuances that are worth unpacking. First, attentive readers will note that CoreWeave is comparing Rubin against a **GB200 NVL72 2025 baseline**. In some ways, comparing performance at the early stages of GB200 NVL72’s lifecycle is fair, since Rubin performance is expected to massively improve from this early stage in its own lifecycle. Our analysis will also use the GB200 NVL72 early performance results from 2025, but we also compare how GB200 NVL72 did by 2026 as well as the most current GPU worth comparing to: **GB300 NVL72**. We will directly compare GB300 NVL72 performance from early in 2026 with Rubin’s comparable early lifecycle performance.
The second nuance in CoreWeave’s performance results is that they are using DeepSeek R1 671B, a model that is not widely used anymore. One would perhaps wish that CoreWeave used a more modern model like GLM5.2, Kimi K2.5, Qwen3.5, or DeepSeek V4. Even better would be Kimi K3 or Qwen3.8, both of which are [coming soon to InferenceX](https://inferencemax.ai/)! At least CoreWeave is [not using GPTOSS 120B in Summer 2026 like AMD is for MI455X UALoE72](https://github.com/ROCm/aiter/pull/3676) performance metrics. We expect that the fog of war created by benchmarking old models will be cleared up once Nvidia starts benchmarking Rubin on more modern model architectures with InferenceX in Q3 CY2026. Oddly enough, CoreWeave’s choice of using DeepSeek R1 671B is theoretically more favourable towards the Blackwell baseline, and not Rubin. Rubin’s main advantages lie in a higher HBM capacity, higher CPU DRAM capacity, and greater HBM bandwidth, meaning that Rubin is more optimized for multi-trillion parameter models like Fable 5, Gemini Pro, Kimi K3, and Qwen3.8 2.4T. The third noteworthy item is that CoreWeave uses only single turn 8k/1k input/output tokens. Theoretically, multi-turn long context workloads like Agentic Coding should do better on Rubin, due to Rubin’s higher HBM capacity and bandwidth, but this would not be captured on a simple single-turn benchmark. [Our upcoming AgentX benchmark scenario created in collaboration with Weka, LMCache, the vLLM/SGLang community, Nvidia, AMD, and many others in the community will provide a realistic agentic workload to benchmark inference performance.](https://inferencex.semianalysis.com/datasets/cc-traces-weka-062126) We encourage everyone to adopt this inference benchmark! Finally, we note that CoreWeave’s testing was done on a pre-production rack without a scale-out fabric. Specifically, CoreWeave used a Dell Engineering Sample (ES) rack. We do believe these results are valuable as they use wide EP and PD disagg, which uses the NVL72 scale-up backplane and proves that it is working well. This backplane faced many reliability challenges during the ramp of GB200 NVL72 Oberon, [as we have noted in our Accelerator model.](https://semianalysis.com/accelerator-hbm-model/)
[Source](https://x.com/CoreWeave/status/2061146723200962763/photo/2) ### Rubin Versus Blackwell Performance per MegaWatt The metric Nvidia chose to lead with was “output tokens per second per all-in utility megawatt”, counting every GPU in the system. To compare apples-to-apples, we renormalize our own InferenceX benchmark data onto the same total-GPU basis. Below, we put VR NVL72 up against our official GB200 and GB300 July 2026 benchmarks, as well as CoreWeave’s 2025 GB200 baseline. The eye-catching multiples in Nvidia’s charts all come from the 2025 baseline. When comparing benchmark data, we believe we should use figures from the same time period, so the July 2026 GB200 and GB300 benchmarks are the more useful comparison. In theory, datacenter PUE can be lower for Vera Rubin [since Vera Rubin can operate with 45 degrees Celsius coolant temperatures in custom datacenters without chillers](https://blogs.nvidia.com/blog/liquid-cooling-ai-factories/). However, for our comparisons, since most datacenters are designed to accommodate a wide variety of systems, we use the same PUE across the DLC cooled chips. The following pareto curves plot output throughput per total-GPU megawatt against interactivity. Each line stops where its recipe’s frontier ends.
Next, we provide the same data in table form. When a cell says "impossible," we mean that the interactivity is past that recipe's frontier, simply not allowing the configuration to serve that workload at that speed.
Here is the same frontier as bars, across the 100 to 300 tok/s/user band. All four recipes have data through 250 tok/s/user, and only Rubin and GB300 make it to 300 tok/s/user.
Let’s first compare Rubin against the July 2026 GB300 NVL72 baseline. Rubin’s lead is smallest at low interactivity and continues growing through the middle of the interactivity curve. Rubin sits at near 2x the throughput of Blackwell up through 100 tok/s/user, then widens to roughly 4x around 200 tok/s/user, where the gap peaks. Then, the gap begins to narrow again. The headline 5.4x performance gain over GB300 at 300 tok/s/user isn’t Rubin pulling further ahead. Rather, it is GB300 running the last, barely viable point on its frontier, which causes the ratio to balloon. GB200 can’t reach 300 tok/s/user at all. Blackwell’s per-GPU throughput drops off fast as the batch shrinks at high interactivity, while Rubin is still on a flatter part of its frontier. Comparing Rubin against the 2025 GB200 NVL72 baseline is different, showing the biggest lead in the middle of the curve. The gap starts at under 3x at low speeds, but increases to about 10x at 150 tok/s/user (the point Nvidia highlights in their chart), before falling back to 6x at around 200 tok/s/user. The data from that line is accurate, but as we have mentioned it uses a software stack that is a year old, not the GB200 you would run today. At the very top of the interactivity range, the Blackwell curves drop off. By 350 tok/s/user, neither GB200 nor GB300 can serve the workloads at all, leaving only Rubin with an actual curve, delivering 96,446 tok/s/MW at 300 tok/s/user and 70,703 at 350 tok/s/user. Clearly, Rubin is going to give us a lot more “fast mode” than Blackwell. ### Rubin Versus Blackwell Performance per TCO Per-megawatt performance only counts performance against power. Cost per million output tokens folds in the hardware’s total cost of ownership (TCO) including IT capital costs as well as electricity and datacenter costs. [Our TCO model breaks this down comprehensively, providing capital costs and operating costs across server generations.](https://semianalysis.com/ai-cloud-tco-model/) Here, we divide each SKU's all-in TCO by that same renormalized output throughput, so lower is better. Rubin carries a higher TCO per GPU than Blackwell, $3.57 per GPU-hour against $1.84 for GB200 and $2.36 for GB300 in the operator ownership scenario (not rental prices). The charts and tables below will show how Rubin’s $ per token lead comes out a little smaller than its per-megawatt lead.
As with the per-MW analysis, the 2025 GB200 baseline produces the largest gains in performance for Rubin, but the July 2026 GB200 and GB300 numbers are the more relevant baseline for comparison for anyone buying capacity today. The following pareto curves plot cost per million output tokens against interactivity. Each line stops where its recipe’s frontier ends.
Next, we provide the same data in table form, with the ratio showing how many times cheaper Rubin is at each interactivity. Again, a cell marked "impossible" is a speed that the recipe's frontier can't reach.
The chart below plots the frontier as bars across the 100 to 300 tok/s/user band. All four recipes have data through 250 tok/s/user, and only Rubin and GB300 reach 300 tok/s/user.
Against July 2026 GB200 & GB300, Rubin is cheaper at every interactivity, and the gap widens as you climb. The gap starts at about 1.5x cheaper than GB200 through 100 tok/s/user, and improves to 3x by 200 tok/s/user through to 250 tok/s/user. The 5x edge over GB300 at 300 tok/s/user is the same as the per-MW view, where GB300 can barely serve tokens and GB200 can’t serve at this interactivity level at all. The 2025 GB200 NVL72 baseline is once again the more dramatic one, cresting in the middle of the curve. Rubin is a little over 2x cheaper at low speeds, peaks near 8x at 150 tok/s/user, and then moves back to 5x by 200 tok/s/user. Same comment as on the per-MW version: the 2025 GB200 baseline measures a year-old software stack, not the GB200 you would run today. At the very top of the range, things are the same. GB200 has no operating point past 250 tok/s/user and GB300 has none past 300 tok/s/user, so by 350 tok/s/user only Rubin can serve at all, delivering a cost of $4.18 per million output tokens. Click to see the full InferenceX dashboard → _The article continues with Rubin's performance compared against the best-known publicly available MI355X distributed inference performance, plus a brief analysis of how the Triton Compiler, PyTorch, vLLM, and Dynamo software will function on Rubin, in the [subscriber edition on the SemiAnalysis newsletter](https://newsletter.semianalysis.com/p/vera-rubin-nvl72-vs-gb200-nvl72-inference)._ {`{ "@context": "https://schema.org", "@type": "FAQPage", "mainEntity": [ { "@type": "Question", "name": "How much faster is Vera Rubin NVL72 than GB200 NVL72 on inference?", "acceptedAnswer": { "@type": "Answer", "text": "Running DeepSeek R1 on a single-turn 8k input / 1k output workload, Vera Rubin NVL72 engineering-sample results deliver 5.4x performance per megawatt and 5x performance per dollar over GB200 NVL72 measured on the July 2026 software stack. Against CoreWeave's 2025 GB200 NVL72 baseline the per-megawatt gap peaks near 10x at 150 tok/s/user, but that baseline measures a year-old software stack. Rubin is still in early bringup, so the gap is expected to keep widening as its software matures." } }, { "@type": "Question", "name": "What is the Rubin LUT-based Tensor Core weight decompression format?", "acceptedAnswer": { "@type": "Answer", "text": "Rubin SM107 adds LUT B, a Tensor Core MMA mode where the weight operand is a compressed matrix of 3-bit indices. Each index selects one of eight E4M3 entries in a lookup table shared by an 8x64 weight block, and the reconstruction happens inside the MMA with the multiply running at FP8, with no separate dequantization pass. The stored footprint is 3.125 bits per weight: 3 bits of index plus 64 bits of codebook spread across 512 weights. For a 2.8T-parameter model like Kimi K3, that is about 1,094 GB of weights versus about 1,487.5 GB in MXFP4, saving roughly 2 of the 6 288 GB HBM4 packages the weights would otherwise need." } }, { "@type": "Question", "name": "What are the caveats in CoreWeave's 10x tokens-per-megawatt Vera Rubin claim?", "acceptedAnswer": { "@type": "Answer", "text": "Four nuances. First, the 10x compares Rubin against a 2025 GB200 NVL72 baseline running a year-old software stack; against July 2026 GB200 and GB300 InferenceX baselines the per-megawatt lead is roughly 2x through 100 tok/s/user, peaking around 4x versus GB300 near 200 tok/s/user. Second, the benchmark uses DeepSeek R1 671B, an older model that theoretically favors the Blackwell baseline since Rubin's advantages are HBM capacity and bandwidth suited to multi-trillion parameter models. Third, it is a single-turn 8k in / 1k out workload, which does not capture multi-turn agentic serving where Rubin should do better. Fourth, testing ran on a pre-production Dell engineering sample rack without a scale-out fabric." } }, { "@type": "Question", "name": "Can Rubin GPUs run existing Blackwell CUDA kernels?", "acceptedAnswer": { "@type": "Answer", "text": "Yes. Rubin (SM107) runs Blackwell SM100-family kernels across the important kernel libraries including DeepGEMM, FlashMLA, and CUTLASS, which makes bringup much smoother than the Hopper-to-Blackwell transition, where kernels had to be rewritten from scratch. Engineers still need to tune and rewrite kernels for speed-of-light performance on Rubin. Feynman (SM140) is a completely new architecture family, so the Rubin-to-Feynman transition will again require rewriting many kernels." } }, { "@type": "Question", "name": "How does Vera Rubin NVL72 compare to Blackwell on cost per million tokens?", "acceptedAnswer": { "@type": "Answer", "text": "Rubin carries a higher TCO at $3.57 per GPU-hour against $1.84 for GB200 and $2.36 for GB300 in the operator ownership scenario, per the SemiAnalysis AI Cloud TCO Model. Even so, against July 2026 GB200 and GB300 baselines on DeepSeek R1, Rubin is cheaper per million output tokens at every interactivity: about 1.5x cheaper than GB200 through 100 tok/s/user, improving to 3x by 200 through 250 tok/s/user, and about 5x cheaper than GB300 at 300 tok/s/user. GB200 has no operating point past 250 tok/s/user and GB300 none past 300, so by 350 tok/s/user only Rubin can serve at all, at $4.18 per million output tokens." } } ] }`} --- # DeepSeekV4 1.6T Day 0 to Day 43 Performance Over Time — Huawei, GB300 NVL72, MI355X, B200 > Day 0 Inference Performance, InferenceX, 100x performance improvement in 26 Days, Cost per Million Tokens, Huawei 950DT Inference Trace Analysis - **Author**: SemiAnalysis - **Date**: 2026-06-09 - **URL**: https://inferencex.semianalysis.com/blog/deepseekv4-16t-day-0-to-day-43-performance - **Tags**: benchmark, gpu, inference, deepseek, nvidia, amd, huawei, gb300, b300, b200, mi355x, h200, sglang, vllm, trtllm, cann - **Reading time**: 29 min _Originally published on the [SemiAnalysis newsletter](https://newsletter.semianalysis.com/p/deepseekv4-16t-day-0-to-day-43-performance) on June 9, 2026._ The release of DeepSeek v4 marks another step forward for the open model community - unsurprisingly, it is the product of a Chinese lab. The evolution of its performance over time is of paramount importance to the AI Ecosystem. [The open-source InferenceX engineering team has pulled multiple all-nighters to measure performance results for this model on Day 0, Day 1, Day 2, and beyond and bring these results to the world.](https://inferencex.semianalysis.com/) In this article, we will highlight DeepSeek v4’s Day 0 performance and explain the significant improvements made in the subsequent weeks following the model’s release. We will also explain core components of DeepSeek v4’s model architecture and discuss how it was co-designed in part for Huawei Ascend inference. In section 2 of our blog post, we do a comprehensive analysis of DeepSeekv4’s inference on Day 0 Huawei Ascend 950DT. This article serves as the first analysis of Ascend 950DT DeepSeekv4 inference and we break down the compute ↔ communication overlap & the different compute streams that Huawei did to optimize performance. A key goal of InferenceX, especially during a model’s Day 0 release window, is to record each SKU’s performance using open-sourced images and recipes across as many frameworks as possible, regardless of how well these images and recipes perform. This enables us to track improvements over time, which we believe best reflects the real, deployable performance of each chip. The video below shows iterative improvements for non-MTP configs from Day 0 onward for vLLM/SGLang, respectively. [visit inference.com to see the MTP configs from day 0 onwards too](https://inferencemax.ai/).