kernels

The pattern that made attention slow recurs across the stack, and once you know what to look for it is easy to find. This lesson applies the diagnosis to normalisation layers, optimizer steps, loss functions and inference decoding, covers why fused attention silently stops applying when a model deviates slightly from standard, and gives the profiling routine that decides where to look first.

The reason most teams never fuse anything is that CUDA asks you to manage threads, shared memory and synchronisation by hand. Triton moves the unit of programming from a thread to a block and hands the rest to a compiler. This lesson covers what that buys, what it still asks of you, how to decide a kernel is worth writing, and how to be sure it is correct.

Softmax normalises over a whole row, which appears to require the whole row before anything can be produced, which appears to require the score matrix to exist. This lesson works through the rescaling identity that removes that obstacle, the running max and sum that make it numerically safe, and the backward pass trick of recomputing the matrix from two saved statistics.

For years attention was optimised by reducing FLOPs, and approximate methods that cut FLOPs kept failing to run faster. The reason is that attention was never compute-bound: it spends its time moving a matrix between GPU memory tiers. This lesson establishes that hierarchy, counts the traffic a standard implementation generates, and shows why an exact algorithm beat every approximation.

How CUDA carves a problem into a grid of blocks of threads. Host vs device code, the __global__ qualifier, the launch syntax, and how every thread figures out which slice of data it owns.