Skip to content
Bud's Wiki

Retire the Abstractions — CUDA DSLs Heading Toward Retirement

TL;DR

Section titled “TL;DR”

A manifesto post from the Hazy Research group (creators of ThunderKittens and the megakernel pattern) arguing that CUDA DSLs are heading toward retirement now that coding agents can do the cognitive-offloading job that C++ template abstractions used to do. The team reports that this year’s MoE megakernel (Cursor’s Mixture of Kittens) was built without a hand-designed intermediate C++ abstraction — the agent worked directly against the complexity, with the abstraction expressed informally in the prompt. Tasks that used to require abstractions became manageable because there is now a “compiler” (the agent) that takes vague instructions and produces code. The provocation: if a prompt can generate every layer of the stack, what is a codebase worth? The post argues codebases will stop being the “source of truth” and become disposable caches of one particular compilation; intent, invariants, tests, and hard-earned domain knowledge stay.

Key claims

Section titled “Key claims”
  • Last year’s Llama-1B megakernel required a hand-designed C++ abstraction layer (github.com/HazyResearch/Megakernels) because the complexity — data structures, inter-thread/SM/GPU synchronization, deeply nested control flow — couldn’t be held in a human head [§“Retire the Abstractions”].
  • This year’s MoE megakernel (Cursor’s Mixture of Kittens) was built without the intermediate abstraction: with agents in the loop, the team could work through the complexity directly and produce target-optimized code from scratch [§“Offloading the offloader”].
  • Two-tier pattern in agent-driven kernel work today: tasks that were previously doable without abstractions (e.g., optimized GEMM kernels) are now “nearly automated given the right prompt” — the team still specifies PTX instructions and warp-specialization design, but the agent gets to near-SOTA quickly; tasks that previously required abstractions (e.g., a megakernel) are not one-shot yet, but the abstraction can now live in an incomplete/messy form inside the prompt rather than as carefully designed C++ templates [§“Offloading the offloader”].
  • Framing claim: an agent acts as a “compiler that can take vague instructions and produce code,” so complexity that was unmanageable last year suddenly becomes manageable — the underlying kernel-design ideas still matter, but it is much easier to express them [§“Offloading the offloader”].
  • Provocation on codebases: a codebase’s value is precision (unambiguous, deterministic on a machine), but that comes with brittleness (tied to language, framework, hardware target, team conventions); a prompt is fuzzy but portable — hand the same intent to a different worker (human or machine) and if the worker is smart enough, you get a correct result without standardizing gaps in advance [§“What next”].
  • The bet is that an intelligent executor reading underspecified instructions replaces the current practice of specifying every gap so a dumb executor doesn’t need judgment; DSLs and frameworks are that specification expressed as a codebase, and lose their ground if the executor stops being dumb [§“What next”].
  • Compiler-precedent argument: opaque intelligent transformation between intent and machine is already accepted (few write the assembly that actually runs; the compiler reorders/inlines/vectorizes), so letting intelligence sit higher in the stack is a difference in degree, not in kind [§“What next”].
  • Three eligibility conditions for retiring an abstraction: (1) an oracle must outlive it — reference implementations, numerics tolerances, intuition about the profiled Gantt chart; if nobody knows what the oracle is, keep the scaffolding [§“Eligibility”]; (2) a shared surface where application/reuse/review attach must be preserved somehow — ten teams generating bespoke megakernels get ten disjoint verification problems that don’t stratify [§“Eligibility”]; (3) an abstraction also transmits knowledge to newcomers, and the post explicitly acknowledges its authors are a biased sample who deleted an abstraction in a domain they knew deeply [§“Eligibility”].
  • What survives the transition: intent, invariants, tests, and hard-earned domain knowledge that currently lives inside ThunderKittens’s abstractions and keeps them correct on hardware. What changes: trust moves up a level — scrutinize the spec and the oracle, not the diff; the implementation becomes a disposable cache of one particular compilation, not the source of truth [§“The pension”].
  • Author-of-record commitment (footnote): the team is still going to maintain ThunderKittens (Vera Rubin kernels coming soon) — this is a directional argument, not an immediate retirement announcement [footnote 1].

Method

Section titled “Method”

The post is an essay, not an experimental writeup. Its “method” is comparing two of the team’s own projects — last year’s Llama-1B megakernel (which required a designed C++ template abstraction to be tractable) and this year’s Cursor Mixture of Kittens MoE megakernel (built with agents, no intermediate abstraction). From that comparison the post extracts a diagnosis (the cognitive-offloader job of abstractions is being taken over by agents), a prediction (CUDA DSLs including ThunderKittens are next on the retirement list, “probably next year, or maybe sooner”), and a set of conditions (oracle, shared surface, knowledge transmission) under which such a retirement is safe.

Results

Section titled “Results”

Not applicable — the post reports no benchmarks. The concrete existence proofs cited are: (a) the team shipped a MoE megakernel this year with agents and no C++ abstraction; (b) agent-driven GEMM kernels are described as reaching SOTA quickly given the right prompt (no numbers). Both align with filed quantitative results elsewhere on the wiki (see below).

Why it’s interesting

Section titled “Why it’s interesting”

This is the manifesto that pairs with two concrete quantitative datapoints already on the wiki. First, CUDA Agent: Large-Scale Agentic RL for High-Performance CUDA Kernel Generation is the “hard evidence” version of the post’s thesis for the simpler GEMM/operator-fusion end: a Seed1.6-based agentic RL system reaches 96.8% faster-than-torch.compile rate on KernelBench with a 128-turn ReAct loop and a torch.compile milestone-based reward, beating Claude Opus 4.5 and Gemini 3 Pro by ~40 points on the hardest split — i.e., the “we still tell the agent which PTX instructions to use, but it gets to SOTA pretty darn quickly” claim as measured. Second, Better MoE model inference with warp decode is Cursor’s public writeup of warp decode, the Blackwell-targeted MoE inference kernel that flips expert-vs-warp parallelism — the codebase-side artifact of the Mixture of Kittens project the post cites as its existence proof for the “no intermediate abstraction” case. The post’s argument is that these two datapoints, taken together, mean the ThunderKittens abstraction layer itself is now on the retirement list — a striking claim from the group that built it. It also sits in productive tension with Kittens virtual machine fuses entire training runs into a single GPU kernel (Ben Spector / Hazy Research) (Ben Spector’s kittens-VM tease, from the same lab) and Loads and Loads of Fluffy Kittens: Compute-Communication Kernels with Multi-GPU ThunderKittens (the ParallelKittens multi-GPU release), both of which double down on the DSL as the load-bearing abstraction; the retire-the-abstractions post is the group publicly reversing on where the compiler layer should sit, mapping neatly onto the AI-for-AI Research cluster’s “trust moves up a level” theme. The three eligibility conditions (oracle, shared surface, knowledge transmission) are also concrete tests worth applying to other filed DSL / framework artifacts on the wiki.

See also

Section titled “See also”