Why a GPU needs so many threads

01 / The jobGive one thread one number to produce

We want C = A × B. Matrix A has four rows and four columns. B has four rows and 32 columns. Their product C therefore has four rows and 32 columns: 128 output numbers. The small dimensions make the bookkeeping visible; this is not a workload large enough to demonstrate real GPU performance.

Take output C[0,2], using zero-based indices. Pair row 0 of A with column 2 of B, multiply corresponding entries, and add. The four products are 1×3, 2×0, 3×1, and 4×1. A thread responsible for this output can keep one private running sum:

No other output is needed to calculate this one. That is the independence we can spend. Assign one thread to each output: 128 threads, each making four multiply-add updates. A fused multiply-add, or FMA, evaluates a × b + accumulator with one final rounding in floating-point arithmetic. Our small integer-valued inputs give exact results in the floating-point format used by the sketch.

// CUDA kernel sketch: A[4][4], B[4][32], C[4][32], row-major.
// Device pointers and a launch of exactly <<<4, 32>>> are assumed.
__global__ void tiny_product(const float* A, const float* B, float* C) {
    int row = blockIdx.x;   // four blocks: one per output row
    int col = threadIdx.x;  // 32 threads in each block
    float acc = 0.0f;
    for (int k = 0; k < 4; ++k) {
        float a = A[row * 4 + k];
        float b = B[k * 32 + col];
        acc = fmaf(a, b, acc);
    }
    C[row * 32 + col] = acc;
}

This is an intentionally simple mapping in NVIDIA’s CUDA programming platform, not an optimized matrix multiplication. No GPU execution or timing is claimed here. The compiler may unroll or rearrange the actual instructions; the scheduling experiment below has its own explicit rules.

02 / The unit of issueGive 32 threads the same instruction

A CUDA thread block is a group of threads assigned together to a streaming multiprocessor, or SM. An SM contains schedulers, registers, and execution resources; it is not just one arithmetic unit. Within a block, NVIDIA groups consecutive thread IDs into warps of 32 threads. Our 32-thread blocks give us exactly four warps, W0 through W3. Each produces one complete output row. [1]

At k = 0, W0’s threads all execute the same multiply-add operation. They have different values of col, so they read different B entries and update different private sums. Here are only four of the warp’s 32 lanes:

This is single-instruction, multiple-thread (SIMT) execution: one issued instruction operates on the participating threads, whose registers remain distinct. It does not give every thread a private instruction-issue engine. Modern hardware also supports independent thread scheduling; do not turn this straight-line example into a general assumption that lanes are implicitly synchronized. [2]

So there are two kinds of multiplicity already. One instruction does work across lanes. And several warps can be available to take turns issuing instructions. The second kind becomes important when the first kind runs out of operands.

03 / The holeFind the hole between load and use

Our thread’s arithmetic is cheap to describe: load A, load B, multiply-add. But issuing a load starts a request; the result need not be available for the next instruction. A dependent multiply-add must wait for both operands. An independent load can still issue while the first load is outstanding.

To see the consequence, construct a small machine with one warp-instruction issue slot per cycle. A load issued at the start of cycle t becomes usable at the start of t + 4. A multiply-add occupies one issue slot, and its result is usable next cycle. These are chosen teaching parameters, not latencies measured on a particular GPU.

Allow only W0 to be resident. At cycle 0 it issues the A load; at cycle 1 it issues the independent B load. Its next instruction is the multiply-add. A will arrive at cycle 4, but B will not arrive until cycle 5:

Three of those six issue slots did useful dispatch work. Three had nothing ready to issue. The other output rows are independent, but if their warps are not resident, they cannot supply an instruction to this scheduler. Work that exists somewhere is not automatically work that can run here now.

Before adding hardware: could a second warp put an instruction into one of those blank slots?

04 / The loomLet another warp use that hole

Keep the arithmetic and the load latency unchanged. Now allow all four warps to remain resident, each with its own register state. Use a round-robin scheduler that skips a warp when its next instruction is not ready. The first twelve cycles become:

W0 issues B at cycle 4, so that operand arrives at 8. While it travels, W1, W2, and W3 issue their B loads. When the scheduler returns to W0, its multiply-add is ready. The load still took four cycles. Other warps filled the time.

At four-cycle load latency, the complete job takes 96 simulated cycles with one resident warp, 64 with two, and 48 with four. In every case it issues the same 32 warp-load instructions and 16 warp multiply-add instructions. The four-resident run removes the 48 idle slots of the one-resident run. It does not double the issue width.

There is a useful sting here. W0 alone finishes its row after 24 cycles. Sharing the issue slot with three other warps makes W0 finish after 45 cycles. Yet the whole job finishes twice as quickly: 48 cycles instead of 96. Aggregate throughput improved even though this particular row waited longer for some of its turns. “Latency hiding” describes what happens to the machine’s otherwise unused opportunities; it does not promise that every individual task finishes sooner.

05 / The storage behind the trickWhere do the waiting threads live?

The scheduler needs more than a queue of names. If W0 has accumulated 6 for C[0,2], that 6 must still exist when W1 runs. Resident warps keep their execution context on-chip, including register state. Choosing another ready warp therefore does not require an operating system to save one software thread’s registers to memory and restore another on every issue. [2]

In our example, when W0 reaches its second multiply-add, its four displayed accumulators become [1,4,3,6]. W1 has a different row of A and therefore different accumulators. These are simultaneously retained values, even though this model issues only one warp’s instruction in a cycle. Storage capacity and execution bandwidth are separate resources.

This distinction makes the hardware’s apparently extravagant thread count intelligible. The machine needs enough stored independent contexts to keep finding ready instructions. A software thread, a resident hardware context, and an arithmetic lane are not interchangeable things to count.

On actual hardware, occupancy is the fraction of the SM’s supported resident warp capacity that is occupied. Registers, shared memory, and block limits can constrain it. More occupancy supplies more possible work, but performance depends on what that work is ready to do. A kernel with enough independent instructions may already hide the relevant waits at lower occupancy; forcing higher occupancy can also trade away registers and cause spills. [3]

The experiment’s “issue slots used” is a different quantity. Its 100% means every modeled issue opportunity contained an instruction, including loads. It does not mean that every arithmetic unit on a real GPU would be fully utilized. Likewise, four resident warps here is a deliberately tiny capacity setting, not an SM specification.

06 / Break one assumptionWhat if the wait outgrows the spare work?

Keep four resident warps, but move load latency to twelve cycles. W0’s B load still issues at cycle 4; now its result is unavailable until cycle 16. By the end of cycle 7, all four warps have issued their B loads. The scheduler has nothing else eligible from cycles 8 through 15. The contexts are occupied. The issue slot is empty.

The full run now takes 80 simulated cycles, including 32 idle slots. The same four warps that completely hid a four-cycle load cannot completely hide a twelve-cycle load under this instruction sequence. The missing resource is ready independent work during those intervals.

There is another limit our toy deliberately removes. Memory cannot accept unlimited requests per second. On a real device, extra warps may create enough outstanding work to keep the memory system busy, after which more requests mainly compete for that capacity. The distinction should feel familiar: recovering unused time and increasing the service capacity are different achievements. Here, however, the waiting contexts and the scheduler are inside the processor.

Return to the original matrix product. Each of the 32 threads in a row asks for the same A value at a given k. The four row warps ask for the same B row. Hardware caches and combined accesses may already service some of this repetition efficiently, so source-level load counts are not off-chip transaction counts. Still, the repeated operands create a new design opportunity: arrange the computation so a value, once fetched, has more useful work to do nearby.

For now, take one testable distinction away: when a GPU stalls, ask whether work is absent, nonresident, resident but unready, or ready but competing for a limited execution resource. Those are different states. Adding threads fixes only some of them.