Every kernel in the step can run near peak bandwidth while the step itself runs at two thirds of it. The answer reads the three rows of the timeline that tell you why, and turns the gaps into milliseconds.
an account raises the free count per topic
← CUDA, Triton & GPU Performance / 10
mediumNewNVIDIAOpenAIAnthropic
Here is an Nsight Systems timeline of one decode step. Find the launch gaps, synchronizations and CPU stalls, and price each one.
Every kernel in the step can run near peak bandwidth while the step itself runs at two thirds of it. The answer reads the three rows of the timeline that tell you why, and turns the gaps into milliseconds.
Updated Oct 2026 · Worked answers on inference mechanisms, with calculations, implementation details, and follow-up questions.
UP NEXT ON YOUR JOURNEY
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?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?
DISCUSSION · 0
Nobody has commented on this answer. Add a correction or another way to work it.