Thread hierarchy (CUDA)

CUDA organises threads in two levels. A launch creates a grid of blocks, and each block holds up to 1024 threads. The thread hierarchy is what every kernel uses to work out which element it is responsible for.

int i = blockIdx.x * blockDim.x + threadIdx.x;

Four built-in variables are available inside a kernel:

• threadIdx is the index of the thread within its block
• blockIdx is the index of the block within the grid
• blockDim is the number of threads per block
• gridDim is the number of blocks in the grid

All four are dim3 structs with .x, .y, and .z fields, so grids and blocks can be 1D, 2D, or 3D. A 2D launch maps naturally onto images and matrices:

dim3 threads(16, 16);
dim3 blocks((width + 15) / 16, (height + 15) / 16);
blur<<<blocks, threads>>>(img, width, height);
 
// inside the kernel:
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x < width && y < height) { /* ... */ }

The block is not just bookkeeping. Threads in the same block run on the same streaming multiprocessor, can share shared memory, and can synchronise with __syncthreads(). Threads in different blocks can do neither. Blocks are scheduled in an unspecified order and may not be resident at the same time, so any algorithm that needs blocks to communicate mid-kernel needs to be split into two kernel launches instead.

Block size should be a multiple of 32 (Warps (CUDA)). 128 or 256 is a reasonable default.