Chapter 5: Memory Access, Autotuning & Benchmarking

Your vector-add kernel is correct — but is it fast? Most simple kernels are memory-bound, so their speed is decided almost entirely by how efficiently they move data, not by the arithmetic. This chapter is about turning a working kernel into a fast one and, just as importantly, measuring whether you actually succeeded.

Three skills come together here. First, coalesced memory access: structuring offsets so a warp's reads hit contiguous addresses and saturate DRAM bandwidth. Second, autotuning: instead of guessing BLOCK_SIZE and num_warps, you hand Triton a list of candidate configurations and @triton.autotune benchmarks them on real hardware and caches the winner per input shape. Third, benchmarking and the roofline: using triton.testing.do_bench to measure honestly (with warmup and synchronization), converting that to effective bandwidth or FLOP/s, and placing the result on the roofline to see how much performance is left.

The theme is intellectual honesty: a kernel is only as fast as your benchmark proves it is, and you only know it's 'good enough' when it's near the roofline ceiling for its arithmetic intensity.

This chapter covers:

• Coalesced Memory Access: hitting peak bandwidth
• Autotuning Kernels: let the compiler pick the config
• Benchmarking with triton.testing: measuring without fooling yourself
• Reading the Roofline: how close to the ceiling are you?
• Debugging Performance: finding the real bottleneck