GPU Architecture
Chapter: 6.2 GPU Architecture Difficulty: Level 5 — Mastery (cross-domain: performance modeling, architecture, and CUDA coding) Time limit: 90 minutes Total marks: 60
Instructions: Answer all three questions. Show all reasoning, derivations, and code. State every assumption explicitly. Calculators permitted.
Question 1 — Roofline Modeling & Memory-Bound Analysis (22 marks)
A GPU has the following specifications:
- Peak FP32 throughput:
- Peak memory bandwidth:
- 108 Streaming Multiprocessors (SMs), each capable of resident threads
- Warp size threads
(a) Define arithmetic intensity and derive the ridge point (in FLOP/byte) at which the roofline transitions from memory-bound to compute-bound. Compute its numeric value. (4)
(b) A kernel performs a SAXPY operation over single-precision elements. Counting reads and writes to global memory, derive its arithmetic intensity and classify it (memory- vs compute-bound) using the ridge point from (a). Then compute the achievable throughput in GFLOP/s. (6)
(c) A dense matrix–matrix multiply with all matrices (FP32, naive global-memory implementation, no tiling) reads and once per output element. Derive as a function of , and find the smallest (power of 2) for which the kernel becomes compute-bound. (6)
(d) Explain, using the occupancy concept, why a memory-bound kernel can still approach peak bandwidth even though each individual memory access has ~500-cycle latency. Give the quantitative condition (Little's Law form) linking required in-flight requests to latency and throughput. (6)
Question 2 — SIMT Execution, Warp Divergence & Coalescing (20 marks)
(a) Consider the CUDA kernel fragment executed by a single warp (32 threads, tid = 0..31):
if (tid % 2 == 0) {
// Branch A: 20 instructions
} else {
// Branch B: 12 instructions
}
// Branch C (common): 8 instructionsUnder classic (pre-Volta) SIMT lock-step execution, derive the total number of instruction-issue cycles for this warp (assume 1 cycle/instruction). Compare against the ideal (no divergence) cost and compute the divergence penalty factor. (6)
(b) Prove that if the branch condition were if (tid < 16) instead, the total instruction count is identical but the effective SIMT utilization during the divergent region differs from part (a). Compute the average lane-utilization (fraction of active lanes averaged over issued cycles) for both conditions and explain the architectural reason they are equal despite different masks. (6)
(c) A kernel accesses global memory as float val = data[stride * tid];. For a warp of 32 threads and a 128-byte cache line (32 floats), derive the number of memory transactions as a function of stride for stride = 1, 2, 32. Explain coalescing and quantify the bandwidth efficiency for each stride. (8)
Question 3 — Shared Memory Bank Conflicts & Tiled GEMM (18 marks)
Shared memory on this GPU has 32 banks, each 4 bytes wide, with successive 4-byte words mapped to successive banks.
(a) A thread block loads a tile into __shared__ float tile[32][32]; and threads in a warp access tile[tid][0] (column access). Derive the number of bank conflicts (the "way" of the conflict) and the resulting slowdown versus conflict-free access. Then show that padding to tile[32][33] eliminates the conflict — prove it via the bank-index arithmetic. (9)
(b) For a tiled GEMM with tile width , each output tile requires loading and sub-tiles from global memory into shared memory. Derive the reduction in global memory traffic relative to the naive version of Q1(c) as a function of , and state the arithmetic-intensity improvement factor. Given shared memory budget of 48 KB per SM and the need to hold two FP32 tiles simultaneously, find the maximum (power of 2) and its resulting intensity gain. (9)
End of paper.
Answer keyMark scheme & solutions
Question 1
(a) Ridge point (4 marks)
Arithmetic intensity (FLOP/byte). (1)
Roofline: attainable performance . The ridge point is where memory bound equals compute bound: (1)
B \cdot I_{\text{ridge}} = P_{\text{fp32}} \;\Rightarrow\; I_{\text{ridge}} = \frac{P_{\text{fp32}}}{B} = \frac{19.5\times10^{12}}{1555\times10^{9}} \approx 12.54\ \text{FLOP/byte}. \tag{2}
(b) SAXPY (6 marks)
Per element: 1 multiply + 1 add = 2 FLOPs. (1) Memory traffic per element: read (4 B), read (4 B), write (4 B) = 12 bytes. (2)
I_{\text{saxpy}} = \frac{2}{12} = 0.1\overline{6}\ \text{FLOP/byte}. \tag{1}
Since , the kernel is strongly memory-bound. (1)
Achievable throughput FLOP/s GFLOP/s (≈1.3% of peak). (1)
(c) Naive GEMM intensity (6 marks)
Total FLOPs (each of outputs needs MACs = FLOPs). (1) Naive traffic: for each of outputs, read a full row of ( floats) and column of ( floats) = reads, plus 1 write. Bytes . (2)
I(n) = \frac{2n^3}{4(2n^3+n^2)} = \frac{2n^3}{8n^3+4n^2} = \frac{n}{4n+2} \xrightarrow{n\to\infty} \frac14 = 0.25\ \text{FLOP/byte}. \tag{2}
The asymptote : naive GEMM never reaches compute-bound. (1) So no finite makes it compute-bound — the correct mastery answer is "none: naive un-tiled GEMM saturates at 0.25 FLOP/byte, far below the ridge point; tiling is required (see Q3)."
(d) Latency hiding & Little's Law (6 marks)
Occupancy = ratio of active resident warps to max supported. With enough resident warps, the SM issues instructions from ready warps while others wait on memory, so the memory pipeline stays saturated. (2)
Little's Law: required concurrency (in-flight bytes) . (2)
Numeric illustration: at ~1.4 GHz, 500 cycles ≈ 357 ns; of in-flight data must be requested to hide latency. Only high occupancy supplies enough independent warps/requests to reach this; low occupancy leaves the memory system idle. (2)
Question 2
(a) Divergence penalty (6 marks)
Even lanes take Branch A (16 threads), odd take Branch B (16 threads). Under lock-step SIMT the two branches serialize: (2)
\text{cycles} = \underbrace{20}_{A} + \underbrace{12}_{B} + \underbrace{8}_{C} = 40\ \text{cycles}. \tag{1}
Ideal (no divergence, all lanes agree): the branch region would cost the larger single path plus common cycles (best case if all took A) — but the meaningful ideal is running each branch's work only once with full lanes. Penalty relative to non-divergent equivalent (28): (1)
\text{penalty factor} = \frac{40}{28} \approx 1.43. \tag{2}
(b) Utilization equivalence (6 marks)
For tid % 2: A active lanes = 16, B active = 16, C = 32. For tid < 16: A = 16, B = 16, C = 32 — same active-lane counts per region. (2)
Total instruction cycles identical: in both. (1)
Average lane utilization (active-lane-cycles / (32 × total cycles)): U = \frac{16\cdot20 + 16\cdot12 + 32\cdot8}{32\cdot40} = \frac{320+192+256}{1280} = \frac{768}{1280} = 0.6 = 60\%. \tag{2}
Architectural reason: SIMT utilization depends only on how many lanes are active per issued instruction (the popcount of the mask), not which lanes — divergence cost is set by mask cardinality per branch, so both partitions (interleaved vs contiguous) yield identical serialization and identical utilization. (1)
(c) Coalescing (8 marks)
A warp requests 32 floats; a 128-byte (=32-float) transaction serves consecutive words. Transactions = distinct 128-byte segments touched. (2)
- stride = 1: addresses → all within one 128-byte segment → 1 transaction. Efficiency = requested/fetched = 128/128 = 100%. (2)
- stride = 2: addresses → span 64 floats = 256 bytes = 2 segments → 2 transactions. Useful 128 B out of 256 B fetched → efficiency 50%. (2)
- stride = 32: each thread hits a different 128-byte segment → 32 transactions; each fetches 128 B but uses only 4 B → efficiency . (2)
Coalescing = hardware merging of a warp's accesses that fall in the same aligned segment into one transaction; strided/scattered access defeats it and wastes bandwidth.
Question 3
(a) Column access bank conflict (9 marks)
tile[32][32] is row-major: element tile[i][0] is at linear index floats. Bank index for all . (3)
All 32 threads (accessing tile[tid][0]) map to bank 0 → a 32-way bank conflict → 32 serialized accesses → 32× slowdown vs conflict-free. (2)
With padding tile[32][33]: tile[i][0] is at linear index . Bank index (since ). (3)
For these give — all distinct banks → conflict-free, 1 cycle. Padding by one column shifts the stride so consecutive rows land in consecutive banks. (1)
(b) Tiled GEMM traffic (9 marks)
Naive traffic (Q1c): bytes ( float reads). With tiles, the block loads each element of and from global memory only once per tile-row/column rather than once per output. Global traffic scales as floats: each output tile reuses loaded data times. (3)
Reduction factor in traffic . (1)
Arithmetic intensity improves by factor : I_{\text{tiled}} \approx \frac{2n^3}{4\cdot(2n^3/T)} = \frac{T}{4}\ \text{FLOP/byte}. \tag{2}
Shared memory budget: hold two FP32 tiles simultaneously (-tile + -tile): 2\times T^2 \times 4\ \text{B} \le 48\times1024\ \text{B} \;\Rightarrow\; T^2 \le 6144 \;\Rightarrow\; T \le 78.4. \tag{2}
Largest power of 2: . Resulting intensity FLOP/byte → the tiled kernel becomes compute-bound. Intensity gain factor over naive baseline (0.25→16). (1)
[
{"claim":"Ridge point ~12.54 FLOP/byte", "code":"P=Rational(195,10)*10**12; B=1555*10**9; result = abs(float(P/B)-12.5401)<0.01"},
{"claim":"SAXPY intensity = 1/6", "code":"result = Rational(2,12)==Rational(1,6)"},
{"claim":"SAXPY achievable throughput ~259 GFLOP/s", "code":"B=1555*10**9; I=Rational(1,6); T=B*I; result = abs(float(T)/1e9-259.17)<1.0"},
{"claim":"Naive GEMM intensity asymptote 0.25", "code":"n=symbols('n',positive=True); I=n/(4*n+2); result = limit(I,n,oo)==Rational(1,4)"},
{"claim":"Divergence penalty factor 40/28", "code":"result = abs(float(Rational(40,28))-1.42857)<0.001"},
{"claim":"SIMT utilization = 0.6", "code":"result = Rational(16*20+16*12+32*8,32*40)==Rational(3,5)"},
{"claim":"Padded bank index distinct: 33 mod 32 = 1", "code":"result = (33 % 32)==1 and len({ (33*i)%32 for i in range(32)})==32"},
{"claim":"Unpadded column all bank 0", "code":"result = len({ (32*i)%32 for i in range(32)})==1"},
{"claim":"Max tile T=64 fits 48KB for two tiles", "code":"result = (2*64**2*4 <= 48*1024) and (2*128**2*4 > 48*1024)"},
{"claim":"Tiled intensity T/4 at T=64 exceeds ridge", "code":"result = Rational(64,4)==16 and 16 > 12.54"}
]