Shared Memory & Warp-Level Programming in CUDA C++

Shared Memory as a Software Cache (Theory)

Shared memory is:

• On-chip (much lower latency than global)
• Per-block (scope is block)
• Manually managed (you load and synchronize)

Typical pattern:

1. Threads cooperatively load a tile from global → shared
2. __syncthreads()
3. Threads compute using shared
4. __syncthreads() before reusing shared for next tile

Code Example 1 — Tiled Matrix Multiply (Shared Memory)

#include <cuda_runtime.h>
#include <cstdio>

#define TILE 16

__global__ void matmul_shared(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++) {
    // Load tile
    int tiledCol = t * TILE + threadIdx.x;
    int tiledRow = t * TILE + threadIdx.y;

    As[threadIdx.y][threadIdx.x] = (row < N && tiledCol < N) ? A[row * N + tiledCol] : 0.0f;
    Bs[threadIdx.y][threadIdx.x] = (tiledRow < N && col < N) ? B[tiledRow * N + col] : 0.0f;

    __syncthreads();

    // Compute
    #pragma unroll
    for (int k = 0; k < TILE; k++) {
      sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
    }

    __syncthreads();
  }

  if (row < N && col < N) C[row * N + col] = sum;
}

Code Example 2 — Warp-Level Reduction with Shuffle

Warp-level reductions avoid:

• shared memory arrays
• block-wide synchronization

Example: sum within each warp.

#include <cuda_runtime.h>

__inline__ __device__ float warp_reduce_sum(float val) {
  // full mask for active threads
  unsigned mask = 0xffffffff;
  for (int offset = 16; offset > 0; offset >>= 1) {
    val += __shfl_down_sync(mask, val, offset);
  }
  return val;
}

__global__ void reduce_warp(const float* x, float* out, int N) {
  int idx = blockIdx.x * blockDim.x + threadIdx.x;
  float v = (idx < N) ? x[idx] : 0.0f;

  // Reduce inside each warp
  v = warp_reduce_sum(v);

  // One lane per warp writes
  int lane = threadIdx.x % warpSize;
  int warpId = threadIdx.x / warpSize;

  if (lane == 0) {
    out[blockIdx.x * (blockDim.x / warpSize) + warpId] = v;
  }
}

Common Gotchas

• Bank conflicts: shared memory is banked; certain access patterns serialize. (Padding can help.)
• Missing __syncthreads(): causes stale/incorrect reads from shared.
• Assuming warp-wide behavior beyond a warp: shuffle works within a warp only.
• Divergence: warp primitives assume participating threads use the provided mask correctly.

Quick Checklist

• Use shared memory for cooperative tiling and reuse
• Synchronize before/after shared memory reuse
• Use warp shuffle for fast warp-local communication
• Benchmark; warp-level code can be faster but also more complex