100 DAYS OF GPU Course preview
← Glossary
CUDA glossaryExecution model
CC 7.5

What is a grid-stride loop in CUDA?

A loop that lets a fixed-size grid cover any N by having each thread step forward by the total thread count.

One thread per element makes the launch configuration an argument your kernel never declared. out[i] = a[i] + b[i] under if (i < n) is right only while the grid holds at least n threads, and there is nowhere in C to write that precondition down. The loop deletes it by turning the bounds check into a loop condition. Each thread starts at its global index and adds blockDim.x * gridDim.x, the thread count of the whole grid. Call that T. Element i is then written by thread i % T on pass i / T, exactly once, for every T of one or more, so the kernel stops caring which grid it got.

Undersize the grid on a monolithic kernel and nothing complains. Every thread that ran passed its bounds check, every write it made was correct, and the elements past the last thread keep whatever the allocator left there. cudaGetLastError reports success, cudaDeviceSynchronize reports success, the process exits 0. That is a different failure from writing out of bounds, which error checking does catch on the next runtime call. It is also why day 8 fills the output with NaN and not zero: against a zero-filled buffer, a kernel that skips element 0 still matches.

Two smaller strides look plausible and both are wrong. Step by blockDim.x and every block walks the array from its own start, so element i is written once per block. A plain map survives that, because each of those threads stores the same value and your test passes while the kernel does the work several times over. Anything that accumulates is wrong by a factor that moves with the grid. The other is writing the product with no cast: blockDim.x and gridDim.x are both unsigned int, so the multiply happens in 32 bits and wraps. The full stride also protects the access pattern, because inside one pass thread t and thread t + 1 are still one element apart, so a warp covers 128 contiguous bytes and coalescing survives the loop.

The repair people reach for instead is one contiguous chunk per thread, [t * C, (t + 1) * C). On a CPU that is correct: each core streams its own region into its own cache. On a GPU the 32 lanes of a warp execute the same k at the same instant, so their addresses sit C elements apart, and you have built a strided access and given it a tidy name. Day 11 measured that pattern at 32.6 GB/s against 232.9 GB/s for consecutive lanes on identical byte counts. Not the worst row in that benchmark, since stride 32 came in at 9.6 GB/s and a warp's chunk stays cache-resident, but seven times off what the loop gives you for nothing.

Measured

Day 8 ran the same vector add three ways at ten sizes on a Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -O3 -arch=sm_75. The grid came from cudaOccupancyMaxActiveBlocksPerMultiprocessor rather than from n: 4 blocks per SM across 40 SMs, so 160 blocks of 256 threads, 40,960 threads, one grid for every row.

The loop was correct at n = 0, 1, 31, 32, 611, 1,024, 1,025, 40,960, 40,961 and 4,194,304, and under <<<1, 1>>>, <<<1, 256>>>, <<<160, 256>>> and <<<10000, 32>>>. One thread per element on that same 40,960-thread grid was correct up to 40,960 and wrong from 40,961, with the first bad element at index 40,960, which is exactly the first element no thread owns. Sizing the grid to n instead moves the failure to the other end: at n = 0 it asks for zero blocks and the driver refuses the launch with invalid configuration argument.

Nothing on this page is a timing. Day 8 measures correctness only, and the two bandwidth figures above belong to day 11's run on the same card.

Diagram

index-tracer, set to a grid of 1,024 threads over 4,096 elements. Monolithic mode fills the first quarter of the strip and leaves the rest hollow with no error anywhere. Grid-stride mode fills all 4,096 in four passes, shaded by pass, with thread 0 owning cells 0, 1024, 2048 and 3072 and lanes still adjacent inside each pass.

Alt text: "A grid of 1,024 threads writes all 4,096 elements when each thread loops and only the first 1,024 when each thread takes one element. Nothing in either run reports the difference."

Code

From code/day08-grid-stride/grid_stride.cu. The cast is load-bearing, not style.

__global__ void addGridStride(const float* a, const float* b, float* out,
                              size_t n) {
    const size_t step = gridDim.x * static_cast<size_t>(blockDim.x);
    for (size_t i = blockIdx.x * static_cast<size_t>(blockDim.x) + threadIdx.x;
         i < n; i += step) {
        out[i] = a[i] + b[i];
    }
}

The diff against the monolithic version deletes if (i < n) { and its brace and puts the for in their place. Nothing else moves.

Related terms

Where you meet this

Sources

Byline

Written by: pending. Reviewed by: pending. Written on: pending. Last checked on: pending. This entry stays a draft until a named author and a different named reviewer sign it.