GPU Fundamentals: Coalescing & Occupancy

Memory Coalescing (Theory)

Coalescing depends on address patterns across threads in the same warp.

Good:

• thread t=0..31 reads base + t * sizeof(T) (contiguous)

Bad:

• thread t=0..31 reads base + t * sizeof(T) * stride (strided), causing many separate transactions

Data layout principle

When storing 2D arrays in 1D buffers:

• Prefer row-major patterns where threads in x dimension move across columns.

Code Example 1 — Coalesced vs Strided Access (CUDA C++)

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

__global__ void coalesced(const float* A, float* out, int N) {
  int i = blockIdx.x * blockDim.x + threadIdx.x; // 0..N-1
  if (i < N) {
    // contiguous across threads -> coalesced
    out[i] = A[i];
  }
}

__global__ void strided(const float* A, float* out, int N) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < N) {
    int stride = 128; // example stride in elements
    // strided addresses -> poor coalescing
    out[i] = A[(i * stride) % N];
  }
}

int main() { return 0; }

Code Example 2 — Fix a Row/Column Traversal Bug

For a 2D matrix M[rows][cols] stored row-major:

• row-major traversal (vary column within a warp) is coalesced
• column-major traversal is usually strided

#include <cuda_runtime.h>

__global__ void row_major_access(const float* M, float* out, int rows, int cols) {
  int row = blockIdx.y * blockDim.y + threadIdx.y;
  int col = blockIdx.x * blockDim.x + threadIdx.x;

  if (row < rows && col < cols) {
    int idx = row * cols + col;  // row-major index
    out[idx] = M[idx];
  }
}

Occupancy (Theory)

Occupancy is influenced by:

• register usage per thread
• shared memory per block
• max blocks/warps per SM (hardware limits)
• block size and resulting resource allocation

Higher occupancy can help:

• hide long-latency operations (especially global memory)
• keep SMs busy with other warps while one warp waits

But occupancy can be limited by resource usage:

• Using too many registers (or huge shared arrays) reduces resident blocks.

Code Example — Using Occupancy Calculator Pattern

This snippet shows the pattern to query recommended block sizing:

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

__global__ void work(const float* x, float* y, int N) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < N) y[i] = x[i] * 2.0f;
}

int main() {
  int minGridSize = 0;
  int blockSize = 0;

  size_t dynamicSMemSize = 0;

  cudaOccupancyMaxPotentialBlockSize(
    &minGridSize,
    &blockSize,
    work,
    dynamicSMemSize,
    0
  );

  printf("Suggested threads/block: %d\n", blockSize);
  // In real use, test blockSize variants and benchmark.
  return 0;
}

Common Gotchas

• Assuming coalescing from intuition: verify your indexing math.
• Over-optimizing occupancy: you can have high occupancy but low bandwidth utilization.
• Ignoring shared memory/register pressure: tuning block size may backfire.

Quick Checklist

• Map warp threads to contiguous addresses
• Prefer row-major access when possible
• Use occupancy APIs (or Nsight) to estimate limits
• Benchmark multiple block sizes; don’t trust a single metric