Skip to content
← CUDA, Triton & GPU Performance / 08
hardNewNVIDIATogether AIFireworks AI

Fused MoE kernels sort tokens by expert before the GEMM. Why, and what does a grouped GEMM do that a loop over experts cannot?

A Python loop over 256 experts is correct and slow for three separate reasons, only one of which is launch count. The answer builds the expert-sorted layout the fused kernel runs on and counts what it costs in padding.

Updated Oct 2026 · Worked answers on inference mechanisms, with calculations, implementation details, and follow-up questions.

A Python loop over 256 experts is correct and slow for three separate reasons, only one of which is launch count. The answer builds the expert-sorted layout the fused kernel runs on and counts what it costs in padding.

an account raises the free count per topic
UP NEXT ON YOUR JOURNEY
Next in this trackIn a W4A16 GEMM, where does the INT4 unpacking happen, and how does the kernel keep that work hidden under the memory loads?Next in this trackHere is an Nsight Systems timeline of one decode step. Find the launch gaps, synchronizations and CPU stalls, and price each one.Next in this trackHopper replaced Ampere's cp.async and mma.sync with TMA and WGMMA. What changed in each, and how does it change the way you design a kernel?
DISCUSSION · 0

Nobody has commented on this answer. Add a correction or another way to work it.