Kittens virtual machine fuses entire training runs into a single GPU kernel (Ben Spector / Hazy Research)

TL;DR

Ben Spector (Hazy Research / Stanford) describes a “kittens” virtual machine built on top of ThunderKittens that takes over the whole GPU and lets the team go past the traditional sequential-kernel model. The headline claim: they can fuse an entire training run into a single kernel — extending the megakernel pattern (one persistent CUDA kernel reading an instruction tensor at runtime) from low-latency inference (ThunderMLA, Llama-1B no-bubbles, the HazyResearch/Megakernels repo) to full forward+backward+optimizer-step training loops. Post 4/5 of a thread; the underlying ThunderKittens/Megakernels work has been the team’s research direction since ThunderMLA (March 2025).

Key claims

The “kittens” virtual machine takes over the whole GPU and allows new kinds of co-optimization, going past the traditional sequential kernel model [tweet body].
An example of this co-optimization is fusing entire training runs into a single kernel [tweet body].
The author teases “even weirder stuff” beyond training-run fusion, without specifying what [tweet body].

Method

Per the broader Hazy Research line of work that this tweet is part of: a megakernel is a single persistent CUDA kernel that occupies the entire GPU and reads an instruction tensor at runtime to decide what work each SM does next. The ThunderKittens framework provides tile primitives (warp / warpgroup / block / grid scopes, with declarative async producer-consumer scheduling) on top of which the megakernel pattern is implemented as an interpreter template — arbitrary “simple kernels” become opcodes that get dispatched inside the megakernel. ThunderMLA (March 2025) instantiated this for MLA attention decode, “Look Ma, No Bubbles” (May 2025) for end-to-end Llama-1B inference, “One Kernel for All Your GPUs” (Sept 2025) for multi-GPU all-reduce / reduce-scatter, and “Loads and Loads of Fluffy Kittens” (Nov 2025) for multi-node MoE megakernels with BF16/MXFP8/NVFP4 support. This tweet extends the line from inference to training. The tweet itself doesn’t ship code; the public artifact closest to this claim is HazyResearch/Megakernels on GitHub.

Results

No quantitative numbers in this tweet. Prior milestones in the line: ThunderMLA reports 20–35% faster than DeepSeek’s FlashMLA on MLA decode (41 µs vs 52 µs at 4-prompt batch on H100); the 8-GPU LLaMA-70B megakernel hides most NVLink communication during decode; the all-reduce kernel beats NCCL on 8× B200 in under 100 lines of CUDA.

Why it’s interesting

This generalizes the megakernel pattern from inference-only primitives to full training loops, which is the load-bearing missing piece in the IO-Aware Kernel Design story so far: every paper there (FlashAttention-4: Algorithm and Kernel Pipelining Co-Design for Asymmetric Hardware Scaling, Gram Newton-Schulz: A Fast, Hardware-Aware Newton-Schulz Algorithm for Muon, Flash-KMeans: Fast and Memory-Efficient Exact K-Means, FlashSampling: Fast and Memory-Efficient Exact Sampling) is a single primitive that fuses one produce/consume pair, but the megakernel-for-training claim says the whole compute graph of forward+backward+optimizer-step can be co-scheduled into one dispatch — which is the natural endpoint of “kernel launches are now the overhead” reasoning. It also pairs with Attn-QAT: 4-Bit Attention With Quantization-Aware Training‘s training-side co-design point: if FA4-backward’s algebraic identities couple forward and backward precision, then a single training kernel makes that coupling structural rather than coincidental. Open question: does the kittens VM expose a programmability cliff (each new model architecture = new megakernel) or is the interpreter template enough abstraction that the cost amortizes across runs.