GPU and Heterogeneous Computing

Why one more processor — and a very different one

Everything in this ladder, from the stored-program model through virtual memory and threads, quietly assumed one kind of engine: a CPU built to run a single instruction stream fast, with deep branch predictors, big caches, and out-of-order execution all spent on making one thread finish sooner. A GPU is the opposite bet. Instead of a few clever cores chasing low latency, it packs thousands of simple arithmetic lanes and bets everything on throughput: do the same operation across a huge array of data at once. The CPU is a sports car for one passenger; the GPU is a freight train. GPU computing is the practice of using that freight train for general computation, not just pixels.

The word for a machine that mixes such different engines is heterogeneous: a single program runs partly on the CPU and partly on one or more accelerators that are good at things the CPU is bad at. This is not exotic anymore — your phone, your laptop, and every cloud GPU node are heterogeneous. The systems-programming consequence is sharp and is the theme of this whole guide: you can no longer write one stream of instructions over one address space and call it a program. You must decide which engine runs which part of the work, and move data between memories that the two engines do not share. The skill is no longer just being close to the metal — it is being close to two very different metals at once.

SIMT: how thousands of threads run one program

You already met SIMD on the CPU: one instruction applies the same operation to a small vector of values — eight floats in one register, say. The GPU takes that idea to an extreme and dresses it differently, in a model called SIMT, single-instruction-multiple-thread. You write a small function — a kernel — that looks like it computes one element: it reads its own index, loads one value, does some math, writes one result. The GPU then launches that same kernel across thousands of lanes at once, each lane with its own index. Where SIMD makes you think in vectors, SIMT lets you write scalar-looking code and have the hardware run a swarm of copies. It is the same data parallelism underneath, with a friendlier face.

// CPU view: a loop over N elements
for (int i = 0; i < N; i++)
    c[i] = a[i] + b[i];

// GPU (SIMT) view: ONE kernel, launched N times,
// each instance handed its own index i
__global__ void add(const float *a, const float *b,
                    float *c, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N)            // guard: last block may overhang
        c[i] = a[i] + b[i];
}
// the loop has vanished; the hardware *is* the loop

The same vector add, two ways. On the CPU the loop is explicit; the GPU kernel computes one element and the runtime launches N copies, each with a distinct index. The if-guard matters because the launch is rounded up to whole blocks, so some lanes have no element to touch.

Here is the honest catch that the friendly face hides. Those lanes are not truly independent: the hardware runs them in lockstep groups (NVIDIA calls a group of 32 a warp). Every lane in a group must execute the same instruction at the same time. So when your kernel hits an if and some lanes take the branch while others do not, the hardware cannot run both sides at once — it runs the taken lanes while the others sit idle, then runs the other side. This is branch divergence, and it is the GPU's version of the branch cost you learned on the CPU, except the penalty is wasted lanes rather than a flushed pipeline. Branchy code on a GPU silently throws away most of your throughput, which is exactly why GPUs love flat, uniform, arithmetic-heavy loops and hate decision trees.

The offload model: data must cross a gap

A GPU does not run your whole program — it runs kernels that the CPU asks it to run. The CPU stays the boss: it sets up the data, hands the GPU a kernel and its arguments, waits for the result, and carries on. This is the offload model, and the single hardest fact about it is that the GPU usually has its own memory. A discrete GPU sits across a bus (PCIe, or a faster on-package link) with gigabytes of its own RAM, and the address space you allocate with malloc() on the CPU means nothing over there. A CPU pointer is not a GPU pointer. Before a kernel can touch your array, the bytes must be physically copied across the bus into GPU memory; after it finishes, the results must be copied back.

Allocate buffers in GPU memory (the GPU's own RAM), separate from your CPU malloc()'d arrays.
Copy the input bytes host-to-device across the bus — the CPU initiates it, but a DMA engine moves the bytes without the CPU babysitting each one.
Launch the kernel, specifying how many lanes (threads, grouped into blocks) to spawn; the call returns immediately and the GPU runs asynchronously.
Synchronize — wait until the GPU reports the kernel finished — then copy the results device-to-host, and only now check them on the CPU.

That bus crossing is where heterogeneous performance lives or dies. The GPU can add a million floats in the blink of an eye, but if you copy the million floats over, add them, and copy them back, the copies dwarf the compute and the whole exercise is slower than the CPU would have been. The rule of thumb is brutal and worth memorizing: offload pays only when the arithmetic per byte is high enough to hide the transfer, or when the data can live on the GPU across many kernels so you copy once and compute many times. This is the same compute-bound versus memory-bound thinking from the performance rung, now stretched across a physical bus. Modern systems blur it with unified memory, where the runtime migrates pages on demand behind a single pointer — but a page still moves physically when first touched, so the cost has not vanished, only hidden.

The programming models: CUDA, SYCL, and compute shaders

How do you actually write a kernel? The dominant answer is CUDA, NVIDIA's toolkit (first released 2007), which extends C++ with a few keywords: __global__ marks a function as a kernel, the angle-bracket launch syntax kernel<<<blocks, threads>>>(args) spawns it, and runtime calls handle the allocate-copy-launch-copy dance. CUDA is mature and fast, but it ties you to NVIDIA hardware — that is its honest cost. The portable alternatives trade some peak performance for running on more than one vendor's silicon. SYCL is a Khronos standard that expresses the same offload in pure, standard C++ (no language extension), so the same source can target NVIDIA, AMD, and Intel GPUs through different back ends. And compute shaders — in Vulkan, Metal, or WebGPU — let you run general SIMT kernels through a graphics API, which is how the browser guide before this one reaches the GPU from WebAssembly.

Underneath the syntax, all three share the architecture you have now seen: a host program, kernels compiled for the device, explicit or migrated data movement, and SIMT execution. They even share a compiler back end story — a CUDA kernel is compiled to an intermediate form called PTX and then finalized to the actual GPU's instructions by the driver at load time, much like a portable bytecode. So the abstraction stack you internalized for the CPU repeats here, one level over: source, to a portable IR, to vendor machine code, just for a different machine. What is genuinely new is not the compiler pipeline; it is that the target is a thousand-lane throughput engine across a bus, and your job as a systems programmer is to feed it without starving it on data transfer.

Where this sits in the frontier

Step back and see why GPUs headline the systems frontier. The CPU stopped getting dramatically faster per-thread roughly two decades ago — clock speeds plateaued, and the free lunch of instruction-level parallelism hitting diminishing returns. The industry's answer was to go wide and go specialized: more cores, then whole different engines for the workloads that dominate now — graphics, simulation, and above all the dense linear algebra at the heart of machine learning. A modern training run is, underneath, billions of the vector-add kernels you saw above, fused and tiled, kept resident in GPU memory to dodge the bus. Understanding the offload model and SIMT is no longer niche; it is how a large fraction of the world's compute actually runs.

And the GPU is only the most visible accelerator. The same heterogeneous pattern — host plus a specialized device across a bus, with its own memory and its own instruction stream — covers tensor units (TPUs), DSPs, FPGAs, and the smart NICs the kernel-bypass networking guide touched. The systems lesson generalizes: the future is not one faster engine but many specialized ones, and the hard, durable skills are deciding what runs where, managing data movement across memories that are not shared, and reasoning about ordering and synchronization between engines that run at the same time. The single-CPU, single-address-space world this ladder began in was always a simplification; heterogeneous computing is the place where that simplification finally, honestly, ends.