Worked examples — Coalesced memory access

This deep-dive lives under the parent topic. There we learned what coalescing is. Here we grind through every kind of access pattern a warp can produce — perfect, offset, strided, mixed, degenerate, and a real word problem — so you never meet an access pattern in the wild that you haven't already solved on paper.

---

Ground rules for this whole page

So that every number is checkable, we fix one model architecture for all examples (exactly the parent's Volta-style model):

Everything reduces to "how many distinct sectors do the 32 addresses land in?" Call that number $N_{\text{tx}}$. That single count decides transactions, bytes fetched, and efficiency.

---

The address-ramp metaphor (our "regime map")

Before the matrix, one picture ties every example together. Plot each thread index $t$ on the horizontal axis and its byte address $A(t)$ on the vertical axis. Every access pattern is just a line (or scatter) of dots — an address ramp. The slope of that ramp is the stride, and the slope alone decides how coalesced you are. This is our stand-in for the "which quadrant / which sign" case-splitting you do with angles: instead of quadrants we have slope regimes, drawn as labelled bands in the figure.

The four slope regimes are exactly our case split (example IDs are introduced fully in the matrix below and the worked sections):

• Slope $=0$ (flat): broadcast — Example 7.
• Slope $=E$ (gentle, unbroken): coalesced — Examples 1, 6; broken by a shift in Example 2.
• Slope strictly between $E$ and $T$ (medium): partial waste — Examples 3, 9.
• Slope $\ge T$ (steep): one transaction per thread, the floor — Examples 4, 5, and the column-access half of Example 8 (Example 8a, read below before relying on this bullet).
• No single slope (scatter): mixed/irregular — Example 8b.

---

Where the numbers come from (T = 32 B is not arbitrary)

---

The scenario matrix

Every access pattern a warp can throw at you falls into one of these cells, listed in cell order A → I. The examples that follow cover every cell. The last two cells are the ones students underestimate: Cell I is the "exam twist" — reading one field of a struct { float x, y, z; } (an array-of-structs, AoS), where a hidden 12-byte stride wrecks coalescing — which we build in full at Example 9. Cell J is a mixed warp with no single stride, built at Example 8b.

Cell	Pattern class	What's special about it	Covered by
A	Perfectly sequential	stride = 1 element, aligned start	Example 1
B	Sequential but misaligned start	consecutive, offset by non-multiple of $T$	Example 2
C	Small stride (2 elements)	partial waste, some sharing	Example 3
D	Stride = one sector exactly	boundary case, zero sharing	Example 4
E	Huge stride (worst case)	one transaction per thread	Example 5
F	Reversed / permuted order	same addresses, shuffled threads	Example 6
G	Degenerate: all threads same address	broadcast, minimum possible	Example 7
H	Real-world word problem (matrix rows vs cols)	why layout choice matters	Example 8a
I	Exam twist: struct of 3 floats (AoS)	element size ≠ 4, hidden stride	Example 9
J	Mixed / irregular warp (no single stride)	half sequential, half scattered	Example 8b

Each row maps onto one slope regime of the ramp picture above: increasing vs decreasing (Cell F), aligned vs shifted (Cell B), zero-slope (Cell G), the steep limiting case (Cell E), and no-single-slope scatter (Cell J).

---

Example 1 — Cell A: Perfect sequential (the baseline)

Forecast first: Guess the sector count before reading on. (Hint: 32 floats = how many bytes? divide by 32.)

1. Write the address function. $A(t)=4t$ bytes (each float is 4 bytes, stride 1). Why this step? Everything downstream is just $\lfloor A(t)/32\rfloor$; we need $A(t)$ explicitly first.
2. Find the sector of each thread. $\text{sector}(t)=\lfloor 4t/32\rfloor =\lfloor t/8\rfloor$. Why this step? This turns 32 addresses into 32 sector numbers — the only thing that matters.
3. Count distinct sectors. As $t$ runs $0\to 31$, $\lfloor t/8\rfloor$ takes values $0,1,2,3$. So $N_{\text{tx}}=4$. Why this step? $N_{\text{tx}}$ is our master number.
4. Efficiency. $\frac{128}{32\times 4}=\frac{128}{128}=1.0=100\%$. Why this step? Confirms zero wasted bytes — the definition of coalesced.

Figure walkthrough: In the map above, the 32 coral bars (thread reads) sit exactly on top of the 4 mint sectors with no gaps and no grey (unused) sectors. That "coral perfectly tiles mint" is what 100% efficiency looks like.

Verify: 4 sectors × 32 bytes = 128 bytes fetched = 128 bytes used. Addresses span 0x00–0x7F = exactly 128 bytes, exactly 4 aligned sectors. Units: bytes/bytes = dimensionless efficiency. ✓

---

Example 2 — Cell B: Sequential but misaligned start

Forecast: Same data size (128 bytes) — but does shifting the window across sector boundaries cost an extra transaction?

1. Address function. $A(t)=8+4t$. Why? The offset 8 is not a multiple of $T=32$, so the 128-byte window straddles one extra sector boundary.
2. Sectors touched. Smallest address $A(0)=8\Rightarrow$ sector $\lfloor 8/32\rfloor =0$. Largest $A(31)=8+124=132\Rightarrow$ sector $\lfloor 132/32\rfloor =4$. Why? The span now reaches from sector 0 up to sector 4.
3. Count distinct sectors. Sectors $0,1,2,3,4$ are all touched $\Rightarrow N_{\text{tx}}=5$. Why? The 128-byte window slid partly into a 5th sector; we must fetch all of it.
4. Efficiency. $\frac{128}{32\times 5}=\frac{128}{160}=0.8=80\%$. Why? We fetched 160 bytes but only used 128 — the misalignment cost us 20%.

Figure walkthrough: Compare with Example 1's map. The coral bars have slid right by 8 bytes, so the first two bars now poke into sector 0 while the last bar spills into a fifth sector (sector 4, drawn grey-mint with mostly unused space). That single extra half-empty sector is the 20% you lost.

Verify: span = 0x08–0x84 (132 = 0x84). $\lfloor 132/32\rfloor -\lfloor 8/32\rfloor +1=4-0+1=5$ sectors. Efficiency $128/160=0.8$. ✓

---

Example 3 — Cell C: Small stride (2 elements)

Forecast: We now skip every other float. We use 128 bytes but they're spread over twice the span. Guess $N_{\text{tx}}$.

1. Address function. $A(t)=4\times 2t=8t$. Why? Stride 2 floats = 8 bytes between consecutive threads.
2. Sectors. $\text{sector}(t)=\lfloor 8t/32\rfloor =\lfloor t/4\rfloor$, running $0,1,\dots ,7$ as $t:0\to 31$. Why? Now only 4 threads share each sector (not 8), so we spill into more sectors.
3. Count. $\lfloor t/4\rfloor$ takes $\mathrm{0..7}\Rightarrow N_{\text{tx}}=8$. Why? Distinct sector count is what sets transactions.
4. Efficiency. $\frac{128}{32\times 8}=\frac{128}{256}=0.5=50\%$. Why? Half of every fetched sector is data we skipped.

Verify: span 0x00–0xF8 (248 bytes); 8 sectors. Used 128 of 256 bytes = 50%. ✓

---

Example 4 — Cell D: Stride = one sector exactly

Forecast: Each thread jumps exactly one sector. Do any two threads ever share a sector?

1. Address function. $A(t)=4\times 8t=32t$. Why? This is the tipping point where stride equals $T$.
2. Sectors. $\text{sector}(t)=\lfloor 32t/32\rfloor =t$ — a different sector for every thread. Why? Because stride equals $T$ exactly, no two threads can ever land in the same sector.
3. Count. Sectors $0,1,\dots ,31\Rightarrow N_{\text{tx}}=32$. Why? 32 distinct sectors is the master number here.
4. Efficiency. $\frac{128}{32\times 32}=\frac{128}{1024}=0.125=12.5\%$. Why? Only 4 of every 32 fetched bytes are used.

Verify: 32 distinct sectors, 1024 bytes fetched, 128 used → 12.5%. This is the parent note's exact "strided" number. ✓

---

Example 5 — Cell E: Huge stride (worst case, the limiting value)

Forecast: Once stride ≥ $T$, adding more stride can't hurt further — you're already at one transaction per thread. Confirm.

1. Address function. $A(t)=4\times 1024t=4096t$. Why? We test whether extra stride beyond $T$ makes things worse.
2. Sectors. $\text{sector}(t)=\lfloor 4096t/32\rfloor =128t$ — sectors $0,128,256,\dots$ all distinct. Why? Any stride $\ge T$ guarantees 32 distinct sectors. Beyond that, extra stride only spreads them out; the count saturates at 32.
3. Count. $N_{\text{tx}}=32$ (the maximum possible for a 32-thread warp). Why? A 32-thread warp can touch at most 32 distinct sectors — this is that ceiling.
4. Efficiency. $\frac{128}{1024}=12.5\%$ — identical to Example 4. Why? Same transaction count → same efficiency; the floor is reached.

Verify: $N_{\text{tx}}=32$, efficiency $=0.125$, matching Example 4 → the floor is confirmed. ✓

---

Example 6 — Cell F: Reversed / permuted order

Forecast: Does the order threads visit addresses matter, or only the set of addresses?

1. Address function. $A(t)=4(31-t)$. Thread 0 reads 0x7C, thread 31 reads 0x00. Why? We want to test whether coalescing cares about permutation.
2. Set of addresses. As $t$ ranges $\mathrm{0..31}$, $31-t$ ranges $\mathrm{31..0}$ — the same set $\{0,4,\dots ,124\}$ as Example 1. Why? Coalescing counts distinct sectors touched by the warp, not who touches them.
3. Sectors. Exactly $\{0,1,2,3\}$, so $N_{\text{tx}}=4$. Why? Same set of sectors as Example 1 → same count.
4. Efficiency. $100\%$ — identical to Example 1. Why? Same $N_{\text{tx}}$ → same efficiency.

Recall Does thread-to-address permutation change coalescing?

Permutation within the warp is free — only the set of sectors matters ::: Correct. Reverse, shuffle, or interleave — as long as the 32 addresses cover the same sectors, $N_{\text{tx}}$ is unchanged. The hardware collects addresses before issuing transactions.

Verify: same sector set $\{0,1,2,3\}$ as Example 1 → $N_{\text{tx}}=4$, 100%. ✓

---

Example 7 — Cell G: Degenerate — all threads read the same address

Forecast: 32 threads, one address. What's the minimum possible $N_{\text{tx}}$?

1. Address function. $A(t)=0$, constant (zero slope — the degenerate "flat" ramp). Why? This is the boundary case where the address function has stride 0.
2. Sectors. $\text{sector}(t)=\lfloor 0/32\rfloor =0$ for all threads → a single sector $\{0\}$. Why? All 32 addresses are identical, so they cannot spread across more than one sector.
3. Count. $N_{\text{tx}}=1$. The hardware fetches sector 0 once and broadcasts the value to all 32 threads. Why? One distinct sector → one transaction, the absolute minimum a warp load can cost.
4. What efficiency means here (the broadcast footnote). Our page-wide numerator "bytes used $=W\cdot E=128$" assumes 32 threads want 32 distinct elements. In a broadcast they want the same element, so there are really only $E=4$ distinct useful bytes, not 128. Plugging the fixed formula literally would give $\frac{128}{32\times 1}=4.0=400\%$ — a nonsense >100% number that only appears because the numerator over-counts identical requests. So for this one degenerate cell we report the honest operational headline instead: $N_{\text{tx}}=1$, the minimum possible transaction count, with the single fetched sector reused by all 32 threads. Broadcast is the best-behaved pattern, not a 400%-efficient one. Why this step? Without this footnote a reader would trust the $400\%$ figure or think the model broke; naming the numerator's assumption (distinct elements) is exactly what resolves the paradox.

Verify: single distinct sector → $N_{\text{tx}}=1$, the minimum. Distinct useful bytes $=4$; literal formula would give the flagged nonsense $128/32=4.0$, which is why we report $N_{\text{tx}}=1$ instead. ✓

---

Example 8a — Cell H: Real-world word problem (row-major matrix)

Forecast: Which layout traversal — down a column or across a row — feeds the warp consecutive addresses?

1. Column access address function. $A(t)=4\times (t\cdot 1024)=4096t$ bytes. Why? Row-major means consecutive rows are 1024 floats = 4096 bytes apart.
2. Column sectors. $\lfloor 4096t/32\rfloor =128t$ — all distinct → $N_{\text{tx}}=32$, efficiency $12.5\%$. This is Example 5 in disguise! Why? A column in row-major storage has stride = row length ≫ sector → worst case. (This is the column-access case referenced in the slope-regime bullets above.)
3. Row access address function. $A(t)=4t$ → identical to Example 1 → $N_{\text{tx}}=4$, $100\%$. Why? The last index varying with the thread gives stride 1 → perfect coalescing.
4. Design conclusion. For a row-major array, make consecutive threads read consecutive columns (the last index varies with threadIdx.x). To process columns efficiently, first transpose in shared memory. Why this step? Steps 2 and 3 proved the row traversal costs $8\times$ fewer transactions than the column traversal; a step that decides what to do with that measured fact is what turns the arithmetic into an engineering rule — otherwise the derivation ends without an actionable takeaway.

Verify: column: $N_{\text{tx}}=32$, 12.5%; row: $N_{\text{tx}}=4$, 100%. Speedup row-over-column $=32/4=8\times$ in transactions. ✓

---

Example 8b — Cell J: Mixed / irregular warp (no single stride)

Forecast: The warp is half nice, half nasty. Do the good 16 rescue the bad 16, or does the total just add up?

1. Address functions for the two halves.
   • Threads $\mathrm{0..15}$: $A(t)=4t$ → bytes 0,4,...,60.
   • Threads $\mathrm{16..31}$: $A(t)=4(1000+64(t-16))=4000+256(t-16)$ → bytes 4000, 4256, 4512, .... Why? A mixed warp has no closed-form single slope, so we must count each sub-pattern separately, then union the sectors.
2. Sectors from the contiguous half. Bytes 0..60 → sectors $\lfloor 0/32\rfloor ..\lfloor 60/32\rfloor =0,1$. That is 2 sectors. Why? 16 floats = 64 bytes span exactly two aligned sectors — the good half is locally coalesced.
3. Sectors from the scattered half. $\text{sector}(t)=\lfloor (4000+256(t-16))/32\rfloor =125+8(t-16)$ for $t-16=\mathrm{0..15}$ → sectors $125,133,141,\dots$ all distinct → 16 sectors. Why? Stride here is 256 bytes $\ge T$, so each of the 16 lands in its own sector (Example 5 behaviour, locally).
4. Union and count. The two sector sets are disjoint (0–1 vs 125–245), so $N_{\text{tx}}=2+16=18$. Why? Total transactions = distinct sectors over the whole warp; the good half does not rescue the bad half, but it does keep its own cost low.
5. Efficiency. $\frac{128}{32\times 18}=\frac{128}{576}\approx 0.222=22.2\%$. Why? Only the 16 contiguous reads were dense; the 16 scattered reads dragged efficiency down but not all the way to the 12.5% floor.

Verify: contiguous half → 2 sectors; scattered half → 16 sectors; disjoint union → 18; efficiency $128/576\approx 0.2222$. Between the 12.5% floor and 100% ceiling, as a mixed pattern must be. ✓

---

Example 9 — Cell I: Exam twist — Array of Structs (hidden stride)

Forecast: The struct is 12 bytes, so the x fields are 12 bytes apart — not 4. Does that pack neatly into 32-byte sectors?

1. Address function. $A(t)=12t$. Threads read at 0, 12, 24, 36, .... Why? The struct size becomes the stride when you touch one field per thread.
2. Sectors touched. $\text{sector}(t)=\lfloor 12t/32\rfloor$. Evaluate the span: $A(31)=372\Rightarrow$ sector $\lfloor 372/32\rfloor =11$. Smallest is sector 0. Are all of $\mathrm{0..11}$ hit? Step $12/32=0.375$ sector per thread, so consecutive threads never skip a whole sector → yes, sectors $0,1,\dots ,11$ all touched. Why? We must confirm no sector in the range is skipped before counting.
3. Count. $N_{\text{tx}}=12$. Why? 12 distinct sectors is the master number.
4. Efficiency. $\frac{128}{32\times 12}=\frac{128}{384}\approx 0.333=33.3\%$. Why? We fetched the y and z fields we never used — 2/3 of every sector wasted.

Verify: span 0–372 bytes → sectors 0..11 = 12 sectors; efficiency $128/384=1/3$. ✓

---

One-glance summary

Recall Fill in

$N_{\text{tx}}$ for each example (our fixed $T=32$B model) Ex1 seq → 4 · Ex2 misaligned → 5 · Ex3 stride2 → 8 · Ex4 stride8 → 32 · Ex5 huge → 32 · Ex6 reversed → 4 · Ex7 broadcast → 1 · Ex8a column → 32 / row → 4 · Ex8b mixed → 18 · Ex9 AoS → 12 ::: Memorise the shape: gentle stride = few transactions; stride ≥ sector = 32 (the floor); mixed = sum of parts.

Recall Why is efficiency capped at 12.5% for any large uniform stride, never lower?

Because a warp has only 32 threads, so at most 32 distinct sectors ::: With one transaction per thread you fetch $32T$ bytes and use $32E$; the ratio $E/T=4/32=12.5\%$ is a hard floor for a single-element-per-thread uniform-stride pattern.

Prerequisites & neighbours: 6.2.01-GPU-execution-model (warps & SIMT), 6.2.09-Memory-hierarchy (where sectors and L2 live, and why caching softens but never fixes a bad pattern), 6.2.04-Shared-memory-and-synchronization (the transpose fix), 6.3.02-Roofline-model (why bandwidth matters), 5.1.05-Cache-coherence-protocols (line/sector granularity origins), and the Hinglish walkthrough 6.2.08 Coalesced memory access (Hinglish).