Foundations — CUDA programming model basics

6.2.13 · D1 · Hardware › GPU Architecture › CUDA programming model basics

Yeh D1 Foundations ka deep-dive hai CUDA programming model basics ke liye. Kisi bhi kernel ko touch karne se pehle, tumhe har woh symbol pata hona chahiye jo parent note mein use hota hai. Hum har ek ko zero se banate hain: simple words → ek picture → kyun yeh topic uske bina kaam nahi kar sakta.

1. Host aur Device — do alag duniyaan

Picture: do islands socho. Host island par kuch clever workers hain. Device island par hazaaron simple workers hain. Ek single patli bridge — PCIe bus — unhe connect karti hai. Ek island par koi bhi doosre island ki library se seedha koi book nahi le sakta; tumhe usse bridge ke paar ship karna padta hai.

Kyun topic ko yeh chahiye: har CUDA program ek story hai data ko us bridge ke paar move karne ki, Device island par kaam karne ki, phir results wapas ship karne ki. Parent note ke steps "Copy Host → Device" aur "Copy Device → Host" tab hi samajh mein aate hain jab tum dekho ki sach mein do memories hain jo ek doosre ko dekh nahi sakti. Device ke apne internal libraries ka full map dekhne ke liye Memory Hierarchy dekho.

2. Thread — sabse chhota worker

Picture: warehouse ke floor par ek single dot, ek slip of paper pakde hue jisme likha hai "Main element 7 handle karta hoon." Usse element 6 ya 8 ke baare mein kuch nahi pata — sirf apna.

Kyun topic ko yeh chahiye: GPU ka poora point loops likhna band karna hai. "Yeh kaam N baar karo, ek ke baad ek" kehne ki bajaye, tum kehte ho "yeh ek element ke liye kaam hai — ab N threads spawn karo aur sab ek saath chalao." Thread usi idea ka atom hai.

3. Block — threads ki ek team

Picture: dots ko 256 workers ke square huddles mein cluster karo. Har huddle ek block hai. Ek huddle ke andar ke workers ek doosre ko saste mein notes pass kar sakte hain; alag huddles ke workers nahi kar sakte.

Kyun topic ko yeh chahiye: cooperation bade scale par mehenga hai. Saare 10,000 workers ko sabse baat karne dena chaos hai. Unhe blocks mein grouping karne ka matlab hai ki sirf woh ~256 workers jo collaborate karna chahte hain ek table share karte hain. Ek poora block ek physical Streaming Multiprocessor (SM) par schedule hota hai.

4. Grid — saari teams ek saath

Picture: zoom out karo. Poora warehouse floor, huddles se tiled, woh grid hai. Ek kernel launch = ek grid.

Kyun topic ko yeh chahiye: grid scalability layer hai. Tum apni problem ko "itne blocks, itne threads each" ke roop mein describe karte ho, aur GPU yeh figure out karta hai ki un blocks ko jitne bhi SMs us par hain unpar kaise daale — ek chhote card par 20 SMs, ek bade par 100+. Tum hardware hardcode nahi karte. Yahi parent note ka logical→physical mapping hai: Grid → GPU, Blocks → SMs, Threads → CUDA cores.

5. Dot-notation symbols: blockIdx, blockDim, threadIdx

Yeh teen built-in variables hain jisse ek thread discover karta hai ki woh kaun hai. Har ek ka ek .x hai (aur 2D ke liye .y). .x ko "horizontal component" padho.

Picture: tum ek worker ho. blockIdx.x = 2 matlab "Main 3rd huddle mein hoon (huddles 0 se count hote hain)." blockDim.x = 256 matlab "Har huddle mein 256 log hain." threadIdx.x = 5 matlab "Main apne huddle mein 6th person hoon."

6. Global index formula — har piece khud kamao

Ab hum parent note ki central equation assemble karte hain. Hum ise build karte hain, sirf quote nahi karte.

YEH kya karta hai: ek thread ke local seat number ko poori grid mein $0$ se $N-1$ tak ek unique global number mein convert karta hai.

KYUN, step by step (figure s03 dekho):

• Har block blockDim.x threads hold karta hai. Toh mere block se pehle, blockIdx.x full blocks the.
• Mere block se pehle aane wale threads ki sankhya = blockIdx.x × blockDim.x. Yahi mere huddle ka starting offset hai — figure mein red bracket.
• Huddle ke andar meri seat, threadIdx.x, add karo aur main apne unique slot par pahunch jaata hoon.

Ek concrete check. Block $2$, block size $256$, seat $5$: $\text{globalIdx}=2\times 256+5=517.$ Thread $517$ C[517] handle karta hai. Koi do threads kabhi collide nahi karte, kyunki har (blockIdx, threadIdx) pair ek alag sum deta hai — bilkul (page × 256) + line ki tarah jo ek book mein ek line ko uniquely locate karta hai.

Multiplication kyun aur kuch fancy kyun nahi? Hume har block ko indices ki ek contiguous, non-overlapping range claim karni hai. Block number ko block size se multiply karna precisely hai "saare earlier blocks ke slots ko skip karo." Koi trigonometry nahi, koi exotic tool nahi — sirf counting.

7. ceil, aur ceiling-division trick

$N$ elements ko size $T$ ke blocks se cover karne ke liye, hum kitne blocks launch karte hain?

Picture: tumhare paas $N$ eggs hain aur cartons hain jo $T$ eggs hold karte hain. Agar aakhri carton sirf partly full bhi ho, tumhe phir bhi woh poora carton chahiye. Neeche rounding karna eggs ko homeless kar dega.

+ T - 1 trick kyun? Computers integer division remainder chop karke karte hain (neeche rounding). Upar rounding force karne ke liye, pehle numerator ko $T-1$ se nudge karo. Agar $N$ evenly divide hota hai, toh woh nudge agale block mein spillover karne ke liye kaafi nahi hai; agar koi remainder hai, toh hai. Example: $N=1000$, $T=256$: $\frac{1000+255}{256}=\frac{1255}{256}=4\text{(integer division)}.$ Chaar blocks. Lekin $4\times 256=1024>1000$ — isliye agle symbol ki zaroorat hai.

8. Boundary check if (idx < N)

Chaar blocks 1024 threads dete hain, lekin humare paas sirf 1000 elements hain. Threads $1000$ se $1023$ tak extra hain — unke paas touch karne ke liye koi valid egg nahi hai.

Picture: aakhri carton mein 24 khaali egg slots hain. Agar ek extra worker array ke slot $1000$ mein haath daale jo $999$ par ruk jaata hai, toh woh woh memory grab karta hai jo hamaari nahi hai — ek crash ya garbage.

Yeh parent note ka Mistake 1 hai. Yeh optional nahi hai — kyunki block sizes fixed hain, numBlocks × T almost always $N$ ko overshoot karta hai.

9. Memory table mein symbols: cycle, KB/MB/GB

Parent ki memory table yeh units use karti hai — inhe define karo taaki table readable ho.

Picture — kyun closer = faster: registers compute core ke andar rehte hain (millimetres of wire → 1 cycle). Global VRAM off-chip hai (centimetres of wire + ek memory controller → hundreds of cycles). Host RAM PCIe bridge ke paar hai (metres of signalling → thousands of cycles). Distance literally time cost karta hai kyunki electrical signals ko travel karne mein time lagta hai. Poori tiered picture Memory Hierarchy mein hai.

Prerequisite map

Baayein taraf har foundation CUDA execution model ko feed karta hai. Agar koi bhi node unclear hai, parent note padhne se pehle uska section dobara dekho. Aage ke related paths: Thread Warps and SIMT, Streaming Multiprocessors, GPU Architecture Overview.

Equipment checklist

Right side cover karo aur zor se jawab do. Agar atko, toh woh section tumhara agla stop hai.