CUDA essentials: memory coalescing

Bytes, words, sectors

• A word is a 4-byte unit. It is a machine's natural register size.
• 8 words = 32 bytes constitute a sector.

A sector is a fundamental unit of data transfer. Even if you want to read or write just a single word, the entire sector the word's memory belongs to is moved across HBM.

Initial setup

HBM

• The figure shows 64 words in contiguous memory.
• Each rectangle shows a word (4 bytes); this is a sector.
• The memory can be split into 8 sectors.

SRAM and warp lanes

• SRAM
• Warp lanes

Reading words from HBM

float x = A[2];    // thread 0 -> sector S0
float y = A[35];    // thread 1 -> sector S4
// Same 2 values needed, but they land in different sectors - the
// controller has to fetch both sectors in full. 2 fetches, not 1.

• 1 thread reads A[3]: one sector fetched
• 2 threads read A[2] and A[3]: still one sector fetched (coalesced)
• 2 threads read A[2] and A[35]: two sectors fetched (not coalesced)
• 8 threads read A[0] through A[7]: just one sector fetch (coalesced)

You may have noticed: coalesced operations are fast and efficient.

Coalesced and strided warp reads of A[tid]

The above examples were just instructional. In reality, all threads in a warp execute the same instruction together. So when that instruction is a memory read, all 32 threads issue their read at once. The fetch happens for all 32 addresses, not just one.

• Thread tid reads A[tid]: coalesced access
• Thread tid reads A[tid * 2]: strided access
• All threads read the same value: broadcasting

A warp's 32 threads need 32 words (32 x 4 bytes). Whether that costs 4 sectors or 32 sectors depends entirely on which addresses the threads ask for.

__global__ void coalesced(float* A, float* out) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    out[tid] = A[tid];   // thread 0 -> A[0], thread 1 -> A[1], ...
}

__global__ void strided(float* A, float* out, int stride) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    out[tid] = A[tid * stride];   // e.g. stride=2 -> A[0], A[2], A[4], ...
}

__global__ void broadcast(float* A, float* out) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    out[tid] = A[0];   // every thread reads the same word
}