GPU architecture: SMs, warps, the memory hierarchy, and Tensor Cores

When an H100 SXM runs a large matrix multiply, its dense BF16 Tensor Core peak is 989.4 TFLOPS; yet when the same card generates tokens one at a time for a large model, it spends most of its time waiting for data. Both are true at once. A GPU is built for total throughput, not for making one thread fast: it keeps many warps resident and switches to another whenever one is waiting on memory, and whether the Tensor Cores stay busy comes down to how fast data can be moved into the SMs.

Draft, pending review. The site owner has not reviewed this note yet; facts, emphasis and wording may change. Check the primary sources linked in the text before relying on it.

Hardware numbers are mostly for the H100 SXM5; final specifications come from the H100 architecture whitepaper V1.04 (which states that it includes final clocks and TFLOPS), and the rest from NVIDIA product pages, architecture blogs and CUDA documentation; cycle counts and matrix entries in the worked examples are teaching assumptions and are marked as such. How to write kernels is in CUDA, the roofline is in Performance, and the cross-generation comparison is in NVIDIA GPU generations.

Two design goals: wait less, or do more at once

The CUDA Programming Guide puts it plainly: a CPU is built to finish a serial sequence of operations as fast as possible and runs a few tens of threads at once; a GPU is built to run thousands of threads in parallel, trading lower single-thread performance for much higher total throughput. So a CPU spends more of its transistors on caching and control logic, and a GPU spends more on compute units. How a CPU shortens the wait of a single thread is covered in CPU architecture.

This is a design tendency, not a hard line: CPUs have SIMD vector units too, and GPU threads can branch; the costs are just different.

The SM: a GPU is a hundred-odd processors

In CUDA's hardware model, a GPU is a set of Streaming Multiprocessors (SMs) grouped into GPCs; each SM has its own register file, an on-chip memory that can be split between L1 and shared memory, and compute units (programming model). The H100 SXM5 has 132 SMs (Hopper architecture blog), and each SM contains:

The "four processing blocks, one warp scheduler each" layout is described in the Volta whitepaper; the Ampere whitepaper says the A100 has one Tensor Core per partition, and the H100 also has 4 per SM. Residency limits come from the compute capability tables; the H100 is compute capability 9.0.

On the software side, a kernel launch creates a grid made of thread blocks. All threads of a block run on the same SM, which is why they can exchange data through shared memory and synchronize; there is no ordering guarantee between blocks. From compute capability 9.0 there are also thread block clusters: the blocks of a cluster are scheduled onto one GPC and can read and write each other's shared memory. The thread hierarchy in detail is in CUDA.

Warps and SIMT: 32 threads share one instruction

An SM does not schedule threads one by one but in groups of 32, called warps. The CUDA guide calls this SIMT: the threads of a warp execute the same code, but each can take its own branch. When they disagree, the threads not on the current path are masked off and get their turn when the other path runs. This is warp divergence.

Work out the cost once. A warp hits an if: when the condition is true it runs path A (10 instructions), otherwise path B (20 instructions). The instruction counts are a teaching assumption.

A branch is not expensive in itself; splitting within one warp is. Different warps taking different branches pay no such cost.

Two common misreadings. First, SIMT does not mean "the 32 threads are always in lockstep". Since compute capability 7.0 (Volta) there is independent thread scheduling: each thread has its own program counter, and threads in a warp can diverge and reconverge at finer granularity; old code that relies on implicit synchronization must use __syncwarp() (advanced kernel programming). Second, the difference between SIMT and SIMD is the interface: SIMD exposes the vector width to software, while SIMT lets you write code for a single thread and the hardware groups threads into warps.

Latency hiding: keep many warps resident

The CUDA guide explains that each warp's execution context (program counter, registers) stays on chip for the warp's whole lifetime, so switching between warps costs nothing; at every issue cycle a warp scheduler picks one of the ready warps (hardware multithreading). If one warp is waiting on memory, another one goes.

A simple model tells you how many warps you need. Suppose each warp repeats two things: issue $W$ independent instructions (taking $W$ issue cycles), then wait $L$ cycles for a memory load. With $N$ warps resident on one scheduler, the fraction of issue slots used is about:

$$ U \approx \min\left(1,\ \frac{N\,W}{W+L}\right). $$

Take $L=300$ and $W=30$ (teaching assumptions, not measured H100 latencies): a warp uses only 30 issue cycles out of every 330, so you need $N\ge 330/30=11$ warps to fill the slots.

This figure needs JavaScript.

How many warps fit is often set by registers. An H100 SM has 65,536 32-bit registers and at most 64 resident warps; split evenly over 4 schedulers, that is at most 16 each.

This ignores register allocation granularity and shared memory limits; the full occupancy calculation is in CUDA.

The last row deserves a word. High-performance matrix multiply kernels often use many registers per thread and run at modest occupancy, hiding the wait with instruction-level parallelism and asynchronous data movement; the CUDA guide also describes async copies as a way to overlap movement and compute beyond what warp switching does (asynchronous execution). Latency can be hidden; the bandwidth ceiling cannot. That is the next section.

The memory hierarchy: closer is smaller and faster

Capacities and official bandwidths for each level of the H100 SXM5 (whitepaper V1.04, Hopper blog, H100 product page, compute capability tables):

Note a discrepancy between sources: the 2022 Hopper architecture blog gives HBM3 bandwidth as 3,000 GB/s, a preliminary figure; whitepaper V1.04 gives 3,352 GB/s, consistent with the product page's 3.35 TB/s. In the V1.04 spec table the memory bandwidth row label still carries an old "Not Finalized for H100" note, but the cover says the final memory clocks are included, so this note uses 3,352 GB/s.

NVIDIA does not publish absolute on-chip bandwidths for the H100, but the public A100 material is enough for an order-of-magnitude check. The A100 whitepaper gives an L2 read bandwidth of 5,120 bytes per clock and a boost clock of 1,410 MHz; the CUDA guide says shared memory has 32 banks of 32 bits per cycle each, i.e. 128 bytes per SM per cycle. The A100 has 108 SMs:

$$ \begin{aligned} \text{shared memory, all SMs} &\approx 108\times128\ \text{B}\times1.41\ \text{GHz}\approx19.5\ \text{TB/s},\\ \text{L2} &\approx 5120\ \text{B}\times1.41\ \text{GHz}\approx7.2\ \text{TB/s},\\ \text{HBM2e (A100 80GB SXM)} &= 2.039\ \text{TB/s}. \end{aligned} $$

The first two lines are theoretical peaks converted from specifications, not measurements; bank conflicts and cache misses lower the real numbers. The HBM figure comes from the A100 datasheet. The ratios already make the point: the closer to the compute units, the higher the bandwidth and the smaller the capacity. Hence the basic recipe of high-performance kernels is tiling: read a block from HBM into shared memory or registers, reuse it many times there, then write back. How to write it is in the tiling section of CUDA; FlashAttention is the same idea applied to attention.

Blackwell adds one more level. Fifth-generation Tensor Cores have a dedicated Tensor Memory, which PTX ISA §9.7.18.1 describes on sm_100a as 128 rows × 512 columns of 32-bit cells per CTA, i.e. 256 KB; the matrix-multiply accumulator $D$ lives there instead of in registers.

Tensor Cores: one instruction, one small matrix multiply-accumulate

A CUDA core's FMA instruction does one multiply-add. A Tensor Core does a small block at once: $D=A\times B+C$. In the Volta whitepaper, each Tensor Core works on 4×4 matrices and performs 64 FMAs per clock.

A 4×4 multiply-accumulate by hand

Take these small matrices (teaching numbers), with $C$ the all-ones matrix:

$$ A=\begin{bmatrix}1&2&0&1\\0&1&1&0\\2&0&1&0\\1&1&1&1\end{bmatrix},\qquad B=\begin{bmatrix}1&0&2&0\\0&1&0&1\\1&1&0&0\\0&2&1&1\end{bmatrix}. $$

Row 1, column 2 of $D$ is the dot product of row 1 of $A$ with column 2 of $B$, plus the matching entry of $C$: $d_{12}=1\cdot0+2\cdot1+0\cdot1+1\cdot2+1=5$. Doing all 16 entries gives:

$$ D=AB+C=\begin{bmatrix}2&5&4&4\\2&3&1&2\\4&2&5&1\\3&5&4&3\end{bmatrix}. $$

Each output entry takes 4 multiply-adds, i.e. 4 FMAs; 16 entries take $16\times4=64$ FMAs, exactly one Volta Tensor Core's work per clock. Counting an FMA as 2 FLOPs, this tile is 128 FLOPs.

On precision, Volta's $A$ and $B$ are FP16, the products are kept at full precision and accumulated in FP32; $C$ and $D$ may be FP16 or FP32. "Low-precision inputs, higher-precision accumulation" carries through to the later FP8 and FP4.

From one tile to the whole card's peak

Peak = number of Tensor Cores × FMAs per Tensor Core per cycle × 2 × clock. Check against NVIDIA's published numbers:

The V100 row comes from the Volta whitepaper and the A100 row from the Ampere architecture blog. The H100 row's clock and official figure come from whitepaper V1.04: a boost clock of 1,830 MHz for Tensor Core operations, 989.4 TFLOPS dense FP16/BF16 (FP32 accumulate), and 1,978.9 TFLOPS sparse. The 512 FMAs per Tensor Core per cycle is this note's inference: the whitepaper says the per-SM MMA rate is 2× the A100's at the same clock; plugging in $528\times512\times2\times1.83\ \text{GHz}$ gives exactly 989.4 TFLOPS.

The "sparse" figures in datasheets come from 2:4 structured sparsity, introduced with Ampere: when at most 2 of every 4 values are non-zero, Tensor Core throughput doubles. Ordinary dense models do not get that factor, so convert to dense before comparing. Every Tensor row on the H100 product page is marked "with sparsity".

Bigger instructions, issued by fewer threads

Programs never write 4×4 directly. At the PTX level you can see how the interface changed with each generation (PTX ISA §9.7.16–9.7.18):

One mma.sync equals 32 of the 4×4×4 tiles above (4,096 / 128). The trend: each instruction does more work, runs more asynchronously, and is issued by fewer threads. With fewer threads and registers spent feeding the Tensor Cores, the remaining warps can move data or compute softmax; FlashAttention-3 on Hopper uses warp specialization to overlap data movement with compute and interleaves block-wise matmul with softmax (see also attention kernels). The actual code is left to CUDA and libraries such as CUTLASS.

Data types: what each Tensor Core generation added

Sources: FP16 from the Volta whitepaper; TF32 and BF16 from the Ampere blog; the two FP8 formats from the Hopper blog; which input types each compute capability supports from the compute capability tables; FP8 mma needing sm_89 and the FP6/FP4 format names from the PTX ISA. The FP16 versus BF16 trade-off is in Precision and memory and is not repeated here.

The narrower the format, the more it depends on scaling. FP4 E2M1 can only represent 0, 0.5, 1, 1.5, 2, 3, 4, 6 and their negatives. NVFP4 gives every 16 values a shared FP8 (E4M3) scale factor and the whole tensor a further FP32 scale, for 4.5 bits per value on average; the MXFP4 format it is compared with shares one E8M0 scale, restricted to powers of two, across every 32 values.

Work a simplified version by hand (only 4 values; ignore the tensor-level scale and the rounding of the scale itself). The block is $[0.3,-1.2,2.4,3.0]$; its largest magnitude is 3.0, so the scale is $3.0/6=0.5$. Dividing by 0.5 gives $[0.6,-2.4,4.8,6.0]$; rounding to the nearest representable values gives $[0.5,-2,4,6]$; multiplying back gives $[0.25,-1.0,2.0,3.0]$. If a 30 lands in the same block, the scale becomes 5, and 0.3 and −1.2 both round to 0 after dividing by 5: one outlier wipes out the small values of the whole block. The smaller the block, the fewer values one outlier drags down, which is the intuition behind NVFP4 using 16 rather than 32.

NVLink and NVSwitch: between cards

The H100 SXM5 has 18 fourth-generation NVLink links at 25 GB/s per direction each, 900 GB/s in total; its PCIe Gen5 link is 128 GB/s (Hopper blog, H100 product page). NVLink is a point-to-point link; NVSwitch is a switch chip that lets every GPU in a group talk to every other at full NVLink bandwidth. A group of GPUs that can reach each other this way is called an NVLink domain.

The last two rows come from the GB200 NVL72 page and the HGX page. Both numbers add the two directions: HGX B200 has 8 NICs at 400 Gb/s (the DGX B200 page lists 8 single-port ConnectX-7), 400 GB/s per direction in total, 0.8 TB/s counting both. Compared on the same basis, NVLink bandwidth in the same Blackwell generation is about 18× the off-machine network (1.8 TB/s versus about 0.1 TB/s per GPU). That is why tensor parallelism, the most communication-heavy split, usually stays inside an NVLink domain, while only lighter-communication splits cross machines; see Tensor parallelism and Distributed training. NVL72 stretches the NVLink domain from one 8-GPU machine to a whole rack.

What this note leaves out

• How to write kernels and use Nsight to read occupancy and bandwidth: see CUDA.
• One roofline number only: the H100 SXM has 989.4 dense BF16 TFLOPS and 3,352 GB/s of HBM, a ratio of about 295 FLOPs per byte. A kernel that cannot do that many operations per byte read from memory hits the bandwidth limit first. How to reason about this is in Performance.
• Per-generation specs and the worked example of how many GPUs a 70B model needs: see NVIDIA GPU generations. Process nodes, HBM and packaging: see How chips are made.