Coalesced memory access

What Is Coalesced Memory Access?

Coalesced memory access occurs when threads in a GPU warp access consecutive memory addresses that can be served by the minimum number of memory transactions. GPUs fetch memory in fixed-size transactions (bursts of a fixed number of bytes). When the threads of a warp request data that falls within the fewest possible transactions, the access is coalesced. When the requests scatter across many transactions, memory efficiency drops.

Why GPUs Care: The Hardware Reality

Memory Transaction Granularity (Architecture-Dependent!)

GPU memory controllers fetch data in fixed-size transactions. Representative values on NVIDIA GPUs:

Level	Fermi/Kepler/Pascal	Volta/Turing/Ampere+
L1 cache line (fill granularity)	128 bytes	128 bytes
L2 cache line	32 bytes (sector) within 128-byte line	32-byte sectors within 128-byte line
Global memory (DRAM) transaction	128-byte transactions (segmented into 32-byte parts)	Can service 32-byte sectors independently

The key architectural distinction:

Older GPUs (Fermi/Kepler/Pascal): a global load through L1 pulls a full 128-byte cache line.
Newer GPUs (Volta+): the memory subsystem is sectored, so it can service 32-byte sectors independently, wasting less bandwidth on partial accesses.

When thread 0 reads address 0x1000, the hardware doesn't just fetch 4 bytes — it fetches the entire transaction/sector containing that address (e.g., a 32-byte sector 0x1000–0x101F, or a full 128-byte line depending on architecture and cache configuration).

WHY? DRAM chips are organized in rows and columns. Reading one byte costs almost the same as reading a whole burst. The memory controller is optimized to fetch bursts, not individual bytes. Coalescing exploits this hardware reality regardless of the exact burst size.

Warp Execution Model

All 32 threads in a warp execute the same instruction simultaneously (SIMT). When that instruction is a memory load, the warp scheduler collects all 32 addresses and issues the minimum number of memory transactions covering them.

Best case: All 32 addresses lie in the smallest possible set of transactions → minimum transactions (e.g., 1 × 128-byte line, or 4 × 32-byte sectors).
Worst case: All 32 addresses are scattered → up to 32 transactions.

Deriving the Performance Impact

Let's quantify the cost. We fix a model architecture for the numbers below: a Volta-style sectored memory subsystem where the global-memory transaction (sector) size is $T = 32$ bytes. (On Fermi/Kepler/Pascal, replace $T=32$ with a 128-byte line and redo the arithmetic — the shape of the result is identical.)

Assume:

Warp size: $W = 32$ threads
Memory transaction (sector) size: $T = 32$ bytes (256 bits)
Element size: $E = 4$ bytes (float32)
Memory bandwidth: $B = 900$ GB/s
Memory latency: $L = 300$ ns

Coalesced Access (Optimal Pattern)

Threads access consecutive floats: arr[threadIdx.x] where threadIdx.x = 0, 1, 2, ..., 31.

Step 1: How many elements per transaction?

$\text{Elements per transaction} = \frac{T}{E} = \frac{32}{4} = 8 \text{ floats}$

Step 2: Transactions needed for the warp?

$\text{Transactions} = \frac{W}{\text{Elements per transaction}} = \frac{32}{8} = 4 \text{ transactions}$

WHY 4? With 32-byte sectors, each transaction fetches 8 consecutive floats. 32 threads need 32 floats, so $32 \div 8 = 4$ transactions. (On a 128-byte-line architecture, those same 128 bytes are 1 transaction — same data, coarser granularity.)

Step 3: Bytes fetched?

$\text{Bytes} = 4 \times 32 = 128 \text{ bytes}$

Step 4: Time to fetch?

Bandwidth-limited time (ignoring latency for sustained transfers):

$t_{\text{coalesced}} = \frac{128 \text{ bytes}}{900 \times 10^9 \text{ bytes/s}} = 0.142 \text{ ns}$

But memory latency dominates for single accesses: $t_{\text{total}} \approx 300$ ns (latency) + $0.142$ ns (transfer) $\approx 300$ ns.

Strided Access (Non-Coalesced)

Threads access with a large stride so each lands in a different transaction, e.g. arr[threadIdx.x * 8] (stride = 32 bytes = one sector) or larger:

Thread 0 → address 0x000 (sector 0)
Thread 1 → address 0x0020 (sector 1)
Thread 2 → address 0x0040 (sector 2)
...

Each thread's address falls in a different transaction/sector.

Transactions needed: 32 (one per thread, each in a different 32-byte sector).

Bytes fetched: $32 \times 32 = 1024$ bytes (but only $32 \times 4 = 128$ bytes are actually used).

$t_{\text{strided}} = 32 \times 300\text{ ns} = 9600 \text{ ns} = 9.6 \, \mu\text{s}$

Speedup from coalescing:

$\text{Speedup} = \frac{9600}{300} = 32\times$

WHY? Non-coalesced access forces 32 separate memory transactions, each with full latency overhead. You fetch 8× more data than needed and pay 32× the latency cost.

Access Patterns: Examples with Derivation

Example 1: Sequential Access (Perfect Coalescing)

__global__ void sequential(float *data) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float val = data[idx];  // Thread i reads data[i]
}

Address analysis:

Thread 0: data[0] → 0x000
Thread 1: data[1] → 0x0004 (4 bytes later)
Thread 31: data[31] → 0x007C (124 bytes from thread 0)

All 32 addresses span exactly 128 contiguous bytes: $[0x0000, 0x007F]$ .

Transactions (architecture-dependent, but always the minimum possible):

128-byte-line architecture (Fermi/Kepler/Pascal): these 128 bytes are one aligned cache line → 1 transaction.
32-byte-sector architecture (Volta+): 128 bytes = four aligned sectors → 4 transactions.

Either way, this is the minimum for reading 128 useful bytes, and 100% of the fetched bytes are used. That is what "coalesced" means — not a fixed transaction count.

WHY this works: Memory addresses are monotonically increasing with stride 4 (size of float). The controller covers all 32 requests with the fewest transactions and zero wasted bytes.

Example 2: Strided Access (Worst Case)

__global__ void strided(float *data) {
    int idx = threadIdx.x * 1024;  // Stride = 1024 floats = 4096 bytes
    float val = data[idx];
}

Address analysis:

Thread 0: 0x0000
Thread 1: 0x1000 (4096 bytes away)
Each thread's address is in a separate transaction/cache line.

Transactions: 32 (one per thread), on any architecture, because no two threads share a transaction.

Bytes fetched (32-byte-sector arch): $32 \times 32 = 1024$ bytes.
Bytes used: $32 \times 4 = 128$ bytes.
Waste: $1024 - 128 = 896$ bytes (87.5% wasted on a sectored GPU; even worse — 96.9% — on a 128-byte-line GPU that must pull a full line per thread).

WHY this fails: Stride exceeds the transaction size. No two threads share a transaction, so every thread triggers a separate one.

Example 3: Structure of Arrays vs. Array of Structures

Array of Structures (AoS) – Non-coalesced:

struct Particle {
    float x, y, z;  // 12 bytes
    float vx, vy, vz;  // 12 bytes
};
__global__ void aos(Particle *particles) {
    int idx = threadIdx.x;
    float x = particles[idx].x;  // Access x-coordinate
}

Address analysis (struct = 24 bytes):

Thread 0: particles[0].x → 0x0000
Thread 1: particles[1].x → 0x0018 (24 bytes later)
Thread 31: particles[31].x → 0x02E8 (744 bytes)

The 32 accessed x-values are spread across $32 \times 24 = 768$ bytes of address space, but only 4 useful bytes are read every 24 bytes.

Transactions (32-byte-sector arch): the touched bytes span $\lceil 768 / 32 \rceil = 24$ sectors, so ~24 transactions, of which most bytes are wasted.

Structure of Arrays (SoA) – Coalesced:

struct Particles {
    float *x;  // All x-coordinates together
    float *y, *z;
    float *vx, *vy, *vz;
};
__global__ void soa(Particles p) {
    int idx = threadIdx.x;
    float x = p.x[idx];  // Consecutive access
}

Address analysis: Same as Example 1 (sequential). Thread $i$ reads p.x[i], stride = 4 bytes.

Transactions: 4 (32-byte sectors) or 1 (128-byte line) — the minimum, 100% efficiency.

WHY SoA wins: Separating fields lets threads access homogeneous data in sequence, with no wasted bytes between the data each thread needs.

Worked example Misaligned Access

__global__ void misaligned(float *data) {
   int idx = threadIdx.x;
   float val = data[idx + 1];  // Offset by 1
}

What happens:

Thread 0: data[1] → 0x0004
Thread 31: data[32] → 0x0080 (128 bytes from base)

Addresses span $[0x0004, 0x0083]$ . On a 128-byte-line architecture this crosses two aligned 128-byte lines (0x0000–0x007F and 0x0080–0x00FF) → 2 transactions instead of 1. On a 32-byte-sector architecture it touches one extra sector → 5 sectors instead of 4.

WHY extra transactions: The offset shifts the pattern so it straddles a transaction boundary. Alignment matters: even a 1-element offset can add a transaction.

Common Mistakes and Steel-Manning

Common mistake Mistake 1: "Coalescing only matters for global memory"

Why it feels right: Documentation emphasizes global memory coalescing because that's where the biggest wins are. Shared memory is on-chip and fast, so students assume access patterns don't matter there.

The reality: Shared memory has memory banks (typically 32 banks on modern GPUs). If multiple threads in a warp access the same bank but different addresses, you get bank conflicts, serializing the access. Sequential access to shared memory often maps to conflict-free patterns, but strided access (stride = multiple of bank count) causes worst-case conflicts.

The fix: Locality principles apply everywhere. For shared memory, aim for conflict-free access: stride-1 or broadcast (all threads access the same address). Pad arrays to avoid bank conflicts.

Common mistake Mistake 2: "If my data isn't naturally sequential, I can't coalesce"

Why it feels right: Many algorithms have inherent random access patterns (hash tables, trees, sparse matrices). If the algorithm dictates the access order, you're stuck, right?

The reality: You can often reorder computation or restructure data to enable coalescing:

Tiling: Break work into blocks where each warp processes a contiguous tile.
Sort: Sort by memory address before accessing (if the access pattern is known).
SoA transformation: As shown above, data layout changes can turn scattered accesses into sequential ones.
Texture cache: For truly random patterns, texture memory (with 2D spatial locality) can outperform raw global memory.

The fix: Coalescing is a target, not a constraint. Profile first. If memory is the bottleneck and coalescing is poor, explore data layout changes or algorithmic tweaks.

Active Recall Flashcards

#flashcards/hardware

What is coalesced memory access in GPU programming?

When threads in a warp access consecutive memory addresses that can be served by the minimum number of memory transactions, exploiting the memory controller's burst-fetching behavior.

Why does the GPU memory controller fetch data in bursts rather than individual bytes?

DRAM architecture makes reading a row nearly as cheap as reading a single cell. Memory controllers are optimized for burst transfers. Fetching bursts amortizes the fixed latency cost over many bytes.

Are GPU transaction/cache-line sizes universal across all NVIDIA GPUs?

No. They depend on compute capability. Fermi/Kepler/Pascal pull full 128-byte L1 lines; Volta/Turing/Ampere use 32-byte sectors and can service partial lines. Always reason with the efficiency ratio, not a fixed count.

If 32 threads access floats with a stride large enough that each lands in a different transaction, how many transactions are needed?

Up to 32 — one per thread. No coalescing occurs because no two threads share a transaction.

What is the memory efficiency formula for GPU access patterns, and why is it architecture-independent?

Efficiency = (Bytes used) / (Bytes fetched) = (Warp size × Element size) / (Num transactions × Transaction size). It's a ratio, so it captures wasted bandwidth regardless of the absolute transaction size.

Why does Structure of Arrays (SoA) typically outperform Array of Structures (AoS) on GPUs?

SoA groups the same field for all elements contiguously, so threads accessing one field read consecutive memory with 100% efficiency. AoS interleaves fields, forcing large strides and wasted bytes per transaction.

How does a 1-element offset (e.g., data[threadIdx.x + 1]) affect coalescing?

It can add a transaction. The offset shifts the access pattern so it straddles a transaction boundary, forcing the warp to touch one extra line/sector.

Connections

6.2.01-GPU-execution-model – Warps execute in lockstep, so memory requests from all 32 threads are collected simultaneously. Coalescing depends on the SIMT model.
6.2.04-Shared-memory-and-synchronization – Locality applies to shared memory via bank conflicts. Strided shared memory access causes serialization.
6.2.09-Memory-hierarchy – Transaction and cache-line sizes (architecture-dependent) determine coalescing granularity. L1 (128 bytes) and L2 (32-byte sectors within 128-byte lines) affect transaction counts.
6.3.02-Roofline-model – Coalescing directly impacts achieved memory bandwidth. Non-coalesced code operates far below the hardware roofline.
5.1.05-Cache-coherence-protocols – GPUs don't use traditional snooping coherence, but L1/L2 caches still exploit spatial locality, which coalescing maximizes.

Recall Explain to a 12-Year-Old

Imagine you and 31 friends are standing in a line at a library, and you all need books. The librarian can carry a big cart with several books at a time. If you all ask for books that are sitting next to each other on the same shelf (book 1, book 2, book 3, ...), the librarian fills the cart in a few quick grabs and hands everyone their book. Easy!

But if each of you asks for a book from a totally different shelf across the library (book 1 from shelf A, book 2000 from shelf Z, book 5000 from shelf Q, ...), the librarian has to make a separate trip for each person — up to 32 trips. That takes way longer, even though you're still only getting one book each.

Coalesced memory access is when your GPU's "librarian" (the memory controller) can fetch all the data your 32 threads need in the fewest trips because you asked for things that are next to each other in memory. When your code asks for scattered data all over the place, the GPU has to make tons of trips, and your program slows down. The trick is to organize your data and access it in order, so the GPU can grab everything in as few trips as possible!

Generated: 2026-07-01 | Context: GPU memory performance optimization

Concept Map

Hinglish (regional understanding)

Intuition Hinglish mein samjho

Hinglish (regional understanding)

Intuition Hinglish mein samjho

Dekho, GPU memory access ka concept thoda warehouse jaisa hai. Socho ki ek warp mein 32 threads hain aur sabko data chahiye. Agar sare threads consecutive memory addresses se data mang rahe hain – matlab ek ke bad ek wale locations se – toh GPU ka memory controller smart hai, woh ek hi transaction (ek 32-byte ya 128-byte chunk) mein saara data utha lata hai. Isko bolte hain coalesced access. Lekin agar har thread kisi alag jagah se data mang raha hai, scattered addresses se, toh memory controller ko har thread ke liye alag transaction karna padta hai. Worst case mein 32 threads = 32 transactions! Yeh 32× zyada slow ho sakta hai. Yeh hardware limitation hai kyunki DRAM rows aur columns mein organized hota hai aur burst fetch karna efficient hai, individual bytes nahi.

Yeh matter kyuarta hai? Kyunki GPU programming mein tumhara bandwidth precious hai. Agar tumne array ko row-major order mein store kiya hai aur threads column-major pattern mein access kar rahe hain, ya phir random indices se scattered read kar rahe hain, toh tumhara kernel 10× ya 30× slow ho jayega, chahe computation simple ho. Real-world impact dekha jaye toh matrix transpose, image processing, particle simulations – sab jagah memory access pattern hi performance decide karta hai. Modern GPUs (Volta onwards) thoda better hain kyunki woh 32-byte sectors ko independently service karakte hain, lekin principle same hai: consecutive addresses = faster, scattered addresses = disaster. Isliye GPU code likhte waqt hamesha socho ki warp ke threads kaise memory touch kar rahe hain. Coalescing optimize karo, performance milega.

Hinglish (regional understanding)

Intuition Hinglish mein samjho

Dekho, GPU memory ka funda samajhna bahut zaroori hai kyunki yahan bottleneck aksar computation nahi, balki memory access hota hai. Core intuition yeh hai ki jab warp ke 32 threads ek saath memory se data maangte hain, toh hardware individual bytes nahi fetch karta — woh ek poora fixed-size chunk (transaction ya sector) laata hai, jaise 32 ya 128 bytes ek saath. Ab agar tumhare 32 threads consecutive addresses maang rahe hain (matlab data aas-paas pada hai), toh sab kuch minimum transactions mein aa jaata hai — yeh hai coalesced access. Lekin agar har thread alag-alag jagah se data maang raha hai, toh hardware ko alag-alag trips lagani padti hain, worst case mein 32 separate transactions. Bus warehouse wala example yaad rakho — same aisle se saman uthana easy hai, poore warehouse mein bhagna waste of time.

Yeh matter isliye karta hai kyunki DRAM ka nature hi aisa hai — ek byte padho ya poora burst, cost lagbhag same hai. Toh agar tum scattered access karte ho, toh tum bandwidth waste kar rahe ho aur speed 32x tak slow ho sakti hai. Matlab tumhara kernel technically sahi answer dega, par bahut dheere chalega. Isiliye jab tum GPU code likhte ho, toh threads ko aise arrange karna chahiye ki consecutive thread consecutive memory address access kare.

Ek important baat yaad rakho — yeh transaction aur cache-line sizes universal nahi hote, architecture pe depend karte hain. Purane GPUs (Fermi, Kepler, Pascal) 128-byte poori line fetch karte the, jabki newer ones (Volta, Ampere) 32-byte sectors independently service kar sakte hain, jisse partial access mein kam bandwidth waste hota hai. Isiliye numbers ko blindly hard-code mat karo — apne actual GPU ko profile karke check karo. Lekin core principle har jagah same rehta hai: memory ko contiguous pattern mein access karo taaki hardware apni burst-fetching capability ka poora fayda utha sake.

What Is Coalesced Memory Access?

Why GPUs Care: The Hardware Reality

Memory Transaction Granularity (Architecture-Dependent!)

Warp Execution Model

Deriving the Performance Impact

Coalesced Access (Optimal Pattern)

Strided Access (Non-Coalesced)

Access Patterns: Examples with Derivation

Example 1: Sequential Access (Perfect Coalescing)

Example 2: Strided Access (Worst Case)

Example 3: Structure of Arrays vs. Array of Structures

Common Mistakes and Steel-Manning

Active Recall Flashcards

Connections

Concept Map

Hinglish (regional understanding)

Hinglish (regional understanding)

Hinglish (regional understanding)

Go deeper — visual, from zero

Test yourself — GPU Architecture

Connections