Coalesced memory access: Question bank — Hardware deep dive

Before we start, three words we will lean on the whole page — each was built in the parent note, but let's re-anchor them so nothing is used before it's clear. (Throughout this vault we wrap key terms in ==double equals==, which Obsidian renders as a highlight and turns into a flashcard cloze — treat every highlighted phrase as "memorise this exact wording".)

Definition The three load-bearing words

Warp — a fixed bundle of 32 threads that execute the same instruction at the same instant. This execution style is called SIMT (Single Instruction, Multiple Threads): one instruction is broadcast, and all 32 threads run it in lockstep on their own data. See 6.2.01-GPU-execution-model.
Transaction (= sector) — the fixed-size burst of bytes the memory controller always fetches, even if you asked for one byte. The two words mean the same thing here: a transaction is the request the controller issues, and on a sectored GPU the unit it fetches is called a sector — so "transaction" and "sector" are interchangeable names for one indivisible burst. On a Volta+ GPU that burst is 32 bytes; on Fermi/Kepler/Pascal it is a 128-byte line.
Coalesced — the warp's 32 addresses fall inside the fewest possible transactions with near-zero wasted bytes. It is a property of the pattern relative to the transaction size, not a fixed transaction count.

True or false — justify

True or false: "Coalesced" means all 32 threads read from exactly one memory transaction.

False. On a 32-byte-sector GPU, 32 floats (128 bytes) span four aligned sectors, yet this is still perfectly coalesced. Coalescing means minimum transactions with no waste, not one transaction.

True or false: If threads read consecutive addresses, the access is always coalesced regardless of the starting address.

False. Alignment matters — if the 128-byte span starts mid-transaction, it straddles an extra sector, so you pay for one more transaction than the ideal. Consecutive but misaligned is worse than consecutive and aligned.

True or false: Reading only 4 useful bytes per thread but fetching 32-byte sectors always wastes bandwidth.

False when coalesced. If neighbouring threads consume the other bytes of the same sector, nothing is wasted — 8 threads share one 32-byte sector and use all 32 bytes. Waste only appears when threads scatter.

True or false: A stride of exactly 2 (data[idx*2]) is just as bad as a stride of 1024.

False, but it's still degraded. Stride-2 uses half of each sector, so efficiency is ~50%, not the ~12.5% of a large stride. The badness scales smoothly with stride until each thread lands in its own transaction.

True or false: Coalescing is a software optimisation you turn on with a compiler flag.

False. It is an emergent property of your access pattern. The hardware coalesces automatically; your job is to write indices so the addresses land contiguously.

True or false: On a modern sectored (Volta+) GPU, non-coalesced access is no longer a problem.

False. Sectoring softens the penalty (32-byte waste per stray thread instead of 128-byte), but a fully scattered warp still issues up to 32 transactions and pays ~32× the latency. It's better, not solved.

True or false: If a kernel is latency-bound, coalescing doesn't matter.

Mostly false. Non-coalescing creates extra latency-bound transactions (up to 32 separate round trips), so poor coalescing is often why you're latency-bound in the first place. See 6.3.02-Roofline-model.

True or false: Coalescing and cache hits are the same thing.

False. Coalescing is about one warp's addresses fitting into the fewest transactions right now. Caching is about reuse across time. A coalesced load can still miss the cache; a cache hit can still come from a scattered pattern.

True or false: Two threads reading the exact same address break coalescing.

False. Identical addresses fall in the same transaction, so the hardware serves both from one fetch (a broadcast). This is cheaper than a stride, not worse.

Spot the error

"data[idx * 4] is coalesced because floats are 4 bytes."

Error: multiplying the index by 4 gives a stride of 4 floats = 16 bytes between threads, leaving 3 of every 4 floats unread per sector. The 4-byte float size is already baked into indexing; you must not multiply by it again.

"To make row-major matrix access coalesced, thread i should read M[i][0]."

Error: M[i][0] steps down a column in row-major storage, so consecutive threads are one full row apart — a huge stride. Coalesced row-major access is thread i reading M[row][i], walking along a row.

"Transaction size is 128 bytes, so coalescing means each warp does exactly 1 transaction."

Error: it conflates a representative number with a universal law. The transaction size is architecture-dependent (128-byte line vs 32-byte sector), and even at 128 bytes a warp of 32 floats is exactly one line only when aligned. See the 6.2.09-Memory-hierarchy.

"My access is strided but I added a __shared__ buffer, so now it's coalesced."

Error: shared memory doesn't make global access coalesced — you use it to reorganise. You coalesce the global load into contiguous addresses, stage it in shared memory, then let threads read the scattered pattern from shared memory (which has different rules — bank conflicts, not coalescing).

"Efficiency = bytes used / bytes fetched, and my kernel uses 128 of 1024 bytes, so efficiency is fine at 12.5%."

Error: 12.5% is terrible, not fine — it means 87.5% of your memory bandwidth is thrown away. The formula is right; the interpretation inverts good and bad.

"Threads reading data[idx] and data[idx+1] overlap, so it's double-counted and slow."

Error: overlapping reads within the same sector are free — the sector is already fetched. Reuse within a transaction is the good case, not a slowdown.

Why questions

Why does the hardware fetch a whole 32-byte sector when a thread asks for 4 bytes?

Because of how DRAM physically works: a read first activates an entire row of the chip into a row buffer (the slow step), then streams out consecutive columns from that buffer in a burst. The costly part — opening the row — is already paid, so pulling 32 contiguous bytes costs almost nothing extra over pulling 4. The controller therefore always transfers a full sector-sized burst rather than a single element; asking for 4 bytes still opens the row and streams the whole burst.

Why does a large stride cause up to 32 transactions instead of just a few?

When the stride exceeds the transaction size, no two threads' addresses share a sector, so every one of the 32 threads forces its own separate fetch — the worst case.

Why is misalignment (offset by a few bytes) enough to cost an extra transaction?

A contiguous 128-byte block that starts in the middle of a sector spills into the next sector, so it now touches one more aligned transaction than a block that starts on a sector boundary.

Why is the speedup often quoted as "up to 32×" and not a fixed number?

32× is the theoretical ceiling (1 transaction vs 32). Real speedup depends on stride, alignment, cache behaviour, and whether you're bandwidth- or latency-bound, so it lands anywhere from ~1× to ~32×.

Why does the Array-of-Structs (AoS) layout hurt coalescing compared to Struct-of-Arrays (SoA)?

In AoS, thread i reading field.x of struct i steps by the whole struct size, scattering the x values across many sectors. SoA stores all x values contiguously, so the same access is a clean stride-1 read.

Why can coalescing matter even when the total data moved is identical between two kernels?

Because scattered access inflates the bytes actually fetched (fetching full sectors to use 4 bytes each). Same useful data, but far more fetched data and more latency-bound transactions.

Why doesn't fixing coalescing help a kernel that is purely compute-bound?

If arithmetic, not memory, is the bottleneck, the memory subsystem already has slack. Improving coalescing frees bandwidth that wasn't the limiter — the roofline tells you which side you're on.

Edge cases

Edge case: A warp where all 32 threads read the same single address.

Fully coalesced (in fact, one transaction with a broadcast). Zero stride is the friendliest pattern, not a degenerate failure.

Edge case: A warp active with only 8 of 32 threads (a partial/predicated warp) reading contiguous floats.

Still coalesced — those 8 addresses (32 bytes) fit one sector. Inactive threads issue no request, so they neither help nor hurt the transaction count.

Edge case: Stride exactly equal to the transaction size (idx * 8 floats = 32 bytes).

Worst case — each thread lands at the start of a distinct sector, so 32 sectors are fetched and only 4 of each 32 bytes is used (≈12.5% efficiency).

Edge case: Reading char (1 byte) instead of float, contiguously.

32 threads read 32 bytes = exactly one 32-byte sector, still fully coalesced. Smaller elements pack more threads per sector, which is fine as long as the pattern stays contiguous.

Edge case: A base pointer misaligned by 8 bytes but otherwise stride-1 access.

The 128-byte span now straddles five 32-byte sectors instead of four, costing one extra transaction. Aligning allocations (e.g. cudaMalloc returns aligned memory) restores the ideal.

Edge case: The array is smaller than one warp's worth of floats (only 5 elements).

The 5 in-bounds threads read one sector coalesced; the out-of-bounds threads must be masked off (bounds check) so they issue no illegal transaction — otherwise you get an out-of-bounds fault, not a coalescing issue.

Recall Fast self-test

Which pattern is coalesced: data[idx], data[idx*2], or data[idx*1024]? ::: Only data[idx] — stride-1 contiguous. The others waste half, then almost all, of each fetched sector. Does a 32-byte-sector GPU make scattered access acceptable? ::: No — it caps the waste per stray thread at 32 bytes, but a fully scattered warp still issues up to 32 latency-heavy transactions. What single property defines "coalesced"? ::: Fewest possible transactions with near-zero wasted bytes for the warp's 32 addresses — not any fixed transaction count.

Related: Coalesced memory access · 6.2.09-Memory-hierarchy · 5.1.05-Cache-coherence-protocols · 6.3.02-Roofline-model