SIMT (single instruction multiple thread): Question bank

6.2.4 · D5 · Hardware › GPU Architecture › SIMT (single instruction multiple thread)

Yeh SIMT (single instruction multiple thread) ka misconception-hunting deep dive hai. Neeche har item ek aisi jagah hai jahan smart log quietly SIMT ko galat samajh lete hain. Question padho, apna jawab zor se bolo, phir reveal karo.

Shuru karne se pehle, woh words aur symbols jo hum baar baar use karte hain. Inhe ek baar padh lo — poora page inhi par based hai:

warp — 32 threads ka fixed group jise hardware ek saath schedule karta hai, ek baar mein ek instruction.
lockstep — ek warp ke saare active threads same cycle mein same instruction execute karte hain.
predication mask — ek 32-bit flag, har thread ke liye ek bit (1 = active, 0 = masked off / paused).
$N$ — ek branch par ek warp ke andar liye gaye distinct control-flow paths ki sankhya. Agar sab threads agree karein, $N = 1$ ; agar saare 32 threads alag destinations chunein, $N = 32$ . Divergence cost $N$ ke saath scale karti hai.
$L$ — kisi operation ki latency, clock cycles mein measure ki gayi (e.g. ek memory load ke liye $L \approx 400$ cycles ho sakte hain data wapas aane se pehle).
$T$ — throughput target, instructions issued per cycle mein measure kiya gaya jo tum SM ko sustain karna chahte ho (often $T = 1$ : har cycle mein ek instruction issue karte raho).
$W$ — ek SM par switch karne ke liye available resident warps ki sankhya.

$L$ , $T$ aur $W$ ke naam rakhne ke baad, neeche wala latency-hiding rule clearly padhta hai: rate $T$ par issue karte rehne ke liye jabki har in-flight op $L$ cycles leta hai, tumhe chahiye $W \geq L \cdot T (warps needed to fill the latency gap).$

Teen pictures jo dimaag mein rakhni hain

Is page ke almost har trap mein actually teen alag pictures ke beech confusion hai. Inhe ek baar study kar lo aur tum aadhe items seedha dekh ke jawab de paoge.

(1) Divergence — ek lane/mask picture. Kaunse lanes kisi given cycle mein active hain?

Figure — SIMT (single instruction multiple thread)

(2) Coalescing — ek address picture. Memory mein 32 lanes kahan point karte hain, aur us se kitne hardware transactions (memory segments) lagte hain.

(3) Occupancy — ek scheduler picture. Kitne warps resident hain taaki scheduler ke paas hamesha ek ready warp ho jabki doosre latency $L$ par wait kar rahe hain.

True ya false — justify karo

Ek warp mein hamesha exactly 32 active threads hote hain.

False. Ek warp mein allocate 32 lanes hote hain, lekin kuch permanently inactive ho sakte hain (100 threads ka ek block apne last warp mein 28 idle lanes chhodta hai) ya divergence ke dauran temporarily masked off ho sakte hain — figure 1 mein greyed lanes dekho.

Ek warp ke saare 32 threads ek program counter share karte hain, isliye woh kabhi alag instructions par nahi ho sakte.

Jo log sochte hain us tarah se False hai. Woh ek scheduled instruction stream share karte hain, lekin divergence ke dauran hardware track karta hai ki kaun se threads kis path par hain aur paths ko ek ek karke issue karta hai — isliye kisi given cycle mein sirf kuch lanes active hote hain (figure 1).

SIMD aur SIMT identical machine behaviour produce karte hain, bas alag alag marketing names hain.

False. SIMD lanes sab ko execute karna padta hai (koi per-lane branching nahi); SIMT lanes ko individually mask kiya ja sakta hai, jisse threads alag control-flow paths le sakten hain. SIMD vs SIMT Comparison dekho — core baat yeh hai ki SIMT ek per-lane mask add karta hai, jo SIMD mein nahi hai.

Agar ek warp do branches mein diverge ho, total time dono branch times ka average hoga.

False. Divergent branches serially chalte hain, isliye time roughly

T_{A} + T_{B}

(sum) hai, na average na max. Figure 1 mein dono phases time mein stacked dikhe hain.

Fully converged code (saare 32 threads same branch lete hain) mein zero divergence penalty hoti hai.

True. Ek hi path lene par, mask poore waqt all-ones rehta hai aur

N = 1

path execute hota hai — koi serialization nahi, full lane utilization.

Zyada occupancy ka matlab hamesha faster kernels hota hai.

False. Occupancy sirf usi point tak help karta hai jahan latency hide ho jaati hai (jab

W \geq L \cdot T

ho jaaye); uske baad, zyada warps same registers/cache ke liye compete karte hain aur hurt kar sakte hain. GPU Occupancy ka core idea yeh hai ki occupancy latency hide karne ka zariya hai, apne aap mein goal nahi — figure 3 dekho.

Warp size (32) CUDA C++ standard se define hoti hai aur tumhe ise hard-code karna chahiye.

False. 32 ek NVIDIA hardware architectural constant hai, koi language guarantee nahi; portable code warpSize query karta hai instead of assume karne ke.

Memory coalescing ek divergence problem hai disguise mein.

False. Yeh dono independent hain. Coalescing ek instruction mein 32 threads ke touch kiye addresses ke baare mein hai (figure 2); divergence is baare mein hai ki kaun se threads ek instruction execute karte hain (figure 1). Ek perfectly converged warp phir bhi badly uncoalesced ho sakta hai. GPU Memory Coalescing dekho.

Kyunki GPUs SIMT use karte hain, unhe 400-cycle memory latency hide karne ke liye out-of-order execution use karna padta hai.

False. GPUs latency doosre ready warps mein switch karke hide karte hain (massive multithreading, figure 3), CPU-style speculation ya out-of-order issue se nahi. Yeh GPU vs CPU Architecture mein headline contrast hai: CPUs latency bade caches aur reordering se hide karte hain, GPUs many warps se hide karte hain.

1024 threads ka ek block ek giant lockstep unit ki tarah run karta hai.

False. 1024 threads = 32 warps; har warp independently schedule hota hai. Sirf ek warp ke andar ke 32 threads lockstep mein hain. CUDA Thread Hierarchy dekho, jahan thread → warp → block → grid key ladder hai.

Galti dhundho

"Saare lanes use karne ke liye, 100 threads launch karo taaki sab busy rahein." — kya galat hai?

100, 32 ka multiple nahi hai. 4th warp mein sirf 4 threads hain (12.5% utilization); 28 lanes waste hain. 32 ka multiple launch karo.

"if (tid % 2 == 0) kaam ko alternate karta hai taaki dono branches parallel mein chalein, speed double ho jaye." — kya galat hai?

Divergence branches ko serially chalata hai, parallel mein nahi (

N = 2

paths). Har phase mein aadhe lanes masked off hote hain, isliye tum slower ho jaate ho (dono paths ka sum), 2× speedup nahi — figure 1.

"Registers per thread badhaana threads ko faster banata hai, isliye hamesha register use maximize karo." — kya galat hai?

Thread per zyada registers reduce karta hai kitne warps

W

ek SM par fit ho sakte hain (registers ek shared budget hain), occupancy girti hai aur latency hiding hurt hoti hai. Yeh ek trade-off hai, free win nahi — register limit GPU Occupancy mein teen occupancy ceilings mein se ek hai.

"data[threadIdx.x * 32] ek nice contiguous read hai kyunki hum warp size se multiply kar rahe hain." — kya galat hai?

Yeh ek strided pattern hai. Har thread alag memory segment par land karta hai → ek packed load ki jagah kaafi alag transactions (figure 2, bottom row). Contiguous ka matlab hai data[threadIdx.x].

"Divergence penalty 32× fixed hai kyunki ek warp mein 32 threads hain." — kya galat hai?

32× worst case hai (har thread ek unique path,

N = 32

). Actual penalty liye gaye distinct paths ki sankhya

N

hai,

N = 1

(koi nahi) se

N = 32

(maximum) tak.

"Alag warps ke threads jo same if branch hit karte hain woh ek doosre ko converge karne mein help karte hain." — kya galat hai?

Convergence ek per-warp property hai. Divergence sirf ek warp ke 32 lanes ke andar matter karta hai; doosre warps independently schedule hote hain aur is warp ke mask se interact nahi karte. Warp Divergence dekho.

"Ek warp scheduler us warp ko issue karta hai jo pehle finish hua, isliye ordering deterministic hai." — kya galat hai?

Scheduler koi bhi ready warp choose karta hai; choice runtime latencies par depend karti hai aur across runs deterministic guaranteed nahi hai, isliye race conditions ko explicitly avoid karna padta hai.

Why questions

32 threads per warp kyun hain, say 30 kyun nahi?

32 power of two hai, isliye tid & 31 se lane index cheaply extract hota hai, memory addresses fixed-size segments mein neatly fall karte hain coalescing ke liye (current NVIDIA GPUs par segment typically 32 ya 128 bytes hota hai — exact size architecture-dependent hai), aur ek fetch/decode unit cleanly 32 parallel ALU lanes par map hota hai.

SIMT programmer ko plain scalar code likhne kyu deta hai instead of explicit vector code ke?

Har thread ke apne logical registers aur program counter hote hain, isliye tum likhte ho "ek thread kya karta hai" aur hardware us instruction ko 32 lanes par broadcast karta hai — unlike SIMD, jahan tumhe haath se vectorize karna padta hai.

Memory latency (

L \approx 400

cycles) poore GPU ko kyun stall nahi karti?

Jab ek warp memory par wait karta hai, scheduler doosre ready warps se issue karta hai (figure 3). Enough resident warps (

W \geq L \cdot T

) ALUs ko busy rakhte hain — latency hide hoti hai, eliminate nahi. GPU Occupancy dekho.

Divergence penalty serialization se kyun aati hai, wasted silicon se nahi?

Dono branch phases run karne padte hain kyunki kuch threads ko har ek ki zaroorat hai; hardware 32 lanes par do alag instructions simultaneously run nahi kar sakta, isliye woh ek ke baad ek chalata hai — cost cycles mein time hai (figure 1).

High occupancy zaruri kyun hai lekin good performance ke liye sufficient kyun nahi?

Occupancy enough warps provide karta hai latency hide karne ke liye (

W \geq L \cdot T

), lekin agar un warps mein divergence, uncoalesced memory, ya bandwidth bottleneck hai, toh woh phir bhi slowly chalenge. GPU Occupancy aur Instruction-Level Parallelism ka takeaway: latency hide karna kaafi levers mein se ek hai, aur independent instructions wala ek single warp bhi apne aap kuch latency hide kar sakta hai.

GPUs many SMs kyun use karte hain instead of ek bahut deep pipeline ke?

Ek single SM ~32–64 resident warps tak cap hota hai, jo deep latency hide karne ke liye zaroori hundreds se bahut kam hai; kaafi SMs mein kaam spread karna real throughput ke liye total warp count deliver karta hai.

__shfl_sync jaise warp-level primitives kuch branches kyun replace kar sakte hain?

Woh data directly lanes ke beech exchange karte hain bina control-flow divergence ke, isliye woh computation jo branch karti (aur

N > 1

phases mein serialize hoti), saare 32 lanes par ek single converged instruction ban jaati hai.

Edge cases

Jab ek block mein sirf 1 thread ho toh warp utilization ka kya hota hai?

Woh ek thread ek full warp occupy karta hai: 1/32 ≈ 3.1% utilization. Baaki 31 lanes poore time masked off rehti hain — bahut hi wasteful (imagine figure 1 mein sirf ek lane lit hai).

Agar ek warp ka har thread alag path par diverge kare (saare 32 alag)?

Warp

N = 32

phases mein serialize hota hai, har phase mein exactly ek active lane — worst-case 32× slowdown aur har phase mein 1/32 utilization.

Agar ek branch condition thread-independent ho (runtime par saare 32 threads ke liye same) toh kya hota hai?

Divergence nahi hoti chahe if ho: saare lanes condition identically evaluate karte hain, saath mein ek path lete hain (

N = 1

), aur mask all-ones rehta hai. Compiler/hardware ek uniform branch dekhta hai.

Jab ek kernel itne registers use kare ki SM par sirf ek warp fit ho, occupancy kya hoti hai?

Occupancy = 1 / (max warps per SM), e.g. 1/64 ≈ 1.6%. Switch karne ke liye almost koi warps nahi (

W

bahut chhota hai), memory latency ab hide nahi ho sakti aur SM stall karta hai.

Kya hota hai jab blockDim 32 ka multiple nahi hota — say 48 threads?

Tumhe 2 warps milte hain: ek full (32) aur ek half-full (16 active, 16 idle). Doosre warp ke idle lanes execution slots waste karte hain har cycle mein jab woh warp run karta hai.

Agar SM ke saare warps ek saath memory par wait kar rahe hoon toh kya hota hai?

Switch karne ke liye koi ready warp nahi hota, isliye SM sach mein stall karta hai jab tak memory response wapas nahi aata — yeh wahi failure mode hai jise high occupancy (

W \geq L \cdot T

) prevent karne ke liye exist karta hai. Warp Divergence aur GPU Occupancy dekho.

Jab stride bahut large ho jaaye (random access) toh coalescing efficiency kahan tak jaati hai?

Yeh apne floor tak pahunchti hai: har thread ek alag memory segment pull karta hai, isliye bytes-transferred balloon karta hai jabki bytes-used ek warp ke data par fixed rehta hai. Exact floor architecture ke segment size par depend karta hai (figure 2, bottom row).

Recall Jaane se pehle quick self-check

Divergence penalty kis cheez ke saath scale karti hai? ::: Ek warp ke andar liye gaye distinct control-flow paths ki sankhya, yaani $N$ , 1 (koi nahi) se 32 (worst case) tak. Latency hiding kis cheez ke saath scale karta hai? ::: SM per resident, ready warps $W$ ki sankhya — $W \geq L \cdot T$ se capture hota hai, jahan $L$ cycles mein latency hai aur $T$ target instructions per cycle hai. Coalescing kis par depend karti hai? ::: Ek single memory instruction mein 32 threads ka address pattern (figure 2), divergence se independent.