09FlashAttention-3 on Hopper: explain warp specialization, TMA and ping-pong scheduling, and which bottleneck each one removes.▼expertNewNVIDIAOpenAIAnthropic○ sign inFlashAttention-2 reaches well under half of an H100's BF16 peak. The answer explains why in one ratio, exponentials against matmuls, and maps each FA3 technique to the stall it hides.Open full answer →
05Why do tensor core GEMMs use tiles like 128 by 128 by 64, and what limits the tile size on Hopper?▼hardNewNVIDIAAMDOpenAIunlockedEvery dimension of a GEMM tile is pinned by a different piece of hardware. The answer derives the reuse that sets the M and N sides, the swizzle row that sets K, and the register and shared-memory budgets that stop the tile growing.Open full answer →
11Hopper 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?▼expertNewNVIDIAOpenAIAnthropic◆ premiumBoth changes move work off the threads: one takes over address generation, the other takes over operand staging. The answer counts the instructions each one removes and shows why a Hopper kernel is organized around barriers instead of thread blocks that do everything.Open full answer →