Foundations — SIMT (single instruction multiple thread)

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

Yeh page har woh word aur symbol build karta hai jis par parent SIMT note tika hua hai. Agar wahan koi term bina explanation ke aaya ho, uski picture yahan milegi. Upar se neeche padho — har block sirf wohi cheezein use karta hai jo uske upar define ho chuki hain.

1. Thread — sabse chhota worker

Parent note baar baar kehta hai "32 threads ek hi step karte hain." Lekin is sentence ka koi matlab bane, pehle ek single worker ka picture banana padega.

Figure 1 dekho: har blue box ek thread hai. Har thread ka apna private scratch space hai (uske registers) aur apna private answer hai, lekin sab ek hi recipe padh rahe hain (same program). Yahi shared recipe poori trick hai — ise yaad rakhna.

2. threadIdx.x — "main kaun sa worker hoon?"

Har thread ko ek naam chahiye taaki woh jaane ki kaunsa data touch karna hai. CUDA mein woh naam hai threadIdx.x.

Topic ko yeh kyun chahiye? Kyunki sab threads ek hi code chalate hain. Thread #7 thread #3 se alag kaam sirf tabhi kar sakta hai jab apna badge number use karke alag array slot pick kare:

int val = data[threadIdx.x];   // thread 7 reads data[7]

Same recipe line, alag plate — yahi hai SIMT code ki ek line mein.

3. Lockstep — "sab saath step 4 par"

Parent note ka key phrase "execute in lockstep" isliye GPU hardware bachata hai. Agar 32 threads guaranteed same instruction par hain, toh sirf ek instruction-fetch unit chahiye jo 32 workers ko feed kare, na ki 32 alag fetch units.

Figure 2 do possible worlds ka contrast karta hai: left par har worker ka apna drummer hai (mehenga, jaise CPU with many cores); right par ek drummer 32 ke liye beat karta hai (sasta — yahi SIMT hai). Is saste version ki jo cost pay karni padti hai woh baad mein aati hai, jab workers alag steps par hona chahein (yahi divergence hai, neeche).

4. Warp — 32 ka bundle

32 kyun, 8 ya 100 kyun nahi? Teen seedhe reasons jo parent list karta hai:

• Ek fetch bahut saaron ko feed karta hai: ek instruction decode 32 arithmetic lanes chalata hai — overhead aur kaam ka achha ratio.
• 32 ek power of two hai ($32=2^5$), toh "main apne warp ke andar kaun sa lane hoon?" ek sasta bit operation hai: threadIdx.x & 31. (& bitwise AND hai, aur 31 se masking low 5 bits rakhti hai — yeh sirf 32 se divide karne ke baad ka remainder hai.)
• Memory hardware 32-wide chunks pasand karta hai (tum ise coalescing ke roop mein miloge).

5. Active vs masked — har lane ka on/off switch

Kyunki 32 threads ek instruction share karte hain, agar ek thread baithna chahiye toh kya hoga? Hardware us worker ke slot ko bundle se skip nahi kar sakta — slot phir bhi tick karta hai. Iski jagah woh lane ko mask karta hai: lane instruction chalata hai lekin result throw away kar deta hai.

Figure 3 mask ko 32 boxes ki strip ki tarah dikhata hai: green = rakha gaya, gray = discard kiya gaya. Yahi mechanism hai parent ke phrase "masked off if diverged" ke peeche. Koi lane kabhi truly remove nahi hota — sirf chup karaya jaata hai.

6. Divergence — do-recipe problem

Figure 4 se socho:

• Phase A: "if" lanes active hain (green), "else" lanes masked hain (gray). Warp branch A chalata hai.
• Phase B: masks flip ho jaate hain. "else" lanes active hain, "if" lanes masked. Warp branch B chalata hai.

Toh warp time A + time B kharach karta hai, chahe har individual thread ko sirf ek hi chahiye tha. Yahi parent ka formula hai:

$T_{\text{diverged}}=T_A+T_B\text{vs}{T}_{\text{converged}}=\max (T_A,T_B)$

Converged mein max kyun? Agar sab threads same branch lete hain, koi masking nahi hoti — warp us ek branch ko ek baar chalata hai, jitna bhi time us branch ko lage. General case mein max ka bas matlab hai "sabse lamba single path", kyunki converged warp sirf ek hi path chalata hai.

7. Latency aur throughput — wait karna vs flow karna

Do words jo parent constantly use karta hai. Inhe alag rakho:

8. Occupancy — kitne warps "oven mein" hain

Topic ko yeh kyun important lagta hai: jab ek warp apni 400-cycle memory latency mein wait karta hai, scheduler doosre warps chalata hai. Jitne zyada warps ready hain ($W$), utna zyada wait hide kiya ja sakta hai. Yahi parent ka rule hai:

$W\ge L\cdot T$

Words mein padho: ready warps ki sankhya kam se kam latency times throughput honi chahiye. $L=400$ cycles hide karne ke liye $T=1$ instruction/cycle par tumhe $W\ge 400$ warps ke barabar independent kaam chahiye. Kyunki ek single SM bahut kam hold karta hai, GPUs bahut saare SMs use karte hain aur high occupancy demand karte hain — isliye occupancy matter karti hai.

9. Streaming Multiprocessor (SM) — workshop

Upar sab kuch — warps, masks, occupancy — ek SM ke andar hota hai. Jab parent kehta hai "warps per SM", SM yahi workshop hai. Iska poora anatomy dekhne ke liye Streaming Multiprocessor jaao.

Pieces topic ko kaise feed karte hain

Har arrow ka matlab hai "target samajhne se pehle source samajhna zaroori hai." Poora graph SIMT model mein funnel hota hai — parent topic mein.

Har foundation aage kahan use hoti hai

• Thread + warp → CUDA Thread Hierarchy (thread → warp → block → grid)
• Masks + divergence → Warp Divergence
• 32-wide access → GPU Memory Coalescing
• Latency vs throughput, occupancy → GPU Occupancy
• Lockstep vs vector lanes → SIMD vs SIMT Comparison
• Kyun bahut saare warps, deep pipelines nahi → GPU vs CPU Architecture aur Instruction-Level Parallelism

Equipment checklist

Khud test karo — right side cover karo aur reveal karne se pehle answer do.