GPU memory bandwidth optimization

6.2.14 · Hardware › GPU Architecture

Overview

GPU memory bandwidth optimization ek aisi practice hai jisme memory access patterns ko is tarah structure kiya jaata hai ki GPU cores aur memory ke beech data transfer rates maximize ho sakein — kyunki memory bandwidth aksar GPU computing ka primary performance bottleneck hoti hai.

Fundamental problem yeh hai: Modern GPUs mein hazaaron cores hote hain lekin limited memory buses hoti hain. Ek single NVIDIA A100 mein 6,912 CUDA cores hain jo ~2 TB/s memory bandwidth share karte hain. Agar har core independently data request kare, toh memory system congested ho jaata hai. Smart access patterns se 10-100× speedups achieve kiye ja sakte hain.

Core Concepts

Yeh kyun exist karta hai: GPU memory controllers data ko bade chunks mein read karte hain (32-128 bytes). Jab thread 0 address A read karta hai, thread 1 A+4 read karta hai, thread 2 A+8 read karta hai, etc., toh hardware 32 alag-alag 4-byte reads ki jagah ek 128-byte block fetch karta hai.

Coalesced bandwidth ki derivation:

$w=32$ threads ke warp ke liye:

• Har thread $b$ bytes request karta hai
• Memory transaction size: $T$ bytes (typically 128B)
• Coalesced case: addresses $\le T$ bytes span karte hain

$\text{Transactions needed}=\lceil \frac{w\cdot b}{T}\rceil$

Best case (fully coalesced): Saare 32 threads consecutive floats access karte hain (4 bytes each): $\text{Span}=32\times 4=128\text{ bytes}=1\text{ transaction}$

Worst case (strided by 32): Thread $i$ address $A+32i\times 4$ access karta hai: $\text{Span}=32\times 32\times 4=4096\text{ bytes}=32\text{ transactions}$

Effective bandwidth ratio: $\text{Efficiency}=\frac{\text{Data used}}{\text{Data transferred}}=\frac{w\cdot b}{\text{Transactions}\times T}$

Coalesced: $\frac{128}{128}=100\%$ | Strided: $\frac{128}{4096}=3.125\%$

$M=B\times \eta$

jahaan: $\eta =\frac{\text{Useful bytes loaded}}{\text{Total bytes transferred from DRAM}}$

Hardware constraints se derivation:

1. Physical bandwidth: $B=\frac{\text{Bus width}\times \text{Clock rate}\times 2\text{ (DDR)}}{\text{8 bits/byte}}$
2. Transaction overhead: Har memory request ka minimum size $T_{min}$ hota hai
3. Warp divergence: Jab threads non-contiguous addresses access karein, toh har subset ko alag transaction chahiye
4. Bank conflicts: Shared memory mein, same bank ka parallel access serialize ho jaata hai

Combined efficiency: ${\eta }_{total}={\eta }_{coalesce}\times {\eta }_{cache}\times {\eta }_{occupancy}$

• ${\eta }_{coalesce}$: Coalescing efficiency (upar derive ki gayi)
• ${\eta }_{cache}$: Cache hit rate (DRAM traffic reduce karta hai)
• ${\eta }_{occupancy}$: Fraction of time jab memory controllers busy hain

Banks kyun exist karte hain: Hardware cost ki wajah se. 32-way parallel access arbitrary addresses ke liye provide karna full crossbar ($32\times 32$ connections) maangega. Banks memory ko partition karte hain taaki har thread ek bank se map ho, jisse simpler hardware ke saath $32\times 1$ connections per bank possible ho.

Bank assignment: Address $A$ aur $N$ banks ke liye: $\text{Bank}=\left(\frac{A}{\text{word size}}\right)\mathrm{mod}N$

Conflict rule: Bank $b$ access karne waale $k$ threads 1 cycle ki jagah $k$ cycles lete hain.

Broadcasting exception: Agar saare threads same address read karein, toh hardware 1 cycle mein broadcast karta hai (koi conflict nahi).

Optimization Techniques

1. Array-of-Structures vs Structure-of-Arrays

Array of Structures (AoS):

struct Particle { float x, y, z, vx, vy, vz; };
Particle particles[N];
 
// Thread i accesses:
float x = particles[i].x;  // Address: base + i*24 + 0
float y = particles[i].y;  // Address: base + i*24 + 4

Yeh slow kyun hai: Thread 0 address A read karta hai, thread 1 A+24 read karta hai, thread 2 A+48 read karta hai. Jo field threads chahte hain (x) woh har 24 bytes par scattered hai. Hardware phir bhi 128-byte lines fetch karta hai, lekin ek warp jo ek float per thread read kar raha hai woh ab $32\times 24=768$ bytes ka address span touch karta hai. Iska matlab hai $\lceil 768/128\rceil =6$ transactions per field — lekin fetch kiye gaye 768 bytes mein se sirf $32\times 4=128$ bytes useful hain (efficiency $\approx 16.7\%$). Toh AoS partially coalesce karta hai; yeh 32 alag reads mein fully serialize nahi hota, lekin ~5/6 bandwidth waste karta hai.

Structure of Arrays (SoA):

struct Particles {
    float x[N], y[N], z[N];
    float vx[N], vy[N], vz[N];
};
Particles particles;
 
// Thread i accesses:
float x = particles.x[i];  // Address: base_x + i*4
float y = particles.y[i];  // Address: base_y + i*4

Yeh fast kyun hai: Thread 0 A read karta hai, thread 1 A+4 read karta hai, thread 2 A+8 read karta hai. Addresses consecutive hain, isliye hardware 1 transaction per warp per field mein combine kar deta hai.

Speedup calculation: 6 fields ke liye, per warp:

• AoS: $6$ fields $\times 6$ transactions/field $=36$ transactions
• SoA: $6$ fields $\times 1$ transaction/field $=6$ transactions

$\text{Reduction}=\frac{36}{6}=6\times \text{ fewer memory transactions}$

(Yeh gap wider structs ke liye barhta hai: ek struct jo per element $S$ bytes span karta hai, AoS ke liye $\lceil 32S/128\rceil$ transactions per field deta hai vs SoA ke liye 1.)

2. Tiling with Shared Memory

Problem: $N\times N$ matrices ke liye, har thread global memory se $2N$ elements load karta hai. Total: $N^2\times 2N=2N^3$ global element reads (bahut saare redundant hain).

Tiled approach:

1. Matrices ko $T\times T$ tiles mein divide karo
2. $A$ ka ek tile aur $B$ ka ek tile shared memory mein ek baar per block load karo
3. Block ke saare threads shared data reuse karte hain
4. Next tile par move karo

Yeh step kyun? Shared memory global memory se ~100× faster hoti hai aur ek block ke andar share hoti hai. Tiles cooperatively load karke, hum $N$ redundant global reads ko 1 global read + $N$ fast shared reads mein convert karte hain.

Speedup ki derivation:

Tiling ke bina: $\text{Global reads per output element}=2N$

$T\times T$ tiles ke saath: $\text{Tiles along the k-dimension}=\frac{N}{T}$ $\text{Global reads per output element}=\frac{2N}{T}\text{ (global)}+2N\text{ (shared)}$

Global memory traffic reduction: $\text{Reduction factor}=\frac{2N}{2N/T}=T$

$T=32$ ke liye: 32× kam global memory traffic.

Arithmetic intensity ki poori derivation:

Operation count: $2N^3$ FLOPs (har multiply-add = 2 FLOPs, $N^3$ baar done).

Global data movement (tiled), bytes count karte hue: har output element $\frac{2N}{T}$ global element reads trigger karta hai, aur har element ek 4-byte float hai: $\text{Bytes moved}=N^2\times \frac{2N}{T}\times 4=\frac{8N^3}{T}\text{ bytes}$

Arithmetic intensity: $I=\frac{\text{FLOPs}}{\text{Bytes}}=\frac{2N^3}{8N^3/T}=\frac{T}{4}\text{ FLOPs/byte}$

$T=32$ ke liye: $I=8$ FLOPs/byte.

Compute-bound vs bandwidth-bound (Roofline): Ek kernel compute-bound hota hai jab uski arithmetic intensity machine balance $I^{\ast }=\frac{\text{peak FLOP/s}}{\text{peak bytes/s}}$ se zyada ho. Ek GPU par jo 20 TFLOP/s peak aur 1 TB/s bandwidth deta hai: $I^{\ast }=\frac{20\times 10^{12}}{1\times 10^{12}}=20\text{ FLOPs/byte}$

• Tiled ($T=32$): $I=8$ FLOPs/byte $<20$ → still bandwidth-bound, lekin naive se 32× kam traffic.
• Larger tiles / register blocking $I$ ko $20$ ke paas aur usse aage push karte hain taaki compute-bound ho jaaye.
• Naive (no tiling, $T=1$): $I=0.25$ FLOPs/byte $\ll 20$ → heavily bandwidth-bound.

3. Avoiding Bank Conflicts in Shared Memory

Naive approach:

__shared__ float tile[32][32];
 
// Load: thread i reads column i → coalesced
tile[threadIdx.y][threadIdx.x] = input[...];
 
// Store: thread i writes row i → each thread different bank
output[...] = tile[threadIdx.x][threadIdx.y];  // CONFLICT!

Conflict kyun hota hai: Jab thread 0 tile[0][0] read karta hai, thread 1 tile[1][0] read karta hai, .., thread 31 tile[31][0] read karta hai. Saare column 0 access karte hain, jo same bank (bank 0) se map hota hai → 32-way conflict → 32× slower.

Bank mapping: Address tile[row][col] = base + (row*32 + col)*4 bytes $\text{Bank}=(row\times 32+col)\mathrm{mod}32=col\mathrm{mod}32$

Column 0 access karne waale saare threads bank 0 hit karte hain.

Solution: Padding

__shared__ float tile[32][33];  // Extra column

Yeh kyun kaam karta hai: Ab address tile[row][col] = base + (row*33 + col)*4 $\text{Bank}=(row\times 33+col)\mathrm{mod}32$

Column 0 ke liye: thread $i$ bank $(i\times 33)\mathrm{mod}32$ access karta hai. Kyunki $\gcd (33,32)=1$, consecutive threads different banks se map hote hain → koi conflict nahi!

Cost: 3% memory overhead (33/32) → 32× speedup transpose mein.

Yeh sahi kyun lagta hai: CPUs par, caches working sets ke relative large hoti hain aur achha automatic optimization provide karti hain.

Reality yeh hai: GPU caches per-thread bahut chhoti hain (128 KB L1 per SM jo 1024+ threads share karte hain = 128 bytes per thread). 32-thread warps lockstep mein execute hone se, agar threads scattered addresses access karein, toh cache data reuse se pehle hi thrash ho jaati hai.

Example: Matrix multiply bina tiling ke. Har thread ko $2N$ elements chahiye, lekin cache sirf ~32 floats per thread hold kar sakti hai. $N=1024$ ke liye, 99.7% accesses cache miss karte hain → bandwidth-bound.

Fix: Deliberate data reuse ke liye explicitly shared memory use karo. Automatic caching par depend mat karo.

Yeh sahi kyun lagta hai: Addresses strided hain (har 24 bytes), toh lagta hai kuch bhi coalesce nahi hota.

Reality yeh hai: Hardware hamesha full 128-byte cache lines fetch karta hai. Ek warp jo ek float per thread 24-byte-stride struct se read kar raha hai woh $32\times 24=768$ bytes span karta hai → $\lceil 768/128\rceil =6$ transactions, 32 nahi. Yeh partially coalesce karta hai lekin ~83% bandwidth waste karta hai (768 bytes mein se sirf 128 bytes useful).

Fix: 1 transaction/field ke liye SoA use karo. AoS ki asli penalty wasted bandwidth hai (efficiency ~17%), "32× more transactions" nahi.

Reality yeh hai: Memory transactions aligned hone chahiye 32/64/128-byte boundaries par. Agar array address 0x102 se start ho (misaligned), toh consecutive accesses bhi boundaries par split ho jaate hain.

Example: Array 0x102 se start hoti hai. Warp 0x102-0x181 (128 bytes) read karta hai. Yeh 0x100 aur 0x180 boundaries span karta hai → 1 ki jagah 2 transactions.

Fix: cudaMalloc (aligned pointers return karta hai) ya __align__ attribute use karo:

__align__(128) float data[N];

Active Recall Practice

GPUs us library jaisi hain. Unke paas hazaaron "students" (threads) hain jo ek saath kaam karte hain. Memory library shelves hai. Agar threads randomly data maangein, toh memory system overwhelm ho jaata hai. Lekin agar nearby threads nearby data maangein, toh memory system sab kuch ek trip mein deliver kar sakti hai.

Kuch tricks hain: (1) data organize karo taaki nearby threads ko nearby data chahiye (jaise students ko unki shelf ke hisaab se arrange karna), (2) ek chhhoti fast "desk" (shared memory) use karo taaki same room ke students ke beech data share ho, (3) ensure karo ki students sab ek saath same drawer se grab na karein (bank conflicts). Yeh tricks GPU programs ko 10-100× faster banati hain!

Connections

• GPU-Warp-Scheduling: coalescing warps ke andar kaam karta hai (32 threads)
• Cache-Hierarchy: L1/L2 caches explicit optimization ko complement karte hain, replace nahi karte
• Roofline-Model: bandwidth optimization programs ko compute-bound regime ki taraf move karta hai
• Parallel-Algorithm-Design: memory access patterns algorithm choice ko influence karte hain
• CUDA-Shared-Memory: banking aur conflicts ke hardware details
• Memory-Latency-Hiding: high occupancy latency tolerance allow karta hai

Summary

GPU memory bandwidth optimization coalescing (threads ko consecutive addresses access karaana), tiling (fast shared memory mein data reuse karna), aur bank conflicts avoid karne (shared memory access patterns structure karna) par center hoti hai. Fundamental trade-off yeh hai: data layout mein chhote changes 10-100× speedups dete hain kyunki memory bandwidth GPU ka scarcest resource hai. Har optimization ek hardware truth se derive hoti hai ki 128 bytes load karne ki cost wahi hai jo 4 bytes load karne ki — isliye threads ki requests ko saath pack karo.

#flashcards/hardware

Concept Map