Cursor Fused MoE Training Into One Kernel and Gave It Away

TL;DR

Cursor open-sourced Mixture-of-Kittens (MoK), the mixture-of-experts training megakernel that powers production training of Composer, its agentic coding model, across tens of thousands of GPUs. MoK fuses every MoE communication and computation step on an NVL72 rack into one fully deterministic kernel. In Cursor's own benchmarks it runs the MXFP8 forward pass up to 2.37x faster than the fastest public baseline, and swapping it into a 512-GPU production run lifted end-to-end throughput from 760.9 to 1,070.2 tokens/sec/GPU. The code is on GitHub under Apache 2.0, with a 28-minute technical writeup explaining the build.

What actually shipped

The MoE layer is the bottleneck of every modern frontier training run: tokens get routed to experts scattered across the rack, which means all-to-all communication is interleaved with grouped matrix multiplies, and the standard way to build that is a parade of separate kernels synchronized by the CPU. Cursor's answer is a megakernel: one kernel launch that holds the entire layer, forward and backward, with some streaming multiprocessors assigned to expert FFN compute and the rest to dispatch and combine traffic, overlapping at the SM task level instead of at kernel launch boundaries.

If that sounds abstract: a conventional MoE stack runs like a kitchen that closes and re-briefs the entire staff between every course. A megakernel keeps every cook at their station for the whole service, grabbing the next ticket the moment they finish the last one. Nobody waits for the room to reset, and the reset was where the time went.

This is not a research demo. The repo is the code Cursor uses for production training of Composer, released under Apache 2.0 with benchmarks, tests, and a functional-layer API: call schedule() once, then forward() and backward().

The numbers

Cursor benchmarked isolated MoE layers on a single NVL72 rack at expert-parallel degree 64, against the strongest public stacks: DeepEP with PyTorch and with TransformerEngine, NCCL with PyTorch, and HybridEP with Megatron, which is NVIDIA's recommended option on this hardware. Against the fastest baseline per configuration:

Layer microbenchmarks are easy to flatter, so the more interesting number is end to end. On a 512-GPU production training run spanning multiple NVL72 racks, replacing the previous DeepEP-based stack with MoK took throughput from 760.9 to 1,070.2 tokens per second per GPU, a 1.41x speedup on the whole run, not just the MoE layer.

The usual caveat applies: these are the vendor's own benchmarks, run on the vendor's own workload. But the benchmark code ships in the repo, and the comparison shapes are not cherry-picked toys. Cursor tested against the expert configurations of Kimi K2.7 Code, GLM-5.2, Qwen3.5-397B-A17B, and DeepSeek-V4-Pro, which is to say, the open-model world's actual production geometries.

How it got there

The writeup is a genuine systems paper, and three tricks carry most of the win. First, communication direction: MoK uses pull-based transfers for forward dispatch, which Cursor measured at up to 29% higher NVLink bandwidth utilization than the push-based equivalent, because it eliminates cross-GPU signaling lanes.

Second, ring token buffers: tokens cycle through a fixed-size ring at minibatch granularity, which fully eliminates CPU-GPU synchronization. The Grace CPUs on a GB300 rack are, in Cursor's dry phrasing, slow relative to the GPUs, so the less they are consulted, the better.

Third, hardware that only exists on Blackwell: MoK leans on Cluster Launch Control, a native work-stealing feature, to keep SMs busy across the compute/communication split. That split is itself tunable, with separate compute-to-comms SM ratios for the forward and backward passes.

One property matters beyond speed: MoK is fully deterministic, producing bitwise-identical output regardless of hardware scheduling. If you have ever burned a week trying to reproduce a loss spike that only happens on Tuesdays, you know why a training team would give up a few percent of peak for that. Here they claim you get the determinism and the speedup.

The catch

The hardware floor is the whole story of who this is for. MoK requires Blackwell SM100 or SM103 silicon, meaning GB200 or GB300 NVL72 racks, plus PyTorch 2.10 and CUDA 13. Benchmarks were run only on GB300 NVL72s. If your homelab has a spare NVL72 rack, congratulations on the electrical service; everyone else reads this one for the ideas.

It also is not fire-and-forget. The config exposes five hyperparameters, and the docs are blunt that minibatch size must be tuned per workload to get the advertised overlap. This is infrastructure for teams already running expert parallelism at rack scale, not a pip install that makes your fine-tune faster.

Why this matters anyway

DeepEP has been the de-facto open standard for MoE expert-parallel communication since DeepSeek released it, and NVIDIA's Megatron stack is the official answer on NVIDIA's own flagship rack. A code-editor company just published something that beats both on that rack, then licensed it Apache 2.0. Frontier-grade training infrastructure keeps commoditizing from the top down, and the direction of travel is consistent: the moat is the data and the recipe, not the plumbing.

The lineage is fun too. MoK is built on ThunderKittens, the Stanford Hazy Research CUDA framework, vendored right there as a submodule, and the megakernel idea traces to Hazy Research's low-latency Llama megakernel work from 2025. What started as an academic bet that kernel launches were the enemy is now training a commercial frontier model, and the kitten branding survived contact with production.

Key Takeaways

• Cursor open-sourced Mixture-of-Kittens, its production MoE training megakernel for NVL72 racks, under Apache 2.0.
• Cursor's benchmarks show up to 2.37x (MXFP8 forward) over the fastest public baseline at the layer level, including DeepEP and NVIDIA's recommended HybridEP + Megatron stack.
• End to end, a 512-GPU Composer training run went from 760.9 to 1,070.2 tokens/sec/GPU, a 1.41x whole-run speedup.
• The kernel is fully deterministic: bitwise-identical results regardless of hardware scheduling, which is rare at this performance tier.
• Hardware floor is steep: Blackwell GB200/GB300 NVL72 only, PyTorch 2.10+, CUDA 13+, and five hyperparameters you are expected to tune.
• Built on Stanford's ThunderKittens; one more layer of frontier training infrastructure just became a free commodity.

Sources: Cursor blog: Mixture-of-Kittens, GitHub: cursor/mixture-of-kittens, AINews (smol.ai), Aug 4 2026