Visual walkthrough — ROCm - OpenCL alternatives

6.2.15 · D2 · Hardware › GPU Architecture › ROCm - OpenCL alternatives

Yeh page parent note ka ek central result pictures mein derive karta hai: ek formula jo GPU pe har single worker ko uska apna unique home address deta hai. Hum abhi woh formula likhne wale nahi hain — uske har symbol ko pehle earn karna hoga. End tak aap dekh paoge ki poori cheez bas "boxes ke andar boxes ginना" hai, aur aap jaanoge ki har edge pe kya hota hai — last box, leftover box, aur empty case.

Pehle parent padho agar nahi padha: ROCm / OpenCL alternatives.

Step 1 — Problem: identical jobs ka ek bada pile

KYA. Soch lo ki hamare paas numbers ki ek giant list hai aur hum har ek pe same chhota sa operation karna chahte hain — jaise, do arrays ko element by element add karna: C[i] = A[i] + B[i]. Maano list mein $16$ elements hain (hum chhota number use karte hain taaki picture mein fit ho; real note mein over a million use hote hain).

KYO. Ek GPU ek fast worker nahi hota. Yeh hazaaron slow-ish workers hota hai. Use karne ke liye, hume har worker ko exactly ek job dena hoga aur ensure karna hoga ki do workers ek hi job na pakdein aur koi job skip na ho. Toh sabse pehla sawaal yeh hai: hum job numbers $0,1,2,\dots ,15$ kaise distribute karein taaki har worker jaane ki unka wala kaunsa hai?

PICTURE. Yeh rahe jobs ka raw pile — numbered boxes ki ek single row.

Step 2 — Pehla symbol: the global ID

KYA. Aap bol sakte ho: "Bas inhe $0$ se $15$ tak ek lambi line mein number karo aur ho gaya." Woh flat position — poore problem mein worker ka address — itna important hai ki isko abhi hi ek naam aur symbol milna chahiye.

KYO. Hum ise abhi symbol dete hain kyunki yahi woh quantity hai jo hum ultimately dhundh rahe hain. Is page ki har baaki cheez sirf $\text{global\_id}$ ko un chhote pieces se reconstruct karne ke liye exist karti hai jo hardware actually deta hai. Is box-number picture ko pakde raho; har baad wala symbol iske against measure kiya jaayega.

PICTURE. Wahi row phir se, ab uske neeche $\text{global\_id}$ naam likha hua.

Step 3 — Hum sabko ek saath number kyun nahi kar sakte

KYA. Toh yahan kyun nahi rukke aur $\text{global\_id}$ directly use karein? Kyunki hardware humein woh flat numbers distribute nahi karne deta. Yeh pehle line ko equal chunks mein kaat'ne par insist karta hai. Har chunk ko work-group kehte hain.

KYO. Real silicon workers ko fixed-size batches mein execute karta hai — NVIDIA inhe $32$ at a time chalaata hai (ek warp), AMD inhe $64$ at a time chalaata hai (ek wavefront). Toh hum force hain ki flat line ko equal groups mein slice karein jo hardware ke batch size se match karein.

PICTURE. Wahi $16$ boxes, lekin ab $4$ ke groups mein slice kiye gaye. (Hum $64$ ki jagah $4$ pick karte hain sirf taaki screen pe fit ho — logic identical hai.)

$\underset{\text{total jobs}=16}{\underbrace{G}}÷\underset{\text{jobs per group}=4}{\underbrace{L}}=\underset{\text{how many groups}}{\underbrace{4}}$

Step 4 — Ek ruler ki jagah do rulers

KYA. Hardware kabhi worker ko uska $\text{global\_id}$ directly nahi batata. Iski jagah yeh har worker ko do chhote addresses deta hai:

• Uska group ID — woh kis box mein rehta hai ($0,1,2,3$).
• Uska local ID — us box ke andar uski seat ($0,1,2,3$).

KYO. Hardware sirf itna jaanta hai "tum group #1 ke andar worker #2 ho". Yeh Step 3 ke batching ka natural consequence hai. Toh humein do chhote rulers diye jaate hain aur hume ek bade ruler ($\text{global\_id}$) ko khud reconstruct karna hota hai.

PICTURE. Wahi strip, do baar label ki gayi: bottom ruler groups count karta hai, top ruler har group ke andar $0$ se restart karta hai.

Notice karein ki top ruler har group ke start pe reset to 0 hota hai. Wahi reset hi wajah hai ki hume formula chahiye: local ID $2$ chaar alag-alag baar appear karta hai, toh yeh akele address nahi ho sakta.

Step 5 — Formula banana: pehle full boxes skip karo, phir andar step karo

KYA. Ab hum finally do rulers se $\text{global\_id}$ assemble karte hain. Worker local ID 2 in group 1 lo. Uska global number find karne ke liye hum do cheezein karte hain:

1. Hamare pehle ke har box ko skip karo. Humare aage $1$ full box hai, jisme $L=4$ workers hain, toh hum $1\times 4=4$ addresses skip karte hain.
2. Apni seat par walk karo. Apna local seat number add karo, $2$.

Yeh deta hai $4+2=6$. Step 4 se check karo: haan, global $6$ label wala box exactly local-2-of-group-1 hai.

KYO. Yeh mixed base mein likh'a number padhne se zyada kuch nahi hai. "Group $1$, seat $2$, with $4$ seats per group" wahi idea hai jaise "$1$ ten aur $2$ ones $=12$", except har place $10$ ki jagah $L$ worth hai.

PICTURE. Red arrow pehle full groups ke upar jumps karta hai (woh multiply hai), phir seat pe steps karta hai (woh add hai).

$\text{global\_id}=1\times 4+2=6✓$

Step 6 — Ulta jaana: global number se box aur seat recover karna

KYA. Kabhi kabhi hamare paas global number hota hai aur hum doosri direction chahte hain: kaun sa box, kaun si seat? Woh bas integer division aur remainder use karke formula ko ulta solve karna hai:

$\text{group\_id}=\lfloor \frac{\text{global\_id}}{L}\rfloor ,\text{local\_id}=\text{global\_id}\mathrm{mod}L$

KYO. Integer division $\lfloor \cdot \rfloor$ jawaab deta hai "mere neeche kitne whole boxes fit hote hain?" aur remainder $\mathrm{mod}$ jawaab deta hai "last full box ke baad main kitna aage hun?" — exactly Step 5 ke skip-then-step ka inverse.

PICTURE. Worker global $6$: $4$ se divide karne pe quotient $1$ milta hai (box) aur remainder $2$ (seat). Dono operations Step 5 ke arrow ke do parts ki taraf point karte hain.

$\lfloor \frac{6}{4}\rfloor =1(\text{group}),6\mathrm{mod}4=2(\text{local})$

Step 7 — Edge case A: sabse pehla aur sabse aakhri worker

KYA. Extremes check karo taaki kuch kabhi undefined na ho.

• Pehla worker: group $0$, local $0$ → $0\times 4+0=0$. Correct: line ka front.
• Aakhri worker: group $3$, local $3$ → $3\times 4+3=15$. Correct: final box, final seat, jo ki $G-1=16-1=15$ hai.

KYO. Hume prove karna hoga ki formula boundaries cover karta hai, kyunki off-by-one errors yahan classic GPU bug hain: ek box zyada door padh lo aur tum us memory ko touch karte ho jo tumhari nahi.

PICTURE. Dono ends lit up — global range exactly $0$ se $15$ tak run karti hai, koi gap nahi, koi overshoot nahi.

Step 8 — Edge case B: jab total evenly divide nahi hota (the "guard" clause)

KYA. Maano hamare paas $G=14$ real jobs hain lekin hardware $L=4$ ke groups force karta hai. Tab $14/4=3.5$, jo whole number nahi hai. Runtime launch ko $4$ groups $=16$ workers tak round up karta hai — lekin sirf $14$ jobs exist karti hain. Workers $14$ aur $15$ ke paas koi real job nahi hai.

KYO. Hardware sirf whole groups launch kar sakta hai. Toh hum hamesha kuch extra workers spawn karte hain aur unhe array ke end ke baad read karne se pehle switch off karna padta hai. Wahi exactly har kernel mein if (i < n) guard hai:

int i = blockIdx.x * blockDim.x + threadIdx.x; // = global_id
if (i < n) y[i] = a * x[i] + y[i];             // <-- the guard

PICTURE. Do extra workers (global $14,15$) ghosts ki tarah draw kiye gaye hain; red guard line unhe drop karti hai taaki woh kuch na karein.

Step 9 — Edge case C: degenerate launches, including the empty one

KYA. Check karne ke liye teen limiting shapes.

• Ek giant group ($L=G=16$): tab group_id hamesha $0$ hoga, aur $\text{global\_id}=0\times 16+\text{local\_id}=\text{local\_id}$. Do rulers ek mein collapse ho jaate hain — formula phir bhi hold karta hai.
• Groups of one ($L=1$): tab local_id hamesha $0$ hoga, aur $\text{global\_id}=\text{group\_id}\times 1+0=\text{group\_id}$. Phir se ek ruler; formula phir bhi hold karta hai.
• Koi jobs hi nahi ($G=0$): line empty hai, $\text{Num Groups}=\lceil 0/L\rceil =0$, toh zero workers launch hote hain. Formula kabhi evaluate nahi hota — koi worker nahi hai jo $\text{global\_id}$ hold kare — aur kuch galat nahi hota. Ek empty workload ek valid, quietly-does-nothing launch hai.

KYO. Ek formula jis par aap trust karte ho use apne extremes survive karne chahiye. $L=G$ aur $L=1$ par "boxes inside boxes" picture ek single ruler mein flat ho jaati hai; $G=0$ par picture blank hai aur Step 8 ka guard (if (i < n) with $n=0$) ka matlab hai koi worker kabhi nahi chalta. Har case mein arithmetic agree karta hai.

PICTURE. Top: ek box jisme saare $16$ hain. Middle: $16$ boxes jisme ek ek hain. Bottom: empty line — koi boxes nahi, koi workers nahi.

Ek-picture summary

Upar sab kuch ek idea hai: jobs ki ek flat line secretly boxes ka ek grid hai, aur address formula woh hai jisse tum dono views ke beech fold karte ho. Full boxes skip karne ke liye multiply karo, apni seat mein step karne ke liye add karo; wapas jaane ke liye divide aur remainder karo.