Exercises — Memory hierarchy (global, shared, registers)

What we do: read the latency column of the parent's summary table and sort ascending.

Answer: registers → shared memory → L2 cache → global memory.

Why this order: speed comes from physical proximity. Registers sit inside the arithmetic unit; global memory is off-chip DRAM (VRAM) that light has to physically travel to.

• Registers → per-thread (private; no other thread can read them).
• Shared memory → per-block (every thread in the same block shares it).
• Global memory → per-device (every thread in every block sees it).

Why: scope and speed go together. Private things are fast (no coordination); shared things need arbitration (banks); global things need a long trip. Widening the audience costs time.

Who manages each level:

• Registers → compiler (automatic; the compiler assigns your local variables to registers without you asking).
• Shared memory → programmer (explicit; you write __shared__ and fill it yourself — that is why it's called a "user-controlled cache").
• Global memory → programmer (explicit; you allocate it and choose every read and write).

The claim is False. The compiler manages registers, not shared memory. Shared and global memory are both programmer-controlled; only registers are automatic.

What / Why: each thread carries away $r$ registers from the fixed pot $R_{total}$, so we divide the pot by the per-thread cost and round down with the floor operator (you cannot run a fractional thread). $N_{threads}=\lfloor \frac{65,536}{48}\rfloor =\lfloor 1365.33\rfloor =1365\text{ threads (idealized budget bound).}$ Warp rounding: hardware allocates in chunks of 32, so $N_{threads}^{\text{hw}}=\lfloor \frac{1365}{32}\rfloor \times 32=42\times 32=1344\text{ threads.}$ Answer: 1365 threads by the plain-floor budget bound, 1344 threads once warp-granularity rounding is applied — the value the SM actually schedules.

Idealized budget bound: $\text{Occupancy}=\frac{1365}{2048}=0.6665\dots \approx 66.7\%.$ Warp-rounded (real hardware): ${\text{Occupancy}}^{\text{hw}}=\frac{1344}{2048}=0.65625=65.6\%.$ Why occupancy matters: GPUs hide the 400–800 cycle global-memory wait by switching to another ready warp (see Memory Latency Hiding). Fewer resident threads = fewer warps to switch to = the SM sits idle waiting. The honest hardware number here is ~65.6%, so roughly two-thirds of the hiding power is available.

What / Why: we first convert the byte address to a word index (divide by width $w$, because banking is per 4-byte word), then take mod $B$ to see which bank cycles around to it. $\text{bank}(260)=\left(\frac{260}{4}\right)\mathrm{mod}32=65\mathrm{mod}32=1.$ Answer: bank 1.

(a) Coalesced (and aligned). Assume $\text{base}$ is a multiple of 128 (aligned). Addresses span $\text{base}+0$ to $\text{base}+124$, a total of $32\times 4=128$ bytes = exactly one 128-byte segment. $N_{transactions}=\lceil 128/128\rceil =1.$ (b) Strided. Each thread's address is 128 bytes apart, so every thread lands in its own segment. $N_{transactions}=32.$ (c) Misaligned base — the hidden edge case. Suppose $\text{base}=64$ (not a multiple of 128). The 128 requested bytes now run from byte 64 to byte 191, straddling the boundary at 128. That touches two segments (0–127 and 128–255): $N_{transactions}=2.$ So even a perfectly consecutive access pays double the transactions if its starting address is misaligned. The lesson: coalescing needs both a consecutive pattern and an aligned base — this is why CUDA allocators return 128/256-byte-aligned pointers.

Analysis: the aligned coalesced-vs-strided ratio is $32/1=32\times$; misalignment quietly turns your best case into a $2\times$ penalty. Same math, same 32 floats wanted — but pattern (b) drags 32 times more bytes across the bus, wasting 31/32 of every fetch. This is the single biggest lever in Memory Bandwidth Optimization.

How to read this figure: the top row of green dots shows the coalesced aligned case — 32 thread requests packed into a single shaded green 128-byte segment (one transaction). The middle row of yellow dots shows the misaligned case — the same consecutive requests straddle the boundary line and spill into a second segment (two transactions). The bottom row of red dots shows the strided case — each request lands in its own separate shaded red segment, so the warp needs many transactions. The horizontal axis is the byte address; the vertical dotted white lines mark 128-byte segment boundaries.

Kernel A: $N=\lfloor 65536/32\rfloor =2048$, capped at $N_{max}=2048$ → occupancy $=2048/2048=100\%$. Kernel B: $N=\lfloor 65536/64\rfloor =1024$ → occupancy $=1024/2048=50\%$.

Analysis: doubling registers halved occupancy. Kernel B has only half as many warps to hide latency with. But — subtlety — B might still win if those extra registers eliminate memory traffic (fewer trips to global). This is the classic register-pressure vs. occupancy tradeoff: you never optimize occupancy in isolation, only alongside the memory traffic it buys you.

Warp Y (thread $t$ → word $t$): words $\mathrm{0..31}$ map to banks $\mathrm{0..31}$ — all distinct → $S=1$. No conflict, full speed.

Warp X (all → word 5): naïvely 32 threads hit one bank, so $\text{bank}(a)=\text{count}=32$, suggesting $S=32$. But shared memory has a broadcast rule: when many threads read the same address, hardware broadcasts it in one shot → $S=1$.

Answer: both have $S=1$, but for different reasons — Y because addresses spread across banks, X because identical addresses trigger broadcast. The dangerous case is in between: several threads hitting the same bank at different addresses (that truly serializes).

(a) $16\times 16=256$ floats $\times 4$ bytes $=1024$ bytes = 1 KB per tile. (b) Two tiles (As and Bs): $2\times 1024=2048$ bytes = 2 KB per block. (c) $\lfloor 96\text{KB}/2\text{KB}\rfloor =\lfloor 98304/2048\rfloor =48$ blocks. Synthesis point: shared-memory footprint is a second occupancy limiter alongside registers. The true block count is the minimum across the register limit, the shared-memory limit, and the hardware block cap.

(a) With [32][32], column $j$ of row $i$ sits at word index $32i+j$. When threads read a column during transpose (varying $i$, fixed $j$), their indices differ by 32 → $(32i+j)\mathrm{mod}32=j$ for all $i$ → all 32 hit the same bank → 32-way conflict. Padding to 33 makes column stride 33, and $\gcd (33,32)=1$, so consecutive column elements land in 32 different banks → conflict-free. (b) Extra column = $32\text{ rows}\times 1\text{ float}\times 4=128$ bytes per block. Synthesis point: a 128-byte "waste" buys a 32× speedup on the transpose step — the cheapest trade on this page.

$N=2048$, tile $=16$, tiles per dimension $=2048/16=128$.

(a) Naive: each of $N^2=2048^2=4,194,304$ output threads reads $2N=4096$ values from global. $4,194,304\times 4096=17,179,869,184\approx 1.72\times 10^{10}\text{ reads.}$ (b) Tiled: each block loads, per k-step, one A-tile + one B-tile $=2\times 256=512$ elements, over $128$ k-steps $=512\times 128=65,536$ reads per block. Number of blocks $=128\times 128=16,384$. $65,536\times 16,384=1,073,741,824\approx 1.07\times 10^9\text{ reads.}$ (c) Reduction factor: $\frac{17,179,869,184}{1,073,741,824}=16\times .$ Explanation: the reduction factor equals the tile edge length (here 16). Why exactly 16? In the naive version each output element re-reads every input value it needs straight from global; in the tiled version, once a tile is loaded into shared memory, all 16 threads sharing that row/column reuse it 16 times before it is evicted. Each loaded byte is therefore amortized over 16 uses, cutting global traffic by that same factor of 16.

Synthesis point: bigger tiles = more reuse (larger reduction factor), but bigger tiles cost more shared memory (Exercise 4.1) and more registers — the recurring three-way tension between reuse, occupancy, and capacity.

(a) $I_{\star }=P/\beta =\frac{10\times 10^{12}}{500\times 10^9}=20$ FLOP/byte. (b) Naive intensity $I=2\text{FLOP}/8\text{B}=0.25$ FLOP/byte. Since $0.25<20=I_{\star }$, the kernel is far to the left of the ridge → memory-bound. It wastes almost all its compute waiting for data. (c) Tiling reuses each loaded byte 16×, so bytes-from-global drop 16× while FLOPs stay the same → intensity rises 16×: $I_{new}=0.25\times 16=4$ FLOP/byte. Still $4<20$, so still memory-bound, but now runs $16\times$ faster on the memory roofline. Mastery point: tiling moves you rightward along the roofline. To cross the ridge to compute-bound you'd need even larger tiles / register blocking to push $I$ past 20.

How to read this figure: the axes are log–log — arithmetic intensity $I$ (FLOP/byte) across the bottom, attainable performance (FLOP/s) up the side. The white line is the roofline: a rising blue-dashed part on the left (memory-bound — performance limited by bandwidth $\beta \cdot I$) and a flat yellow-dashed ceiling on the right (compute-bound — performance capped at $P$). They meet at the dotted vertical line, the ridge $I_{\star }=20$. The red dot ($I=0.25$, naive) sits far left under the sloped part; the green dot ($I=4$, tiled) has slid rightward toward the ridge but is still under the slope — still memory-bound, just faster.

Pick (ii), shared-memory tiling. The kernel is memory-bound with already-coalesced access, so the bus is used efficiently but too much data is being fetched — the cure is reuse, not more coalescing. Lever (i) risks lowering occupancy (Exercise 3.2) and only helps if data is reused within a single thread; lever (iii) doesn't reduce total bytes moved at all. Tiling raises arithmetic intensity (Exercise 5.1), sliding you rightward on the roofline — the only lever here that attacks total global traffic, which is the actual bottleneck.

(a) $1,048,576\times 4=4,194,304$ bytes $\approx 4$ MB of useful data. (b) $t=\frac{4,194,304\text{B}}{500\times 10^9\text{B/s}}=8.388608\times 10^{-6}\text{s}\approx 8.39\mu \text{s}.$ (c) Strided access fetches 32× more bytes (whole segments for single floats), so effective transferred bytes $\approx 32\times$ and time scales up: $8.39\mu \text{s}\times 32\approx 268.4\mu \text{s}$. Mastery point: the coalesced floor is a hard physical limit — you cannot go faster than bandwidth allows. The 32× penalty is entirely self-inflicted by the access pattern, which is why coalescing is checked first in any optimization pass.