Exercises — GPU memory bandwidth optimization
6.2.14 · D4· Hardware › GPU Architecture › GPU memory bandwidth optimization
Is page mein woh sab kuch test hota hai jo parent note 6.2.14 ne banaya tha: coalescing, banks, AoS vs SoA, tiling, arithmetic intensity, aur Roofline picture. Har problem ko uski solution kholne se pehle solve karo.
Shuru karne se pehle, ek shared vocabulary reminder plain words mein, taaki koi bhi symbol bina explanation ke use na ho:
Related tools jinhe tum use kar sakte ho: Roofline-Model, CUDA-Shared-Memory, Cache-Hierarchy, GPU-Warp-Scheduling, Memory-Latency-Hiding, Parallel-Algorithm-Design.
Level 1 — Recognition
Exercise 1.1 — Transactions count karo (coalesced)
32 threads ke ek warp mein se har thread ek 4-byte float padhta hai consecutive addresses se, jo address A se shuru hote hain. Transaction size bytes hai. Warp ko kitne transactions chahiye, aur efficiency kya hai?
Recall Solution
Addresses kitna span karte hain: thread reads , ke liye. Lowest byte touched hai , highest hai . Toh poora warp bytes se lekar tak touch karta hai — exactly 128 bytes.
Transactions: transaction.
Efficiency: useful bytes ; hauled bytes . Yeh ideal "fully coalesced" case hai — ek truck, ek full crate, kuch bhi waste nahi.
Exercise 1.2 — Kaun sa bank?
Shared memory mein banks hain aur word size 4 bytes hai. Ek thread byte address A = 260 access karta hai. Yeh kis bank mein padega?
Recall Solution
Pehle word size se kyun divide karein: banks word ke hisaab se assign hote hain, byte ke hisaab se nahi. Word index . Bank rule apply karo : Word 65 = 32 banks ke 2 full sweeps (words 0–63) plus 1, jo bank 1 par land karta hai.
Level 2 — Application
Exercise 2.1 — Strided access
32-thread warp mein thread ek 4-byte float padhta hai address par — yaani threads ke beech 32 floats ka stride. . Transactions aur efficiency nikalo.
Recall Solution
Span: thread 0 at , thread 31 at ; 4 bytes read add karo bytes … . Har thread ka float apne khud ke 128-byte crate mein hai (spacing B = exactly ek crate ka antar). Transactions: 32 threads mein se har ek ek alag crate force karta hai transactions. Efficiency: Yeh parent note ka worst case hai: haule gaye 96.875% bytes phek diye jaate hain.
Exercise 2.2 — AoS field read
Ek particle struct hai {float x,y,z,vx,vy,vz;} = 24 bytes per element. Ek warp .x field padhta hai: thread reads base + 24i. . Har field ke liye kitne transactions, aur ?
Recall Solution
Chaahiye gaye bytes ka span: thread 0 offset 0 par, thread 31 offset par; +4 bytes 748 bytes ka span. Us span ko cover karne wale crates: . Parent ke count ke saath align karte hue, address range bytes of stride span karta hai: Efficiency: useful B; hauled B. AoS partially coalesce karta hai — 6 trips, 32 nahi — lekin bandwidth ka 5/6 waste karta hai.
Level 3 — Analysis
Exercise 3.1 — AoS vs SoA saare six fields ke liye
Ex 2.2 use karte hue, ek kernel saare 6 fields (x,y,z,vx,vy,vz) ek baar per warp padhta hai. AoS aur SoA ke liye total transactions compute karo, aur reduction factor bhi.
Recall Solution
AoS: har field ko 6 transactions lagte hain (Ex 2.2) transactions.
SoA: har field apna contiguous array hai, thread reads base + 4i 1 transaction/field (Ex 1.1) transactions.
Reduction:
Exercise 3.2 — Tiling traffic aur arithmetic intensity
matrix multiply ke liye tiles se tiled (, 4-byte floats), derive karo (a) naive ke muqable global-memory reduction factor, (b) bytes moved, (c) arithmetic intensity .
Recall Solution
(a) Reduction factor. Naive: global element-reads per output element. Tiled: har element shared memory mein ek baar per tile enter karta hai, aur ke saath tiles hain, jo global reads per output element deta hai. (b) Bytes moved (tiled). outputs reads bytes: (c) Arithmetic intensity. FLOPs (ek multiply + ek add per inner term). Neeche di gayi picture mein dekho ki yeh point roofline par kahan baithta hai.

Level 4 — Synthesis
Exercise 4.1 — Roofline verdict
Machine mein 20 TFLOP/s peak compute aur 1 TB/s peak bandwidth hai. Machine balance nikalo, phir classify karo: naive multiply (), tiled (), aur ek register-blocked kernel jisme hai.
Recall Solution
Machine balance — woh intensity jahan dono ceilings cross karti hain: Rule: bandwidth-bound (memory ceiling hai); compute-bound.
- Naive → bandwidth-bound (heavily).
- Tiled → abhi bhi bandwidth-bound, lekin naive se 32× kam traffic ke saath.
- Register-blocked → compute-bound — memory ab bottleneck nahi hai. Full graph ke liye Roofline-Model dekho.
Exercise 4.2 — Attainable throughput
Is machine par tiled kernel () ke liye attainable performance (FLOP/s) kya hai? Register-blocked wale () ke liye?
Recall Solution
Roofline formula: attainable . Tiled (, bandwidth-bound): Kyunki , memory slope jeet jaata hai — 8 TFLOP/s, yaani peak ka 40%. Register-blocked (): , lekin peak se cap hoga TFLOP/s (full peak).
Level 5 — Mastery
Exercise 5.1 — Bank conflict diagnose aur fix karo
Ek tile __shared__ float tile[32][32] store ki gayi hai, phir column-wise padhi jaati hai: thread (ek poora warp row vary karta hua) tile[i][col] padhta hai fixed col ke liye. Word size 4 B, 32 banks. (a) Thread kis bank ko hit karta hai? (b) Warp kitne cycles leta hai? (c) Ek one-line fix suggest karo aur naaya cycle count do.
Recall Solution
(a) tile[i][col] ka bank. [32][32] array mein linear word index hai.
Har thread ek hi bank = hit karta hai, kyunki term 32 ka multiple hai aur modulo ke under vanish ho jaata hai.
(b) Cycles. 32 threads, ek bank, sab alag addresses (alag rows) full 32-way conflict cycles (1 ideal ke muqable). Yeh broadcast nahi hai: broadcast ke liye same address chahiye, yahan rows differ karti hain.
(c) Fix — row ko pad karo. tile[32][33] declare karo. Ab linear index hai:
jo har ke liye distinct hai 32 alag banks cycle. Ek wasted padding column 32× serialization hata deta hai.

Exercise 5.2 — End-to-end throughput budget
Machine: TB/s, peak 20 TFLOP/s. Ek kernel fully coalesced hai (), cache hit rate , occupancy . (a) Effective throughput . (b) Agar uski arithmetic intensity FLOP/byte hai, toh woh kitna compute rate support karta hai, aur kya yeh compute- ya bandwidth-bound hai?
Recall Solution
(a) Efficiencies combine karo (yeh multiply hoti hain, kyunki har ek independently usable bandwidth trim karti hai): (b) Us memory feed se compute rate = TFLOP/s. Peak 20 TFLOP/s se compare karo: bandwidth-bound — memory feed, ALUs nahi, hame limit kar rahi hai.
Recall Self-test one-liners
32 floats ke fully coalesced warp ko kitne transactions chahiye? ::: 1 AoS 24-byte struct, ek field padhna, transactions per warp? ::: 6 (efficiency 16.67%) tiling se global traffic mein reduction? ::: ka factor Tiled matmul ki arithmetic intensity, , floats? ::: 8 FLOPs/byte 20 TFLOP/s aur 1 TB/s ke liye machine balance? ::: 20 FLOPs/byte 32-way shared-memory column conflict ka fix? ::: Tile row ko width 33 tak pad karo