SIMD, GPUs and coprocessors

The cores of the previous chapter find parallelism between instructions: while one add waits, another independent one runs. That's instruction-level parallelism, and it tops out at a handful of instructions per cycle. Many programs have a far simpler kind of parallelism available: the same operation applied to thousands or millions of values. Brightening an image adds the same number to every pixel. Training a neural network multiplies large matrices. This chapter is about hardware built for that data parallelism: SIMD instructions inside a normal core, GPUs, and other coprocessors.

Instruction-level parallelism also includes VLIW processors, which the VLIW and Itanium chapter covers. A classic VLIW example is Philips's TriMedia, a media processor from the early 2000s that no longer exists as a product line; its ideas live on in digital signal processors such as Qualcomm's Hexagon, a VLIW design found in Snapdragon phone chips. What the TriMedia's multimedia instructions did (work on four bytes packed in one register) is now in every CPU, as SIMD.

One instruction, many data

A SIMD instruction (single instruction, multiple data) treats a wide register as a small array of lanes and does the same operation on every lane at once. A 128-bit register holds sixteen 8-bit values, eight 16-bit, four 32-bit or two 64-bit ones. One paddd on x86, or one add v0.4s, v1.4s, v2.4s on ARM64, adds four pairs of 32-bit integers in the time of one.

You can get the idea without any special hardware. The demo below adds sixteen bytes two ways: one at a time, then four at a time packed into 32-bit ints, a trick called SWAR (SIMD within a register). The only difficulty is the carries: a carry out of one byte must not leak into the next lane. So add4 clears each lane's top bit before adding, which makes a carry across lanes impossible, and puts the top bits back with an XOR.

All 16 lanes match. Look at the last one: 200 + 100 gives 44, because each lane wraps around modulo 256, exactly as x86's paddb does. For pixels, wrapping from bright to dark is wrong, so SIMD instruction sets also have saturating adds (paddusb on x86, uqadd on ARM64) that stop at 255 instead. The TriMedia's multimedia units already had saturating arithmetic.

In the simulator, which compiles without optimization, the byte-by-byte loop took about 400 instructions and the packed loop about 240, function calls included. Real SIMD hardware removes the masking entirely: the adder itself is split at lane boundaries, and one instruction does all sixteen bytes. (The simulator doesn't implement SSE, which is why the demo does it by hand.)

SIMD registers, from SSE to SVE

Every mainstream instruction set has grown SIMD extensions, each wider than the last:

The x86 story is one of repeated widening, and each new width needed new instructions and recompiled software. AVX-512 has also had a bumpy life: Intel dropped it from its 12th-generation desktop chips, while AMD added it with Zen 4. The M2 this chapter was written on has NEON only: sysctl hw.optional lists FEAT_AES, FEAT_SHA256, FEAT_I8MM and FEAT_BF16, but no SVE.

ARM's SVE and RISC-V's vector extension take a different approach: the register width isn't fixed by the ISA. Each chip picks a width, and code is written to work with any of them, which the next section shows.

A loop, vectorized

Here is SAXPY (single-precision a·x plus y), a classic kernel:

void saxpy(int n, float a, const float *restrict x, float *restrict y) {
    for (int i = 0; i < n; i++)
        y[i] = a * x[i] + y[i];
}

Compiled by clang 21 with -O2 for ARM64, the heart of the loop is:

.loop:
    ldp   q1, q4, [x10, #-16]      // load 8 floats of y (two 128-bit registers)
    subs  x12, x12, #8
    ldp   q2, q3, [x11, #-16]      // load 8 floats of x
    add   x11, x11, #32
    fmla  v1.4s, v2.4s, v0.s[0]    // y[0..3] += x[0..3] * a
    fmla  v4.4s, v3.4s, v0.s[0]    // y[4..7] += x[4..7] * a
    stp   q1, q4, [x10, #-16]      // store 8 floats
    add   x10, x10, #32
    b.ne  .loop

The compiler did this on its own, with no special code: it's the auto-vectorizer. Each iteration handles eight floats with two fmla (fused multiply-add) instructions of four lanes each. The compiler also emits a plain scalar loop, one fmadd s1, s0, s1, s2 per element, for the last few elements when n isn't a multiple of 8.

The restrict keywords matter. They promise that x and y don't overlap. Without them, clang has to allow for y being partly x, and adds a run-time overlap check before it can use the vector loop.

For x86-64 the same loop depends on which extensions the compiler may use:

Baseline x86-64 only guarantees SSE2, which has no fused multiply-add, so the default build multiplies and adds separately. That's the catch of x86 SIMD: a binary built for the baseline doesn't use the wide units of the chip it runs on, unless the program checks the CPU at run time and picks a different code path, as libraries like glibc's memcpy do.

SVE avoids the problem. Compiled for Fujitsu's A64FX (-mcpu=a64fx), the loop uses scalable z registers and predicates:

    ptrue  p0.s                          // predicate: all lanes active
    rdvl   x12, #4                       // x12 = 4 × the vector length, in bytes
.loop:
    ldr    z2, [x13]
    ldr    z6, [x14]
    fmad   z2.s, p0/m, z1.s, z6.s        // multiply-add, on every active lane
    ...
    add    x13, x13, x12                 // advance by however wide the vectors are

Nowhere does the code say how many floats a register holds. It asks the hardware with rdvl (read vector length), so the same binary runs 4 lanes at a time on a 128-bit SVE chip and 16 on the A64FX.

What SIMD buys on the M2

To measure it, the same file was compiled twice with clang -O2: once normally, once with -fno-vectorize -fno-slp-vectorize, which turns the vectorizer off (the scalar build contains no vector instructions at all). Each kernel then ran on one M2 Ultra performance core, on a small array that fits in the L1 cache (4,096 elements) and on a large one that doesn't (16 million elements, 64 MB per array). Best of three runs, repeated twice with the same results:

In cache, the vector code is five to eight times faster. The integer sum gains more than four lanes would suggest because the vectorizer also splits the running total into several independent accumulators: the scalar loop is one dependent chain at one add per cycle, and the vector loop isn't.

Out of cache, the gain shrinks. The large SAXPY moves 12 bytes per element (read x, read y, write y); at 0.13 ns per element that's about 90 GB/s from a single core, and the arithmetic is no longer the bottleneck. SIMD makes computation cheap; it doesn't make memory faster.

GPUs: SIMD at scale

A GPU takes the SIMD idea as far as it goes. A classic example is NVIDIA's Fermi architecture (2010), used in the GeForce GTX 580: 16 streaming multiprocessors (SMs) of 32 simple "CUDA cores" each, 512 in all, with a shared 768 KB L2 cache.

The programming model is different from SIMD instructions. A GPU program, a kernel, is written for a single thread ("compute y[i] for my i") and launched over millions of threads. The hardware groups threads into warps of 32 (NVIDIA's term; Apple calls them SIMD groups, AMD wavefronts). A warp runs in lockstep: one instruction is fetched and decoded, and all 32 threads execute it, each on its own data. NVIDIA calls this SIMT, single instruction, multiple threads. It's SIMD in which each lane looks like a thread to the programmer.

Two consequences follow:

• Branch divergence. If threads of the same warp take different sides of an if, the warp runs both sides one after the other, each with the other lanes switched off. Code full of data-dependent branches wastes most of the machine.
• Latency hiding by multithreading. A GPU doesn't have big caches or out-of-order execution to hide memory latency. Instead, each SM keeps dozens of warps resident, each with its own registers, and switches to another ready warp every cycle one of them waits for memory. Network processors' packet engines use the same trick; the chapter on multicore looks at it in CPUs.

Since Fermi, the numbers have grown by more than an order of magnitude:

The H100 is a data-center part with less graphics hardware and a lot more: faster double-precision math, and tensor cores, units that multiply small matrices of 8- or 16-bit numbers in one operation, delivering far more than the FP32 figure for machine learning. Around 2010, GPUs used for general computing were still called "GPGPUs" and were a niche. Today most of the world's new supercomputing and nearly all AI training run on GPUs, as the chapter on clusters and supercomputers shows.

A GPU you can measure

The M2 Ultra has its own GPU: system_profiler SPDisplaysDataType reports 60 cores on this machine. Apple's GPU isn't a separate card: it shares the same memory as the CPU cores, so no data has to be copied over a bus first.

Here is SAXPY as a Metal compute kernel, written for one thread:

kernel void saxpy(device const float *x [[buffer(0)]],
                  device float *y       [[buffer(1)]],
                  constant float &a     [[buffer(2)]],
                  uint i [[thread_position_in_grid]]) {
    y[i] = a * x[i] + y[i];
}

Launched over 67 million threads (2²⁶ elements, 256 MB per array), 20 times in a row, from a small Swift program. Metal reported a threadExecutionWidth of 32 for this kernel, its SIMD-group size. Each 20-launch batch took 25.5 ms of GPU time over five runs: 0.019 ns per element, or 630 GB/s of memory traffic. Apple quotes 800 GB/s as the chip's peak memory bandwidth. One CPU core, running the NEON version above, reached about 90 GB/s. For a job that only streams through memory, the GPU wins because it can keep far more memory requests in flight.

Coprocessors

More generally, a main CPU can hand work to specialized helpers called coprocessors. Three kinds have evolved differently.

Network processors. A programmable network processor puts many small RISC engines on a card, handling packets at wire speed while the main CPU does other work. Around 2010, 40-gigabit Ethernet was the next step; 400 Gb/s ports are now common in data centers and 800 Gb/s ones are shipping. Every network card now offloads checksums and splitting large sends into packets, and data centers use SmartNICs or DPUs (data processing units) that run networking, storage and security functions on their own ARM cores, as network processors foreshadowed.

Cryptography. Crypto coprocessors used to come on plug-in cards. Today the most common crypto hardware is inside the CPU itself. Intel added AES-NI instructions in 2010 (aesenc performs one round of AES), and ARMv8 has aese and aesmc; SHA hashing has instructions too. The difference is dramatic. OpenSSL 3.6 encrypting 16 KB blocks with AES-128-CTR on one M2 core ran at 13.1 GB/s using the AES instructions. Forcing OpenSSL's portable C code with OPENSSL_armcap=0, it ran at 0.32 GB/s: about 40 times slower. Specialized hardware for a common, fixed task is the best case for a coprocessor, even when the "coprocessor" is a few instructions.

Graphics and machine learning. The GPU started as the graphics coprocessor and became a general parallel processor. The newest specialized unit is the neural processing unit (NPU). The M2 Ultra's 32-core Neural Engine is rated by Apple at 31.6 trillion operations per second, on low-precision numbers for neural-network inference. Unlike the GPU, it's not programmable directly: applications reach it through Apple's Core ML framework, which decides what runs where. Intel, AMD and Qualcomm now put similar NPUs in laptop chips.

The pattern repeats across these examples. A task becomes common enough, and regular enough, to justify dedicated hardware. It starts as a separate chip or card, and ends up either as instructions in the CPU (floating point, SIMD, AES) or as a unit beside the CPU cores on the same die (GPU, NPU).

Takeaways

• SIMD instructions do the same operation on every lane of a wide register: 128 bits for SSE and NEON, 256 for AVX2, 512 for AVX-512, and a chip-chosen width for SVE and RISC-V's vectors.
• Compilers auto-vectorize simple loops: clang turned SAXPY into fmla v1.4s on ARM64 and mulps/addps, vfmadd…ymm or …zmm on x86-64 depending on the target flags.
• On one M2 core, vectorization made in-cache kernels 5–9× faster, but only 3.5–5× out of cache, where memory bandwidth (~90 GB/s per core here) is the limit.
• GPUs run thousands of threads in warps of 32 that execute in lockstep (SIMT), hide memory latency by switching warps, and suffer from branch divergence. The M2 Ultra's GPU streamed SAXPY at about 630 GB/s.
• Coprocessors for networks, cryptography, graphics and neural networks tend to migrate onto the CPU die. AES instructions made encryption about 40× faster than portable code on the M2.