gpu

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.

Sharding distributes copies; tensor and pipeline parallelism split the computation. This lesson covers the column-then-row trick that lets a transformer block communicate only twice per layer, sequence parallelism for the parts tensor parallelism cannot reach, the pipeline bubble and why it is (p-1)/m, what 1F1B and interleaving actually fix, and how the three axes compose into a 3D layout.

Data parallelism keeps N identical copies of everything. ZeRO removes that redundancy in three stages, and stage 3 takes a 70B model from 1120 GB per GPU to 17.5 GB across 64. This lesson covers what each stage shards, the communication each costs, how FSDP implements it with prefetch and wrapping policy, when offload is worth it, and why stage 3 is not automatically right.

Training memory is dominated by things that are not the model. Mixed-precision Adam costs 16 bytes per parameter, so a 70B model needs 1120 GB of state before one activation is stored. This lesson works through where every byte goes, why data parallelism helps throughput but not memory, how accumulation and recomputation trade compute for space, and what each axis of parallelism addresses.

Decoding is memory-bandwidth-bound, so fewer bits per weight means more tokens per second, not fewer. This lesson builds the mechanism from the arithmetic up: the affine mapping, a worked example done by hand, why group size costs fractional bits, the difference between W4A16 and W8A8, and which formats the hardware actually accelerates.

Rasterisation walks the geometry and asks which pixels each triangle covers. Two ideas make that fast enough for real time: a coverage test that is three linear functions, and a depth buffer that resolves visibility without sorting anything. This lesson builds both and shows why the design maps onto parallel hardware.

Correct gradients are useless if computing them exhausts memory. Radiative backpropagation, path replay backpropagation's constant-memory trick, and why differentiable renderers needed their own compiler.

A model that runs is not a model that serves. This lesson covers what an inference server does that a naive loop cannot, continuous batching and why it dominates throughput, the memory arithmetic that decides which hardware you need, quantization for serving, and how to size a deployment from a traffic estimate.

An AI data center is not a warehouse of ordinary computers; it is a single supercomputer built from tens of thousands of specialized chips wired together. Learn why GPUs beat regular processors for AI, why the network between chips matters as much as the chips, why one company dominates the market, and the cost structure that makes these buildings so expensive.

Modern AI got better mainly by getting bigger, and bigger means more computation, which is why the story of AI is now a story about hardware, power, and money. Learn what compute actually is, why scaling laws made more of it pay off so reliably, the difference between training and inference demand, and why this buildout is unlike previous technology booms.

A working CUDA kernel is the start, not the finish. How to measure occupancy, spot uncoalesced loads and warp divergence, and read the three numbers in Nsight Compute that actually matter.

Why naive matmul on a GPU is bandwidth-starved, and how tiling with __shared__ memory reduces global memory traffic by a factor of the tile size. The classic optimisation, with the kernel that demonstrates it.

The hello-world of CUDA, done properly. Allocate device memory, copy inputs, launch a kernel, copy results back, free, and check every return code. The full driver + kernel in one runnable file.

The five memory spaces a CUDA kernel can see and why they have wildly different speeds. Global vs shared vs constant vs local vs registers, coalesced access, bank conflicts, and a cheat-sheet table you'll actually reference.

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.

Why GPUs eat CPUs for breakfast on some workloads and choke on others. Serial vs parallel execution, Amdahl's law, the latency-vs-throughput trade-off baked into silicon, and a rule of thumb for when to reach for a GPU.