GPU Computing

Definition

A GPU is not a CPU with more cores. It is a chip that spent roughly the same transistor budget on a different bet: arithmetic units, instead of the cache and control logic that make a single thread fast. That bet pays enormously on work which is identical across thousands of data items, and loses badly on work that branches — and the whole of GPU computing, from why deep learning happened on graphics hardware to why your ported code is somehow slower, follows from it.

GPU computing is the practice of running general-purpose calculations on that machine rather than on a CPU. The reason neural networks ended up there is not luck or marketing. A network's forward pass is a stack of matrix multiplications: the same multiply-add applied to millions of numbers, in an order fixed before any data arrives, with no if statement anywhere in the hot loop. That is the exact shape of work a graphics pipeline was built to do to pixels, and it is the only shape of work at which this architecture is overwhelming.

How It Works

The same transistor budget, two different bets

The clearest way to see the trade is to compare two chips launched in the same year, so that process technology and transistor count are held roughly fixed and only the design choice varies. NVIDIA's H100 SXM5 and AMD's EPYC 9654 "Genoa" both shipped in 2022 on TSMC's 5 nm-class nodes. NVIDIA says the H100 carries 80 billion transistors; AMD says fourth-generation EPYC carries more than 90 billion. Call it the same budget, within about 15%. Here is what each spent it on:

The H100 figures are NVIDIA's published specification. The EPYC FP32 number is derived: AMD quotes 5.376 TFLOPS of peak FP64 for the 9654, and Zen 4's two 256-bit FMA pipes do FP32 at exactly twice the FP64 rate, so 10.75 TFLOPS FP32.

Two ratios fall out, and they point in opposite directions. The GPU gets about 6.2× the FP32 throughput from the same transistors (67 ÷ 10.8), despite running at roughly half the clock. And the CPU carries 7.7× the last-level cache (384 ÷ 50). Normalise the cache by the arithmetic it feeds and the gap is starker still: 384 MB per 10.8 TFLOPS is 35.6 MB of cache per TFLOPS, against 50 MB per 67 TFLOPS, or 0.75 MB per TFLOPS — a 48-fold difference in how much memory each design keeps close to its arithmetic.

That cache is not a luxury on the CPU side; it is the entire mechanism. A working set that fits in 384 MB never makes the trip to DRAM at all, so the CPU avoids memory latency. A GPU holding a 70 GB model has no such option. It has to survive the latency instead — which is where the thousands of threads come from.

Latency you cannot avoid, so you outrun it

Reading from off-chip memory is slow on every machine. Measured on a Hopper-generation part, a global-memory access costs 478.8 cycles, against 29 cycles for the on-chip scratchpad (Luo et al., 2024, Table IV). The H100 SXM's shipping clock follows from its own datasheet — 67 TFLOPS ÷ (16,896 lanes × 2 FLOP) ≈ 1.98 GHz — so those 479 cycles are about 242 nanoseconds of waiting.

Now apply Little's law, which says that to sustain a given rate through a pipe of a given delay, the amount in flight must be rate × delay:

811 KB  =  3.35 TB/s  ×  242 ns

Roughly 811 KB of memory requests must be outstanding at every instant, or the memory system sits idle. A CPU cannot come close — a handful of cores each tracking a few dozen outstanding misses gets you into the tens of kilobytes — which is exactly why it buys cache instead. The GPU's answer is to keep an absurd number of threads resident and switch between them for free. An H100 SM holds up to 2,048 threads — 64 warps — and there are 132 SMs:

132 SMs × 2,048 threads = 270,336 resident threads
270,336 × 4 bytes       = 1.08 MB of loads in flight

1.08 MB against the 811 KB required: a margin of 33%, with one four-byte load per thread. That is the answer to "why does a GPU need a quarter of a million threads?" It is not that the problem has that much parallelism to give. It is that this is the minimum concurrency at which the memory system can be kept busy at all, and the machine was sized to hit it exactly.

Switching between those threads has to be free, which is why an SM carries a 256 KB register file — 65,536 32-bit registers, so that every thread's state can live in registers permanently with nothing to save or restore. Divide it up and you get the constraint that quietly governs GPU performance: 65,536 ÷ 2,048 = 32 registers per thread at full occupancy. Ask the compiler for 64 registers per thread and only 1,024 threads fit per SM. Occupancy halves, loads in flight halve to 541 KB, and you drop below the 811 KB the memory system needs — at which point HBM runs at about two-thirds speed and nothing in your code has changed, errored or warned.

How much bandwidth there is to saturate, and what happens when a workload cannot supply enough arithmetic per byte, is the memory wall's subject. This is the other half: the latency of each access, and the fact that a GPU's answer to it is concurrency rather than caching.

One instruction, thirty-two threads

The arithmetic bet only pays if one instruction can drive many lanes, because instruction fetch and decode are exactly the control logic the GPU declined to buy. NVIDIA hardware therefore groups threads into warps of 32, which execute a single instruction together. This is called SIMT — single instruction, multiple threads — and it differs from a CPU's SIMD in the contract rather than the silicon. AVX-512 puts the vector width in the instruction, so the programmer sees it and the compiler must vectorize explicitly. SIMT lets you write ordinary scalar code per thread and has the hardware mask off lanes that should not be participating. Convenient, and quietly expensive.

When the threads of a warp disagree about a branch, the warp does not split. It executes both paths in sequence, with the non-participating lanes masked off each time. Two paths cost twice. Thirty-two paths cost thirty-two times:

switch (data[i] % 32) { ... }     // 32 bodies, 10 instructions each

  all 32 threads agree   →  10 instruction issues
  all 32 threads differ  → 320 instruction issues   (32× slower)

Up to 32× is the price of a fully divergent branch, and it is the single most useful number a newcomer to GPU performance can hold. Note the shape of the fix it implies: divergence costs nothing when a branch is taken uniformly within a warp. A kernel that sorts or bins its work so that threads 0–31 all take the same path pays nothing for the branch, while the same kernel on unsorted input pays up to 32×. Same code, same data, same total work.

The programming model that exposes warps, blocks and the memory hierarchy to you is CUDA; this page is about why the hardware underneath has that shape.

Amdahl's law, or the antidote to "just use a GPU"

None of the above helps the part of your program that cannot be parallelised, and Gene Amdahl's 1967 observation puts a hard ceiling on what that means. If a fraction p of the work parallelises across N lanes and the rest does not, total speedup is:

S(N) = 1 / ( (1 - p) + p/N )

Run it against the H100's 16,896 lanes:

A workload that is 5% serial cannot exceed about 20×, and 16,896 arithmetic lanes buy you 19.98 of them. The other 16,876 do nothing. This is why the honest answer to "should I port this to a GPU?" starts by asking what fraction of the wall-clock is the parallel part — and why the serial 5% is so often not in the model at all but in data loading, Python glue, and the host-to-device copy across a PCIe 5.0 x16 link that moves about 64 GB/s against HBM3's 3,350 GB/s, a factor of 52.

Real-World Applications

Transformer training and inference are the workload this machine was accidentally built for. Every layer of a transformer is a fixed sequence of large matrix multiplies over a batch, and the sequence does not depend on the values in the data. The same instructions run for every token, so warps never diverge, and the batch dimension supplies concurrency far beyond the 270,336 threads needed to hide latency. When people say GPUs are "good at AI", this is the whole of the claim — and it stops being true the moment control flow depends on the data.

Mixture-of-experts is where the branching comes back, and the fix is the warp fix at a larger scale. A mixture-of-experts model routes each token to a small subset of expert weight matrices, so different tokens in the same batch need different matrices — divergence, at the level of the whole layer. No competent implementation branches per token. Instead it sorts the tokens by assigned expert, then runs one dense matrix multiply per expert over its contiguous group. That is precisely the "align the branch to warp boundaries" move, and it is why MoE kernels are dominated by permutation and gather logic rather than by arithmetic.

Ray tracing is the case where NVIDIA gave up and built separate hardware. Traversing a bounding volume hierarchy is a pointer chase whose next step depends on the last comparison, so the rays in a warp diverge after a few levels and stay diverged; the literature describes BVH traversal as exhibiting "frequent control-flow and memory divergences" (Zhu, 2022). NVIDIA's answer in the 2018 Turing generation was to add RT cores — fixed function units that traverse the hierarchy autonomously while the SM goes off and does something else. When a workload's divergence cannot be sorted away, the SIMT machine does not get better at it; it gets a co-processor.

Bitcoin mining shows the other end of the same spectrum. GPUs displaced CPUs at it almost immediately, then were themselves displaced by ASICs within a couple of years, because the workload is not merely branch-free but fixed: one hash function, forever. Against that, even the GPU's residual generality is wasted silicon. A TPU or an NPU is a partial version of the same argument applied to neural networks, and tensor cores are that argument made inside the GPU.

Key Concepts

• SIMT is a contract, not a speed: the hardware executes 32 lanes per instruction either way; SIMT differs from SIMD only in hiding that width behind per-thread code. The performance model does not hide it, which is why code that looks scalar can behave like a vector unit with 31 lanes switched off.
• Occupancy is a latency budget, not a utilisation score: resident warps matter because outstanding loads are the only thing keeping HBM busy, and the H100 has just a 33% margin at full occupancy. Trading occupancy for registers is usually a loss.
• Coalescing costs nothing when you get it right and up to 32× when you do not: 32 threads reading 32 adjacent addresses become a few memory transactions; 32 scattered addresses become up to 32 separate ones. It is warp divergence's twin in the memory system.
• Utilisation counts residency, not work: nvidia-smi reporting 100% means a kernel was resident during the sampling window, not that the arithmetic units did anything. A fully stalled kernel and a perfect one report the same number.
• Perf-per-watt is part of the trade: 67 TFLOPS at 700 W is 96 GFLOPS/W against the EPYC's 10.8 TFLOPS at 360 W, or 30 GFLOPS/W — 3.2×. The arithmetic bet buys efficiency, not just speed.

Challenges

The failure mode is silence. Every mistake described above produces a program that is correct and slow, with nothing to indicate which. The compiler allocating 64 registers instead of 32 halves your occupancy without a warning. An unsorted branch costs 32× without a warning. Scattered memory access costs an order of magnitude without a warning. And the profiler, in the default reading, reports the GPU as fully utilised throughout, because a warp stalled on a memory load still counts as resident. Debugging GPU performance is largely the practice of distrusting the one number that is easy to read.

Many algorithms simply do not have the shape. Graph traversal, sparse iterative solvers, branch-and-bound search, discrete-event simulation, tree-structured recursion: these are built out of data-dependent control flow and irregular memory access, the exact intersection of the two things this architecture is worst at. Graph neural networks land awkwardly in the middle — the per-node transforms are dense matrix work the GPU loves, the neighbourhood gather between them is a scatter it hates. For a genuinely irregular workload a GPU port may be slower than good CPU code, and a rewrite that regularises the algorithm usually beats one that merely parallelises it.

The serial fraction migrates rather than disappearing, and so does the latency problem. Once the model is on the GPU, Amdahl's residue reappears as host-side orchestration: thousands of small kernels, each a few microseconds of launch overhead for a few microseconds of work, with a Python interpreter in the loop. Kernel fusion and CUDA graphs exist to attack exactly this, which is why torch.compile can be worth a large factor on a small model and nothing on a large one. And having spent its area on lanes rather than on making one thread fast, the GPU has no answer at all to a long serial dependency — which is precisely what generating tokens one at a time is.

Future Trends

The trade is being pushed further in the same direction, inside the chip. Tensor cores are already a fixed-function matrix engine embedded in a programmable machine: less generality, more arithmetic per transistor. Every generation moves more of the die into units that do one thing, and the endpoint of that direction is a TPU or a domain-specific ASIC. The live question is not whether GPUs beat ASICs but how much programmability the market will pay for — and that answer moves every time model architectures do.

Compilers are taking the warp away from the programmer. Triton, torch.compile and the XLA lineage let you describe work in terms of tiles and leave the mapping onto warps, registers and occupancy to the toolchain. The arithmetic on this page — registers per thread, loads in flight — is increasingly reasoned about by a compiler rather than a human, against exactly the numbers above.

Workloads, meanwhile, are getting less uniform at exactly the wrong moment. Mixture-of-experts routing, speculative decoding with its accept/reject step, dynamic sparsity, early exit and variable-length sequences are all forms of data-dependent control flow, arriving just as the hardware has bet even harder on uniformity. The resolution so far has always been the same one: do not branch, sort. Whether that keeps working as models get more conditional is the open architectural question of the next few years.

Code Example

Two calculations, both small enough to run, that make the arithmetic above concrete. First, the cost of divergence — the two kernels do identical total work and differ only in whether the branch aligns to warp boundaries:

__global__ void divergent(const int* in, float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= n) return;
    // in[] is unsorted, so neighbouring threads take different paths:
    // the warp serialises every path any of its 32 threads needs.
    switch (in[i] & 31) {
        case 0:  out[i] = work0(i);  break;
        case 1:  out[i] = work1(i);  break;
        // ... 32 cases
    }
}

__global__ void converged(const int* in_sorted, float* out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= n) return;
    // in_sorted[] is bucketed by (value & 31), so all 32 threads of a
    // warp hit the same case and exactly one path executes.
    switch (in_sorted[i] & 31) {
        case 0:  out[i] = work0(i);  break;
        case 1:  out[i] = work1(i);  break;
        // ... 32 cases
    }
}

Second, the occupancy calculation that decides whether your kernel can keep HBM busy at all:

# H100 SXM5, from NVIDIA's published specification
SMS            = 132
REGS_PER_SM    = 65536          # 256 KB register file / 4 bytes
MAX_THREADS_SM = 2048
BANDWIDTH      = 3.35e12        # bytes/s
LATENCY        = 242e-9         # seconds (479 cycles at ~1.98 GHz)

required = BANDWIDTH * LATENCY  # Little's law: bytes that must be in flight

for regs_per_thread in (32, 48, 64, 96, 128):
    threads_per_sm = min(MAX_THREADS_SM, REGS_PER_SM // regs_per_thread)
    in_flight = SMS * threads_per_sm * 4    # one 4-byte load per thread
    print(f"{regs_per_thread:4} regs -> {threads_per_sm:5} threads/SM, "
          f"{in_flight/1e3:7.0f} KB in flight, "
          f"{'saturates' if in_flight >= required else 'STARVES'} "
          f"({in_flight/required:.2f}x of {required/1e3:.0f} KB needed)")

#   32 regs ->  2048 threads/SM,    1081 KB in flight, saturates (1.33x of 811 KB needed)
#   48 regs ->  1365 threads/SM,     721 KB in flight, STARVES (0.89x of 811 KB needed)
#   64 regs ->  1024 threads/SM,     541 KB in flight, STARVES (0.67x of 811 KB needed)
#   96 regs ->   682 threads/SM,     360 KB in flight, STARVES (0.44x of 811 KB needed)
#  128 regs ->   512 threads/SM,     270 KB in flight, STARVES (0.33x of 811 KB needed)

The model is deliberately crude — real kernels issue several independent loads per thread, which is the other way to buy memory-level parallelism, and real access patterns are wider than four bytes. But the shape is right, and it is the shape that matters: a few extra registers per thread can take a correct kernel below the concurrency its own memory system requires, and nothing anywhere will tell you.