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

TL;DR: A bandwidth-bound kernel only reaches peak if the bytes in flight cover bandwidth times memory latency: at 3.35 TB/s and an assumed 600 ns that is about 2 MB across the GPU, roughly 15 KB per SM. Decode GEMVs miss it in three common ways: too few thread blocks (Llama 8B's down_proj tiled 64 rows per block launches 64 programs on 132 SMs), loads too narrow or too few per thread, and matrices so small that launch gaps rival the 2.5 microseconds the transfer takes. Split-K, 128-bit vectorized loads with unrolling, and fused projections under CUDA graphs fix them in that order.

How to approach it

Establish that the kernel should be bandwidth-bound first: at batch 1 a GEMV does 1 FLOP per byte in BF16, so compute is irrelevant and the only question is bytes per second. Then ask for the shapes, because the failure modes depend on them, and for the profiler evidence: achieved DRAM throughput, achieved occupancy, the grid size, and the timeline gaps between kernels. Ask whether the 50% is per kernel or for the whole decode step; those have different causes.

A strong answer

The governing rule is Little's law applied to memory: sustained bandwidth equals bytes in flight divided by latency. To sustain the datasheet figure, the kernel must keep this many bytes outstanding at once:

bytes in flight = 3.35e12 B/s x 600e-9 s (assumed loaded latency) = 2.0 MB
per SM          = 2.0 MB / 132 SMs                                = about 15 KB
as 16-byte loads                                                  = about 950 per SM

The latency figure is an assumption, and under load it is higher than an idle-pointer-chase measurement, but the conclusion survives any reasonable value: every SM needs on the order of a thousand 16-byte loads outstanding. Anything that reduces that count caps bandwidth proportionally. The causes, in the order to check them:

1. Too few thread blocks. Decode shapes are skinny. In Llama 3.1 8B, down_proj is 4,096 output rows by 14,336 inputs. Tile 64 rows per block and the grid is 64 programs on a 132-SM GPU, so 68 SMs never receive work and the kernel tops out near half of peak whatever else is right. The fix is split-K: give each program a slice of the reduction dimension as well, and combine partial sums. Four splits make 256 programs, about two full waves.
2. Not enough bytes per thread in flight. Two-byte scalar loads need eight times the threads of 16-byte vectorized loads to cover the same bytes, and register-heavy kernels cannot run that many threads. Use 128-bit loads (8 BF16 values), unroll so each thread issues several independent loads before using any, and keep register use low enough for decent occupancy. Profiler signature: low achieved occupancy alongside long-scoreboard stalls.
3. Shapes too small to amortize a launch. The same model's k_proj is 1,024 by 4,096, 8.4 MB, which streams in 2.5 microseconds at peak. A kernel launch and its ramp-up cost a few microseconds on their own, so separate Q, K and V GEMVs can lose a large fraction of each transfer to gaps. Fuse them into one 6,144-row projection (and gate with up), and capture the decode step in a CUDA graph so the CPU is not between kernels.

A split-K GEMV in Triton, checked against an FP32 reference on five shapes including 7 by 129 tails (maximum relative error under $3 \times 10^{-7}$):

import triton, triton.language as tl

@triton.jit
def gemv_splitk(W, x, y, N, K, stride_wn,
                BLOCK_N: tl.constexpr, BLOCK_K: tl.constexpr,
                K_PER_SPLIT: tl.constexpr):
    pid_n, pid_k = tl.program_id(0), tl.program_id(1)
    rows = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)
    acc = tl.zeros([BLOCK_N], dtype=tl.float32)           # accumulate in FP32
    k0 = pid_k * K_PER_SPLIT
    for k in range(k0, k0 + K_PER_SPLIT, BLOCK_K):
        cols = k + tl.arange(0, BLOCK_K)
        w = tl.load(W + rows[:, None] * stride_wn + cols[None, :],
                    mask=(rows[:, None] < N) & (cols[None, :] < K), other=0.0)
        xv = tl.load(x + cols, mask=cols < K, other=0.0)
        acc += tl.sum(w.to(tl.float32) * xv.to(tl.float32)[None, :], axis=1)
    tl.atomic_add(y + rows, acc, mask=rows < N)           # combine the K splits

Rows are contiguous in memory, so each program's loads along cols coalesce, and the compiler vectorizes them. The grid is (cdiv(N, BLOCK_N), number of K splits).

The reversal condition: split-K is not free. The partial sums cost extra writes and an atomic or a second reduction pass, which is waste when the grid already covers the GPU several times over. The LM head (128,256 rows, over 2,000 programs) needs no split at all. Choose the split count from the shape, which is what tuned libraries do with a lookup per shape.

What interviewers probe next

"Your split-K output differs in the last bit between runs. Why?" Floating-point addition is not associative and the atomics land in a different order each run. For bitwise reproducibility, write partials to a buffer and reduce them in a fixed order.

"Now the weights are INT4." Bytes fall 4x, so the same kernel needs proportionally fewer bytes in flight but now does dequantization arithmetic per byte. The kernel must unpack in registers while the next loads are outstanding, which is the design of Marlin-style kernels.

"Batch goes from 1 to 16." Each weight byte now feeds 16 multiply-adds, still far below the ridge of 295, so it remains bandwidth-bound, but the operation is a skinny GEMM and tensor cores become worth using for the arithmetic.

How it's scored

Weak: proposes shared-memory tiling or tensor cores for a batch-1 GEMV, or says "increase occupancy" without counting what is needed.

Hire: derives the bytes-in-flight requirement from Little's law, identifies grid underfill on skinny shapes, and proposes split-K plus vectorized loads.

Strong Hire: quantifies grid coverage against the SM count for real projection shapes, connects small-kernel launch gaps to fusion and CUDA graphs, and knows when split-K costs more than it saves and why its atomics break determinism.

Common mistakes

• Tiling the reduction through shared memory as if it were a GEMM. A GEMV has no reuse of W to exploit; staging adds traffic and latency.
• Reading "50% of bandwidth" from a whole-step timeline and blaming the kernel, when the missing time is gaps between kernels.
• Measuring against a matrix small enough to sit in the 50 MB L2, then reporting bandwidth above the HBM datasheet.
• Splitting K on every shape, including those that already fill the machine many times over.