Grid-Stride Loop: Vector Add

The launch intentionally uses a small grid. The starter kernel only handles one element per thread, so most of the array is left unwritten. Rewrite the body with a grid-stride loop so the same kernel works for every tested N.

Languages: CUDA C++, Python, Rust, HIP C++

Lesson

A CUDA kernel should not depend on launching at least one thread per element. A grid-stride loop lets each thread process multiple elements separated by the total thread count in the grid, making the kernel correct across different input sizes and launch shapes.

Keep the starter grid of 64 blocks and 128 threads while learning the loop. Increasing the grid would avoid the intended indexing exercise, even if it passes correctness.

What you will learn

Convert a one-element-per-thread kernel into a grid-stride loop.
Understand why Grid X can be smaller than ceilDiv(N, blockX) when the kernel loops.
Keep the boundary check inside the loop so tiny and tail-sized inputs stay correct.

Starter code

    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        C[i] = A[i] + B[i];
    }