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 9.

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 11 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 3. 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 5 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 9 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 10, 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 8 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, whatever it computes: copy the bulk of the data across once, keep it there, and run as many kernels on it as you can before copying anything back. A physics simulation copies its grid over once and steps it thousands of times on the device. A video encoder keeps the frames on the card while it works on them. A game keeps its textures and geometry there for the whole session.

Running a language model is the case this series cares about, so take it as the worked example. When the model loads, its weights and biases, tens of gigabytes for a large one, are copied to the card once and stay there for as long as the model is in use. But that is not the end of the traffic across the link, only the end of the heavy traffic. Every request still sends something in and gets something back: the prompt's tokens go across as a few kilobytes of integers, and each generated token comes back as one integer, or a few thousand numbers if the program wants the probabilities as well. Between those two small copies, the device does all the work. Each token generated is hundreds of kernel launches, one for each layer's matrix multiply, attention, normalisation, and so on, all working on the weights and on intermediate results that never leave the device. So the link carries a huge copy once, then a trickle of small ones for the life of the model, and that is the shape you want: the ratio of bytes crossing the link to arithmetic done on the card is tiny. 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.

Who moves the bytes

A question that the copy lines raise, and that deserves a plain answer: does the GPU fetch its data by itself, or does everything have to go through the CPU? In a conventional PC the answer has two halves, and they are different.

The first half is that the GPU cannot reach a file or a network on its own. It has no file system and no drivers for the disk or the network card. A model's weights start life on an SSD or on the far side of a network, and getting them into the machine's memory is the operating system's job, which is to say the CPU's. So the data always arrives in host RAM first, and from there it has to cross PCI Express to reach the card. That is the sense in which everything goes through the CPU.

The second half is that "through the CPU" does not mean the CPU carries the bytes. When the program calls cudaMemcpy, the CPU does not sit in a loop reading numbers from its memory and writing them to the card. It writes a short description of the transfer, the source address, the destination address and the length, into a queue, and a copy engine on the GPU does the rest: a small piece of hardware, of which every Nvidia GPU has a few, that reads host RAM across the PCIe link and writes into the card's own memory by itself, at the link's full speed, while the CPU carries on with something else. Engineers call this direct memory access, DMA. The one wrinkle is that the copy engine can only read host memory that the operating system has promised not to move, called pinned memory. Ordinary memory is not pinned, so for an ordinary buffer the driver first copies it into a pinned staging area, which does cost CPU time, and then hands it to the copy engine. Programs that move a lot of data allocate pinned buffers from the start.

Once the data is in the card's memory, the CPU is out of the picture. The SMs read and write that memory directly at its full bandwidth, a terabyte a second and up, with no CPU involved, and that is the normal state of affairs for the whole of a kernel's life. The CPU's only remaining job is to queue up the next kernel.

There are three ways of bending this picture, and all three are about avoiding the trip through host RAM rather than avoiding the CPU's role as organiser. The SMs can read host RAM directly across PCIe during a kernel, which CUDA calls zero-copy and which is fine for a few kilobytes of parameters and terrible for anything large, since every access pays the link's latency and its 32 to 64 GB/s. With unified memory, the program allocates one buffer that both sides can address, and when a kernel touches a page that is not on the card the hardware raises a page fault and the driver migrates that page across, so the GPU does in a sense fetch data by itself, page by page, at the cost of a stall of tens of microseconds each time. Be clear about what this does and does not save. On a card in a slot the bytes still cross PCIe, moved by the same copy engines. What is skipped is the explicit copy call and the bookkeeping of two buffers, not the trip. The other meaning of the phrase, the physically shared pool of a Mac or a DGX Spark from Part 6, is a different thing: there the bytes move nowhere because there is only one memory. And in the data centre, GPUDirect lets the copy engines of an SSD or a network card write straight into the GPU's memory over PCIe, skipping host RAM altogether: the CPU still sets the transfer up, but the bytes never touch its memory. Part 13 adds the fourth path, NVLink, over which one GPU reads another's memory directly.

All three of these are the same machinery seen from different angles. In each case a DMA engine moves the bytes across PCIe, the CPU's part is to set it up, and what differs is who asks for the move, how big a piece moves at a time, and whether host RAM is on the route.

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 11. 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 9 through 11.

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.