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

__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. 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).