Mixture-of-Kittens: our open-source MoE megakernel for NVL72s · Cursor
cursor.com · 4,694 words · saved by 3 readers
We're open-sourcing Mixture-of-Kittens, a deterministic MoE training megakernel for NVL72s that fuses communication and computation into a single kernel.
Today, we're open-sourcing Mixture-of-Kittens (MoK), our production MoE training megakernel for NVL72s. As we have scaled the training and inference of Composer, our agentic coding model, the mixture-of-experts layer has consistently remained the major bottleneck. Depending on the workload and training configuration, it can consume more than half of end-to-end training time. MoK addresses that bottleneck by fusing all MoE communication and computation into a single, fully deterministic kernel. It now powers Composer training across tens of thousands of GPUs. You can try MoK and explore…
saved by
related reading
- Mixture of Experts Quantile Balancing: Validated at 32B-A5B (1e22 FLOPs) Scaleopenathena.ai
- Better MoE model inference with warp decode · Cursorcursor.com
- DeepSeek-V3: A Large-Scale MoE Pretraining Benchmark for MLPerf Training v6.0mlcommons.org
- [2607.16051] Loop the Loopies!arxiv.org
- 1.5x faster MoE training with custom MXFP8 kernels · Cursorcursor.com
- How to Optimize a CUDA Matmul Kernel for cuBLAS-like Performance: a Worklogsiboehm.com
- The 4-bitter Lesson | humans&humansand.ai
- Inside NVIDIA GPUs: Anatomy of high performance matmul kernels - Aleksa Gordićaleksagordic.com
- Compiling Models to Megakernels - by Luminal and Joe Fiotiblog.luminal.com
- Keep the Tokens Flowing: Lessons from 16 Open-Source RL Librarieshuggingface.co
- Grouped GEMM for Imbalanced Experts on Blackwell: A WIP Workloggauravjain.bearblog.dev
- [2410.20399] ThunderKittens: Simple, Fast, and Adorable AI Kernelsarxiv.org