# Synchronization (CUDA)
`__syncthreads()` is a barrier across all threads in a block. **Synchronization** in CUDA exists at three levels, and each has a different scope: within a block, within a [[warps-cuda|warp]], and between the host and the device.

```c
__shared__ float tile[256];

tile[threadIdx.x] = in[i];
__syncthreads();                  // every thread has written before any thread reads
float left = tile[threadIdx.x - 1];
```

Without the barrier, a thread could read a slot its neighbour has not filled yet. This is the single most common correctness bug in [[shared-memory-cuda|shared memory]] kernels.

`__syncthreads()` must be reached by every thread in the block. Putting it inside a divergent branch is undefined behaviour, and in practice it hangs:

```c
if (threadIdx.x < 32)
    __syncthreads();   // wrong: the other threads never arrive
```

There is no barrier across blocks. Blocks are scheduled in an arbitrary order and may not be resident at the same time, so a grid-wide barrier is impossible in the normal programming model. When an algorithm needs one, the answer is to end the kernel and launch another, since consecutive launches on the same stream are ordered.

Host-side synchronisation is separate again:

 - `cudaDeviceSynchronize()` blocks the host until all device work is done
 - `cudaStreamSynchronize(stream)` blocks until one [[streams-cuda|stream]] drains
 - `cudaEventSynchronize(event)` blocks until a recorded [[events-cuda|event]] is reached

Within a warp, threads no longer advance in lockstep automatically on Volta and later. Code that relied on implicit warp synchrony needs `__syncwarp()` or the explicit-mask shuffle intrinsics instead.