Noesis
6.2.5 · D5 · HinglishGPU Architecture

Question bank — Warps and warp scheduling

1,933 words9 min read↑ Read in English

6.2.5 · D5 · Hardware › GPU Architecture › Warps and warp scheduling

Yeh ek rapid-fire misconception buster hai. Har line ek Question ::: Answer reveal hai — question padho, apna answer zor se bolo, phir reveal karo. Agar tumhari reasoning justification se match nahi ki, toh woh ek trap tha jisme tumhe girna tha.

Yahan sab kuch Warps and warp scheduling pe build karta hai. Shuru karne se pehle, make sure karo ki tum ek hi saanth mein ye teen words explain kar sako:

  • warp — 32 consecutive threads ka ek fixed group jo ek instruction fetch share karta hai aur program mein ek instruction at a time saath chalta hai.
  • lockstep — sabhi 32 threads hamesha same instruction point kar rahe hote hain; agar kuch threads ko woh run nahi karni chahiye, toh unhe masked off kar diya jata hai (present hain par kuch nahi kar rahe), skip nahi kiya jata.
  • stall — ek warp jo is cycle mein progress nahi kar sakta kyunki woh kisi cheez ka wait kar raha hai (memory data, pehle ka result, ya koi barrier).

Related deep dives jo tumhe open rakhne chahiye: SIMT-vs-SIMD, Branch-Divergence-Patterns, Occupancy-vs-Performance, Register-Pressure, Thread-Blocks-and-Grids, GPU-Memory-Hierarchy.

True or false — justify

A warp always contains exactly 32 threads on every GPU ever made.
False. NVIDIA GPUs pe ek warp 32 threads ka hota hai, lekin AMD ke "wavefronts" 64 ke hote hain. Number ko hamesha vendor se jodo; portable reasoning mein 32 hard-code kabhi mat karo.
If a thread block has 40 threads, it uses exactly 2 warps.
Count mein True hai lekin waste ke saath. 40 threads round up hokar 2 warps (64 slots) ban jaate hain; warp 1 mein sirf 8 active threads hain aur 24 masked-off — SM phir bhi ek full warp ke resources reserve karta hai. Dekho Occupancy-vs-Performance.
Two threads in the same warp can execute two different instructions in the same cycle.
False. Yahi lockstep ki defining property hai — ek warp ke liye ek cycle mein ek hi instruction fetch hoti hai. Divergent threads ko paths sequentially masks ke saath chalaakar handle kiya jaata hai, simultaneously nahi.
Threads in the same warp always finish their kernel at the same time.
False. Agar woh diverge karein, toh kuch threads masked-off baithte hain jab doosre ek branch chalate hain; warp tabhi "reconverge" karta hai jab paths dobara milte hain, isliye effective per-thread finish times alag hoti hain, chahe woh program counter share karein.
Higher occupancy always means higher performance.
False. Occupancy sirf latency hiding ko enable karti hai — scheduler ko switch karne ke liye zyada warps deke. Ek kernel memory-bandwidth-bound ho sakta hai ya 50% occupancy pe already latency hide kar raha hota hai, toh occupancy aur badhane se kuch nahi hoga. Dekho Occupancy-vs-Performance.
Switching from a stalled warp to a ready warp costs the GPU a context-save like a CPU thread switch.
False. Har warp ke registers aur program counter SM mein resident rehte hain, toh switching sirf ek alag pointer choose karna hai — zero-overhead, same cycle. Yahi puri wajah hai ki GPUs itni badi memory latency tolerate karte hain.
A warp scheduler can only run one warp per SM at a time.
False. Modern SMs mein multiple schedulers hote hain (aksar 4), har ek apne warps ke pool se issue karta hai, toh same cycle mein kai warps progress karte hain.
__syncthreads() synchronises threads within a warp.
Misleading hai. Ek warp ke threads already lockstep mein hote hain; __syncthreads() ek block-wide barrier hai jo block ke saare warps ko synchronise karta hai. Warp ke andar coordination ke liye warp-level primitives use karo, dekho Cooperative-Groups.
SIMT and SIMD are the same thing with different names.
False. SIMD ek wide vector register expose karta hai jise programmer explicitly pack karta hai; SIMT har thread ko apne registers deta hai aur woh independent appear karta hai, jabki hardware unhe ek instruction mein 32 ki tarah gang karta hai. Divergence SIMD mein invisible hai lekin SIMT mein ek real cost hai — dekho SIMT-vs-SIMD.
Using more registers per thread can lower the number of resident warps.
True. Register file per SM fixed hoti hai; thread ke liye zyada registers matlab kam threads (aur warps) ek saath fit honge, occupancy kam ho jaati hai. Yeh register pressure hai — dekho Register-Pressure.

Spot the error

"To avoid divergence, I made blockDim = (30, 1, 1) so each warp is nearly full."
Error yeh hai: 30, 32 ka multiple nahi hai, toh warp 0 mein 30 active aur 2 masked threads hain — tumne permanent waste add kar liya, aur iska divergence se koi lena-dena nahi. Divergence branches ke baare mein hai, block size ke nahi; 32 lo.
"My kernel branches on threadIdx.x < 16, so half the warp diverges — unavoidable."
Error yeh hai: divergence is baat pe depend karta hai ki branch har 32-thread warp ke andar kaise split hoti hai. Kyunki threads 0–15 aur 16–31 same warp mein hain, haan yeh diverge karega. Branch kisi aisi cheez pe karo jo warp boundaries ke saath aligned ho (jaise warpId % 2) aur koi divergence nahi hoga. Dekho Branch-Divergence-Patterns.
"A memory load takes 400 cycles, so my program is 400× slower."
Error yeh hai: latency hidden hoti hai, serially har warp ko pay nahi karni padti. Jab warp 0, 400 cycles wait karta hai, scheduler doosre ready warps chalata hai, toh agar occupancy sufficient hai toh throughput barely girta hai.
"Divergence with all 32 threads taking the if branch still costs 2× because there's an if-else."
Error yeh hai: agar koi bhi thread else branch nahi leta, toh warp masked-out else ko puri tarah skip kar deta hai — koi divergence nahi, koi penalty nahi. Cost tabhi aati hai jab dono branches mein kam se kam ek active thread ho.
"Warp ID equals threadIdx.x / 32."
Error yeh hai: yeh sirf 1D blocks ke liye sahi hai. Generally pehle teeno dimensions use karke linear index flatten karo, phir 32 se divide karo. x-only formula 2D/3D blocks ke liye silently break ho jaata hai — dekho Thread-Blocks-and-Grids.
"Two warps from different blocks can never run on the same SM."
Error yeh hai: ek SM aksar ek saath kai blocks hold karta hai (registers aur shared memory se limited), aur unke warps ek scheduling pool mein mix karta hai — yahi variety latency hiding ko feed karti hai.

Why questions

Why does the scheduler operate at warp granularity instead of per thread?
Kyunki instruction fetch/decode 32 threads mein amortise hoti hai: ek fetch 32 executions drive karta hai, toh control-hardware cost 32 baar ki jagah ek baar 32 lanes ke liye pay hoti hai. Individual threads schedule karna woh saving kho deta.
Why do block dimensions being multiples of 32 matter?
Koi bhi partial final warp phir bhi ek full warp ke slots aur resources reserve karta hai lekin lanes masked-off chhod deta hai, jo execution slots waste karta hai aur effective occupancy bina kisi fayde ke gira deta hai.
Why does memory latency hiding need many warps rather than a few?
Tumhe poora stall window fill karne ke liye kafi independent work chahiye: agar ek load cycles costa hai aur ek warp sirf cycles useful work karta hai loads ke beech, toh tumhe roughly warps per scheduler chahiye taaki koi na koi hamesha ready rahe. Bahut kam hone pe SM idle ho jaata hai.
Why can high occupancy still leave the SM idle?
Agar har resident warp ek hi dependency pe ek hi time mein stall ho jaaye (poor "warp diversity"), toh scheduler ke paas kai warps hain lekin koi eligible nahi. Occupancy resident warps count karta hai; sirf ready warps latency hide karte hain.
Why does branch divergence hurt even when both branch bodies are short?
Kyunki dono paths sequentially complementary masks ke under run hote hain; warp time mein T_if + T_else pay karta hai chahe zyada se zyada 32 threads ka kaam hua ho. Har masked pass mein idle lanes pure waste hain. Dekho Branch-Divergence-Patterns.
Why is per-warp register partitioning what makes zero-overhead switching possible?
Har warp ke registers SM ke register file mein physically resident rehte hain, toh switch ke time kuch save ya reload nahi karna padta — scheduler sirf apni issue logic ko ek alag warp ki taraf point karta hai. Dekho Register-Pressure.

Edge cases

A block of exactly 32 threads — how many warps and any waste?
Exactly 1 warp, fully packed, zero waste. Yeh sabse chhota "clean" block size hai aur ek safe default hai.
A block of 1 thread — what happens to the warp?
Woh phir bhi ek full warp occupy karta hai; 1 lane active, 31 masked-off. Enormously wasteful, isliye single-thread blocks almost kabhi sensible nahi hote.
What if only ONE thread in a warp takes the else branch?
Poora warp phir bhi dono paths serialize karta hai — 31 threads us single thread ke liye else pass ke dauran idle baithte hain. Divergence cost is baat se set hoti hai ki kya dono paths liye gaye, nahi ki kitne threads ne kaunsa liya. Dekho Branch-Divergence-Patterns.
What if every resident warp is stalled on global memory at once?
Scheduler ke paas zero eligible warps hain, toh SM kuch issue nahi karta aur pehla memory reply aane tak idle raha ta hai — yeh insufficient warp diversity ka failure mode hai, low occupancy ka nahi.
A block with 1024 threads and a warp size of 32 — how many warps, and is that automatically good?
32 warps. Woh count satisfy karta hai, lekin agar har thread registers hog kare, toh 1024 saath mein resident nahi ho sakte; bade blocks register/shared-memory limits ki wajah se occupancy reduce kar sakte hain. Dekho Register-Pressure aur Occupancy-vs-Performance.
A warp where all 32 threads read consecutive global addresses vs. scattered addresses — same latency?
Nahi. Consecutive (coalesced) accesses kam memory transactions mein collapse ho jaate hain; scattered accesses kai issue karte hain, effective latency multiply ho jaati hai. Warp scheduling latency hide karta hai; access pattern decide karta hai ki kitni latency hai. Dekho GPU-Memory-Hierarchy.
Recall Fast self-test before you close this page

Name the one situation where an if-else inside a kernel costs nothing extra. ::: Jab warp ke saare active threads ek hi branch lein, toh koi bhi path serialize nahi hota. State the difference between "resident" and "eligible" warps in one sentence. ::: Resident = SM pe loaded, registers reserved; eligible = resident aur stalled nahi, yaani is cycle mein actually schedulable.