Grid-Stride Loops and Launch Size Independence

The natural kernel launches one thread per element: the grid size is derived from the buffer length, each thread computes one linear index, and the kernel body is the operation applied once. It is the simplest arrangement and it has a structural problem. The launch size is fixed at launch time, so the host must know the problem size, and the number of threads that can be resident at once is bounded by the hardware. Between those two facts sits a decision that the one-thread-per-element form does not let the programmer make: how much work each resident thread performs before it retires.

The grid-stride loop moves that decision back into the kernel. The launch covers a fixed number of threads — usually a small multiple of what the device can hold — and each thread walks the problem in steps equal to the total thread count.

The loop and the constant it depends on

__global__ void saxpy(int n, float a, float* y, const float* x) {
    // stride is the grid's total thread count, which the launch fixes.
    const 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];
    }
}

Two properties follow from that single line of arithmetic. First, the traversal is complete: every index in $[0, n)$ is visited exactly once, because the indices form a set of residue classes modulo stride and the loop covers each class up to $n$. Second, the order in which a thread visits its elements is increasing but not adjacent — consecutive visits by one thread are stride elements apart. Consecutive threads are still adjacent to each other within each iteration, so a warp’s accesses within one trip are contiguous and the memory system sees the same coalesced pattern it would in the one-element form.

That second property is why the loop is not a performance compromise by construction. The warp still loads a contiguous run of 32 elements per trip; it merely performs several such runs, spaced by the grid width, before terminating. The coalescing analysis that applies to a block’s single trip applies unchanged within each trip of the loop, because the warp’s membership is the same in every iteration.

Why the stride is the grid, not the block

A loop that stepped by blockDim.x instead would be a block-local loop: each block would cover a contiguous window of the problem and blocks would not overlap. That version is also correct, and it is what a tiled kernel does for its own tile. The difference appears when the launch is smaller than the number of tiles. With a block-local loop, a grid of 64 blocks can only cover the window of the first 64 blocks’ worth of work, and the rest of the problem is simply never processed. With a grid-stride loop, block 0 processes elements at offsets $0$, $1 \cdot \text{stride}$, $2 \cdot \text{stride}$ and so on, so the grid’s coverage is independent of the grid’s size. Any grid of at least one block covers the whole problem.

This is what makes the launch size a tuning parameter rather than a correctness parameter. The same kernel can be launched with a grid sized to the device, with a grid one tenth that size, or with a grid of exactly one block, and it produces the same result. That property is what allows the host to cap the launch at a number derived from the device’s capabilities, to subsume small problems into a single block, and to change the launch size without touching the kernel.

Choosing the launch size

The number of threads that can be resident on a device is the product of the multiprocessor count, the maximum resident blocks per multiprocessor, and the block size, subject to register and shared-memory limits. A launch of that size fills the machine for one wave. A launch of several times that size queues multiple waves, which is what the hardware would do anyway with a one-thread-per-element launch.

There are two reasons to cap the grid rather than let it grow with $n$. Each block has a launch cost — the scheduler must dispatch it, and its threads must be issued their first instructions, which includes reading the kernel parameters and computing the initial index. For a large problem those costs are amortised over a lot of work per block; for a small problem with one element per thread, the block does one element and retires, and the dispatch cost is a meaningful fraction of the total. The grid-stride loop with a capped grid makes each block do many elements, so the fixed per-block cost is paid once.

The other reason is that the number of blocks in the grid is a 32-bit quantity per axis, and the total number of threads in a one-element-per-element launch is only bounded by the buffer size. Capping the grid keeps the launch parameters small and keeps the grid size independent of buffer growth.

The counter-argument is load balance. A grid-stride loop that fills exactly one wave and then loops gives every resident thread a number of trips that differs by at most one, which is near-optimal. A grid that is not a multiple of the machine’s capacity fills a partial final wave, and in the worst case the final wave contains one block, so for a short interval one block runs while the rest of the device is idle. The same is true of the one-element form, so the grid-stride loop does not make this worse — but it does mean that a cap chosen as “some multiple of the device’s capacity” has a quantisation effect, and the optimum shifts with problem size.

The dirty detail: trip counts and index width

The index in the loop above is int, and n is too, so every comparison is 32-bit and the index wraps silently once a thread’s position passes $2^{31}$. Declaring n as 64-bit does not fix it on its own: the comparison promotes i to 64 bits, but the increment i += stride still happens in 32 bits, so the index wraps between iterations where the test cannot see it. The safe form uses a 64-bit index and a 64-bit stride:

const size_t stride = (size_t)blockDim.x * gridDim.x;
for (size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x; i < n; i += stride)

The product blockDim.x * gridDim.x is computed in 32 bits unless one operand is widened first — the same trap described in thread hierarchy and index arithmetic. A block of 1024 threads and a grid of 2 million blocks is $2 \times 10^9$ threads and still fits; a larger grid does not, and the resulting stride wraps to a value that silently skips elements.

There is a subtler issue with the trip count itself. If a thread’s iterations are dependent — an accumulator that must be combined across trips — the loop introduces a linear dependence chain of length $\lceil n / \text{stride} \rceil$. For an accumulate-and-store kernel this is irrelevant, because the dependency is on memory which was already serialised. For a compute-heavy body, the chain limits instruction-level parallelism: the next trip cannot begin until the previous one’s arithmetic retires. The one-element form has no such chain; the price of skipping the chain is a launch that is as large as the problem.

Instruction overhead

Every trip executes the loop test and the stride addition, which is a handful of instructions. For a body of a single multiply-add, that overhead is a large fraction of the body — the loop is slower than the unrolled one-element form for exactly the small, linear kernels where the launch cost per block matters most. For a body of tens of instructions the overhead is negligible.

The compiler cannot generally unroll the loop by a fixed factor, because the trip count is not known at compile time and the body may be predicated. A common manual remedy is to process several consecutive elements per trip with the loop still stepping by the grid stride:

const size_t tid = (size_t)blockIdx.x * blockDim.x + threadIdx.x;
const size_t stride = (size_t)blockDim.x * gridDim.x;

// Four elements per trip, each one stride apart and each coalesced
// across the warp; the loop control is paid once for all four.
for (size_t base = tid; base < n; base += stride * 4) {
    for (int k = 0; k < 4; ++k) {
        const size_t i = base + k * stride;
        if (i < n) body(i);
    }
}

Each of the four inner accesses is still coalesced within a warp, and the loop overhead is paid once per four elements. The inner bounds tests remain, so the gain is the loop control only, not the checks. This is a vectorisation-of-the-loop rather than vectorisation of the data, and the two are unrelated: the four accesses are 32-bit loads at addresses separated by stride elements, not one wide load.

When the loop is the wrong shape

A grid-stride loop assumes that iteration order is irrelevant. Kernels where the order of writes matters — a compaction that needs to know how many earlier threads matched, a scan whose output at $i$ depends on $i-1$, or any transform with a sequential dependency — cannot be expressed by distributing iterations round-robin, because the per-thread state would not respect the order. Those kernels either restructure the algorithm so the dependency is explicit, as in the single-pass scan discussed in single-pass prefix scan with decoupled look-back, or they keep one element per thread and pay for a large launch.

Kernels that use shared memory per tile also do not fit the loop directly. The tile is loaded cooperatively and consumed after a barrier, and a grid-stride iteration would need the tile reloaded for each iteration, with a barrier inside the loop. That is legal, and it is what a persistent-thread formulation of a tiled algorithm does, but it brings the divergence and barrier constraints described in block barriers and divergent synchronization into the loop body, so the loop is no longer a free abstraction.

What it buys in the end

The grid-stride loop is a standing constraint on the kernel — every thread must be correct at any index, and any per-block resource must be sized for the trip count, not the element count. What it returns is a kernel whose launch is a tuning decision and whose correctness holds for a one-block launch of a billion-element problem. For a library kernel that must run on unknown hardware with unknown problem sizes, that independence is usually worth more than the loop’s instruction overhead.