Open-Source MoK Megakernel Bypasses CPU Bottlenecks on NVL72

Cursor released Mixture-of-Kittens (MoK), an Apache 2.0-licensed Mixture-of-Experts megakernel built specifically for NVIDIA Blackwell NVL72 systems. By fusing computation and communication into a single deterministic kernel, MoK bypasses the CPU synchronization delays that typically stall large-scale MoE training.

Fusing Compute and Networking

Standard MoE architectures rely on external communication libraries like NCCL to manage inter-GPU routing. On GB200 and GB300 hardware, this creates bottlenecks when the integrated Grace CPUs lag behind the much faster Blackwell GPUs, leading to wasted GPU idle time.

MoK eliminates this dependency during the MoE phase. Built with pure CUDA and PTX using the ThunderKittens tile-primitive library, the kernel overlaps network transfers with computation at a configurable granularity. Instead of waiting for host-device handshakes, MoK uses ring token buffers to handle routing, effectively removing all CPU-GPU synchronization requirements within a 72-GPU NVLink domain.

The release represents a unified evolution of Cursor’s earlier warp decode optimizations, transitioning from custom MXFP8 work into a comprehensive megakernel.

Benchmark Results

Tested against leading public baselines, including implementations using NCCL, DeepEP, and TransformerEngine, MoK demonstrated substantial throughput improvements across both execution directions.

Precision	Forward Throughput Gain	Backward Pass Speedup
MXFP8	2.37x	1.78x
BF16	1.92x	1.58x

In practical application, Cursor reports that integrating MoK into their internal production stack yielded a 1.41x increase in end-to-end training throughput for their Composer and Grok models.

Hardware and Software Constraints

Because the kernel is heavily optimized for a specific network topology, deployment is highly constrained. MoK strictly targets SM100 and SM103 architectures, meaning it cannot be executed on older Hopper or Ampere hardware.

The required software stack includes Python 3.12+, PyTorch 2.10+, and CUDA Toolkit 13.0+.

If you manage training clusters for large-scale Mixture-of-Experts architectures, adopting a fused megakernel approach fundamentally alters your hardware utilization. By abstracting away host-side orchestration and executing the entire MoE layer directly on-device, MoK allows high-cost Blackwell capacity to spend significantly less time waiting on CPU synchronization and network routing.