Memory-Level Parallelism and Latency Tolerance

A memory system has a fixed latency and a fixed bandwidth, and the two are connected by the number of requests that can be outstanding at once. A kernel that issues one dependent load at a time can use only as much bandwidth as one request per latency allows, which is a small fraction of the interface. Raising that count is the subject, and it is not the same as raising occupancy.

The concurrency requirement restated

If the interface can deliver $B$ bytes per second at a latency of $L$ seconds and transfers $S$ bytes per request, then saturating it requires

$$ C = \frac{B \cdot L}{S} $$

outstanding requests. Nothing about the kernel enters this expression. It is a property of the memory system, and it sets a target the kernel must meet by whatever means available.

With a bandwidth near 1 TB/s, a latency near 400 ns, and 128-byte requests, $C$ is roughly 3,100 requests for the whole device. A device with 80 multiprocessors holding 64 warps each holds about 5,100 warps, so the requirement is under one outstanding request per warp. The same figure follows from a per-multiprocessor calculation: an even share of the bandwidth is about 12.5 GB/s, which over 400 ns at 128 bytes works out to roughly 39 outstanding requests per multiprocessor, against 64 resident warps.

The requirement is therefore modest in aggregate and unforgiving in distribution. It assumes every resident warp keeps a request in flight at all times, which a dependent chain prevents: a warp executing load; use; load; use has one request outstanding and spends the latency of each load stalled, so the warps that are stalled contribute nothing while the interface waits. A kernel that reaches the aggregate figure in principle can still miss it by a wide margin if half its warps are between requests at any moment.

A warp that executes four loads and then four uses has four outstanding and hides four times the latency. The difference is entirely in the instruction order, and it costs nothing in register terms if the four values are live simultaneously.

Where the outstanding requests are counted

The hardware limits outstanding requests per warp and per multiprocessor. A warp cannot have arbitrarily many loads in flight: the scoreboard tracks a fixed number of pending operations, and once exhausted, further loads stall. Per-multiprocessor limits on miss-status holding registers impose a second cap, and this is often the real ceiling for kernels that attempt deep unrolling.

The practical consequence is that per-thread independence has diminishing returns. Going from one outstanding load to four raises the count substantially. Going from four to sixteen may hit the per-warp limit, at which point additional unrolling produces code that the hardware executes serially and register pressure that reduces occupancy. The optimum is usually in the range of four to eight independent accesses per thread, though the exact figure is architecture-dependent and should be found empirically rather than assumed.

Independent loads per thread	Registers consumed	Effect
1	Minimal	Latency exposed, bandwidth underused
2–4	Small	Most of the latency hidden for typical kernels
4–8	Moderate	Near saturation for streaming kernels
8–16	Large	Diminishing returns, occupancy loss likely
>16	Severe	Per-warp limits reached, spilling possible

Unrolling that actually reorders

A compiler unroll will not necessarily produce independent loads. If the loop body is acc += a[i] * b[i], unrolling by four gives four loads and four multiply-adds, and the compiler may still emit them in an order that consumes each value before the next load issues, because it has no reason to believe the loads are independent unless it can prove no aliasing.

Two techniques change this. The first is to separate the load phase from the compute phase explicitly, using an array of intermediate values:

// Four loads issued before any of them is consumed.
float v0 = x[i+0], v1 = x[i+1], v2 = x[i+2], v3 = x[i+3];
float w0 = y[i+0], w1 = y[i+1], w2 = y[i+2], w3 = y[i+3];
acc += v0*w0 + v1*w1 + v2*w2 + v3*w3;

The compiler emits the eight loads before the first multiply because the values are needed together. The register cost is eight floats plus the accumulator, which is affordable at typical register budgets.

The second is to make the independence provable. Pointer arguments are assumed to alias unless the compiler is told otherwise, so a[i] and b[i] may be the same memory and the compiler must order the accesses conservatively. The __restrict__ qualifier on pointer parameters tells the compiler the ranges do not overlap, which permits reordering. This is not a micro-optimization; on some kernels it is the difference between a serialized load chain and a parallel one.

__global__ void axpy4(int n, float a, const float * __restrict__ x, float * __restrict__ y) {
    int i = 4 * (blockIdx.x * blockDim.x + threadIdx.x);
    if (i + 3 < n) {
        float x0 = x[i], x1 = x[i+1], x2 = x[i+2], x3 = x[i+3];
        y[i]   = a * x0;
        y[i+1] = a * x1;
        y[i+2] = a * x2;
        y[i+3] = a * x3;
    } else {
        for (int j = i; j < n; ++j)   // at most three leftover elements
            y[j] = a * x[j];
    }
}

The four-element unrolling does not by itself improve coalescing: each thread’s four elements are sixteen bytes apart from the next thread’s, so an individual load instruction is strided across the warp and touches more sectors than a single fully coalesced load of the same size. The total sectors touched over the four instructions are the same as before, but the transaction count per instruction is higher. Loading a float4 per thread instead restores per-instruction coalescing, since one instruction then covers 16 contiguous bytes per thread and 512 contiguous bytes per warp. The coalescing side is treated in sector counting.

The interaction with occupancy

Occupancy and per-thread parallelism multiply to give total concurrency:

$$ C_{total} = W_{resident} \times C_{per_warp} $$

The product is what matters, and the two factors trade against each other through register pressure. A kernel using 32 registers per thread can be at full occupancy with one load in flight per warp. A kernel using 64 registers per thread halves occupancy but can have four loads in flight per warp, for a net doubling of total concurrency. This is the standard reason a low-occupancy kernel can saturate bandwidth, and the reason occupancy targets derived from a different kernel’s behaviour are unreliable.

The trade has a limit. Register pressure beyond the point where the compiler spills converts a latency problem into a traffic problem, since spilled values go to local memory and every spill is a store followed by a load. A kernel that spills is adding memory traffic to solve a memory latency problem, which is a losing exchange in any kernel that is already bandwidth bound.

The diagnostic is straightforward: compare the reported register count against the occupancy it yields, and check whether the compiler reports spills. A kernel with spills and low occupancy is in the regime where the trade has gone wrong, and the answer is to reduce live values rather than to force more unrolling.

Why this is not the same as prefetching on a CPU

A CPU core hides memory latency with out-of-order execution, which finds independent instructions within a window and issues them ahead of stalled ones. A GPU warp is in-order: it issues instructions in program order and stalls at the first one whose operands are not ready. There is no window to search, so the independence must be visible in the instruction stream and the scheduler cannot manufacture it.

This makes the GPU more sensitive to code structure and less sensitive to compiler scheduling than a CPU. It also means that source-level transformations which are irrelevant on a CPU, such as declaring pointers non-aliasing or manually staging loads into temporaries, can change GPU performance by a large factor.

The other side of the coin is that GPU warps are cheap to switch. Where a CPU stalls a thread, a GPU stalls a warp and issues from another, so a kernel with low per-warp concurrency but high occupancy can still do well. The two mechanisms are substitutes, and the optimum mix depends on which resource is scarcer, which is exactly what the occupancy calculation in occupancy as a latency-hiding budget determines.

Measuring whether concurrency is the limit

Three signals indicate that memory-level parallelism, rather than bandwidth or issue rate, is the binding constraint.

Achieved bandwidth rises when the grid is enlarged with the same total work, meaning the kernel was not saturating with the smaller grid. This is the cleanest signal and the easiest experiment.

Achieved bandwidth rises when per-thread unrolling increases without a change in occupancy. This isolates the per-warp factor.

The stall histogram is dominated by long scoreboard while the achieved bandwidth is well below the practical ceiling, and occupancy is already at its maximum. Then the only remaining lever is per-warp concurrency.

The third signal is the one that most often gets missed, because it requires holding two figures in view at once. A kernel at maximum occupancy with a long-scoreboard stall profile and a low achieved rate has no more warps to add, so the conclusion “increase occupancy” is unavailable and the correct one is “increase independence”. Conversely, a kernel with low occupancy but high per-warp concurrency and a high achieved rate is already at its ceiling and should not be tuned further on either axis.

A fourth, negative signal is worth stating: if enlarging the grid, increasing unrolling, and raising occupancy all fail to move the achieved rate, the kernel is at the memory system’s practical ceiling and the remaining lever is to move fewer bytes. That is a data-layout and fusion question, taken up in kernel fusion and launch boundaries, not a concurrency question.