Streaming multiprocessors (SM)

6.2.2 · Hardware › GPU Architecture

Overview

Streaming Multiprocessor (SM) ek GPU ke andar ka fundamental execution unit hai. Ise ek mini-processor samjho jo parallel execution ke liye optimize hua hai – jahan ek CPU mein shayad 4-16 cores hon, ek modern GPU mein dozens of SMs hote hain, aur har SM ek saath hundreds of threads execute kar sakta hai.

Yeh hardware-level parallelism hai: kai SMs bilkul independently chalte hain, har ek hundreds of threads juggle karta hua. GPU scheduler thread blocks ko SMs mein distribute karta hai – ek SM shayad pixels 0-1023 render kare jab doosra 1024-2047 render kare.

Architecture of a Single SM

SM ke Andar Kya Hota Hai?

Yeh structure kyun? Har component ek specific bottleneck ko target karta hai:

• CUDA cores → raw arithmetic throughput
• SFUs → expensive math offload karo (CPU par 20-100 cycles, SFU par 1-2 cycles)
• Registers → memory latency khatam karo (threads apne variables registers mein rakhte hain)
• Shared memory → fast inter-thread communication (global memory se 100x faster)

Hierarchical Execution Model

SM threads ko ek strict hierarchy mein execute karta hai:

Warp mein 32 threads kyun? Yeh SIMT hai (Single Instruction Multiple Thread) – warp scheduler ek instruction issue karta hai jo saare 32 threads ek saath execute karte hain. Yeh SIMD (vector processing) jaisa hai lekin zyada flexibility ke saath: threads alag branches le sakte hain, halanki isse performance cost aati hai (divergence).

Thread scheduling ki derivation:

Har cycle mein, warp scheduler:

1. Ek eligible warp chunti hai (threads jo memory/dependencies par stall nahi hain)
2. Us warp ke 32 threads ko ek instruction issue karta hai
3. CUDA cores instruction ko parallel mein execute karte hain

Agar ek warp stall ho jaye (memory ka wait kar raha ho), scheduler turant doosre warp par switch kar leta hai – isliye GPUs ko hazaron threads chahiye: latency hide karne ke liye.

Occupancy calculation:

Occupancy kyun matter karta hai: Zyada occupancy → zyada warps available → better latency hiding. Lekin occupancy in chezon se limit hoti hai:

• Registers per thread × threads per block ≤ Total registers
• Shared memory per block ≤ Total shared memory
• Blocks per SM ≤ Max blocks (usually 16-32)

Example: Agar aapka kernel 64 registers/thread use kare aur SM mein 65,536 registers hon, max threads = 65,536 / 64 = 1024. Lekin agar max warps = 64 hai, toh aap 64 × 32 = 2048 threads par cap ho jaoge waise bhi.

Work Kaise Distribute Hota Hai

Scenario: Ek kernel launch karo 10,000 threads ke saath jo 40 blocks of 256 threads each mein organized hain. GPU mein 80 SMs hain.

Kya hota hai:

1. Block assignment: GPU ka Grid Distribution Unit blocks ko SMs par assign karta hai. Har block exactly ek SM par start se finish tak chalta hai (koi migration nahi).

   • 80 SMs aur 40 blocks ke saath, har SM ko pehle ⌊40/80⌋ = 0 ya 1 block milta hai, phir jaise blocks complete hote hain, zyada milte hain.

2. Warp creation: Har 256-thread block, 256/32 = 8 warps ban jaata hai. SM ka warp scheduler ab in 8 warps ko juggle karta hai.

3. Instruction issue: Har cycle, scheduler ek eligible warp chunti hai aur uska next instruction CUDA cores ko issue karta hai.

   • Agar Warp 0 memory se fetch kar raha hai (300 cycles), scheduler doosre eligible warps chalaata hai jab Warp 0 ka data in flight hai.

Yeh step kyun? (Latency-hiding math sahi se) Maano ki ek global-memory access ek warp ko 300 cycles ke liye stall kare, aur ek scheduler 1 instruction per cycle issue kare. Us stall ko poori tarah hide karne ke liye aapko roughly 300 independent in-flight warps chahiye (latency ke har ek cycle ke liye ek) — ek single block ke ek SM ke 8 warps akele 300-cycle stall cover karne ke liye kaafi nahi hain. Practice mein hum iske kareeb jaate hain ek saath kai blocks resident rakh ke: real SM cap of ~48-64 warps ke saath, plus kai schedulers jo har cycle issue karte hain, hum hundreds of cycles hide karne ke kareeb pahunch jaate hain. Lesson yeh hai: resident warps maximize karo, kyunki 8 warps sirf ~8 cycles latency hi cover karte hain per scheduler pass.

Problem: Aapka kernel 128 registers per thread use karta hai. SM mein 65,536 registers hain. Block size = 256 threads.

Calculation:

• Registers needed per block: 256 × 128 = 32,768
• Max blocks per SM: 65,536 / 32,768 = 2 blocks
• Warps per block: 256/32 = 8 → Total warps = 2 × 8 = 16 warps
• Agar max warps/SM = 64, occupancy = 16/64 = 25%

Yeh step kyun? High register usage SM ko parallelism se starve kar deta hai. Sirf 16 warps ke saath, aur ek scheduler ~1 instruction/warp/cycle issue karte hue, aap ek 400-800 cycle global-memory stall ka sirf ek chhota sa hissa hi cover karte ho — isliye ek memory-bound kernel stall karega. Hundreds of cycles hide karne ke liye aapko dozens of resident warps chahiye, jo yahan register pressure rok deta hai.

Fix: Registers kam karo (algorithmic change), resident blocks/warps badhao, ya lower occupancy accept karo agar compute-bound ho.

Memory Hierarchy aur SM

Har SM mein kai memory tiers hain, har ek ki distinct latency characteristics hain:

Note: Shared memory aur L1 data cache aksar usi physical SRAM ko share karte hain lekin yeh same cheez nahi hain — shared memory explicitly managed hoti hai (aap directly isme index karte ho, koi tag lookup nahi), jabki L1 hardware-managed hai with tag checks aur eviction. Isliye unki latencies aur behavior architecture ke hisaab se alag hoti hain.

Shared memory kyun? Ek block ke threads shared memory ke zariye cooperate kar sakte hain:

__shared__ float cache[256];  // Visible to all threads in block
cache[threadIdx.x] = data[i]; // Each thread writes
__syncthreads();              // Wait for all
float sum = cache[0] + cache[1] + ..; // Read others' data

Shared memory global memory se ~100x faster hai aur tiled matrix multiply, reductions, aur prefix sums jaisi algorithms ko enable karti hai.

Bank conflicts (sirf shared memory mein): Shared memory 32 banks mein divided hoti hai. Agar ek warp ke kai threads same bank mein alag-alag addresses access karen, toh accesses serialize ho jaate hain. Access patterns is tarah design karo ki thread i, bank i % 32 access kare. (Registers per-thread private hote hain aur is tarah ke bank conflicts nahi suffer karte.)

Warp Divergence

Yeh sahi kyun lagta hai: GPU marketing kehti hai "thousands of threads in parallel!"

Reality: Ek warp ke threads same instruction execute karte hain. Agar threads alag branches len:

if (threadIdx.x < 16) {
    A;  // Threads 0-15
} else {
    B;  // Threads 16-31
}

Warp ko dono paths serially execute karne padte hain:

1. A execute karo with threads 0-15 active (16-31 masked)
2. B execute karo with threads 16-31 active (0-15 masked)

Performance cost: 2× instructions. Complex branches ke liye aur bhi bura.

Steel-man: GPU flexibility ke liye divergence allow karta hai (pure SIMD ke unlike). Lekin yeh free nahi hai – ise sparingly use karo.

Fix: Data reorganize karo taaki ek warp ke threads same branch len. Example: processing se pehle data ko category ke hisaab se sort karo.

Tum 10,000 squares ko 256 ke groups (blocks) mein divide karte ho aur har helper ko ek group dete ho. Helper apne 256 squares ko aur 32 ke 8 mini-groups mein split karta hai (warps). Har second, wo ek mini-group chunte hain aur same crayon stroke se saare 32 squares color karte hain.

Agar ek mini-group naye crayon box ka wait kar raha ho (memory fetch), helper turant doosre mini-group par switch kar leta hai – koi time waste nahi! Lekin ek crayon box aane mein bahut time lag sakta hai (jaise 300 seconds), isliye tumhe actually bahut saare mini-groups chahiye waiting mein taaki helper ke paas hamesha kuch color karne ko ho. Isliye hazaron threads chahiye: lambe waits ko poori tarah cover karne ke liye.

Aur yahan ek trick hai: agar instructions kehte hain "squares 0-15 red karo aur 16-31 blue karo," toh helper ko do strokes karne padte hain (red, phir blue) ek ki jagah. Yeh slower hai! Isliye best hota hai jab ek mini-group ke saare 32 squares ko same color chahiye.

Ya: Streaming Machine – threads ki streams isme continuously flow karti hain.

Connections

• CUDA-cores – Har SM ke andar ke arithmetic units
• Warp-scheduling – SM kaise decide karta hai ki har cycle mein kaun se threads chalenge
• Shared-memory – Ek SM ke andar fast inter-thread communication
• Thread-blocks – Ek SM ko assign hone wala work unit
• Occupancy-optimization – Active warps maximize karne ke liye resource usage tune karna
• GPU-memory-hierarchy – SMs kaise L1/L2/global memory access karte hain
• Tensor-cores – Modern SMs mein specialized matrix units

#flashcards/hardware

Concept Map