← Back to examples

CUDA Thread Indexing: grid → block → thread

When you launch a kernel with kernel<<<blocks, threads>>>(), CUDA creates a grid of blocks, and each block holds a team of threads. Every thread computes one global index to find its own piece of data: i = blockIdx.x * blockDim.x + threadIdx.x. Click any thread below to see exactly how it maps onto the data array — and what the bounds check if (i < n) is for.

Threads per block (blockDim.x): Number of blocks (gridDim.x): Data size (n):

The grid of blocks (each square is one thread — click one)

The data array in GPU memory (float d_x[n])

// click a thread above to compute its global index

selected thread data element it writes extra thread (i ≥ n) — bounds check skips it

The takeaway: you launch more threads than data elements whenever n isn't an exact multiple of the block size — here the last block has threads that run off the end of the array. That is why every CUDA kernel starts with if (i < n) return; (or wraps its work in if (i < n)): the extra threads must do nothing, or they corrupt memory past the array.