Coalesced memory access: Foundations — Hardware deep dive

This page assumes nothing. Before you can understand why one memory pattern is 32× faster than another, you need to know what a thread, a warp, an address, a byte, a transaction, and a stride actually are — and see the picture behind each. We build them in order, each one leaning on the one before.

1. Byte — the unit everything is measured in

Picture a row of lockers. Each locker is one byte. The number written on the locker door is its position in the row — locker 0, locker 1, locker 2, and so on.

Why we need it: every quantity in this topic — how big a float is, how big a transaction is, where a thread reads — is counted in bytes. It is the ruler.

2. Address — the number on the locker door

You will see addresses written in hexadecimal, like 0x0004. That is just another way to write a normal number:

Hex	Plain number
`0x0000`	0
`0x0004`	4
`0x0020`	32
`0x0080`	128
`0x1000`	4096

Why hex? Because memory sizes are powers of two (32, 128, 4096…), and those come out as round, easy-to-spot numbers in base 16. 0x1000 = 4096 is tidy; "4096" in base 10 hides the roundness.

Question — what plain number is 0x0020? ::: 32 Question — why do GPU engineers write addresses in hex? ::: Memory sizes are powers of two, which look round and easy to align in base 16.

3. Element and its size $E$ — a value spread over several bytes

One byte only holds 0–255, which is tiny. To store a real decimal number (a float32) the computer glues 4 bytes together into one element.

So the address of data[i] is:

$address of data[i] = base + i \times E$

where base is the address of the very first element data[0].

Why we need $E$ : to turn an array index (data[7]) into an address (byte 28), which is the language the memory hardware actually speaks.

4. Thread — one worker running the code

Picture 32 identical workers, each holding a slip of paper with their number: 0, 1, 2, …, 31.

Why we need it: coalescing is entirely about which threads read which addresses at the same time. No threads, no story. (This page treats a thread simply as "one worker with an ID"; that is all you need here.)

5. Warp $W$ — 32 threads that move in perfect lock-step

This is the heart of the whole topic.

The consequence that makes coalescing possible: when the shared instruction is "read from memory," all 32 threads hit memory simultaneously. The hardware sees 32 addresses arriving together and gets to decide how to serve them cleverly — as one batch, not 32 separate errands.

Question — what does SIMT force all 32 threads to do? ::: Execute the same instruction at the same time, so their memory requests arrive together.

6. Transaction size $T$ (a.k.a. "sector") — memory only sells in fixed-size bundles

Memory hardware never hands you a single byte. It always delivers a fixed-size burst.

Aligned means the burst always starts at an address that is a whole multiple of $T$ . With $T = 32$ : transaction 0 covers bytes 0x00–0x1F, transaction 1 covers 0x20–0x3F, and so on. You cannot get a burst that starts in the middle.

Why $T$ exists: DRAM chips read whole rows at once (see 6.2.09-Memory-hierarchy). Grabbing one byte costs almost the same as grabbing a full burst, so the controller is built to move bursts. This single fact is the entire reason coalescing matters.

The base-alignment edge case

Because transactions start only at multiples of $T$ , where the array itself begins matters.

The picture: a perfectly packed run of bytes still spills into an extra burst if it does not start on a transaction boundary. This is why libraries return memory aligned to 128 or 256 bytes — to guarantee the base sits on a boundary and the minimum transaction count is achievable.

Question — why can a perfectly sequential warp still cost 5 transactions instead of 4? ::: If the array base is not a multiple of T, the 128 useful bytes straddle an extra transaction boundary.

7. Stride — the gap between one thread's address and the next

Stride is the knob that decides everything:

Small stride (= element size): consecutive threads read neighbouring lockers, so many of them land inside the same transaction. → coalesced, few bursts.
Large stride (≥ transaction size): every thread lands in a different transaction. → non-coalesced, up to 32 bursts.

Every case, spelled out (assume base aligned to $T$ , $T = 32$ , $E = 4$ , so 8 floats per transaction):

Stride	What happens	Transactions for the warp
$= E$ (4 B)	8 floats fit per transaction; 32 floats span 128 B	4 (the minimum)
$= 2 E$ (8 B)	threads step 2 floats apart; 4 threads per transaction; warp spans 256 B	8
$= 4 E$ (16 B)	2 threads per transaction; warp spans 512 B	16
$= T$ (32 B)	exactly 1 thread per transaction; each starts a fresh one	32 (worst case)
$> T$ (e.g. 4096 B)	threads even further apart, still one transaction each	32 (worst case)
stride 0 (all read same address)	one transaction serves everyone (broadcast)	1

The pattern for intermediate strides ( $E \leq stride \leq T$ , base aligned): the number of threads that share one transaction is $⌈ T / stride ⌉$ — you round up, because a thread whose element even partially falls inside a transaction still counts as landing there. Equivalently, threads-per-transaction $= ⌊ (T - 1) / stride ⌋ + 1$ . So transactions $\approx W \div ⌈ T / stride ⌉$ . As the stride grows from $E$ toward $T$ , transactions climb steadily from 4 up to 32 — coalescing degrades gradually, it is not an all-or-nothing cliff.

Notice the worst case saturates at 32: you can never need more than one transaction per thread, because there are only 32 threads.

Question — what stride gives perfect coalescing for float32? ::: Stride equal to the element size, 4 bytes (each thread reads the next element). Question — with stride 8 bytes (2 floats), how many transactions does the warp need? ::: 8 — four threads share each 32-byte transaction, so 32÷4 = 8. Question — how many threads share a transaction for stride s (T-byte transactions)? ::: ceil(T/s) — round up, because a partially-overlapping element still counts. Question — why can't a warp ever need more than 32 transactions? ::: There are only 32 threads; worst case is one transaction each.

8. Coalesced vs non-coalesced — the definition, stated plainly

Now that transactions and strides exist, we can define the term the whole topic is named after.

Note this is relative to the minimum, not a fixed number: 32 consecutive floats are coalesced whether that costs 1 transaction (128-byte-line GPU) or 4 transactions (32-byte-sector GPU) — in both cases it is the fewest bursts possible for 128 useful bytes.

Question — define coalesced access in one line. ::: The warp's addresses are served by the minimum possible number of transactions, so almost every fetched byte is used. Question — is "coalesced" a fixed transaction count? ::: No — it means the minimum for that architecture; the count differs between a 128-byte-line and a 32-byte-sector GPU.

9. Bandwidth $B$ and Latency $L$ — the two costs of a memory trip

Two different clocks tick when you touch memory. Confusing them is the classic beginner error.

This is exactly the tension the 6.3.02-Roofline-model captures.

10. $N_{transactions}$ and Efficiency — the scorecard

First, name the counter we have been using all along:

Now we can grade any pattern. Build the fraction piece by piece:

Bytes actually used — each of the $W$ threads genuinely wants one element of $E$ bytes, and no more. So useful work $= W \times E$ . With $W = 32$ , $E = 4$ that is $128$ bytes.
Bytes actually fetched — the hardware pulls $N_{transactions}$ whole bursts, each $T$ bytes, whether or not every byte is wanted. So fetched $= N_{transactions} \times T$ .

Every symbol in that fraction was built above: $W$ (§5), $E$ (§3), $T$ (§6), $N_{transactions}$ (§10). Nothing here is new — you now own the whole formula.

Question — why is "bytes used" equal to W·E? ::: Each of the W threads wants exactly one E-byte element, and nothing more. Question — why is "bytes fetched" equal to N_transactions·T? ::: The hardware always pulls whole T-byte bursts, one per transaction, wanted or not. Question — what is the efficiency of stride-32-byte access on a 32-byte-transaction GPU? ::: 12.5% — 128 useful bytes out of 1024 fetched.

How the pieces feed the topic

The map below traces the exact build order of this page: bytes give us addresses; addresses plus the element size locate each data[i]; threads grouped into a warp all read at once; the stride between their addresses and the transaction size together decide $N_{transactions}$ ; that count defines whether access is coalesced, and fed with latency/bandwidth it gives the time cost and, with used-vs-fetched bytes, the efficiency. Read each arrow as "is needed before you can understand."

Related prerequisites in the vault: 6.2.04-Shared-memory-and-synchronization (an alternative when coalescing is impossible), 6.2.09-Memory-hierarchy (where transactions live), and 5.1.05-Cache-coherence-protocols (cache-line ideas that echo here).

Equipment checklist

Cover the right side and answer out loud. If any line surprises you, reread its section before tackling the parent note.

A byte is... ::: one memory box holding a number 0–255; memory is a numbered row of them. An address is... ::: the number identifying which byte you mean. 0x0020 in plain digits is... ::: 32. An element is... ::: one logical data value the program treats as a single number, made of E bytes. Element size E for float32 is... ::: 4 bytes; data[i] starts at base + i×4. A thread is... ::: one independent worker with its own ID, threadIdx.x. A warp (size W) is... ::: 32 threads forced to run the same instruction at the same instant (SIMT). A transaction / sector (size T) is... ::: the smallest fixed, aligned burst the memory controller fetches (T=32 B on Volta-style); the two words mean the same thing here. "Aligned" means... ::: the burst starts at a multiple of T; you can't fetch one starting mid-burst. A misaligned array base can... ::: push a packed warp into one extra transaction, because the useful bytes straddle a boundary. Stride is... ::: the fixed address gap between consecutive threads' reads. Stride for perfect coalescing is... ::: equal to the element size (4 B), so threads read neighbouring lockers. Threads sharing one transaction for stride s is... ::: ceil(T/s) — round up, not floor. Coalesced access is... ::: a warp served by the minimum possible number of transactions, so nearly every fetched byte is used. Non-coalesced access is... ::: scattered addresses needing more than the minimum transactions, wasting most fetched bytes. N_transactions is... ::: how many T-byte bursts the hardware issues for one warp's read. Latency L is... ::: the wait before the first byte arrives (a representative ~300 ns; real GPUs 400–800 ns+), paid once per transaction. Bandwidth B is... ::: how fast bytes stream once flowing (~900 GB/s); does not include the wait. Non-coalescing is slow mainly because... ::: you pay latency L up to 32 times, once per scattered transaction. Efficiency equals... ::: bytes used over bytes fetched = W·E / (N_transactions·T).