# Shared memory (CUDA)
**Shared memory** is a small, fast scratchpad private to each thread block, declared with `__shared__`. It sits on the multiprocessor itself, roughly an order of magnitude faster than [[global-memory-cuda|global memory]], and it is the main tool for cutting memory traffic when threads reuse the same data.

Unlike a CPU cache, it is managed by hand. You decide what goes in, when, and when it leaves.

```c
__global__ void stencil(const float *in, float *out, int n) {
    __shared__ float tile[BLOCK + 2];
    int g = blockIdx.x * blockDim.x + threadIdx.x;
    int t = threadIdx.x + 1;

    tile[t] = in[g];                                  // each element loaded once
    if (threadIdx.x == 0)          tile[0] = in[g - 1];
    if (threadIdx.x == blockDim.x - 1) tile[t + 1] = in[g + 1];

    __syncthreads();                                  // wait for the whole tile

    out[g] = 0.25f * tile[t - 1] + 0.5f * tile[t] + 0.25f * tile[t + 1];
}
```

Without the tile, each element of `in` would be read three times from global memory, once by each neighbouring thread. With it, each element is read once and then shared. The `__syncthreads()` is mandatory, since a thread must not read a slot another thread has not filled yet ([[synchronization-cuda]]).

Shared memory is a limited resource, typically 48 KB to 164 KB per multiprocessor depending on architecture. A block that asks for a lot of it reduces how many blocks fit on a multiprocessor at once, which lowers [[occupancy-cuda|occupancy]]. That trade is often worth making, but it should be a deliberate choice rather than an accident.

The size can also be set at launch instead of compile time, with `extern __shared__ float tile[];` in the kernel and the byte count as the third launch parameter.

Shared memory is banked into 32 four-byte-wide banks, and access patterns that collide across banks serialise. See [[bank-conflicts-cuda]].