# 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.

```c
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:

```c
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-cuda|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.