One Token, 2,400 Kernel Launches: Megakernels Bring 1.5–2x Decoding Speedups to vLLM

Start with a counterintuitive phenomenon. Feed a 500-word document to a large language model: the prompt is swallowed instantly with the GPU fully saturated. Then the model starts emitting the answer token by token—and slows down. Compute hasn't changed, so where does the slowness come from?

The bottleneck is data movement. Every time the model generates one token, it must run through dozens of layers from top to bottom, and on every pass it moves all weights from VRAM into the compute units. When you're the only user (low concurrency), each token serves only your request, leaving compute largely idle. The real choke point is memory bandwidth—data can't be moved fast enough, so the compute units sit waiting for material.

Bandwidth itself is already fast. H100's memory moves 3.35 TB per second. The problem is that the movement is chopped into fragments.

2,400 launches per token

Take GLM-5.2 as an example (78 layers: 3 dense plus 75 mixture-of-experts). Running on eight GPUs under vLLM, measurements from the official repo show: generating one token requires 2,459 consecutive micro-kernel launches. Each task finishes and exits, then the next one starts.

Each launch costs one to two microseconds. That sounds trivial, but multiply by 2,459, then add another layer of waste: launches happen in strict order—the next kernel can only begin after every work block of the previous one is done. One kernel splits into 512 work blocks, while a B200 GPU has only 148 SMs, so the last round inevitably leaves 80 SMs idle. Together, these two costs leave only about half of memory bandwidth utilized during token-by-token generation.

The solution was written down in May 2025. Stanford's Hazy Research team published a blog post titled "Look Ma, No Bubbles." Using their own ThunderKittens framework, they welded all of the model's computation into a single never-exiting task—the so-called megakernel. Intermediate results stay in registers and on-chip caches instead of being written back to VRAM, and the kernel never exits, so there are no launch-stop cycles. Measured on H100, memory bandwidth utilization hit 78%, more than 1.5x faster than the inference frameworks of the time. It was validated on NVIDIA hardware.

AMD moved first

What's interesting is the sequence of events that followed. On October 10, the first complete megakernel decoding implementation was merged into vLLM upstream. Not on the NVIDIA side—AMD: built for DeepSeek-V4.1-Flash, running on MI355X, 9,100 lines of code, disabled by default and only active when manually enabled. The backbone of thirty standard layers is split into two halves, each welded into one persistent mega-task, with the attention front-end, attention back-end, mixture-of-experts, and inter-GPU communication all sealed inside.

Three days later, SemiAnalysis issued an alert: AMD had brought megakernels into vLLM, ROCm-only, cutting micro-kernel launches by 10x.

Then came the part the video bloggers missed. On October 16, NVIDIA's own implementation was merged: dedicated decoding for Kimi K3, written in NVIDIA's CuTeDSL, targeting GB300 GPUs, 12,800 lines, with the same official figure of 1.5–2x. Six days behind AMD.

So the claim that "NVIDIA cards still don't have this" was a snapshot from October 13, and it has since expired. The real score is one to one—each vendor has one model merged into upstream. AMD's bench is deeper: megakernels for GLM-5 and GLM-5.2 (8,000 and 9,500 lines) are still in the review queue, with the accompanying standard proposal already posted, and documentation authoring guidelines opened a new entry this morning. The first move is real, but it's hardly a rout.

Why it only works at low concurrency

The performance numbers themselves draw the boundary. AMD's server-side measurements: token latency drops 40.9% at concurrency 4, and 27.4% at concurrency 8; stretch the context to 60,000 tokens and those shrink to 18.4% and 8.1%. At high concurrency, one weight movement is amortized across many requests, the compute units fill up, memory movement stops being the bottleneck, and the fusion gains shrink accordingly. Under high concurrency this switch automatically falls back to the ordinary path.

It treats the "single user" disease. And the most typical single user is an agent. An agent's work is one long chain where every step is low-concurrency token generation, each step waiting on the last—the gaps between tokens simply add up into total task time. AMD's measurement: in agent scenarios, the token interval dropped from 15.0 ms to 6.5 ms, a 2.28x improvement. In the era of long-chain reasoning, this number keeps getting more valuable.

Are the numbers right? Accuracy didn't degrade—GSM8K ran three times scoring 1228, 1234, and 1235, versus 1240 on the ordinary path, with paired-test p-values of 0.22 to 0.63—statistically indistinguishable. Genuinely faster, same answers.

The moat faces its first counterattack

Zoom out. When Stanford wrote the megakernel on H100 in May 2025, it was a research blog validating "where the limit lies." A year and a half later, it has become an engineering standard that two chip vendors race to implement inside the same open-source repository: AMD merged six days earlier, NVIDIA followed six days after, and both share the interface standard.

NVIDIA's software moat has not collapsed. The ecosystem, the operator libraries, the installed base—all still there. What changed is that the defender must, for the first time, respond on its own home turf—previously everyone else scrambled for compatibility with it; now it too has to play on someone else's tempo. For model users, this is purely good news: the same code, with a free 1–2x token generation speedup.

One thing you can watch yourself: the 9,500-line megakernel for GLM-5.2 (PR #60889 in the vLLM repository on GitHub) is still sitting in the review queue.

---

*Sources: vLLM repo PR #60397 (merged 2026-10-10, DeepSeek-V4.1-Flash mono decode, MI355X), #61867 (merged 2026-10-16, Kimi K3 CuTeDSL, GB300), RFC #60904 (performance tables and GLM-5.2 data), SemiAnalysis tweet 2026-10-13, Stanford Hazy Research blog "Look Ma, No Bubbles!" (2025-05-27, H100 measurements).*