Skip to content
← CUDA, Triton & GPU Performance / 11
expertNewNVIDIAOpenAIAnthropic

Hopper 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?

Both 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.

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

Both 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.

10 answers per topic instead of 5, plus saved progress and bookmarks · costs nothingor unlock all 136 remaining answers · ₹2,000 / $25

Every worked answer and derivation in all 12 topics, every concept page, and the full practice-test bank.

6 months · One payment · No auto-renewal

UP NEXT ON YOUR JOURNEY
Next in this trackBlackwell added tensor memory and 2-CTA MMA. What problem does each one solve, and what changes in a GEMM kernel written for it?Next in this trackCould you keep the KV cache of hot sequences resident in the H100's 50 MB L2 during decode, and would it be worth it?Next in this trackYou are about to report one number for your new kernel in a design review. How do you measure it so the number survives scrutiny?
DISCUSSION · 0

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