Question 1

Your decode GEMV kernel reaches only about half of HBM bandwidth on an H100. Find the reasons and get it above 80%.

Accepted Answer

A matrix-vector product has one job: keep enough bytes in flight. The answer counts how many that is per SM, then finds the three places a decode GEMV quietly drops below it.

Question 2

In a decode GEMV over a row-major weight matrix, which thread-to-element mapping coalesces, and why does the transposed one not?

Accepted Answer

Two GEMV kernels do the same arithmetic on the same matrix and differ only in which thread reads which element. One streams at HBM speed; the other turns every warp load into 32 separate transactions.

Question 3

A kernel uses 128 registers per thread and 48 KB of shared memory per block. How many warps are active per SM on an H100?

Accepted Answer

Occupancy is the minimum of four separate budgets, and candidates usually compute one of them. The answer works all four for an H100, finds which one binds, and then asks whether raising occupancy would help at all.

Question 4

Your shared-memory transpose tile has bank conflicts. Do the bank arithmetic, apply the padding fix, and say when padding is the wrong fix.

Accepted Answer

The one-column pad is the textbook fix, and it works for FP32 scalars. The answer derives why, then shows what happens with BF16 and with 16-byte vector stores, where the pad stops working and kernels swizzle instead.

Question 5

Why do tensor core GEMMs use tiles like 128 by 128 by 64, and what limits the tile size on Hopper?

Accepted Answer

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

Question 6

A GEMM launches 140 output tiles on a 132-SM H100. What does that cost, and how do you fix it?

Accepted Answer

Eight tiles too many turns a one-wave GEMM into two. The answer prices the waste in microseconds for a real prefill shape, then compares the three fixes, including the one that lives in the scheduler, not the kernel.

Question 7

Write a fused RMSNorm kernel in Triton and count its memory traffic against the unfused PyTorch version.

Accepted Answer

RMSNorm does almost no arithmetic, so its cost is bytes, and the eager PyTorch version moves nine times more of them than it needs to. The answer writes the one-pass kernel, tests it on awkward hidden sizes, and says where fusion stops paying.

Question 8

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

Accepted Answer

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.

Question 9

In a W4A16 GEMM, where does the INT4 unpacking happen, and how does the kernel keep that work hidden under the memory loads?

Accepted Answer

A 4-bit weight is only cheaper if turning it back into FP16 costs less than the bytes it saved. The answer counts the instruction budget per weight at full HBM bandwidth and shows the bit trick that makes the conversion fit.

Question 10

Here is an Nsight Systems timeline of one decode step. Find the launch gaps, synchronizations and CPU stalls, and price each one.

Accepted Answer

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.

Question 11

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?

Accepted Answer

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.

Question 12

Blackwell added tensor memory and 2-CTA MMA. What problem does each one solve, and what changes in a GEMM kernel written for it?

Accepted Answer

Hopper's biggest GEMM tiles spend half the register file on the accumulator and feed every operand byte through one SM's shared memory. Blackwell attacks both limits, and the answer counts the bytes each change gives back.

Question 13

Could you keep the KV cache of hot sequences resident in the H100's 50 MB L2 during decode, and would it be worth it?

Accepted Answer

CUDA lets you reserve part of L2 for data you choose, and the KV cache is re-read every step. The answer measures what fits against what a decode step streams, finds the one case where pinning pays, and the kernel change that pays more.

Question 14

You are about to report one number for your new kernel in a design review. How do you measure it so the number survives scrutiny?

Accepted Answer

A kernel time is a distribution measured under conditions, and a reviewer will ask about both. The answer builds the timing loop, says what each line defends against, and adds the one sanity check that catches most bad numbers.

Question 15

A kernel runs 20% slower on an H200 than on an H100, even though the H200 has 43% more memory bandwidth. How can that happen?

Accepted Answer

On paper an H200 SXM is an H100 with more and faster memory, so nothing should get slower. The answer works through the four ways it can anyway, from the spec sheet you did not read to the tuning table that does not know the new device's name.

CUDA, Triton & GPU Performance

GPU Performance Interview Questions for Inference