Loading…
Topic
MXFP8
3 posts about MXFP8. Every summary links to the original.
- 3 posts
- 1 company
The Good Stuff
Paid picks we use ourselves. Links carry a referral code — a subscription through one earns us a commission at no extra cost to you.
Newest first
Posts
Mixture-of-Kittens: our open-source MoE megakernel for NVL72s
Mixture-of-Kittens (MoK) is an open-source production mixture-of-experts (MoE) training megakernel designed for GB300 NVL72 systems, where the MoE layer had become a major training bottleneck and could consume more than half of end-to-end training time. It redesigns the layer around a single fully deterministic kernel that fuses communication and computation, minimizing CPU work and CPU-GPU synchronization while overlapping operations across the 72-GPU NVLink domain. The post describes target-specific scheduling, push-versus-pull communication choices, pipelined expert computation, ring token buffers, megakernel structure, determinism, and MXFP8 support for DeepSeek-V3-style layers. Across four model shapes, MoK is reported as up to 2.37x faster for MXFP8 forward and 1.78x for backward, while production training on 512 GPUs reached 1.41x tokens per second versus a DeepEP-based implementation. MoK is fully open-sourced and intended to remain modifiable for other platforms.
Better MoE model inference with warp decode
MoE inference commonly organizes decode around experts, but the post describes “warp decode,” an output-centric kernel for small-batch generation on NVIDIA Blackwell GPUs. Instead of staging tokens into expert-major buffers, each warp owns one output scalar, streams the required weight rows, accumulates contributions from all eight routed experts in FP32, and uses warp-level reductions; the gate/up and down projections are compressed into two kernels. This removes padding, scatter, combine, activation gather, and per-expert output buffers, while keeping BF16 activations and avoiding intermediate quantization. On B200s, the approach delivered 1.84x higher decode throughput, sustained 3.95 TB/s at batch size 32, and produced outputs 1.4x closer to a full FP32 reference.
1.5x faster MoE training with custom MXFP8 kernels
An upgrade from Hopper H100s to Blackwell B200s exposed the Mixture-of-Experts (MoE) layer, implemented with MegaBlocks, as a major training bottleneck: it consumed nearly 53% of forward-pass time and 27% of backward-pass time. The team rewrote the layer at the GPU-kernel level using CUDA, PTX, and some ThunderKittens, while moving from BF16 to microscaled FP8 (MXFP8) with nearly zero loss in training quality. The implementation addresses Blackwell-specific tensor memory constraints and quantization overhead, including custom MXFP8 quantization that sustains more than 6.2 TB/s and emits scale matrices directly compatible with tcgen05.mma. On an internal model, MoE execution became 3.5x faster in both passes, producing a 1.5x end-to-end speedup on Blackwell and 2x versus the original Hopper setup, measured in tokens per second per GPU.
Related topics
All topics →Weekly digest
The best engineering reads, once a week.