GPU memory bandwidth optimization: Exercises — Hardware deep dive

This page tests everything the parent note 6.2.14 built: coalescing, banks, AoS vs SoA, tiling, arithmetic intensity, and the Roofline picture. Work each problem before opening its solution.

Before we begin, one shared vocabulary reminder in plain words, so no symbol is used unexplained:

Related tools you may lean on: Roofline-Model, CUDA-Shared-Memory, Cache-Hierarchy, GPU-Warp-Scheduling, Memory-Latency-Hiding, Parallel-Algorithm-Design.

Level 1 — Recognition

Exercise 1.1 — Count the transactions (coalesced)

A warp of 32 threads each reads one 4-byte float from consecutive addresses starting at address A. Transaction size $T = 128$ bytes. How many transactions does the warp need, and what is the efficiency $η$ ?

Recall Solution

What the addresses span: thread $i$ reads $A + 4 i$ , for $i = 0 \dots 31$ . The lowest byte touched is $A$ , the highest is $A + 4 \cdot 31 + 3 = A + 127$ . So the whole warp touches bytes $A$ through $A + 127$ — exactly 128 bytes.

Transactions: $⌈ \frac{32 \times 4}{128} ⌉ = ⌈ \frac{128}{128} ⌉ === 1 ==$ transaction.

Efficiency: useful bytes $= 32 \times 4 = 128$ ; hauled bytes $= 1 \times 128 = 128$ . $η = \frac{128}{128} === 100% == .$ This is the ideal "fully coalesced" case — one truck, one full crate, nothing wasted.

Exercise 1.2 — Which bank?

Shared memory has $N = 32$ banks and a word size of 4 bytes. A thread accesses byte address A = 260. Which bank does it land in?

Recall Solution

Why divide by word size first: banks are assigned per word, not per byte. Word index $= A /4 = 260/4 = 65$ . Apply the bank rule $Bank = (word index) mod N$ : $Bank = 65 mod 32 === 1 == .$ Word 65 = 2 full sweeps of 32 banks (words 0–63) plus 1, landing on bank 1.

Level 2 — Application

Exercise 2.1 — Strided access

Thread $i$ in a 32-thread warp reads a 4-byte float at address $A + 4 \cdot (32 i)$ — i.e. a stride of 32 floats between threads. $T = 128$ . Find transactions and efficiency.

Recall Solution

Span: thread 0 at $A$ , thread 31 at $A + 4 \cdot 32 \cdot 31 = A + 3968$ ; add the 4 bytes read $\Rightarrow$ bytes $A$ … $A + 3971$ . Each thread's float sits in its own 128-byte crate (spacing $4 \times 32 = 128$ B = exactly one crate apart). Transactions: each of the 32 threads forces a separate crate $\Rightarrow== 32 ==$ transactions. Efficiency: $η = \frac{32 \times 4}{32 \times 128} = \frac{128}{4096} === 3.125% == .$ This is the worst case in the parent note: 96.875% of hauled bytes are thrown away.

Exercise 2.2 — AoS field read

A particle struct is {float x,y,z,vx,vy,vz;} = 24 bytes per element. A warp reads the .x field: thread $i$ reads base + 24i. $T = 128$ . How many transactions per field, and $η$ ?

Recall Solution

Span of the wanted bytes: thread 0 at offset 0, thread 31 at offset $24 \times 31 = 744$ ; +4 bytes $\Rightarrow$ 748 bytes of span. Crates covering that: $⌈ 748/128 ⌉$ . Aligning to the parent's count, the address range spans $32 \times 24 = 768$ bytes of stride: $transactions = ⌈ \frac{768}{128} ⌉ === 6 == .$ Efficiency: useful $= 128$ B; hauled $= 6 \times 128 = 768$ B. $η = \frac{128}{768} === 16.67% == .$ AoS partially coalesces — 6 trips, not 32 — but wastes 5/6 of the bandwidth.

Level 3 — Analysis

Exercise 3.1 — AoS vs SoA over all six fields

Using Ex 2.2, a kernel reads all 6 fields (x,y,z,vx,vy,vz) once per warp. Compute total transactions for AoS and for SoA, and the reduction factor.

Recall Solution

AoS: each field costs 6 transactions (Ex 2.2) $\Rightarrow 6 fields \times 6 === 36 ==$ transactions. SoA: each field is its own contiguous array, thread $i$ reads base + 4i $\Rightarrow$ 1 transaction/field (Ex 1.1) $\Rightarrow 6 \times 1 === 6 ==$ transactions. Reduction: $\frac{36}{6} === 6 \times == fewer transactions .$

Exercise 3.2 — Tiling traffic and arithmetic intensity

For $N \times N$ matrix multiply tiled with $T \times T$ tiles ( $T = 32$ , 4-byte floats), derive (a) global-memory reduction factor vs naive, (b) bytes moved, (c) arithmetic intensity $I$ .

Recall Solution

(a) Reduction factor. Naive: $2 N$ global element-reads per output element. Tiled: each element enters shared memory once per tile, and there are $N / T$ tiles along $k$ , giving $2 N / T$ global reads per output element. $reduction = \frac{2 N}{2 N / T} = T === 32 \times == .$ (b) Bytes moved (tiled). $N^{2}$ outputs $\times \frac{2 N}{T}$ reads $\times 4$ bytes: $Bytes = \frac{8 N ^{3}}{T} = \frac{8 N ^{3}}{32} = \frac{N ^{3}}{4} .$ (c) Arithmetic intensity. FLOPs $= 2 N^{3}$ (one multiply + one add per inner term). $I = \frac{2 N ^{3}}{8 N ^{3} / T} = \frac{T}{4} = \frac{32}{4} === 8 == FLOPs/byte .$ See the picture below for how this point sits on the roofline.

Figure — GPU memory bandwidth optimization

Level 4 — Synthesis

Exercise 4.1 — Roofline verdict

The machine has 20 TFLOP/s peak compute and 1 TB/s peak bandwidth. Find the machine balance $I^{*}$ , then classify: naive multiply ( $I = 0.25$ ), tiled $T = 32$ ( $I = 8$ ), and a register-blocked kernel with $I = 25$ .

Recall Solution

Machine balance — the intensity where the two ceilings cross: $I^{*} = \frac{20 \times 1 0 ^{12} FLOP/s}{1 \times 1 0 ^{12} byte/s} === 20 == FLOPs/byte .$ Rule: $I < I^{*} \Rightarrow$ bandwidth-bound (memory is the ceiling); $I > I^{*} \Rightarrow$ compute-bound.

Naive $I = 0.25 < 20$ → bandwidth-bound (heavily).
Tiled $I = 8 < 20$ → still bandwidth-bound, but with 32× less traffic than naive.
Register-blocked $I = 25 > 20$ → compute-bound — memory is no longer the bottleneck. See Roofline-Model for the full graph.

Exercise 4.2 — Attainable throughput

For the tiled kernel ( $I = 8$ ) on this machine, what is the attainable performance (FLOP/s)? For the register-blocked one ( $I = 25$ )?

Recall Solution

The roofline formula: attainable $= min (peak FLOP/s, B \times I)$ . Tiled ( $I = 8$ , bandwidth-bound): $B \times I = 1 \times 1 0^{12} \times 8 === 8 \times 1 0^{12} === 8 TFLOP/s .$ Since $8 < 20$ , the memory slope wins — 8 TFLOP/s, i.e. 40% of peak. Register-blocked ( $I = 25$ ): $B \times I = 25 \times 1 0^{12}$ , but capped by peak $\Rightarrow min (20, 25) \times 1 0^{12} === 20 \times 1 0^{12} === 20$ TFLOP/s (full peak).

Level 5 — Mastery

Exercise 5.1 — Diagnose and fix a bank conflict

A tile __shared__ float tile[32][32] is stored, then read column-wise: thread $i$ (a whole warp varying the row) reads tile[i][col] for fixed col. Word size 4 B, 32 banks. (a) Which bank does thread $i$ hit? (b) How many cycles does the warp take? (c) Propose a one-line fix and give the new cycle count.

Recall Solution

(a) Bank of tile[i][col]. In a [32][32] array the linear word index is $32 i + col$ . $Bank = (32 i + col) mod 32 = col mod 32.$ Every thread $i$ hits the same bank = $col mod 32$ , because the $32 i$ term is a multiple of 32 and vanishes under the modulo. (b) Cycles. 32 threads, one bank, all different addresses (different rows) $\Rightarrow$ a full 32-way conflict $\Rightarrow== 32 ==$ cycles (vs 1 ideal). Not a broadcast: broadcast needs the same address, here rows differ. (c) Fix — pad the row. Declare tile[32][33]. Now the linear index is $33 i + col$ : $Bank = (33 i + col) mod 32 = (i + col) mod 32,$ which is distinct for each $i$ $\Rightarrow$ 32 different banks $\Rightarrow== 1 ==$ cycle. One wasted padding column removes the 32× serialization.

Exercise 5.2 — End-to-end throughput budget

Machine: $B = 2$ TB/s, peak 20 TFLOP/s. A kernel is fully coalesced ( $η_{co a l esce} = 1$ ), cache hit rate $η_{c a c h e} = 0.9$ , occupancy $η_{occ u p an cy} = 0.8$ . (a) Effective throughput $M$ . (b) If its arithmetic intensity is $I = 5$ FLOP/byte, what compute rate does that $M$ support, and is it compute- or bandwidth-bound?

Recall Solution

(a) Combine the efficiencies (they multiply, since each independently trims usable bandwidth): $η_{t o t a l} = 1 \times 0.9 \times 0.8 = 0.72.$ $M = B \times η_{t o t a l} = 2 \times 1 0^{12} \times 0.72 === 1.44 \times 1 0^{12} == byte/s = 1.44 TB/s .$ (b) Compute rate from that memory feed = $M \times I = 1.44 \times 1 0^{12} \times 5 === 7.2 \times 1 0^{12} === 7.2$ TFLOP/s. Compare to peak 20 TFLOP/s: $7.2 < 20 \Rightarrow$ bandwidth-bound — the memory feed, not the ALUs, limits us.

Recall Self-test one-liners

Fully coalesced warp of 32 floats needs how many transactions? ::: 1 AoS 24-byte struct, reading one field, transactions per warp? ::: 6 (efficiency 16.67%) Reduction in global traffic from $T \times T$ tiling? ::: A factor of $T$ Arithmetic intensity of tiled matmul, $T = 32$ , floats? ::: 8 FLOPs/byte Machine balance for 20 TFLOP/s and 1 TB/s? ::: 20 FLOPs/byte Fix for a 32-way shared-memory column conflict? ::: Pad the tile row to width 33