---
title: "MXFP8"
description: "3 posts about MXFP8, summarised, each linking to the original."
---

# MXFP8
> 3 posts about MXFP8, summarised, each linking to the original.

## Articles

### [Mixture-of-Kittens: our open-source MoE megakernel for NVL72s](https://yomu.fyi/post/mixture-of-kittens-our-open-source-moe-megakernel-for-nvl72s.md)
- Company: [Cursor](https://yomu.fyi/company/cursor.md)
- Author: Stuart Sul, Nash Brown, Henry Wildermuth, William Lin & Federico Cassano
- Published: Aug 4, 2026
- Type: Problem & solution

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](https://yomu.fyi/post/better-moe-model-inference-with-warp-decode.md)
- Company: [Cursor](https://yomu.fyi/company/cursor.md)
- Author: Less Wright, Federico Cassano & Zhiyuan Zhang
- Published: Apr 6, 2026
- Type: Benchmark

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](https://yomu.fyi/post/1-5x-faster-moe-training-with-custom-mxfp8-kernels.md)
- Company: [Cursor](https://yomu.fyi/company/cursor.md)
- Author: Stuart Sul
- Published: Aug 29, 2025
- Type: Benchmark

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.