CUDA Memory Hierarchy & Optimization

Why Memory Dominates GPU Performance (Theory)

CUDA kernels often become memory-bound: even if you have many compute units, slow memory access stalls execution.

Key pieces:

• Global memory (device DRAM): large but high latency.
• Shared memory (per-block, on-chip): low latency, manually managed by you.
• Registers (per-thread): fastest, but limited.
• L2 cache (device-wide): helps reuse, but you shouldn’t rely on it for correctness.

Optimization mantra:

1. Move data closer (global → shared → registers)
2. Reuse data (tiling)
3. Make accesses coalesced (warp-friendly layout)
4. Avoid wasting bandwidth (don’t reload the same values)

Code Example 1 — Coalesced vs Non-Coalesced Access

Consider reading a 2D array stored row-major in a flat 1D buffer.

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

__global__ void read_rows(const float* A, float* out, int N) {
  // Each thread reads A[i][j] where i varies slowly and j varies with threadIdx
  int j = blockIdx.x * blockDim.x + threadIdx.x; // contiguous across threads
  if (j < N) out[j] = A[j * N + j]; // a simple diagonal pattern
}

__global__ void read_columns(const float* A, float* out, int N) {
  // Threads read addresses that are far apart (strided) => poor coalescing
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < N) out[i] = A[i * N + (i % 32)]; // strided behavior
}

int main() {
  int N = 1 << 20;
  // In real code you’d benchmark and measure with nvprof/nsys.
  // This file is illustrative.
  return 0;
}

What to take away: if threads in the same warp read scattered addresses, memory transactions increase and bandwidth drops. Even “correct” code may run far slower.

Code Example 2 — Tiled Matrix Multiply Using Shared Memory

A classic optimization: compute C = A * B using tiling:

• Each block computes a tile of C
• Each thread loads elements into shared memory
• Compute uses shared memory, reducing global loads

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

#define TILE 16

__global__ void matmul_tiled(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;

  // Loop over tiles of the K dimension
  for (int t = 0; t < (N + TILE - 1) / TILE; t++) {
    int tiledCol = t * TILE + threadIdx.x;
    int tiledRow = t * TILE + threadIdx.y;

    // Load one tile of A and B into shared memory
    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(); // ensure As/Bs are fully loaded

    // Multiply the shared-memory tiles
    #pragma unroll
    for (int k = 0; k < TILE; k++) {
      sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
    }

    __syncthreads(); // safe to reuse shared memory for next tile
  }

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

Common Gotchas

• Missing __syncthreads(): shared memory loads are not guaranteed visible before computation.
• Bank conflicts: shared memory is banked; certain access patterns can serialize reads.
• Occupancy vs shared memory usage: using large shared arrays can reduce the number of concurrent blocks.
• Indexing mistakes: off-by-one in bounds checks can silently corrupt memory.

Quick Checklist

• Ensure warp threads access contiguous addresses
• Use tiling when you repeatedly reuse data
• Keep shared memory arrays sized reasonably
• Use bounds checks for non-multiple tile sizes