Global Memory & Coalescing

When the 32 threads of a warp access global memory, the hardware combines (coalesces) their requests into as few 32-byte sectors / 128-byte transactions as possible.

10.1 Coalescing

When the 32 threads of a warp access global memory, the hardware combines (coalesces) their requests into as few 32-byte sectors / 128-byte transactions as possible.

// COALESCED: consecutive threads access consecutive addresses
float v = a[blockIdx.x * blockDim.x + threadIdx.x];

// STRIDED: thread i accesses a[i * stride] → many transactions, wasted bandwidth
float w = a[(blockIdx.x * blockDim.x + threadIdx.x) * stride];

10.2 Array of Structures vs Structure of Arrays

Prefer SoA on GPUs (or vectorized loads like float4 for AoS with 16-byte elements).

struct ParticleAoS { float x, y, z, mass; };     // AoS: thread i reads p[i].x → stride of 16 bytes
ParticleAoS* p;

struct ParticlesSoA { float *x, *y, *z, *mass; }; // SoA: thread i reads x[i] → fully coalesced

10.3 Vectorized loads

Requires aligned data (cudaMalloc returns ≥256-byte aligned pointers).

__global__ void copy_vec4(const float4* __restrict__ in, float4* __restrict__ out, int n4) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n4) out[i] = in[i];                  // 16-byte loads/stores: fewer instructions
}

10.4 Matrix transpose — a coalescing case study

Naive transpose reads rows (coalesced) but writes columns (strided). The fix: stage a tile in shared memory (§11.4).

10.5 __restrict__ and read-only loads

__global__ void k(const float* __restrict__ in, float* __restrict__ out);   // no aliasing → better optimization
float v = __ldg(&in[i]);                                                     // read-only data cache load

10.6 Memory bandwidth math

Effective bandwidth = (bytes read + bytes written) / time. Compare with the GPU's peak (e.g., H100 SXM ≈ 3.35 TB/s). Memory-bound kernels should reach 70–90% of peak.