<<<CUDA C++ Coursegrid · block · warp · lane
EN中文
Stage 2 · Execution model and memory06 / 15about 22 mincuda/07_transpose.cu

Memory coalescing: the same kernel, a 5x gap

How a warp issues memory transactions, what stride access is, and the AoS vs SoA choice

After this lesson you can
  • See how a warp's loads are coalesced into 32-byte sectors
  • Spot the common patterns that break coalescing
  • Apply AoS → SoA, usually the highest-payoff layout change
  • Use shared memory to repair coalescing that you cannot avoid breaking

A GPU does not fetch device memory by the byte. The unit is a 32-byte sector. When 32 threads in a warp issue loads together, hardware packs their addresses into as few sector requests as it can. Ideal case: 32 threads read 32 consecutive floats (128 bytes) = 4 sectors, one trip, nothing wasted. Worst case: 32 threads each read a far-apart float = 32 sectors = 1024 bytes moved, of which you use 128. 87.5% of the bandwidth is thrown away.

Coalesced    c[i] = a[i]            stride-8 access  c[i] = a[i * 8]

addr: 0  4  8 ... 124                addr: 0  32  64 ... 992
thrd: t0 t1 t2 ... t31               thrd: t0 t1  t2 ... t31

 ┌────┬────┬────┬────┐                 ┌────┐    ┌────┐    ┌────┐
 │sec0│sec1│sec2│sec3│                 │sec0│... │sec8│... │sec31│
 └────┴────┴────┴────┘                 └────┘    └────┘    └────┘
   all hit, no waste                      32 sectors, each uses 4/32 B

 moved 128 B / used 128 B = 100%       moved 1024 B / used 128 B = 12.5%
Coalesced vs strided: both read 32 floats
▸Open the coalescing simulatorTune stride, offset, and element width, and watch how many sector transactions a warp produces and what the bandwidth utilization is. Drag stride from 1 to 2 to 32. The drop in efficiency is obvious.

Four patterns that break coalescing

1. Strided access
1// Wrong: each thread skips ahead. The warp spans 32 * 8 * 4 = 1024 bytes2__global__ void strided(const float* in, float* out, int stride) {3    int i = blockIdx.x * blockDim.x + threadIdx.x;4    out[i] = in[i * stride];5}6 7// Correct: if stride comes from multi-channel data, put channel on blockIdx and element on threadIdx8__global__ void byChannel(const float* in, float* out, int n) {9    int i = blockIdx.x * blockDim.x + threadIdx.x;   // element dim, contiguous10    int c = blockIdx.y;                              // channel dim, constant inside the warp11    out[c * n + i] = in[c * n + i];12}
2. Row and column swapped (the usual 2D indexing bug)
1// Wrong: threadIdx.x is the row. The 32 threads in a warp each stride a full row2int row = blockIdx.x * blockDim.x + threadIdx.x;3int col = blockIdx.y * blockDim.y + threadIdx.y;4out[row * cols + col] = in[row * cols + col];5 6// Correct: threadIdx.x is the column. The warp covers 128 contiguous bytes7int col = blockIdx.x * blockDim.x + threadIdx.x;8int row = blockIdx.y * blockDim.y + threadIdx.y;9out[row * cols + col] = in[row * cols + col];

3. AoS vs SoA: the classic particle-system case

Array of structures (AoS): natural to write, terrible for coalescing
1struct Particle { float x, y, z, vx, vy, vz; };   // 24 bytes2 3__global__ void stepAoS(Particle* p, float dt, int n) {4    int i = blockIdx.x * blockDim.x + threadIdx.x;5    if (i >= n) return;6    p[i].x += p[i].vx * dt;      // x of neighboring threads is 24 bytes apart7    p[i].y += p[i].vy * dt;8    p[i].z += p[i].vz * dt;9}10// A warp reading x spans 32 * 24 = 768 bytes,11// and 24 is not a factor of 32, so sectors also misalign and efficiency drops further.
Structure of arrays (SoA): fully coalesced
1struct Particles {          // each field is its own contiguous array2    float *x, *y, *z;3    float *vx, *vy, *vz;4};5 6__global__ void stepSoA(Particles p, float dt, int n) {7    int i = blockIdx.x * blockDim.x + threadIdx.x;8    if (i >= n) return;9    p.x[i] += p.vx[i] * dt;      // warp reads x[0..31] = 128 contiguous bytes10    p.y[i] += p.vy[i] * dt;11    p.z[i] += p.vz[i] * dt;12}

4. When non-coalesced access is unavoidable: stage through shared memory

Matrix transpose is the classic case: of the read and the write, you can keep at most one coalesced. The fix is to keep both global loads and stores coalesced, and do the transpose shuffle inside shared memory, where random access is far cheaper than device memory.

07_transpose.cu: shared-memory transpose
1#define TILE 322__global__ void transposeShared(const float* in, float* out, int w, int h) {3    __shared__ float tile[TILE][TILE + 1];   // +1 kills the bank conflict; next lesson covers why4 5    int x = blockIdx.x * TILE + threadIdx.x;6    int y = blockIdx.y * TILE + threadIdx.y;7 8    // Load: warp walks x contiguously → coalesced9    if (x < w && y < h) {10        tile[threadIdx.y][threadIdx.x] = in[y * w + x];11    }12    __syncthreads();13 14    // Store: swap the block indices so the output-side warp also walks contiguously → coalesced too15    x = blockIdx.y * TILE + threadIdx.x;16    y = blockIdx.x * TILE + threadIdx.y;17    if (x < h && y < w) {18        out[y * h + x] = tile[threadIdx.x][threadIdx.y];   // the transpose happens in shared memory19    }20}

Self-check

Answer in your head first, then open the solution. If you cannot, reread that section.