Shared Memory & Tiling

11.1 Shared memory basics

On-chip, per-block memory, ~100× lower latency than global memory.

Used for data reuse, inter-thread communication, and coalescing transformations.

__global__ void reverseBlock(int* d) {
    __shared__ int s[256];                 // static shared memory
    int t = threadIdx.x;
    s[t] = d[blockIdx.x * 256 + t];
    __syncthreads();                       // wait until all threads have written
    d[blockIdx.x * 256 + t] = s[255 - t];
}

// Dynamic shared memory (size given at launch)
__global__ void dyn(float* d) {
    extern __shared__ float buf[];
    ...
}
dyn<<<blocks, threads, threads * sizeof(float)>>>(d_ptr);
// > 48 KB requires opt-in:
cudaFuncSetAttribute(dyn, cudaFuncAttributeMaxDynamicSharedMemorySize, 100 * 1024);

11.2 Naive matrix multiplication

// C = A × B, all N×N, row-major
__global__ void matmulNaive(const float* A, const float* B, float* C, int N) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < N && col < N) {
        float sum = 0.0f;
        for (int k = 0; k < N; ++k)
            sum += A[row * N + k] * B[k * N + col];   // each element read N times from global memory
        C[row * N + col] = sum;
    }
}

11.3 Tiled matrix multiplication with shared memory

Global loads drop by a factor of TILE. Further optimizations: register tiling (each thread computes a 4×4 or 8×8 sub-block), vectorized loads, double buffering, Tensor Cores — or just use cuBLAS (§30).

#define TILE 32

__global__ void matmulTiled(const float* A, const float* B, float* C, int N) {
    __shared__ float As[TILE][TILE];
    __shared__ float Bs[TILE][TILE];

    int row = blockIdx.y * TILE + threadIdx.y;
    int col = blockIdx.x * TILE + threadIdx.x;
    float sum = 0.0f;

    for (int t = 0; t < (N + TILE - 1) / TILE; ++t) {
        // Cooperative load of one tile of A and B (coalesced)
        int aCol = t * TILE + threadIdx.x;
        int bRow = t * TILE + threadIdx.y;
        As[threadIdx.y][threadIdx.x] = (row < N && aCol < N) ? A[row * N + aCol] : 0.0f;
        Bs[threadIdx.y][threadIdx.x] = (bRow < N && col < N) ? B[bRow * N + col] : 0.0f;
        __syncthreads();                                     // tile fully loaded

        #pragma unroll
        for (int k = 0; k < TILE; ++k)
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];  // reuse from shared memory
        __syncthreads();                                     // done with this tile
    }
    if (row < N && col < N) C[row * N + col] = sum;
}

dim3 block(TILE, TILE);
dim3 grid((N + TILE - 1) / TILE, (N + TILE - 1) / TILE);
matmulTiled<<<grid, block>>>(d_A, d_B, d_C, N);

11.4 Bank conflicts

Shared memory is divided into 32 banks (4 bytes wide). If multiple threads of a warp access different addresses in the same bank, the accesses serialize.

Transpose with shared memory and padding:

Broadcast (all threads reading the same address) is conflict-free.

__shared__ float tile[32][32];      // column access tile[threadIdx.x][c] → all threads hit the same bank!
__shared__ float tileP[32][33];     // padding by 1 shifts each row → conflict-free column access

__global__ void transpose(const float* in, float* out, int w, int h) {
    __shared__ float tile[32][33];
    int x = blockIdx.x * 32 + threadIdx.x;
    int y = blockIdx.y * 32 + threadIdx.y;
    if (x < w && y < h) tile[threadIdx.y][threadIdx.x] = in[y * w + x];   // coalesced read
    __syncthreads();
    x = blockIdx.y * 32 + threadIdx.x;
    y = blockIdx.x * 32 + threadIdx.y;
    if (x < h && y < w) out[y * h + x] = tile[threadIdx.x][threadIdx.y];  // coalesced write
}

11.5 L1/shared memory carveout

cudaFuncSetAttribute(kernel, cudaFuncAttributePreferredSharedMemoryCarveout, cudaSharedmemCarveoutMaxShared);