The Machine: threads, warps, and SMs

Almost every performance mistake in a GPU kernel comes from one of two wrong beliefs: that a thread is like a CPU thread, or that "more parallelism" is automatically better. Both are wrong in specific, fixable ways. This part replaces them.

We are going to build the model bottom-up, and every claim will be tied to a number you can look up. By the end you should be able to answer, for any kernel you write: how many warps are resident, what are they waiting on, and what is the hardware doing while they wait?

A single CUDA thread is slow. On an H100 the clock is around 1.98 GHz — slower than a laptop CPU — and there is no deep out-of-order machinery, no aggressive branch prediction, no big private cache to hide behind. A dependent chain of fused multiply-adds retires at roughly one every 4 cycles, same as it did a decade ago.

What the GPU has instead is thousands of independent instruction streams that cost nothing to switch between. When one stalls on memory, the hardware issues from another in the very next cycle. The CUDA Programming Guide is blunt about why this is cheap:

A CPU hides latency with caches and speculation, and it does so per instruction stream. A GPU hides latency with other instruction streams. That single sentence explains most of the design: the enormous register file, the small caches, the wide memory bus, and the fact that occupancy is a thing you have to think about at all.

The consequence for you: your job as a kernel author is almost never to make one thread faster. It is to make sure that at every cycle, some warp on every SM has work it can issue.

Keep the actual latencies in view, because every design decision later in the series is a negotiation with this table. These are measured cycle counts from published microbenchmarks, not marketing numbers:

Source: Luo et al., Benchmarking and Dissecting the NVIDIA Hopper GPU Architecture, Table IV. The Hopper column is an H800 PCIe measurement; the H100/H200 time-domain extrapolation used below is labelled separately rather than presented as a measurement of those cards.

An HBM access costs roughly 120× a register access. If a warp issues a load and then immediately needs the result, that warp is dead for ~479 cycles. Whether the SM is dead too depends entirely on how many other warps are resident — which is what the rest of this part is about.

CUDA gives you five levels of grouping. They are not equally real: some are hardware, some are pure software convention, and confusing the two is a classic source of wrong intuitions.

1.2.1 The warp is the only unit that matters for correctness intuition

Threads do not execute independently. A warp fetches one instruction and applies it across 32 lanes. That is why the warp partitioning rule is worth memorising exactly:

Warp 0 always contains threads 0–31 of the block, in order. This means a 2D block of (32, 8) gives you warps that are rows — consecutive x, fixed y. A block of (8, 32) gives you warps that straddle 4 rows of y each. That choice silently determines whether your global loads coalesce, which is the single biggest factor in Part 2.

1.2.2 Block size, and the warps you waste

Because warps are cut at 32, a block size that is not a multiple of 32 permanently wastes lanes. A block of 100 threads occupies 4 warps — 128 lanes — and 28 of them are inactive for the entire lifetime of the block. The hardware will not reclaim them.

An H100 has 132 Streaming Multiprocessors. Every one of them, on every architecture from Maxwell to Blackwell, is built the same way at the top level: four sub-partitions, each with its own warp scheduler, its own slice of the register file, and its own execution units — including memory load/store units. The four sub-partitions converge on shared SM-wide resources such as the instruction cache and the L1/shared-memory array.

That last clause matters: a warp is pinned to one sub-partition. Its registers come from that sub-partition's 64 KB slice, and its instructions issue from that sub-partition's scheduler at a maximum of one per cycle. On Hopper, four schedulers × 1 instruction gives a peak of 4 warp instructions per clock per SM.

1.3.1 Why the FP32 lane count is a 2× story

An A100 sub-partition has 16 FP32 lanes; a warp is 32 threads. So a single FP32 multiply occupies the datapath for two cycles. Hopper doubled it to 32 lanes, so the same instruction retires in one. This is directly readable from the Programming Guide's per-clock throughput table:

Two rows deserve a bookmark. INT32 throughput is half of FP32 on Hopper. Address arithmetic, loop counters and predicate math all run on that path, and in a heavily unrolled kernel they can become the bottleneck — which is exactly why Hopper added TMA to move address generation into fixed-function hardware (Part 6).

And the special-function unit does 16 results per clock per SM. On an H100 that works out to about 3.9 Tops/s of exp against 989 TFLOP/s of tensor-core matmul — FlashAttention-3 puts the gap at 256×, computed at a 1.83 GHz sustained clock rather than the 1.98 GHz boost. Hold that number; it is the entire reason FlashAttention-3 needs a ping-pong schedule (Part 7).

When threads in a warp take different paths through an if, the warp does not split into two warps. It executes both sides, one after the other, with the non-participating lanes masked off.

The cost model follows mechanically: a two-way branch where both sides are taken costs the sum of both sides, not the max. A 32-way switch where every lane picks a different case costs 32×.

1.4.1 Divergence that costs nothing

Divergence is only expensive when it happens inside a warp. This is fine:

// Every warp takes exactly one side: warp_id is uniform within a warp.
int warp_id = threadIdx.x / 32;
if (warp_id < 4) { produce(); }   // warps 0-3
else            { consume(); }   // warps 4-7

And this is not:

// Lanes within one warp disagree -> both arms run, masked.
if (x[threadIdx.x] > 0.0f) y[threadIdx.x] = fast_path();
else                      y[threadIdx.x] = slow_path();

A useful rule: ask whether the predicate is constant across each group of 32 consecutive threads. If yes, the branch is free. If no, budget for both arms.

1.4.2 Independent thread scheduling, and why __syncwarp() exists

Before Volta, all 32 threads shared one program counter and an active mask. Volta gave each thread its own PC and call stack, which fixed a class of deadlocks (a lane holding a lock while its warp-mates spin) but broke a very common optimization: warp-synchronous programming, where people omitted synchronization inside a warp because "warps are lockstep anyway."

They are no longer guaranteed to be. The convergence optimizer may group lanes back together, but you cannot rely on it. Any intra-warp data exchange must name its participants explicitly:

// WRONG on Volta and later: assumes implicit lockstep.
sdata[tid] += sdata[tid + 16];
sdata[tid] += sdata[tid +  8];

// RIGHT: explicit mask, explicit sync.
float v = sdata[tid];
for (int off = 16; off > 0; off >>= 1)
    v += __shfl_down_sync(0xffffffff, v, off);   // mask names all 32 lanes

Now we can put the pieces together. Consider the simplest possible memory-bound kernel: load a value, do a little arithmetic, store it. One warp's timeline is mostly dead air. The question is how many warps you need before the dead air is completely covered.

The governing relation is Little's Law, which Vasily Volkov put at the centre of GPU performance analysis in his 2010 GTC talk:

For memory: to saturate HBM you must keep bandwidth × latency bytes in flight at all times. Plug in real numbers:

Luo et al. measured 466 cycles on A100 (331 ns at 1.41 GHz) and 479 cycles on H800 PCIe (273 ns at 1.755 GHz). The starred H100/H200 entries reuse that H800 time as an explicit same-GH100-family estimate; they are neither direct measurements nor conversions at those cards' own clocks. Bytes in flight = bandwidth × latency, with decimal SI units: 914,550 B = 915 kB = 893 KiB. B200 latency is unpublished; the original measurement behind the directional row finds it higher than H100's.

Read the last column carefully. On an H100 at full occupancy, one 4-byte load per thread is roughly enough to saturate HBM. On a B200 it is not — you need something closer to a 16-byte vector load per thread in flight. Each generation raises the amount of memory-level parallelism a kernel must expose. That is the structural reason modern kernels use float4 loads, deep software pipelines, and TMA.

Occupancy is resident warps per SM divided by the architectural maximum (64 on Hopper and datacenter Blackwell; 48 on consumer Blackwell and Ampere GA10x). It is an upper bound on how much latency you can hide, and it is decided at launch by three independent constraints:

1. Registers. Each SM has 65,536 32-bit registers. A thread using R registers means a warp uses 32R, so the SM fits ⌊65536 / (32R)⌋ warps. At 32 registers/thread you get all 64 warps; at 64 you get 32; at the 255 maximum you get 8.
2. Shared memory. An SM has up to 228 KB on Hopper, but CUDA reserves 1 KB per resident block. Two blocks asking for 113 KB each consume 228 KB including those reservations; two asking for 114 KB would need 230 KB and do not fit. A block asking for the 227 KB user maximum fits once. A high-performance GEMM or attention kernel often chooses to sit at 1–2 blocks per SM.
3. Block granularity. At most 32 blocks per SM — the same ceiling on datacenter and consumer Blackwell, per NVIDIA's Blackwell Tuning Guide — and at most 2048 threads (1536 on consumer parts). So 64-thread blocks cap you at 32 × 2 = 64 warps only in the best case, and any register or SMEM pressure drops you below that immediately.

The binding constraint is whichever of the three is smallest. Try it:

Two practical notes. First, registers are allocated in granules, so the real allocation rounds up; use cudaOccupancyMaxActiveBlocksPerMultiprocessor rather than doing it by hand when it matters. Second, CUDA reserves 1 KB of shared memory per block, which is why the per-block maximum is 227 KB while the per-SM maximum is 228 KB. A block that needs more than 48 KB must put the excess in dynamic shared memory and explicitly opt in:

// Anything above 48 KB of shared memory must be dynamic AND opted into.
cudaFuncSetAttribute(my_kernel,
    cudaFuncAttributeMaxDynamicSharedMemorySize, 200 * 1024);
my_kernel<<<grid, block, 200 * 1024>>>(...);

// Ask the runtime what occupancy you actually got.
int blocks;
cudaOccupancyMaxActiveBlocksPerMultiprocessor(&blocks, my_kernel, 256, 200*1024);

1.6.1 Why the calculator disagrees with your profiler

The figure above computes theoretical occupancy: the share of warp slots that could be filled if every selected block stayed resident for the whole kernel. Nsight Compute reports a second number next to it — achieved occupancy, actual resident warps averaged over time — and on perfectly healthy kernels the two differ by 20–30%. That gap has three ordinary causes:

• The compiler knows things the calculator does not. Registers are allocated per warp in 256-register granules and rounded up, and nvcc routinely reserves more than the source implies, to keep spills off the critical path. __launch_bounds__ trades that headroom back for residency, and cudaOccupancyMaxActiveBlocksPerMultiprocessor reads the post-compilation truth rather than your guess.
• Blocks do not retire in lockstep. Occupancy touches its ceiling only while every scheduled block is alive. Blocks finish at different times and replacements trickle in, so averaged over the whole kernel the resident count sits below the peak — worst near the tail of the grid.
• Resident does not mean issuing. A block parked at __syncthreads() keeps all of its warps resident and none of them useful. Barrier-heavy kernels show high theoretical occupancy and low issue rates, and no occupancy number will tell you which is happening.

So treat the calculator as a bound on the possible, not a prediction. Whether those resident warps have anything useful to issue is a different question — the subject of the rest of this part.

Here is where a lot of tuning advice goes wrong. Little's Law says you need parallelism, and warps are only one source of it. The other is instruction-level parallelism: independent work within a single thread. A thread that has eight independent loads in flight contributes as much memory-level parallelism as eight threads with one each — while using a fraction of the warp slots.

Volkov's measurements are still the cleanest demonstration:

Vasily Volkov, Better Performance at Lower Occupancy, GTC 2010. Measured on a GTX 480; the principle has aged better than the hardware.

The honest framing: occupancy is a budget for hiding latency you cannot otherwise avoid. If your kernel has enough independent work per thread, you need less of it. If it is a pointer-chasing, low-ILP kernel, you need a lot. Measure the stall reasons, not the occupancy number.

The naive launch shape — one thread per element — is fine for teaching and rarely optimal in practice. The idiom you will see in real code is the grid-stride loop:

__global__ void saxpy(int n, float a, const float* x, float* y) {
    int stride = blockDim.x * gridDim.x;
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n; i += stride)
        y[i] = a * x[i] + y[i];
}
// Launch with a grid sized to the DEVICE, not to the data:
int blocks; cudaOccupancyMaxActiveBlocksPerMultiprocessor(&blocks, saxpy, 256, 0);
saxpy<<<blocks * numSMs, 256>>>(n, a, x, y);

import triton, triton.language as tl

@triton.jit
def saxpy(x_ptr, y_ptr, a, n, BLOCK: tl.constexpr):
    pid  = tl.program_id(0)
    offs = pid * BLOCK + tl.arange(0, BLOCK)
    mask = offs < n
    x = tl.load(x_ptr + offs, mask=mask)
    y = tl.load(y_ptr + offs, mask=mask)
    tl.store(y_ptr + offs, a * x + y, mask=mask)

grid = lambda m: (triton.cdiv(n, m['BLOCK']),)
saxpy[grid](x, y, a, n, BLOCK=1024)

# Ground truth to check against.
y = a * x + y

Grid-stride loops decouple the grid size from the problem size, which lets the launch be sized to the machine. That matters because of wave quantization: blocks execute in waves of roughly SMs × blocks_per_SM, and a grid of 133 blocks on 132 SMs takes two waves — the second one 99% empty.

Everything above turns into three questions you will ask of every kernel in this series:

1. How many warps are resident, and what are they waiting for? Waiting on HBM is a bandwidth problem (Parts 2–3). Waiting on SMEM is a bank-conflict or layout problem (Parts 2, 6). Waiting on the SFU is a softmax problem (Parts 4, 7). Waiting on nothing is a launch problem.
2. Is the per-thread work independent? If each thread does one dependent chain, you need occupancy. If each thread owns an 8×8 accumulator tile, you do not — and you should spend the registers instead.
3. Does the grid fill the machine, in whole waves? This is the difference between 76% and 99% utilization on a decode-shaped matmul, and no amount of inner-loop tuning fixes it.

Next: where the data actually lives, why a 4-byte load can cost you 8× the bandwidth you meant to use, and the bank-conflict pattern that every tiled kernel has to solve.