Noesis
6.2.2 · D5 · HinglishGPU Architecture

Question bank — Streaming multiprocessors (SM)

1,988 words9 min read↑ Read in English

6.2.2 · D5 · Hardware › GPU Architecture › Streaming multiprocessors (SM)

Yeh ek misconception hunt hai. Neeche har line ek claim, question, ya edge case hai Streaming Multiprocessor ke baare mein. Left side padho, apna jawab pehle decide karo — phir right side dekho. Har jawab tumhe reasoning deta hai, sirf verdict nahi — wahi reasoning hai jo tumhe yaad rakhni chahiye.

Shuru karne se pehle, do words jinhe hum baar baar use karenge, seedhi bhaasha mein:

Agar inme se koi bhi fuzzy lage, pehle Warp-scheduling aur Occupancy-optimization revisit karo — neeche ke traps assume karte hain ki tum dono ko picture kar sakte ho.

True or false — justify karo

SM mein zyada cores hone se hamesha proportionally faster execution hoti hai.
False — cores tab kaam ki nahi hote jab warps memory par stall ho jaayein; agar bahut kam resident warps hon toh cores idle rahte hain chahe kitne bhi hon. Throughput wahan limit hoti hai jahan pehle kuch khatam ho: work, warps, ya bandwidth.
Ek block, ek baar SM ko assign ho jaane ke baad, mid-flight mein kisi kam busy SM par migrate kar sakta hai.
False — ek block shuru se ant tak exactly ek SM par hi chalta hai; uske registers aur shared memory physically usi SM par rehte hain, isliye migration ka matlab hoga private state copy karna, jo hardware kabhi nahi karta.
100% occupancy peak performance guarantee karta hai.
False — occupancy sirf yeh measure karta hai kitne warps chal sakte hain, yeh nahi ki unhe zaroori hai ya nahi. 50% occupancy wala compute-bound kernel 100% wale se better perform kar sakta hai; extra warps sirf tab help karte hain jab tumhare paas actually latency ho jo hide karni ho.
Ek hi warp ke do threads ek hi cycle mein alag alag instructions execute kar sakte hain.
False — yahi core SIMT rule hai: ek instruction ek cycle mein saare 32 lanes ko issue hoti hai. Alag branches ko lanes ko mask off karke paths ko ek ke baad ek chalane se handle kiya jaata hai, simultaneously nahi.
Threads per block badhane se hamesha occupancy badhti hai.
False — bade blocks zyada registers aur shared memory per block consume karte hain, jo resident blocks ki sankhya kum kar sakta hai aur is tarah total resident warps kum ho jaate hain. Occupancy ek resource tug-of-war hai, koi monotonic dial nahi.
Shared memory aur L1 cache ek hi cheez ke do naam hain.
False — yeh often ek hi physical SRAM par hote hain, lekin shared memory explicitly tumhare dwara index ki jaati hai bina kisi tag lookup ke, jabki L1 hardware-managed hai tags aur eviction ke saath. Ek hi bricks, alag rulebook aur alag latency.
Registers mein shared memory ki tarah bank conflicts ho sakte hain.
False — registers har thread ke liye private hote hain, isliye koi bhi do threads kabhi ek hi register slot ke liye contend nahi karte. Bank conflicts ek shared-memory phenomenon hai, jahan warp ke 32 lanes 32 banks ko hit karte hain aur collisions serialize ho jaate hain.
Ek stalled warp wait karte waqt SM ke cycles waste karta hai.
False (yahi toh GPUs ka poora point hai) — jab ek warp stall hota hai, scheduler instantly doosre eligible warp se issue karta hai, toh SM kaam karta rehta hai. Stall sirf tab waste hota hai jab switch karne ke liye koi aur eligible warp na ho.
Agar kernel mein kabhi branch nahi hai, toh warp divergence nahi ho sakta.
True — divergence strictly ek warp ke lanes ke alag alag control-flow paths lene se hoti hai. Koi data-dependent branch nahi matlab saare 32 lanes hamesha agree karte hain, toh koi serialized paths nahi.
Thread per thread register count double karne se maximum resident warps aadhe ho jaate hain.
Zyaatar true — register file fixed size ki hoti hai, isliye warps-per-SM roughly registers-per-thread ke inverse mein scale karta hai jab tak koi doosra cap (max warps, max blocks) pehle binding limit na ban jaaye.

Error dhundho

"Mere block mein 40 threads hain, isliye exactly ek warp banta hai jisme 8 threads unused hain."
Error count mein hai: 40 threads ko do warps chahiye (⌈40/32⌉ = 2). Pehla warp full hai (32), doosre mein 8 active lanes hain aur 24 permanently masked-off lanes slots waste kar rahe hain — yahi wajah hai ki block sizes ko 32 ke multiples mein rakhna chahiye.
"Mere block mein 8 warps hain, toh main 300-cycle memory stall puri tarah hide kar sakta hoon."
8 warps sirf lagbhag 8 cycles ki latency cover karte hain per scheduler pass, 300 nahi. 300-cycle stall hide karne ke liye dozens of resident warps chahiye kai blocks mein — ek block ke 8 warps kaafi nahi hain.
"Maine __shared__ use kiya toh threads automatically ek doosre ki writes immediately dekh lete hain."
Barrier missing hai: writes ke baad __syncthreads() ke bina, ek thread apne neighbor ka slot us se pehle read kar sakta hai jab neighbor ne likha hi na ho. Shared visibility possible hai, lekin ordering tumhari zimmedari hai.
"Bank conflicts se bachne ke liye maine thread i ko address i * 32 access karaya."
Yeh worst case hai: 32 (bank count) se multiply karne par har lane same bank par map ho jaati hai, forcing a 32-way serialized access. Tumhe chahiye ki i % 32 banks mein spread ho, yaani consecutive addresses, strided-by-32 nahi.
"Mera kernel compute-bound hai, toh bhi mujhe occupancy maximum tak badha leni chahiye."
Agar tum compute-bound ho, extra warps help nahi karte — cores already har cycle busy hain. Occupancy push karna hurt bhi kar sakta hai registers per thread kum karke aur slower memory par spills force karke.
"Ek SM ek waqt ek block chalaata hai, phir agla load karta hai."
Ek SM multiple resident blocks simultaneously rakhta hai (jitne registers, shared memory, aur block-count caps allow karein). Blocks ke beech concurrency exactly isi liye hoti hai taaki latency hide karne ke liye kaafi warps ikattha ho sakein.
"Agar ek if-else dono branches chalti hain, toh if...else ki cost double ho jaati hai chahe saare threads ek hi taraf jaayein."
Agar warp ke saare 32 lanes ek hi branch lete hain, sirf wahi ek path execute hoti hai — koi divergence nahi, koi doubling nahi. 2× cost sirf tab aati hai jab ek warp ke andar lanes branches mein split ho jaayein.

Why questions

Warp size 32 kyun hai, say 40 kyun nahi?
Yeh ek hardware constant hai jo SIMT datapath width aur scheduler design se tied hai; poori scheduling machinery, masking, aur shuffle instructions 32-lane warps assume karte hain. Isliye block sizes 32 ke multiples mein honi chahiye taaki dead lanes na hon.
GPUs ko thousands of threads ki zaroorat kyun hai jab ek SM ek cycle mein sirf kuch instructions issue karta hai?
Kyunki inका kaam hai latency hiding: 400–800 cycle memory stall ke saath, tumhe ready warps ka ek deep pool chahiye switch karne ke liye jab data in flight ho. Excess threads pipeline ko busy rakhne ka fuel hain, sab ko ek saath issue karne ke liye nahi.
Variables ko memory ki jagah registers mein kyun rakhein?
Registers ~1-cycle access ke hain, jabki global memory ke liye hundreds of cycles lagte hain. Thread ki working set registers mein rakhne se latency bilkul khatam ho jaati hai — trade-off yeh hai ki heavy register use occupancy kum karta hai.
Shared memory tiled matrix multiply jaise algorithms ko kyun enable karta hai?
Yeh block ke threads ko cooperatively ek tile ek baar load karne aur use ~100× global-memory speed par kai baar reuse karne deta hai, kai slow global reads ko ek shared read plus fast on-chip reuse mein convert karta hai — yahi wajah hai ki Tensor-cores pipelines shared memory se feed karte hain.
Block size badhane se performance kyun kum ho sakti hai register limit se neeche bhi?
Bade blocks scheduler ke liye granularity kum karte hain aur shared-memory-per-block ya max-blocks cap hit kar sakte hain, isliye kam blocks fit hote hain aur SM mein overall kam resident warps ho jaate hain. Bada matlab automatically zyada bhara nahi hota.
SM mein CUDA cores se alag Special Function Units kyun hain?
Transcendentals (sin, exp, sqrt) ek general ALU par dozens of cycles lete; ek dedicated SFU pipeline unhe 1–2 cycles mein karta hai, main cores ko ordinary arithmetic ke liye free karta hai instead of unhe clog karne ke.

Edge cases

Ek SM jo zero warps chala raha hai uski occupancy kya hai?
Occupancy hai — ek valid lekin useless state, matlab SM idle hai. Yeh kernel launches ke beech ya jab koi block assign nahi hua ho tab hota hai.
Ek block itni shared memory request karta hai jo SM ke paas physically hai hi nahi — kya hota hai?
Kernel launch fail ho jaata hai (launch-time error), kyunki SM ek bhi aisa block host nahi kar sakta. Resource requests per-SM limits ke against validate hoti hain pehle kisi bhi thread ke run karne se.
Ek warp jahan saare 32 lanes if lete hain aur koi bhi else nahi leta — kya divergence hai?
Nahi — divergence ke liye warp ke andar aapas mein disagreement zaroori hai. Ek unanimous branch single path ko full efficiency se execute karta hai, bilkul waise jaise branch tha hi nahi.
32 mein se sirf 1 thread ek warp mein active hai (baaki masked hain) — lane utilization kitni efficient hai?
Lagbhag — warp phir bhi ek poora 32-lane issue slot occupy karta hai lekin ek lane ka kaam karta hai. Yeh divergence ki extreme cost hai aur block sizes ka 32 se badly misaligned hone ki bhi.
Itne zyada blocks launch karna jitne SMs hain se bhi zyada — kya kuch tootta hai?
Kuch nahi toottta; Grid Distribution Unit surplus blocks queue karta hai aur unhe SMs free hone par dispatch karta hai. Blocks deliberately independent design kiye jaate hain taaki yeh kisi bhi order ya wave mein chal sakein.
Ek memory-bound kernel jisme high occupancy hai lekin tiny arithmetic intensity hai — kya occupancy ise save karegi?
Sirf us point tak jahan memory bandwidth saturate hoti hai; uske baad, zyada warps sirf lambi queues mein wait karte hain. Occupancy latency hide karta hai, lekin jo DRAM deliver nahi kar sakti uski bandwidth manufacture nahi kar sakta.
Recall Band karne se pehle quick self-test

Agar ek warp stall ho aur SM phir bhi full throughput par chale, toh kaunsi condition ne yeh possible banaya? ::: Kum se kum ek aur eligible warp tha scheduler ke switch karne ke liye — latency hiding ke liye ek spare, ready warp chahiye. Us ek resource ka naam batao jiska over-use ek typical kernel mein occupancy sabse directly crush karta hai. ::: Registers per thread — register file fixed hoti hai, isliye heavy per-thread register use SM ko resident warps se bhookha rakhta hai.