Beyond Peak TFLOPs: Why Communication-Compute Overlap Dictates ML Productivity Goodput

In distributed machine learning, peak theoretical TFLOPs is often a vanity metric. What determines the pace, cost, and feasibility of frontier model pretraining is productivity goodput: the sustained rate of mathematically valid, bitwise-reproducible tokens processed per second per dollar of infrastructure.

As model architectures shift from dense transformers to sparse Mixture-of-Experts (MoE), calculating goodput becomes less about raw matrix math and more about managing distributed communication. In typical expert-parallel (EP) setups, the MoE layer can consume more than 50% of the total iteration wall-clock time. This overhead is driven by:

• Cross-accelerator token dispatch and combine collectives over the network fabric.
• Frequent kernel launch boundaries and stream synchronizations.
• Host CPU-GPU dispatch jitter, where fast accelerators drain their execution queues while waiting on host CPU scheduling loops.

Recently, Cursor Research open-sourced Mixture-of-Kittens (MoK), a deterministic MoE training megakernel targeting NVIDIA Blackwell GB200/GB300 NVL72 racks. MoK collapses the entire distributed MoE pipeline into a single persistent kernel, overlapping communication and compute at fine granularity.

Examining MoK's architectural choices—alongside modern Google TPU MoE implementations—reveals how communication-compute overlap directly dictates ML productivity goodput.

1. The Overlap Paradox: Granularity and Wave Saturation

The standard approach to hiding network latency is pipelining: slicing activations into smaller minibatches ($T$) so that chunk $k$ transmits over the network while chunk $k-1$ evaluates on matrix compute engines.

However, overlap is not free. It presents an architectural trade-off:

• Overly fine-grained minibatches (e.g., $T \le 256$ tokens) cause Tensor Cores or systolic arrays to stall at memory synchronization barriers, failing to saturate heavily pipelined matrix-multiply-accumulate (MMA) units and degrading arithmetic intensity.
• Overly coarse-grained minibatches (e.g., $T \ge 10{,}000$ tokens) introduce pipeline bubbles, leaving compute units idle during initial tile fetches and final write-backs.

To balance this trade-off, MoK derives an analytical lower bound for the minibatch token size ($T$) based on Blackwell GPU Streaming Multiprocessor (SM) count ($C = 148$), hidden dimension ($H$), and intermediate expert dimension ($I$):

$$T \ge \frac{2C \cdot 128 \cdot 256}{\min(2I, H)}$$

For a model shape like Kimi 2.5 ($H = 7168$, $I = 2048$), this establishes an architectural saturation floor of $T \ge 2368$ tokens. Empirical profiling confirms this dynamic:

Key Takeaway: Fine-grained communication overlap only translates into productivity goodput when the transfer slice respects the hardware's wave-quantization floor.

2. Eliminating the Invisible Bottlenecks: Direction, Scheduling, and the Host CPU

Hiding network transfer time behind GEMM tiles is only half the battle. Peak throughput often degrades due to software coordination overheads. MoK introduces three mechanisms to address these issues:

A. Asymmetric Networking: Pull-Based Dispatch vs. Push-Based Collectives

Most distributed collective libraries (like NCCL and DeepEP) rely on push-based transfers, where sending GPUs push activations to destination ranks. Under MoE routing imbalance—where specific popular experts receive a disproportionate number of tokens—push-based communication creates severe destination hotspots, saturating inbound links while outbound links sit idle. Furthermore, push-based dispatch requires cluster-wide memory fences, as destination GPUs must confirm that all 71 remote peers have completed writing.

MoK switches to pull-based forward dispatch (paired with push-based combine):

• Signaling Latency Drops from 103 µs to 18 µs (5.8× reduction): Destination SMs issue loads directly from source ranks and track arrivals locally via memory counters, eliminating cross-GPU memory fences.
• Bandwidth Saturation Under Skew Improves by up to 29%: Distributing metadata and payload requests bidirectionally utilizes both send and receive NVLink lanes evenly.
• Zero-Overhead Scheduling Tables: The pull schedule table requires only two columns ({src_rank, src_idx}), allowing schedule calculation to run device-side in less than 3% of layer runtime. Reinterpreting this single table serves all four forward and backward phases.

B. Decoupling the Host CPU via Macrobatch Ring Buffers

On GB200/GB300 NVL72 platforms, high-speed Blackwell GPUs are paired with energy-efficient NVIDIA Grace ARM CPUs. Because the accelerators execute operations rapidly, standard training loops that rely on the CPU to inspect expert token routing, dynamically size destination buffers, and launch kernels cause the GPU to stall, waiting for host dispatch.

MoK removes the host CPU from the execution loop by allocating a static, circular macrobatch ring buffer (a few hundred megabytes) in symmetric memory. Minibatches cycle continuously through this buffer without dynamic allocations, host-device synchronizations, or token dropping.

By reversing the traversal order of the ring during the forward pass, MoK ensures that forward activations needed first by the backward pass remain immediately available, eliminating activation recalculation.

C. Hiding Pipeline Startup and Drain Latencies

Even in an overlapping pipeline, the initial transfer step and final combine step risk leaving Tensor Cores idle. MoK hides these boundaries by interleaving non-dependent computations:

• Initial Dispatch Step: Compute SMs execute the dense shared expert FFN locally over resident tokens while communication SMs pull the first routed minibatch across NVLink.
• Backward Execution Step: MoK defers parameter weight gradients (wgrads) until all activation gradients (dgrads) finish, overlapping wgrads with the final reverse-dispatch communication rounds while accumulating across the token dimension in one pass for bitwise reproducibility.

3. The Result: Real-World Productivity Goodput

By combining inter-SM partitioning (allocating 4 to 52 dedicated communication SMs alongside compute SMs) with fabric-aware pipelining, MoK demonstrates substantial gains over standard pipelines:

• Isolated MoE Layer Throughput: Up to 2.37× faster in MXFP8 forward and 1.78× faster in MXFP8 backward over NVIDIA's optimized reference implementation (HybridEP + Megatron) on GB300 NVL72 racks.
• Cluster-Scale Production Goodput: Deployed across 512 Blackwell GPUs in production for pretraining Cursor's Composer models, end-to-end throughput rose from 760.9 to 1,070.2 tokens/sec/GPU—a sustained 1.41× (41%) goodput gain.

In operational terms, a 41% goodput improvement reduces a three-week training cluster run to approximately two weeks, translating directly into lower compute costs and faster iteration cycles.

4. Cross-Architecture Perspective: Blackwell vs. Google TPU

While Mixture-of-Kittens is written specifically in CUDA/PTX using ThunderKittens tile primitives for Blackwell NVL72 hardware, the engineering principles it implements are broadly applicable to distributed deep learning.

Comparing MoK on Blackwell with modern MoE implementations on Google TPUs (such as TPU v7x Ironwood running Pallas/OpenXLA) highlights how different architectures address the same communication-compute bottleneck:

Key Architectural Takeaways

1. Software SM Partitioning vs. Dedicated Offload Engines: On Blackwell, general-purpose SMs must be shared between compute and communication, requiring developers to dedicate compute resources to drive NVLink. TPUs feature decoupled, dedicated SparseCore hardware units that run collective operations (like All-to-All and Reduce-Scatter) asynchronously, keeping the Matrix Multiply Units (MXUs) focused entirely on compute.
2. Scratchpad Residency Beats Host Dynamic Allocation: Whether via MoK’s macrobatch ring buffers in symmetric GPU memory or Pallas Fused MoE V2 maintaining token residency in TPU VMEM, the performance requirement is identical: eliminating dynamic memory allocation and host round-trips is mandatory to prevent accelerator starvation.
3. Topology Shapes Transfer Strategy: MoK relies on a flat, single-hop 72-device NVLink crossbar where any rank can address any peer with uniform latency. On TPU multi-hop Torus fabrics, communication kernels must account for physical network hops and link contention, using hierarchical intra-chip and inter-chip schedules to maintain throughput.

Strategic Summary for ML Systems Engineers

Maximizing productivity goodput in the MoE scaling era requires a fundamental shift in systems engineering:

• The Host CPU Cannot Mediate Distributed Execution: Accelerators operate too quickly to wait on host operating system queues, dynamic allocations, or CPU-side routing calculations. Distributed routing, buffer recycling, and completion tracking must run autonomously on the accelerator.
• Wave-Aware Pipelining Is Critical: Attempting to overlap communication and compute without calculating the arithmetic intensity and wave-saturation thresholds of the underlying execution units degrades throughput. Overlap must be aligned with hardware execution widths.
• Bespoke Co-Design Delivers Scaling Gains: The 1.41× production speedup achieved by Mixture-of-Kittens demonstrates that generalized software abstractions increasingly leave performance on the table. Maximizing real-world ML goodput requires vertically integrated, fabric-aware kernel designs that unify networking, memory hierarchy, and compute into a cohesive execution flow.