Making it real: profiling, autotuning, shipping

The single most common failure in kernel work is optimizing something that was never the bottleneck. Parts 1–8 gave you mechanisms; this part is about pointing them at the right thing, and about knowing when to stop.

"This kernel takes 400 µs" is not information. "This kernel moves 268 MB, so it cannot beat 80 µs, and it takes 400" is a bug report with a number attached to it. From Part 3, the procedure is:

1. Count the FLOPs and the compulsory bytes.
2. Compute both ceilings: bytes / bandwidth and FLOPs / peak.
3. The larger is your ideal runtime. If measured throughput reaches roughly 85% of the relevant bandwidth or compute ceiling — equivalently, runtime is no more than ideal time / 0.85 — stop.
4. If you are not, the gap has a cause, and Nsight will name it.

Nsight Compute exposes thousands of counters. About a dozen matter for the kernels in this series, and each maps to something you have already seen.

Two of them deserve special attention because they are the direct measurements of Part 2's two rules:

• Sectors per request. A naturally aligned, coalesced 32-lane, 4-byte load touches 4 32-byte sectors; misalignment can add another. If this instruction reports 32, it generated 8× as many sector transactions as the aligned ideal. That is strong evidence of a lane-to-address problem, not a promise of an 8× whole-kernel speedup or a one-line fix.
• Shared-memory bank conflicts. Zero is the ideal for ordinary loads and stores. A non-zero count is evidence to inspect the stride and instruction layout; padding, swizzling, or a different access mapping are possible remedies, and the right one depends on the instruction.

Two or three profiling passes cover the whole list:

# Full section set on one launch of the kernel, after warm-up — the first
# launches are distorted by autotuning, caching and clock ramp-up:

ncu --set full -k regex:sgemm --launch-skip 2 --launch-count 1 ./app

# Targeted pass when you already know what you are looking for:
# one replay instead of dozens, seconds instead of minutes.

ncu --metrics gpu__dram_throughput.avg.pct_of_peak_sustained_elapsed,\
l1tex__average_t_sectors_per_request_pipe_lsu_mem_global_op_ld.ratio ./app

This is the highest-information view in the profiler. The Warp State Statistics section counts, per issued instruction, how many cycles the warps spent unable to issue — attributed by reason. Each reason supplies a hypothesis; correlate it with the responsible instruction and the memory or pipeline metrics before choosing a remedy.

9.3.1 Reading a breakdown: Part 5's naive matmul

Profile the first kernel of Part 5 — 4092³, 309 GFLOP/s, 1.3% of cuBLAS — with --set full. Speed Of Light shows both headline throughputs in low single digits: nothing is saturated, so this is not yet a bandwidth or arithmetic problem. Achieved occupancy sits close to its theoretical limit, so warp supply is not it either (Part 1). Then Warp State Statistics names the disease: Long Scoreboard dwarfs every other reason, and Sectors per Request reads 32, not 4.

For this kernel, source correlation gives the pair one clear reading: warps wait on the global loads whose column-major walk down B generates eight times the ideal sector traffic. Remap threads so lanes read consecutive columns. The profiler confirms the diagnosis — sectors drop to 4, and the same kernel now measures 1,986 GFLOP/s. Re-profile before celebrating: Long Scoreboard no longer stands alone, while Short Scoreboard and Wait suggest inspecting the MIO dependencies and instruction-level parallelism; the next measured rung stages tiles through shared memory. Form a hypothesis from the top reason, confirm it against the responsible instruction and companion counters, change one thing, and re-measure.

Part 5 had the data point: the best GEMM configuration on an A6000 is BM=BN=128, BK=16, TM=TN=8; on an A100 it is BM=BN=64, BK=16, TM=TN=4, and using the A6000 numbers there costs 6%. The search space is a few hundred valid combinations and the objective is not smooth — occupancy is a staircase, wave quantization is a sawtooth, and shared memory is a cliff.

So you search. In Triton this is one decorator:

@triton.autotune(
    configs=[
        triton.Config({'BLOCK_M':128,'BLOCK_N':256,'BLOCK_K':64,'GROUP_M':8}, num_stages=3, num_warps=8),
        triton.Config({'BLOCK_M':64, 'BLOCK_N':256,'BLOCK_K':32,'GROUP_M':8}, num_stages=4, num_warps=4),
        triton.Config({'BLOCK_M':128,'BLOCK_N':128,'BLOCK_K':32,'GROUP_M':8}, num_stages=4, num_warps=4),
        # ... a few dozen more
    ],
    key=['M', 'N', 'K'],      # re-tune when the shape changes
)
@triton.jit
def matmul_kernel(...):
    ...

# The knobs, and what each one trades off:

BLOCK_M, BLOCK_N   # arithmetic intensity (Part 3) vs registers and
                   # wave quantization. Square often balances reuse for
                   # square GEMMs; rectangular shapes can favor rectangles.
BLOCK_K            # shared-memory footprint per stage → pipeline depth.
                   # Does NOT affect intensity — it cancels out.
GROUP_M            # L2 rasterization order (Part 5, §5.6.1)
num_warps          # threads per block → occupancy staircase (Part 1)
num_stages         # circular buffer depth (Part 2, §2.7).
                   # Fit the device's per-block limit and target residency.

Three practical warnings. First, key matters: if your shapes vary continuously, you will re-tune constantly. Bucket them. Second, autotuning finds the best point in the space you gave it — if every config in your list has the same num_stages, no amount of searching will discover pipelining. Third, Triton runs the kernel repeatedly for every configuration: if it mutates an input, use reset_to_zero, restore_value, or hooks so trials do not accumulate state. And when the space is CUTLASS-sized rather than decorator-sized, use the tool that ships with it: cutlass_profiler sweeps CUTLASS's whole GEMM zoo over your shapes and reports the winner, which is how you pick a tile to pin in code.

Part 3's third regime is invisible on a roofline and dominates small-model, small-batch inference. Every kernel launch costs a few microseconds of CPU and driver work. Twenty small ops per transformer layer, each running for 3 µs and costing 5 µs to launch, spends most of its wall clock in Python and the driver — and the GPU timeline is mostly gaps.

The three tools, in the order you should reach for them:

1. torch.compile — fuses adjacent elementwise work automatically, which removes launches and DRAM round trips (Part 3). Usually the largest single win on an unoptimized model, for the least effort.
2. CUDA Graphs — record a whole sequence of launches once and replay it as a single submission. Removes per-launch CPU cost for a fixed sequence of kernels with fixed shapes. This is what makes small-batch decode viable; it is also why serving frameworks pad batch sizes to a fixed set of buckets.
3. Hand fusion — write the fused kernel. Reach for this when the fusion crosses a boundary the compiler will not cross, such as into a GEMM epilogue.

An honest list, because the answer is "usually don't":

• Plain GEMM. cuBLAS and CUTLASS have absorbed more engineering than you can match. Boehm's 94% took six weekends and was a learning exercise, not a replacement.
• Standard attention. FlashAttention-4 is the current implementation for supported Hopper and Blackwell systems; FA3's historical 75% of H100 peak explains the engineering in Part 7. Use the maintained implementation that supports your hardware and dtype.
• Anything torch.compile will fuse. Try it first; it is one line and it handles the common cases.

The cases where writing one is right:

• A fusion that crosses a library boundary — quantize-and-store in a GEMM epilogue, rotary embedding fused into the attention prologue, an optimizer step fused into the gradient reduction.
• A shape or mask the library does not have — unusual attention biases, MoE routing, sparsity patterns, custom score modifications.
• A schedule that is the optimization — a specific producer/consumer split, a particular TMEM allocation. That is the CUTLASS-level work in Part 6, and it is where the remaining performance actually lives.
• Understanding. Which is most of why you read this.

And the practical middle ground: write it in Triton first. For simple, coalesced streaming kernels, Triton often approaches the same bandwidth ceiling as hand-written CUDA. Memory-bound does not mean every implementation reaches that ceiling: layout, cache behavior, irregular accesses and scheduling still matter. Drop to CUDA when measurement shows that lower-level schedule or layout control is the missing piece.

Nine parts, compressed into the question each one answers.

Everything else was a consequence.

Where to go next

• The PTX ISA. Long, dry, and the only place several things in Part 6 are written down.
• CUTLASS and CuTe, plus Colfax's tutorials, which are the best available writing on Hopper-era kernel structure.
• Programming Massively Parallel Processors (Hwu, Kirk, El Hajj) — the textbook this series compresses. Parts 1–5 are its early chapters with measurements attached.
• Yifan Yang's blog — the swizzle post that prompted this series, and several others in the same style.
• Simon Boehm's matmul walkthrough, with the code for every rung of Part 5's ladder.
• The Triton tutorials — read 03 (matmul) and 06 (attention) against Parts 5 and 7.
• Read a real kernel end to end. FlashAttention's Hopper source, or DeepGEMM, which is small enough to finish.
• FlashInfer's blog and vLLM's — production notes on decode attention, KV-cache layout, and the serving tricks of Part 7, written by the people maintaining them.