1.5x faster MoE training with custom MXFP8 kernels · Cursor
We want to build the world’s best AI coding models, but training large language models can be expensive. For instance, our largest internal models can take weeks to train on tens of thousands of GPUs. This is not only computationally expensive, but it also slows the pace at which improvements reach our users. We recently began upgrading from Hopper GPUs (H100s) to Blackwell GPUs (B200s) and saw this as an opportunity to deeply optimize our training workloads. Profiling revealed that the main bottleneck was the Mixture-of-Experts (MoE) layer, implemented with MegaBlocks, which accounted for nearly 53% of forward-pass time and 27% of backward-pass time. That’s why, over the past few weeks, we rewrote the entire MoE layer from scratch at the GPU kernel level with zero dependencies on any CUDA libraries. Instead, we used pure, good old CUDA and PTX, with a few bits of ThunderKittens sprinkled in. As a result, we achieved a 3.5x improvement in MoE layer performance for both the forward and
Blog / research We want to build the world’s best AI coding models, but training large language models can be expensive. For instance, our largest internal models can take weeks to train on tens of thousands of GPUs. This is not only computationally expensive, but it also slows the pace at which improvements reach our users. We recently began upgrading from Hopper GPUs (H100s) to Blackwell GPUs (B200s) and saw this as an opportunity to deeply optimize our training workloads. Profiling revealed that the main bottleneck was the Mixture-of-Experts (MoE) layer, implemented with MegaBlocks , which
Explore this link on the map →saved by
related reading
- Best practices to accelerate inference for large-scale production workloadstogether.ai
- Inside NVIDIA GPUs: Anatomy of high performance matmul kernels - Aleksa Gordićaleksagordic.com
- How to Think About TPUs | How To Scale Your Modeljax-ml.github.io
- All About Rooflines | How To Scale Your Modeljax-ml.github.io
- How To Scale Your Modeljax-ml.github.io
- How to Optimize a CUDA Matmul Kernel for cuBLAS-like Performance: a Worklogsiboehm.com
- Look Ma, No Bubbles! Designing a Low-Latency Megakernel for Llama-1B · Hazy Researchhazyresearch.stanford.edu
- Making Deep Learning go Brrrr From First Principleshorace.io
- Transformer Inference Arithmetic | kipply's blogkipp.ly
- The Best GPUs for Deep Learning in 2023 — An In-depth Analysistimdettmers.com
- CVPR2023_eff_tutorial_molchanov.pdfnvlabs.github.io
- Reiner Pope – The math behind how LLMs are trained and serveddwarkesh.com