Coalesced memory access: Exercises — Hardware deep dive

Before we start, one shared model so every number is reproducible. This is the same model architecture the parent note fixed:

Definition Our fixed model (unless a problem says otherwise)

Warp size $W = 32$ threads (all execute the same instruction at the same time, the SIMT rule from 6.2.01-GPU-execution-model).
Transaction (sector) size $T = 32$ bytes — the smallest chunk DRAM hands back.
Element size $E = 4$ bytes (a float32).
Memory latency $L = 300$ ns per transaction; bandwidth $B = 900$ GB/s.

Two formulas we will reuse. What they say in words: the first counts distinct sectors the warp forces the controller to fetch; the second is the fraction of fetched bytes you actually keep. $N_{transactions} = {distinct T -byte aligned sectors touched by any thread of the warp}$ $Efficiency = \frac{bytes used}{bytes fetched} = \frac{W \times E}{N _{transactions} \times T}$ Read $N_{transactions}$ literally: take every thread's byte address, round it down to a multiple of $T$ (that names its sector), and count how many different sector-names appear. That count — not the number of threads — is what the memory controller must fetch.

Reading the two-extremes figure

What the picture shows. On the left (coalesced), eight black threads all point up into a single red 32-byte sector — the red box is the one object the controller fetches, and all 8 useful floats live inside it, so no byte is wasted. On the right (strided), four black threads each point into their own red sector: four boxes, one thread apiece, so most of each fetched sector is thrown away. Use this as your mental test: coalesced = many arrows converging on few red boxes; strided = one arrow per red box.

L1 — Recognition

Problem 1.1

A warp runs float v = data[threadIdx.x];. Thread $i$ reads element $i$ . Is this access coalesced or non-coalesced?

Recall Solution

What we check: do consecutive threads hit consecutive addresses? Thread $i$ reads data[i], whose byte address is $i \times E = 4 i$ . So thread 0 → byte 0, thread 1 → byte 4, …, thread 31 → byte 124. The 32 addresses fill the contiguous block $[0, 127]$ with no gaps. Round each address down to a multiple of $T = 32$ : they name sectors ${0, 32, 64, 96}$ — exactly 4 distinct sectors, all fully used. By the definition, efficiency $= 1$ → coalesced (perfectly).

Problem 1.2

Another warp runs float v = data[threadIdx.x * 1024];. Is this coalesced?

Recall Solution

Thread $i$ reads data[1024 i], byte address $4096 i$ . Thread 0 → byte 0, thread 1 → byte 4096. Rounding down to multiples of 32 gives 32 different sector-names (0, 4096, 8192, …). $N_{transactions} = 32$ , one per thread. Efficiency $= 128/ (32 \times 32) = 12.5%$ → non-coalesced (worst case).

L2 — Application

Problem 2.1

For the coalesced pattern data[threadIdx.x] (32 floats), compute (a) elements per transaction, (b) number of transactions, (c) bytes fetched.

Recall Solution

(a) A sector is $T = 32$ bytes and holds $T / E = 32/4 = 8$ floats — this is why one transaction covers 8 threads' worth of contiguous data. (b) The warp needs 32 contiguous floats; packed 8 to a sector they tile into $N = W \div (T / E) = 32/8 = 4$ sectors (the division is valid because the run is contiguous and starts aligned — see Problem 2.4 for the misaligned case). (c) Bytes fetched $= N \times T = 4 \times 32 = 128$ bytes. Every one of those 128 bytes is used, so nothing is wasted.

Problem 2.2

Compute the memory efficiency for the strided pattern data[threadIdx.x * 8] (byte stride $= 8 \times 4 = 32$ bytes $= T$ ).

Recall Solution

Where each thread lands: thread $i$ → byte $32 i$ . Rounding down to multiples of 32 gives sector-names ${0, 32, 64, \dots}$ — 32 distinct names. So $N = 32$ transactions. Bytes used $= W \times E = 32 \times 4 = 128$ . Bytes fetched $= N \times T = 32 \times 32 = 1024$ . $Efficiency = \frac{128}{1024} = 0.125 = 12.5%.$ Only 1 of every 8 fetched bytes is kept.

Problem 2.3

Using latency $L = 300$ ns per transaction, estimate the fetch time for (a) the coalesced warp (4 transactions issued together) and (b) the strided warp (32 transactions each in a different sector). Then give the speedup. State the hardware mechanism that lets the coalesced transactions overlap.

Recall Solution

(a) Coalesced: the 4 sectors are the four aligned quarters of one 128-byte line requested by a single warp-load instruction. The memory controller has multiple outstanding-request slots (miss-status/handling registers), so it launches all 4 sector fetches back-to-back without waiting for the first to return. Their 300 ns latencies therefore overlap into one $\approx L = 300$ ns bubble rather than adding up: $t_{coal} \approx 300$ ns. The overlap works precisely because the requests are issued simultaneously by one instruction and there are enough outstanding slots to hold them all. (b) Strided: 32 sectors in 32 unrelated locations. In the worst case they exceed the outstanding-slot budget and serialize, each paying full latency: $t_{strided} = 32 \times 300 = 9600$ ns $= 9.6 μ s$ . Speedup $= 9600/300 = 32 \times$ .

Problem 2.4 — Misaligned base address

The same contiguous pattern data[threadIdx.x], but the array's first byte is at address 0x0010 (16 bytes into a sector), not on a 32-byte boundary. How many transactions now, and what is the efficiency?

Recall Solution

The 32 floats occupy bytes $[16, 143]$ . Round each address down to a multiple of 32: the touched sector-names are ${0, 32, 64, 96, 128}$ — 5 sectors, because the block now straddles a boundary at both ends (it dips into sector 0's upper half and into sector 128's lower half). $N = 5$ , bytes used $= 128$ , bytes fetched $= 5 \times 32 = 160$ . $Efficiency = \frac{128}{160} = 0.8 = 80%.$ Lesson: misalignment costs one extra transaction and drops efficiency below 100% even though the access is contiguous. This is why CUDA allocators return sector-aligned pointers, and why cudaMalloc guarantees alignment — align your base address and the naive division $N = W / (T / E)$ is exact again.

L3 — Analysis

Problem 3.1

A kernel reads a matrix stored row-major as A[row * N + col]. Compare two mappings:

Mapping X: threadIdx.x → col (threads sweep across a row).
Mapping Y: threadIdx.x → row (threads sweep down a column).

Which mapping coalesces, and why?

Recall Solution

Row-major means element A[row * N + col] sits at byte $(r o w \cdot N + co l) \cdot E$ . Neighbouring columns are 4 bytes apart; neighbouring rows are $N \cdot 4$ bytes apart.

Mapping X (thread $i$ → col = i, fixed row): addresses $4 i$ apart → contiguous → distinct sectors ${0, 32, 64, 96}$ → coalesced, 4 transactions.
Mapping Y (thread $i$ → row = i, fixed col): addresses $4 N i$ apart. For any realistic $N$ this stride $≫ T$ → 32 distinct sector-names → non-coalesced, 32 transactions.

The rule: map threadIdx.x to the axis that is contiguous in memory (the last index in row-major).

Problem 3.2

On a 128-byte-line GPU (Fermi/Kepler/Pascal), the coalesced 32-float read is 1 transaction; on a 32-byte-sector GPU (Volta+) it is 4 transactions. Both read the same 128 bytes with 100% efficiency. Which architecture wastes fewer bytes on the strided data[threadIdx.x * 1024] pattern, and why does sectoring help there?

Recall Solution

Strided, 128-byte-line GPU: each thread pulls a full 128-byte line it barely uses. Fetched $= 32 \times 128 = 4096$ bytes, used $= 128$ . Efficiency $= 128/4096 = 3.125%$ (waste 96.9%). Strided, 32-byte-sector GPU: each thread pulls only its 32-byte sector. Fetched $= 32 \times 32 = 1024$ , used $= 128$ . Efficiency $= 128/1024 = 12.5%$ (waste 87.5%). Sectoring wins because it fetches at a finer granularity — a scattered access drags in only 32 bytes, not 128. See 6.2.09-Memory-hierarchy for where sectors live.

L4 — Synthesis

Problem 4.1 — Array-of-Structs vs Struct-of-Arrays

A particle has struct P { float x, y, z, w; } (16 bytes). A warp reads only the x field of 32 particles.

AoS: P *p; float xi = p[threadIdx.x].x; — the 32 x values are 16 bytes apart.
SoA: float *px; float xi = px[threadIdx.x]; — the 32 x values are 4 bytes apart.

Compute efficiency for each (32-byte-sector model). Which layout do you choose?

Recall Solution

SoA: neighbouring x are 4 bytes apart → contiguous 128-byte block → distinct sectors ${0, 32, 64, 96}$ → $N = 4$ . Efficiency $= (32 \times 4) / (4 \times 32) = 128/128 = 100%$ . AoS: thread $i$ 's x sits at byte $16 i$ . Round each down to a multiple of 32 to name its sector: $thread i 012345 byte 16 i 01632486480 sector ⌊ 16 i /32 ⌋ \cdot 32 0032326464$ The pattern is clear: two consecutive threads share one sector (bytes $0$ – $15$ and $16$ – $31$ both live in sector 0), so the 32 threads name $32/2 = 16$ distinct sectors → $N = 16$ . Efficiency $= (32 \times 4) / (16 \times 32) = 128/512 = 25%$ . Choose SoA — it turns a strided field-access into a contiguous stream. This is the canonical GPU data-layout transform.

Problem 4.2

Continuing 4.1, if the kernel instead reads all four fields x, y, z, w of 32 particles, does AoS suddenly become fine? Reason it through.

Recall Solution

Reading all four fields means the warp touches all $32 \times 16 = 512$ contiguous bytes of p[0..31]. That is $512/32 = 16$ sectors, all fully used: efficiency $= 512/ (16 \times 32) = 512/512 = 100%$ . So AoS is efficient exactly when you consume the whole struct; it is wasteful only when you cherry-pick one field. The layout choice depends on the access pattern, not on AoS/SoA in the abstract.

L5 — Mastery

Problem 5.1 — Transpose via shared memory

A naive matrix transpose B[col*N+row] = A[row*N+col] has a coalesced read of A but a strided write to B (neighbouring threads write $N \cdot 4$ bytes apart). Explain how a shared-memory tile (from 6.2.04-Shared-memory-and-synchronization) fixes the write, and state the efficiency of the global write before and after.

Recall Solution

Before: the write B[col*N+row] has stride $4 N$ per thread → 32 distinct sectors, one per thread → efficiency $128/ (32 \times 32) = 12.5%$ . The fix: each block reads a tile of A coalesced (row-contiguous) into a shared-memory buffer tile[ ][ ], __syncthreads(), then writes the tile back to B reading tile transposed but writing global memory row-contiguous. Because the transpose now happens inside on-chip shared memory (which has no coalescing penalty), both the global read and the global write become contiguous. After: the global write is contiguous → 4 sectors, all used → efficiency $128/128 = 100%$ . The strided cost is paid in fast shared memory, not slow DRAM.

Problem 5.2 — Latency-bound vs bandwidth-bound

You measure a kernel at 12.5% efficiency and 32 transactions per warp. You refactor to a coalesced layout achieving 4 transactions per warp at 100% efficiency. If each transaction is latency-bound at 300 ns and the kernel issues 1,000,000 warp-loads, estimate the wall-clock memory time before and after, and the speedup. Then explain when this latency model stops applying.

Recall Solution

Before (latency-bound, serialized): $32 transactions \times 300 ns \times 1 0^{6} = 9.6 \times 1 0^{9}$ ns $= 9.6$ s. After: the 4 sectors pipeline into one latency $\approx 300$ ns per warp: $300 ns \times 1 0^{6} = 3.0 \times 1 0^{8}$ ns $= 0.3$ s. Speedup $= 9.6/0.3 = 32 \times$ . When the latency model stops applying: the two regimes meet when the data volume is large enough that transfer time overtakes latency. Compare the two times: latency-bound time is $N_{transactions} \times L$ ; bandwidth-bound time is $(bytes fetched) / B$ . The access is bandwidth-bound once $\frac{bytes fetched}{B} ≳ N_{transactions} \times L \times (concurrency hiding factor) .$ In practice, launching many concurrent warps lets the scheduler hide the 300 ns latency behind other warps' work; once enough warps are in flight the GPU saturates $B$ and you should model with bandwidth, not per-transaction latency. Coalescing helps in both regimes: it lowers $N_{transactions}$ (latency-bound) and raises the fraction of $B$ spent on useful bytes (bandwidth-bound). Locate your kernel on the 6.3.02-Roofline-model to know which term dominates.

Recall Quick self-check clozes

A warp of 32 threads reading data[threadIdx.x] (aligned) needs 4 transactions on a 32-byte-sector GPU. $N_{transactions}$ equals the number of distinct 32-byte aligned sectors touched by any thread. Coalescing is judged by memory efficiency (bytes used / bytes fetched), not an absolute transaction count. A contiguous read starting at a misaligned base address needs one extra transaction. The layout transform that makes single-field reads coalesced is Array-of-Structs → Struct-of-Arrays (SoA).

Read this in Hinglish: 6.2.08 Coalesced memory access (Hinglish).