Latency Hiding and the Arithmetic of Warp Counts

The reason a GPU holds tens of warps per SM rather than a handful of threads is that memory latency is enormous relative to the work a warp can do between accesses. A global load takes on the order of hundreds of cycles to return; a warp can issue perhaps a few dozen instructions in that window. Hiding the difference requires other warps ready to issue while the first waits, and the number required is computable from two quantities: how long the stall lasts and how many instructions the warp can issue between stalls.

The computation is a Little’s law argument, and like all such arguments it is exact in the steady state and misleading at the edges. Knowing where it misleads is most of the practical value.

The basic ratio

Take a scheduler that issues one instruction per cycle when a ready warp exists. A warp whose instruction stream contains a stall of L cycles every k issued instructions can occupy the scheduler for k/(k+L) of the time. With W warps sharing the scheduler and assuming their stalls are independent and uniformly distributed, the fraction of cycles the scheduler finds at least one ready warp rises quickly with W: the scheduler is idle only when all W warps are simultaneously stalled, which for independent warps happens with probability on the order of (L/(k+L))^W.

The practical criterion is not probability but sufficiency: the scheduler is fully occupied once W · k ≥ L approximately, because W warps each supply k instructions during a stall window. Rearranged, the warps needed to cover a latency are

W_needed ≈ L / k

where k is the number of independent instructions the warp can issue between the stall and the next dependent use. With a 400-cycle global load latency and 16 instructions of independent work, a scheduler needs 25 warps assigned to it. With the same latency and 4 instructions, it needs 100, more warps than an SM has to give, which is the definition of a latency-bound kernel.

The figure k is the quantity the programmer controls, and it is controlled by restructuring, not by occupancy. The number of warps the hardware can hold is fixed; the number needed is set by the code. Every technique that increases instruction-level parallelism between a load and its use, unrolling to expose independent loads, prefetching, or computing several output elements per thread, reduces W_needed and therefore reduces the pressure on occupancy.

What the ratio leaves out

Four corrections make the simple model conservative in one direction and optimistic in the other.

The first is that instructions are not one cycle each. A warp issuing an arithmetic instruction occupies the scheduler for that cycle, but a warp issuing a load occupies it for one cycle and then stalls for hundreds. The model counts issue slots correctly but treats all issues as equal, which is right for scheduler occupancy and wrong for datapath occupancy. A kernel whose issue stream is dominated by memory instructions keeps the scheduler busy while leaving the ALUs idle; the warps are resident and issuing, and the throughput is still memory-bound.

The second is that stalls are not independent. Warps in the same block that reach the same barrier stall together, and warps accessing the same cache line contend for the same transaction. When stalls correlate, the probability that all warps are stalled is much higher than the independent model suggests, and the effective W is lower than the resident count. This is the main reason block-level synchronization changes the arithmetic: a barrier converts independent warps into a single synchronized group whose slowest member sets the pace.

The third is that latency is not one number. A hit in the last-level cache costs tens of cycles, a miss to local memory costs hundreds, and the same instruction can experience either depending on the access pattern. The L in the formula is the weighted average over the kernel’s accesses, and a kernel with a 95 percent hit rate in a cache with a 30-cycle hit and a 400-cycle miss has an effective latency near 50 cycles, not 400. That is why the working set size matters more than the nominal memory latency: the achievable occupancy is fixed, so the way to reduce W_needed is to reduce L by making the accesses hit.

The fourth is that occupancy is capped by resources other than warps. The W available is whatever the register and shared memory limits allow, as described in register allocation and the occupancy ceiling. A kernel that needs 25 warps and can hold 16 is latency-bound regardless of what the code does, and the fix is either fewer registers or fewer cycles of latency.

Making a warp’s stall productive

The most direct way to raise k is to give each warp more independent work per stall. A loop that loads one element, uses it, and iterates has k near zero: the load’s result is needed by the very next instruction, so the warp stalls immediately. A loop that loads four elements into four registers before using any of them has k equal to the three intervening loads plus the address arithmetic, and the four loads overlap in the memory system.

/* Four independent loads per iteration: the first three cover part of the
 * fourth's latency, and the arithmetic covers the rest. */
__global__ void scale4(float *out, const float *in, int n)
{
    int base = (blockIdx.x * blockDim.x + threadIdx.x) * 4;
    if (base + 3 >= n) return;       /* partial tail handled by another path */
    float a = in[base + 0];
    float b = in[base + 1];
    float c = in[base + 2];
    float d = in[base + 3];          /* four loads issued back to back */
    out[base + 0] = a * 2.0f;        /* first use comes after all four */
    out[base + 1] = b * 2.0f;
    out[base + 2] = c * 2.0f;
    out[base + 3] = d * 2.0f;
}

The cost is four registers per thread for the temporaries plus the address registers, which raises the register count and lowers the achievable W. The two effects are opposite and both are large, which is why the balance point is kernel-specific. Loading four elements per thread costs roughly four registers and buys a factor of four in k; loading sixteen costs sixteen registers and buys sixteen, but may cross an allocation boundary that halves the warp count.

A cheaper variant is to keep k large without holding many values: prefetch the next iteration’s data into registers while computing the current iteration. That is the classic software pipeline, and it needs only one extra set of live registers rather than a full unroll factor. The loop becomes load-next, compute-current, store-current, rotate, so that one load is in flight during each compute phase.

Prefetch and the limits of interleaving

Prefetching raises k without raising the register count much, but it cannot exceed the memory system’s capacity to hold outstanding requests. Each SM has a finite number of outstanding memory requests it can track, and once that limit is reached, further loads queue rather than overlap. The practical consequence is that the useful k saturates: past a certain number of in-flight loads per warp, additional independent loads add latency rather than hiding it, because they are waiting behind earlier ones in the same queue.

This saturation is the reason that raising occupancy and raising per-thread parallelism are not additive. A machine with 32 resident warps each having 16 outstanding loads wants 512 concurrent requests; the memory system will not sustain that, so the loads serialize at the queue. The right configuration depends on where the bottleneck actually sits, the queue depth, the transaction throughput, or the issue rate, and the three have different fixes. A kernel bound by outstanding-request capacity needs fewer, larger requests; a kernel bound by transaction throughput needs better coalescing; a kernel bound by issue rate needs fewer instructions.

The memory hierarchy’s side of this is outside the scheduling model, but the interaction is not: the scheduling model says how many warps are needed to cover a latency, and the memory model says what latency is actually available at a given concurrency. A divergent region compounds both, because the phases described in divergence amplification issue their loads separately and each phase’s latency must be covered independently. Both are needed to predict what a change will do, and neither alone is sufficient.

Reading the result

The model’s practical output is a target occupancy derived from the kernel’s access pattern, not a target of 100 percent. A kernel with L/k of 8 needs eight warps per scheduler and gains nothing from more; a kernel with L/k of 40 asks for more warps per scheduler than the SM can be partitioned into, and is bound by something that restructuring must address.

The honest summary is that latency hiding is a ratio, and both terms of the ratio are design variables. Latency can be reduced by improving locality, which is a memory-layout question. Work per stall can be increased by restructuring the loop, which is a register-pressure question. Occupancy is the third variable, and it is the one most often measured and least often the binding constraint. When it is the binding constraint, the diagnosis is a mismatch between the kernel’s per-thread parallelism and its resource footprint, and the resolution is one of the two trades above rather than a change to the launch configuration.