Talking to the GPU

Everything so far has been about what a GPU is. This part is about how you get it to do anything, and it starts from a fact that surprises people: a GPU cannot run a program on its own. It has no operating system, no way to read a file, no keyboard, and no notion of "starting". It sits in a slot on the motherboard waiting for the CPU to hand it work, do the work, and hand back the answer. In the vocabulary of GPU programming, the CPU and its memory are the host, the GPU and its memory are the device, and every GPU program is a conversation between the two.

For the first decade of GPUs, that conversation could only be about pictures. If you wanted to use the chip for arithmetic, you had to disguise your numbers as a texture, write your calculation as a pixel shader, and read the answer back out of an image. People did this. In 2007 Nvidia released CUDA (Compute Unified Device Architecture), which let you write ordinary C and run it on the lanes, and that decision, more than any piece of hardware, is why Nvidia rather than anyone else ended up at the centre of AI. This part shows you what that C looks like, and how exactly it fits the machine from Part 7.

A function that runs once per thread

The central idea of CUDA is the kernel: a function that you write once, as if for a single thread, and that the GPU then runs simultaneously on thousands of threads, each with a different index. Here is our running calculation as a kernel, applied to a whole vector at once: every element gets its own w, its own x, and the same b. (A real layer of neurons sums each output over all its inputs, which is the matrix multiply of Part 9 and is left to a library. The element-wise version shows the mechanism with nothing in the way.)

__global__ void neuron(const float* w, const float* x, float b, float* y, int n)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;   // which thread am I?
    if (i < n) {
        y[i] = w[i] * x[i] + b;                       // my one multiply-add
    }
}

The __global__ marker says "this runs on the device". The body is one multiply-add, the same w × x + b we have followed since Part 2. The only unusual line is the first, which works out this thread's index i from three built-in values: which block the thread belongs to, how big blocks are, and the thread's position within its block. Thread number 5 in block number 3, with blocks of 256, computes element 773. Every thread runs the same code and gets a different i, so together they compute every element of y. The if guards the last block, which may have more threads than there are elements left.

The host side is where the conversation happens:

float *d_w, *d_x, *d_y;
cudaMalloc(&d_w, n * sizeof(float));                            // memory on the card
cudaMalloc(&d_x, n * sizeof(float));
cudaMalloc(&d_y, n * sizeof(float));

cudaMemcpy(d_w, w, n * sizeof(float), cudaMemcpyHostToDevice);  // copy w over PCIe
cudaMemcpy(d_x, x, n * sizeof(float), cudaMemcpyHostToDevice);  // copy x over PCIe

int threadsPerBlock = 256;
int blocks = (n + threadsPerBlock - 1) / threadsPerBlock;
neuron<<<blocks, threadsPerBlock>>>(d_w, d_x, b, d_y, n);       // launch

cudaMemcpy(y, d_y, n * sizeof(float), cudaMemcpyDeviceToHost);  // copy the answer back

Read it top to bottom and it is the file clerk of Part 4 at a larger scale. Allocate space on the device. Copy the inputs across. Launch the kernel, telling it how many blocks of how many threads. Copy the result back. The odd-looking triple angle brackets are the launch syntax, and the two numbers inside them are the whole of the parallelism: blocks of threadsPerBlock threads each, a grid.

The model is the hardware

Here is why the programming model is worth a part of its own. Every level of it corresponds exactly to a level of the hardware, and once you see the correspondence, both halves of this series click together.

A thread runs on one lane. Thirty-two consecutive threads of a block form a warp, which is what the scheduler actually issues instructions to, and which is why block sizes are always multiples of 32 in practice. A block is assigned to one SM and stays there until it finishes. The threads of a block can share that SM's shared memory and wait for each other at a barrier, and that is the only way threads cooperate: threads in different blocks cannot talk, because they may be on different SMs or may not be running at the same time at all. A grid is all the blocks of one launch, and the GigaThread engine from Part 7 deals them out to whichever SMs have room, refilling as blocks complete. When a program uses shared memory to hold a tile of a matrix, as in Part 8, it is one block doing it, on one SM, for the threads of that block.

Play with the launch parameters below and watch the mapping.

Two lessons fall out. A vector of a few thousand elements does not come close to filling a GPU, which is the crossover from Part 6 seen from the other side, and a big reason real inference engines batch requests together. And a block size that is not a multiple of 32 quietly throws lanes away. Nothing in the language stops you. The hardware just does less.

The slow road between host and device

Look again at the cudaMemcpy lines, because they hide the single most important performance fact about using a GPU. The host and the device are connected by PCIe (peripheral component interconnect express), the same slot standard as a network card or an SSD. An RTX 4090 uses PCIe generation 4 with 16 lanes, which is about 32 GB/s in each direction. An H100 uses generation 5, about 64 GB/s. Now compare with the device's own memory: 1 TB/s on the 4090, 3.35 TB/s on the H100. The road between host and device is thirty to fifty times narrower than the road between the device and its own memory.

The consequence shapes every GPU program ever written: copy the data across once, keep it there, and run as many kernels on it as you can before copying anything back. A language model's weights are copied to the card when the model loads and never again. Each token generated is hundreds of kernel launches, one for each layer's matrix multiply, attention, normalisation, and so on, all working on data that never leaves the device. The host's role is to queue the launches, which it does asynchronously, throwing them onto a stream and moving on while the GPU works through the backlog. A kernel launch costs a few microseconds of overhead, and for small kernels that overhead can exceed the work, which is why frameworks go to some lengths to fuse small operations into fewer, larger kernels.

From C to the machine

The kernel above is compiled by Nvidia's compiler into an intermediate assembly language called PTX, which is portable across generations. When the program runs, the driver translates PTX into the actual machine code of whichever GPU is present, called SASS, tuned to that chip's SM. This two-stage arrangement is why a CUDA program written for Volta still runs on Blackwell. Each generation has a compute capability number that tells the compiler what it can use: 8.9 for Ada, 9.0 for Hopper, 10.0 for Blackwell. A kernel that uses Hopper's thread block clusters will not compile for capability 8.9, and a spec sheet that lists the number is telling you which features the chip has.

The moat

Almost nobody who uses a GPU for AI writes a kernel. They write torch.matmul(a, b), and a chain of software underneath does the rest: PyTorch calls cuBLAS, Nvidia's matrix library, which contains thousands of hand-tuned kernels, one for each combination of shape, precision, and SM generation, choosing tile sizes and shared memory layouts that squeeze the last per cent out of the tensor cores of Part 9. Neural-network-specific operations go to cuDNN. Attention goes to kernels that have been the subject of research papers. Multi-GPU communication goes to another library we meet in the next part.

That stack is nearly two decades deep, and every machine-learning framework in the world is built on top of it. When people say Nvidia's advantage is software rather than hardware, this is what they mean. A competitor can build a chip with more tensor cores. What they cannot quickly build is eighteen years of libraries, the tools that debug and profile them, and the habit of a million developers who know that if they write CUDA it will work. Alternatives exist, AMD's ROCm, the Triton language from OpenAI, various compilers that target several vendors, and they are improving. But the reason a GPU is the default machine for AI is as much the answer to "what happens when I type torch.matmul" as it is anything in Parts 7 through 9.

Where this leaves us

To use a GPU you write a function for one thread and launch it for a million, and the hardware runs them in warps of 32, in blocks that live on one SM, in a grid that fills the chip. The model and the machine are the same shape, which is why code that respects the machine, multiples of 32, data kept on the device, reuse through shared memory, runs fast, and code that does not runs mysteriously slowly.

One thing the model takes for granted is that the device is one GPU with one memory. For the models that matter now, that stopped being true years ago. A large language model does not fit on one card, and the moment it is split across several, the cards have to talk to each other at a speed PCIe cannot begin to provide. That is the next part, and it is where "the GPU" turns into a rack.

Next: Many GPUs as One: NVLink, NVSwitch, the superchip, the 72-GPU rack that Nvidia now sells as a single machine, and the ladder of bandwidths from HBM down to the network.