Warps and warp scheduling — Hinglish notes

6.2.5 · Hardware › GPU Architecture

Overview

Warps GPU execution ki fundamental unit hain NVIDIA GPUs mein, jo ek aisi thread group ko represent karti hain jo same instruction ko lockstep mein execute karti hain. Warp scheduling ko samajhna GPU performance optimize karne ke liye critical hai kyunki GPU scheduler warp granularity par operate karta hai, individual threads par nahi.

Core Concepts

Warps Thread Blocks Se Kaise Map Hote Hain

Thread indices se derivation:

Ek 3D thread block ke liye jiska dimensions (blockDim.x, blockDim.y, blockDim.z) hai, threads ko warps mein row-major order mein assign kiya jaata hai:

Linear Thread Index = threadIdx.z \times (blockDim.x \times blockDim.y) + threadIdx.y \times blockDim.x + threadIdx.x

$Warp ID within Block = ⌊ \frac{Linear Thread Index}{32} ⌋$

Yeh kyun matter karta hai: Agar blockDim.x = 64 hai, toh threads 0-31 warp 0 banate hain, threads 32-63 warp 1 banate hain. Agar blockDim.x = 17 hai, toh threads 0-31 warp 0 banate hain, lekin sirf 17 threads active hain—baaki 15 slots waste hote hain. Isliye block dimensions 32 ke multiples hone chahiye.

Warp Scheduling Mechanics

The Scheduling Problem

Ek SM multiple thread blocks simultaneously hold kar sakta hai (registers, shared memory, aur max blocks/SM se limited). Har block multiple warps contribute karta hai. Agar ek SM 2048 threads hold kar sakta hai aur uske paas 64 warps resident hain, toh scheduler ko decide karna hota hai ki aage kaun sa warp execute hoga.

Zero-overhead switching kyun kaam karta hai, iska derivation:

Jab warp W₁ memory par stall hota hai, scheduler usi cycle mein warp W₂ par immediately switch kar leta hai. Yeh isliye possible hai kyunki:

Register files partitioned hain: Har warp ka dedicated register space hota hai (koi context save/restore nahi chahiye)
Program counters per-warp hain: Warp switch karna = ek pointer update karna (1 cycle)
Instruction buffers alag hain: Har warp ki next instruction pre-fetch ho jaati hai

Latency hiding capacity:

Agar ek memory operation $L$ cycles leti hai aur har warp $W$ instructions mein ek memory operation karta hai, toh aapko zaroori hai:

N_{warps needed} = ⌈ \frac{L}{W} ⌉

Kyun? Agar $L = 400$ cycles aur $W = 10$ instructions hain memory ops ke beech, toh SM ko busy rakhne ke liye 400/10 = 40 warps chahiye jab ek warp wait kar raha ho.

Worked example Warp Scheduling Example: Latency Hiding

Scenario: SM with 4 warp schedulers, 48 resident warps, global memory latency 400 cycles.

Timeline:

Cycle 0: Warp 0 global memory se load issue karta hai → stalls
Cycle 1: Scheduler Warp 1 par switch karta hai (koi overhead nahi)
Cycle 2: Warp 1 arithmetic execute karta hai
Cycle 3: Scheduler Warp 2 par switch karta hai
...
Cycle 400: Warp 0 ka memory data aa jaata hai → Warp 0 ready ho jaata hai Yeh kyun kaam karta hai: 48 warps aur 4 schedulers ke saath jo instructions issue kar rahe hain, SM 4 instructions per cycle issue kar sakta hai. Agar warps ke beech memory ops ke darmiyaan kaafi compute hai, toh jab tak hum bahut saare warps mein cycle karte hain, pehle warp ki memory ready ho jaati hai.

Yeh step kyun? Har scheduler ~12 warps (48/4) mein se choose kar sakta hai. Jab tak har cycle mein har scheduler ke liye 1 warp ready hai, hum full throughput maintain kar sakte hain.

Scheduling Policies

Modern GPUs in policies ke variations use karte hain:

Round-robin with priority: Eligible warps mein cycle karo, long-stalled warps ko prioritize karo
Greedy-then-oldest (GTO): Ek warp ko tab tak execute karo jab tak woh stall na ho, phir sabse purana stalled warp pick karo
Two-level scheduling: Un warps ko prioritize karo jo cache mein hit karenge

Formula Occupancy aur Scheduling Efficiency

Occupancy active warps ka maximum warps per SM se ratio hai:

Occupancy = \frac{Active Warps per SM}{Max Warps per SM}

Occupancy kyun matter karti hai: Zyada occupancy → zyada warps available → better latency hiding. Lekin occupancy akele performance guarantee nahi karti; tumhe warp diversity bhi chahiye (warps alag-alag instruction stages par).

Optimal occupancy derive karna:

Maano $T_{issue}$ = har scheduler ke liye instruction issues ke beech cycles, $L$ = average stall latency, $N$ = schedulers ki sankhya.

Sab schedulers ko busy rakhne ke liye:

Warps needed \geq N \times \frac{L}{T _{issue}}

$N = 4$ schedulers, $L = 400$ cycle memory latency, $T_{issue} = 1$ ke liye (ek instruction per cycle scheduler):

Warps needed \geq 4 \times \frac{400}{1} = 1600 instructions’ worth

Agar har warp stall hone se pehle 20 independent instructions karta hai: 1600/20 = 80 warps. Lekin max typically 64 hota hai, isliye tumhe warps ke andar instruction-level parallelism bhi chahiye.

Warp Divergence

Divergence penalty ka derivation:

Ek if-else consider karo jahan $n_{if}$ threads if branch lete hain aur $n_{else} = 32 - n_{if}$ else branch lete hain.

Divergence ke bina (sab threads same path lete hain):

T_{normal} = T_{if}

Divergence ke saath (maano $n_{if} > 0$ aur $n_{else} > 0$ ): $T_{diverged} = T_{if} + T_{else}$

Worst case: Dono branches ki equal cost $T$ hai aur equally populated hain:

Slowdown = \frac{T + T}{T} = 2 \times

Yeh kyun hota hai: Warp if branch ko ek mask ke saath execute karta hai (jo threads else lete hain unhe disable karta hai), phir else branch ko opposite mask ke saath execute karta hai. Dono paths sequentially run karte hain chahe sirf 1 thread ko hi koi path chahiye ho.

Worked example Divergence Example: Boundary Handling

__global__ void kernel(float* data, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {  // Boundary check
        data[idx] = data[idx] * 2.0f;
    }

Case 1: N = 128, blockDim.x = 128, 4 warps

Warps 0-3: Sab 32 threads idx < N pass karte hain → koi divergence nahi, full speed Case 2: N = 100, blockDim.x = 128, 4 warps
Warp 0 (threads 0-31): Sab pass karte hain → koi divergence nahi
Warp 1 (threads 32-63): Sab pass karte hain → koi divergence nahi
Warp 2 (threads 64-95): Sab pass karte hain → koi divergence nahi
Warp 3 (threads 96-127): Sirf threads 96-99 pass karte hain (4/32) → divergence, 28 threads disabled

Performance: Warp 3 if body execute karta hai (4 active threads) lekin phir bhi ek full instruction ke liye execution pipeline occupy karta hai. Agar body 10 cycles leti hai, toh warp 3 4 threads ke liye 10 cycles leta hai 32 ki jagah → 87.5% waste

Yeh step kyun? Hardware inactive threads ke instruction slots skip nahi kar sakta; woh instruction execute karta hai jisme zyaadatar threads masked off hote hain.

Common mistake Common Divergence Mistake: Thread-ID-Based Branching

Galat approach:

if (threadIdx.x % 2 == 0) {
    // Even threads
} else {
    // Odd threads
}

Yeh sahi kyun lagta hai: "Main sirf threads ko evenly split kar raha hoon, yeh balanced hai!"

Yeh galat kyun hai: Har warp mein 16 even aur 16 odd threads hote hain (threadIdx.x 0-31 tak jaata hai). Har warp diverge karta hai aur dono branches serially run karta hai → poore kernel mein 2× slowdown.

Fix: Kaam ko is tarah organize karo ki poore warps same path lein:

int warpId = threadIdx.x / 32;
if (warpId % 2 == 0) {
    // Even warps
} else {
    // Odd warps
}

Ab warp 0 ek branch execute karta hai, warp 1 doosra—koi intra-warp divergence nahi.

Steel-man the mistake: Even/odd split sahi lagta hai kyunki tum kaam evenly distribute kar rahe ho. Hidden cost yeh hai ki warps SIMT mein execute karte hain—execution ki unit warp hai, thread nahi. Threads ki jagah warps mein sochne ki aadat na hona is galati ka kaaran hai.

Warp Scheduling ke Liye Optimize Karna

1. Occupancy Maximize Karo (Reason ke Saath)

Per SM constraints:

Max threads: 2048 (architecture ke hisaab se alag)
Max blocks: 32 (alag ho sakta hai)
Max warps: 64 (compute capability 7.0+)
Register file size: jaise 65,536 registers
Shared memory: jaise 64 KB (L1 cache ke saath configurable split)

Occupancy formula:

\frac{\text{Threads per Block} \times \text{Blocks per SM}}{\text{Max Threads per SM}}, \frac{\text{Blocks per SM}}{\text{Max Blocks per SM}}, \frac{\text{Warps per SM}}{\text{Max Warps per SM}} \right)$$ Subject to:

\text{Registers per Block} \times \text{Blocks per SM} \leq \text{Total Registers}

Shared Memory per Block \times Blocks per SM \leq Total Shared Memory

2. Divergence Minimize Karo

Warp-wide operations use karo: __ballot_sync(), __shfl_sync()
Conditionals is tarah structure karo ki poore warps same path lein
Data structures ko warp boundaries par pad karo

3. Instruction Mix Balance Karo

Instruction throughput type ke hisaab se alag hota hai:

FP32: typically 1 op/cycle/SM
FP64: 1/2 rate (ya consumer GPUs par 1/32)
Special functions (sin, cos): 1/4 rate

Agar sab warps ko FP64 chahiye, toh schedulers limited FP64 units ke liye compete karte hain → FP32 units ka underutilization.

Recall Ek 12-Saal-Ke Bachche Ko Samjhao

Socho tumhara teacher chahta hai ki class math problems solve kare. Har student ko individually call karne ki jagah (jo slow hai!), woh tumhe 32 ke teams mein group karta hai. Har team ko same problem milta hai, aur sab ek saath apni copy par kaam karte hain. Wahi warp hai!

Teacher (warp scheduler) ghoom-ghoom kar check karta hai kaun si teams ready hain. Agar Team A library se book ka wait kar rahi hai (memory access), woh immediately Team B ki help karne jaata hai. Jab tak woh kai teams mein cycle karta hai, Team A ki book aa jaati hai, isliye koi time waste nahi hota.

Lekin yahan ek tricky part hai: agar Team A ke problem mein likha hai "even IDs wale addition karo, odd IDs wale subtraction karo," toh team ko split hona padta hai. Pehle even-ID students solve karte hain jab odd-ID students idle baithe hain. Phir woh switch karte hain. Team ko double time lagta hai kyunki woh saath mein kaam nahi kar sakte. Wahi divergence hai!

Seekh: apni poori warp team ko ek saath same cheez karte rakhne se maximum speed milti hai.

Connections

SIMT-vs-SIMD - GPU SIMT CPU SIMD se kaise alag hai
Thread-Blocks-and-Grids - Warps larger hierarchy mein kaise fit hote hain
GPU-Memory-Hierarchy - Memory latency hide karne ke liye warp scheduling kyun matter karti hai
Branch-Divergence-Patterns - Common divergence scenarios aur solutions
Occupancy-vs-Performance - 100% occupancy hamesha optimal kyun nahi hoti
Cooperative-Groups - Modern warp-synchronous programming primitives
Register-Pressure - Register usage occupancy aur warp count kaise limit karta hai

#flashcards/hardware

GPU architecture mein warp kya hota hai? :: 32 consecutive threads ka ek group jo SIMT fashion mein ek saath same instruction execute karta hai. Warps NVIDIA GPUs mein fundamental scheduling unit hain.

Warp ID thread indices se kaise calculate hota hai?

Linear thread index ko 32 se divide karo (floor division). Linear index = threadIdx.z * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x, phir warpId = linearIdx / 32.

Thread block dimensions 32 ke multiples kyun hone chahiye?

Wasted warp slots se bachne ke liye. Agar blockDim 32 ka multiple nahi hai, toh last warp mein inactive threads honge jo resources consume karte hain lekin koi kaam nahi karte.

Warp divergence kya hoti hai?

Jab ek hi warp ke threads alag-alag execution paths lete hain (alag branches). GPU ko divergent paths serialize karne padte hain, har path ko kuch threads masked off karke execute karna padta hai, jisse performance kam hoti hai.

Warp divergence ki worst-case performance penalty kya hai?

2× slowdown jab dono branches ki equally costly aur equally populated hon. Warp dono branches sequentially execute karta hai ek path ki jagah.

Warp scheduling memory latency kaise hide karta hai?

Jab ek warp memory par stall hota hai (100s of cycles), scheduler immediately kisi doosre ready warp par zero overhead ke saath switch kar leta hai. Kaafi resident warps ke saath, SM busy rehta hai jab stalled warps wait karte hain.

GPU occupancy kya hoti hai?

Active warps ka maximum warps per SM se ratio. Formula: Active Warps / Max Warps per SM. Zyada occupancy latency hiding ke liye zyada warps provide karti hai.

Occupancy ko kya factors limit karte hain?

Registers per thread, shared memory per block, threads per block, max blocks per SM, aur max warps per SM. Sabse restrictive limit actual occupancy determine karti hai.

Warp context switching ka zero overhead kyun hota hai?

Register files per warp partitioned hain, program counters per-warp hain, aur instructions pre-fetch ho jaati hain. Warps switch karna sirf ek pointer update karta hai, koi context save/restore nahi chahiye.

Warp schedulers aur cores mein kya farq hai?

Warp schedulers decide karte hain ki aage kaun sa warp execute hoga. Cores (CUDA cores) actually instructions execute karte hain. Modern SMs mein multiple schedulers hote hain (jaise 4) jo har cycle mein multiple warps se instructions issue karte hain.

256-thread block mein kitne warps hote hain?

256 / 32 = 8 warps. Har block ke threads ko 32 consecutive threads ke groups mein divide kiya jaata hai.

Ek warp scheduling ke liye eligible kya banata hai?

Warp ko stalled nahi hona chahiye (koi memory wait, dependency wait, ya synchronization barrier nahi) AUR resources available hone chahiye (execution units, memory ports).

threadIdx.x % 2 branching performance ke liye kyun buri hai?

Yeh har warp mein divergence cause karta hai kyunki har warp mein dono even (16) aur odd (16) threadIdx values hoti hain. Har warp ko dono branches serially execute karni padti hain.

Latency L hide karne ke liye warps ki formula kya hai?

Warps needed ≥ (Number of schedulers × L) / Instructions between stalls. 4 schedulers aur 400-cycle latency ke liye 10 instructions ke saath memory ops ke beech: 4 × 400 / 10 = 160 warps worth of work.

Warp-synchronous operation kya hoti hai?

Operations jo ek warp ke sab threads mein simultaneously act karti hain, jaise __shfl_sync() data exchange ke liye ya __ballot_sync() voting ke liye, explicit synchronization ki zaroorat ke bina.