The Classic Index Formula

The whole point of a CUDA kernel is that every thread handles one piece of data. To do that, each thread needs a unique global index.

Inside a block, threadIdx.x only counts 0 up to blockDim minus one. Many blocks reuse those same small numbers, so it cannot be your final index.

Picture the grid as blocks laid end to end. Each block owns a contiguous slice of the array, and blockIdx.x tells you which slice you are in.

blockDim.x is the number of threads per block. It is the width of each slice, so it scales your block offset to the right spot.

Combine the three: skip past earlier blocks, then add your position inside this block. That single line gives every thread a unique index.

With 4 threads per block, block 0 covers 0 to 3, block 1 covers 4 to 7, block 2 covers 8 to 11. The offsets never overlap.

Thread 2 in block 3 with blockDim 256 lands at 3 times 256 plus 2, which is 770. That is its global position in the data.

Once you have i, you treat it as an array subscript. Each thread reads and writes only its own element, with no overlap.

Launch n threads and the formula produces every value from 0 to n minus 1 exactly once. That is a perfect one-to-one cover of the array.

Always multiply before you add. blockIdx.x * blockDim.x is the start of your slice, and threadIdx.x is the step inside it.

The same idea extends to 2D and 3D using the .y and .z members, but for flat arrays the .x formula is all you need. 🚀

Compute one thread's global index.

You learned the formula every kernel uses: blockIdx.x * blockDim.x + threadIdx.x. It hands each thread one unique slot in your array. 🎉