GPU Kernel Memory Access Optimization: Coalescing, Shared Memory, and Bank Conflicts

Note: This is a review article (no first-hand experimental data), based on public practice and the CUDA programming model.

Experiment Objective

Many GPU kernels run slow not because "compute can't keep up" but because "fetching data is slow" — memory bandwidth is not saturated. GPU parallelism relies on "many threads accessing memory simultaneously," and access efficiency is governed by three mechanisms: global-memory coalescing, correct shared-memory use, and bank-conflict avoidance. Goal: dissect the three mechanisms with reproducible code — why contiguous access is fast, how to use shared memory, how to avoid bank conflicts, and vectorized loads. After reading you can answer: where this kernel's access efficiency is lost and how to fix it.

Code Walkthrough

Part 1: Coalesced vs non-coalesced access

// Bad: threads fetch by column (non-coalesced, bandwidth wasted 8-32×)
__global__ void bad(float* a, float* b, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    b[i] = a[i * 32];        // thread i reads a[i*32]: stride 32 elements
}

// Good: threads fetch contiguously (coalesced, one transaction fills)
__global__ void good(float* a, float* b, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    b[i] = a[i];             // adjacent threads read adjacent addresses: perfect
}

Key points:

1. Coalescing defined: a warp (32 threads) accessing contiguous addresses simultaneously — the hardware merges one warp access into minimal memory transactions (128B each). Non-contiguous (strided/transposed) access uses only part of each transaction; bandwidth utilization drops to 1/8-1/32.
2. 90% of access problems are layout/thread-mapping mismatches: design array layout by "thread i reads a[i]," not by "logical index" — transposes, column-wise traversal, and arrays-of-structs (AoS) are the classic pitfalls.
3. Verification: Nsight Compute's memory workload analysis (ch08's nsight article) gives "coalescing/bandwidth utilization" directly — measure before changing.

Part 2: Shared memory and bank conflicts

// Shared memory: block-shared, ~100× faster than global, organized by banks
__global__ void smem(float* in, float* out, int n) {
    __shared__ float tile[32][33];   // note: 2nd dim deliberately +1 (33, not 32)
    int i = threadIdx.x;
    tile[i][0] = in[blockIdx.x * 32 + i];   // write
    __syncthreads();
    // read tile[i][i]: if 2nd dim is 32, same-bank conflict; 33 staggers them
    out[blockIdx.x * 32 + i] = tile[i][i];
}

Key points:

• Shared memory is bank-organized: 32 banks (4B each); a warp hitting different addresses in the same bank = bank conflict (serialized, 2-32× slower); different banks = parallel.
• tile[32][33]'s +1 is the classic trick: 2D arrays are row-major; with row width 32, the i-th row's i-th column all land in the same bank (diagonal access conflicts); width +1 (33) staggers each row, so diagonal access hits different banks — "padding one column" is the most common bank-conflict fix.
• __syncthreads() is shared-memory discipline: always synchronize after writes before reads, or threads read unwritten data.

Principles

Three-level memory hierarchy and optimization order

Global memory (hundreds of GB, slow) → shared memory (tens of KB/SM, 100× faster)
  → registers (fastest, but private)

Optimization order: ① ensure global coalescing first (data layout) → ② move to shared memory when data is reused → ③ use registers for temporaries in loops → ④ handle bank conflicts last. Shared memory pays only when the data is reused by multiple threads in the block (matrix tiling, conv sliding windows) — one-read data in shared memory is a loss.

Vectorized loads (float4)

// float4 reads 16B at once: fewer instructions, better bandwidth (aligned)
float4 v = reinterpret_cast<float4*>(a)[i];

• Vectorization cooperates with coalescing: each thread reads float4 (16B); a warp reads 512B per transaction — 4× fewer instructions than per-float reads, fuller bandwidth.
• Alignment requirement: float4 needs 16B alignment — allocate layouts 16B-aligned, or illegal address.

Access-pattern vs compute-pattern matching

Compute pattern	Access pattern	Optimization
element-wise (add/mul)	contiguous stream	coalescing + vectorization
tiled (matmul/conv)	tile reuse	shared memory + padding
reduction (sum/extrema)	cross-thread	staged shared-memory reduction

Core insight: memory optimization is not "adding cache" but matching access patterns to the hardware's parallel access mechanisms (coalescing, banks, vectorization) — design data layout by compute pattern, then pick the optimization by access pattern.

Results

(Estimated, not measured; bandwidth-bound kernels)

Optimization	Bandwidth utilization gain	Cost
coalescing (layout fix)	2-8×	change layout/indexing
shared memory + padding	1.5-4× (reuse scenarios)	code complexity
float4 vectorization	1.3-2×	alignment constraint

Engineering Takeaways

1. Layout before optimization: design data layout by "thread i reads a[i]"; coalescing is the foundation — wrong layout makes everything else moot.
2. Shared memory only on "reuse": matrix tiling/conv reuse scenarios justify moving to shared memory; one-read scenarios stay with global reads.
3. Pad one column to cure bank conflicts: for diagonal/strided conflicts, width +1 is the fastest fix.
4. Verify coalescing with Nsight: memory workload metrics (ch08's nsight article) quantify "coalescing/bandwidth utilization"; accept with before/after data.

Next article preview: memory-access optimization done; ch08 and ch04 both topped up this round; the next batch could return to ch11 training or ch13 application, or pivot to SEO wrap-up. Your call.