Visual walkthrough — CUDA programming model basics

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

Is page mein parent note se zyada gehraai hai. Agar koi word naya lage, hum use use karne se pehle define karenge. Jo prerequisites hain woh Thread Warps and SIMT aur Streaming Multiprocessors mein milenge.

Step 1 — Ek worker ek job ke liye, ek line mein

KYA. Socho tumhe do lists ke numbers position-by-position add karne hain: C[i] = A[i] + B[i]. Ek hi insaan list ke neeche chalke kaam karne ki jagah, tum ek position ke liye ek worker hire karte ho. Position 0 ke liye worker, position 1 ke liye worker, aur aage bhi.

KYUN. Yahi GPU ka poora point hai: iske paas hazaaron chhote workers hote hain. Agar har worker exactly ek array slot handle kare, toh poora addition N sequential steps ki jagah ek parallel step mein khatam ho jaata hai. Lekin yeh tabhi kaam karta hai jab har worker ko apna slot number pata ho. Wahi number hai jo hume compute karna hai.

TASVEER. Neeche ki row memory mein array hai: boxes pe A[0], A[1], A[2], … likha hai. Har box ke upar ek worker khada hai. Abhi har worker ek hi sawaal pooch raha hai: "mera box kaun sa hai?"

Step 2 — Kyun hum saare workers ko ek flat line mein nahi rakh sakte

KYA. Tum soch sakte ho: bas workers ko 0, 1, 2, … , N-1 ek lambi line mein number kar do aur kaam khatam. GPU aisa nahi karne deta. Woh tumhe workers ko pehle fixed-size groups mein pack karne par majboor karta hai.

KYUN. Hardware (Streaming Multiprocessors) ek baar mein limited workers hi manage kar sakta hai, aur woh unhe bundles mein schedule karta hai. Isliye CUDA tumhe workers ko equal-size blocks mein organize karne par majboor karta hai (ek block mein 256 workers ho sakte hain). Blocks woh unit hai jo machine distribute karti hai. Yeh ek hardware constraint hai, koi choice nahi — aur isliye hi naive "ek flat line" numbering available nahi hai, aur flat number reconstruct karne ke liye ek formula chahiye.

TASVEER. Step 1 ki ek lambi row ab barabar chunks mein kaat di gayi hai. Har chunk ek block hai. Block ke andar workers ki numbering phir se 0 se shuru hoti hai — toh worker numbering restart hoti hai har block boundary pe.

.x ka matlab sirf "pehla dimension" hai — ek plain list ke liye bas yahi chahiye. Step 7 mein hum .y se milenge.

Step 3 — Do numbers jo har worker actually jaanta hai

KYA. Ek CUDA thread apni flat global position directly nahi jaanta. Hardware usse sirf do local facts batata hai:

• blockIdx.x — us box ka label jisme uska block baitha hai.
• threadIdx.x — woh us block ke andar kitna aage khada hai.

KYUN. Hardware saste mein har block ko ek ID stamp kar sakta hai aur har worker ko ek within-block offset de sakta hai. Ek million threads mein se har ek ko ek pre-computed global number dena wasteful hoga; do chhote local numbers dena aur arithmetic se baaki kaam karwana kaafi sasta hai. Toh formula ka kaam hai: (block label, position-in-block) ko ek flat address mein badlo.

TASVEER. Block number 2 pe zoom karo, jahan blockDim.x = 4 hai. Chaar workers ke tags hain threadIdx.x = 0,1,2,3. Block pe hi stamp hai: blockIdx.x = 2. Flat address jo hume chahiye woh har worker ke neeche halke se likha hai — aur woh unke tags se same nahi hai.

Step 4 — Un workers ko count karo jo tumse pehle aaye

KYA. Apna flat address dhundhne ke liye, pehle left mein saare earlier blocks ke har ek ko count karo. Agar har block mein blockDim.x workers hain, aur tum block number blockIdx.x mein ho, toh tumse pehle ke saare blocks mein workers ki sankhya hai: $\underset{\text{mere se pehle kitne poore blocks hain}}{\underbrace{\text{blockIdx.x}}}\times \underset{\text{har block mein workers}}{\underbrace{\text{blockDim.x}}}$

KYUN. Multiplication exactly "ek hi amount ko kai baar add karna" hai. Har earlier block ne blockDim.x workers diye. Block 0 ne ek block ka contribution diya, block 1 ne ek aur, … block blockIdx.x − 1 tak. Yeh blockIdx.x blocks hain, har ek ka size blockDim.x hai. Yah product tumhare block ka starting flat address hai.

TASVEER. blockDim.x = 4 ke saath: block 0 flat slots 0–3 cover karta hai, block 1 cover karta hai 4–7, block 2 shuru hota hai $2\times 4=8$ se. Pale-yellow bracket block 2 se pehle ke 8 workers span karta hai; arrow block 2 ke pehle slot pe jaata hai.

Step 5 — Apna step block ke andar add karo

KYA. Ab apne block ke andar tum personally kitni door khade ho woh add karo — woh hai threadIdx.x. Tumhara flat address block ka start plus tumhara local offset hai: $\text{globalIdx}=\underset{\text{mere block ka start}}{\underbrace{\text{blockIdx.x}\times \text{blockDim.x}}}+\underset{\text{block ke andar mera step}}{\underbrace{\text{threadIdx.x}}}$

KYUN. Block ka starting slot us worker ka hai jiska threadIdx.x = 0 hai. Agla worker (threadIdx.x = 1) ek slot aur right mein hai, aur aage bhi. Toh hum within-block position ko block ke start mein add karte hain. Woh ek sum hi flat global address hai — aur yeh unique hai, kyunki koi bhi do workers ek hi block aur ek hi within-block position share nahi karte.

TASVEER. Wahi block 2. blockStart = 8 (yellow). Worker jiska threadIdx.x = 3 hai use milta hai 8 + 3 = 11. Blue arrow: block start tak jump karo. Pink arrow: 3 aur aage step karo. Landing slot 11 light up hota hai.

Step 6 — Edge case: end ke baad bacha hua workers

KYA. Blocks fixed sizes mein aate hain, isliye workers ki total sankhya almost kabhi exactly N nahi hoti. Agar N = 1000 aur blockDim.x = 256 hai, toh humein chahiye $\text{numBlocks}=\lceil \frac{1000}{256}\rceil =4$ blocks, jisse milte hain $4\times 256=1024$ workers — lekin array mein sirf 1000 slots hain. Aakhri 24 workers (global indices 1000–1023) array ke end se aage point karte hain.

KYUN. Tum fractional block launch nahi kar sakte. Block count ko upar ki taraf round karna padta hai (ceiling ⌈ ⌉), isliye tum hamesha kam se kam enough workers launch karte ho — matlab kabhi kabhi zyada bhi. Woh extra workers, agar C[idx] = ... run karein, toh woh aisi memory read aur write karenge jo tumhari nahi hai: crash ya silent garbage. Fix hai ek simple guard line:

if (idx < N) { C[idx] = A[idx] + B[idx]; }

Jin workers ka globalIdx 1000 ya zyada hai woh simply kuch nahi karte.

TASVEER. 1024 workers ki line. Pehle 1000 real array boxes ke upar khade hain (blue). Aakhri 24 array ke end ke baad khaali jagah pe laTke hain (pink, hatched "koi box nahi"). Ek guard rail — if (idx < N) line — unhe andar jaane se rokti hai.

Step 7 — 2D case: images ke liye rows aur columns

KYA. Ek image pixels ka grid hai, line nahi. Ab har worker ko do pairs of numbers milte hain — ek horizontal direction ke liye (.x) aur ek vertical ke liye (.y) — aur hum wahi Step-5 formula do baar run karte hain: $\text{col}=\text{blockIdx.x}\times \text{blockDim.x}+\text{threadIdx.x}$ $\text{row}=\text{blockIdx.y}\times \text{blockDim.y}+\text{threadIdx.y}$

KYUN. Memory phir bhi ek flat line hai — width-wide image ka pixel (row, col) flat slot row × width + col pe rehta hai. Toh hum column compute karte hain Step 5 wali exact reasoning se (.x numbers use karke), row usi tarah compute karte hain (.y numbers use karke), phir 2D address ko 1D memory address mein fold karte hain: $\text{globalIdx}=\underset{\text{mere upar ke poore rows skip karo}}{\underbrace{\text{row}\times \text{width}}}+\underset{\text{apni row mein step karo}}{\underbrace{\text{col}}}$ width se multiply karna "tumhare upar ki har poori row skip karna" hai — bilkul wahi logic jaisa Step 4 mein blockDim.x se multiply karna tha, bas pixel-row scale pe.

TASVEER. Blocks ka checkerboard ek image tile kar raha hai. Ek worker highlighted hai: uske (col, row) arrows right aur down point karte hain; flat address row × width + col trace hota hai poori rows count karke (yellow) phir across step karke (pink).

Ek tasveer mein summary

Upar sab kuch ek single map mein collapse ho jaata hai: do local numbers andar, ek unique global address bahar, aur overshoot ke liye ek guard rail.

Related reading: Memory Hierarchy (jahan yeh indexed reads jaate hain), Parallel Algorithm Design (block sizes choose karna), GPU Architecture Overview, aur Deep Learning Frameworks jo exactly yeh launches tumhare liye generate karte hain.