ROCm - OpenCL alternatives: Worked examples — Hinglish deep dive

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

Yeh page ek drill hai. Parent topic ne tumhe bataya tha kya hai work-item model aur HIP porting. Yahan hum unhe har case class se guzarte hain — clean divisions, ragged leftovers, degenerate 1-element launches, 2D aur 3D grids, aur woh exam-twist jahan naive formula galat answer deta hai.

Shuru karne se pehle, ek promise: har symbol kamaya hua hai. Agar tumne kabhi get_global_id, "work-group", ya index formula nahi dekha, to pehle do callouts padho — woh poori vocabulary ek single picture se build karte hain.

Definition Teen IDs, boxes ki ek single row se

Socho tumhe do arrays of numbers add karni hain. Tumhare paas ek badi army of tiny workers hai, ek worker per output number. Har worker ko jaanna chahiye: "mera output slot kaunsa hai?"

Global ID worker ka seat number hai jo bilkul shuru se count hota hai: $0, 1, 2, \dots$ poori army mein. Likha jaata hai get_global_id(0).
Army ko equal work-groups (squads) mein divide kiya jaata hai kyunki hardware physically threads ko fixed-size batches mein run karta hai (NVIDIA inhe warps of 32 kehta hai, AMD inhe wavefronts of 64 kehta hai). Local ID worker ka seat hai apne squad ke andar: yeh har squad ke start pe $0$ se reset hota hai. Likha jaata hai get_local_id(0).
Group ID squad ka number hai: $0, 1, 2, \dots$ . Likha jaata hai get_group_id(0).

Figure 1 neeche workers ki ek single row par teeno IDs dikhata hai (squad size 4). Green box worker global_id 9 = group 2 × 4 + local 1 hai.

Scenario matrix

Har launch jo tum ever configure karoge woh in cells mein se kisi ek mein aata hai. Neeche ke examples har ek ko kam se kam ek baar cover karte hain.

#	Case class	Kya special hai	Example
A	Clean 1D divide	$G_{x}$ , $L_{x}$ ka exact multiple hai	Ex 1
B	Ragged 1D divide	$L_{x}$ , $G_{x}$ ko divide nahi karta — padding + guard chahiye	Ex 2
C	Degenerate input	$G_{x} = 1$ (single element) ya empty $G_{x} = 0$	Ex 3
D	2D NDRange	rows × columns par index math	Ex 4
E	3D NDRange	volume, poora $(x, y, z)$ formula	Ex 5
F	Vendor mapping	same code, AMD wavefront 64 vs NVIDIA warp 32	Ex 6
G	Real-world word problem	image blur, launch config choose karo	Ex 7
H	Exam twist	naive `G/L` answer galat hai; ceiling division chahiye	Ex 8

Worked examples

Throughout, hum kernel code ke andar do short names use karte hain: i us worker ka global id hai (yani int i = get_global_id(0);) aur n real data elements ki sankhya hai (yani n = G_data, sahi array length, kisi bhi padding se pehle). Output array C hai (upar defined).

Worked example Ex 1 — Cell A: clean 1D divide

Statement. Tum ek vector add launch karte ho $G_{x} = 1, 048, 576$ elements ke saath, local size $L_{x} = 64$ (AMD wavefront). Dhundho (a) work-groups ki sankhya, (b) worker get_global_id(0) = 4160 kis group aur local seat mein hai.

Forecast: group count ka andaza lagao — kya yeh sixteen thousand ke paas hoga ya million ke paas? Aur andaza lagao: kya worker 4160 group ke start ke paas hai ya end ke?

Groups ki sankhya $= G_{x} / L_{x} = 1, 048, 576/64 = 16, 384$ . Yeh step kyun? Kyunki $L_{x} = 64$ , $G_{x}$ ko exactly divide karta hai (Cell A), poori army equal squads mein split hoti hai kuch bhi leftover nahi — plain division kaam karta hai.
Worker 4160 locate karne ke liye tie-together formula ko invert karo. Hamare paas hai $global_id = group_id \times 64 + local_id$ . Toh $group_id = ⌊ 4160/64 ⌋ = 65$ (floor = round down) aur $local_id = 4160 mod 64 = 0$ (mod = remainder). Yeh step kyun? Squad size se divide karne par pata chalta hai kitne full squads tum jump kar chuke ho (quotient = group), aur remainder batata hai squad ke andar tum kahan land kiye (local seat). Yahi "kaunsi shelf, kaunsa slot" idea hai.

Answer: 16,384 groups; worker 4160 = group 65, local seat 0 (squad 65 ka bilkul pehla worker).

Verify: global id rebuild karo: $65 \times 64 + 0 = 4160$ . ✓ Aur $16, 384 \times 64 = 1, 048, 576$ = total. ✓ Units: pure counts, dimensionless. ✓

Worked example Ex 2 — Cell B: ragged divide (padding + guard)

Statement. Ab $G_{data} = 1, 000, 000$ real elements hain, local size $L_{x} = 256$ . Lekin $256$ , $1, 000, 000$ ko divide nahi karta. Kitne work-groups launch karne chahiye, aur padded NDRange kitna bada hoga? Kernel ko guard if (i < n) (yaad karo i = get_global_id(0), n = G_data) kyun chahiye?

Forecast: kya tum fewer groups launch karoge (leftovers throw away karo) ya ek extra group (kuch workers waste karo)? Padhne se pehle andaza lagao.

Tumhe round UP karna hai, DOWN nahi. Groups $= ⌈ 1, 000, 000/256 ⌉$ . Plain division $3906.25$ deta hai; ceiling $3907$ deta hai. Yeh step kyun? Agar tum down round karke 3906 groups launch karo toh $3906 \times 256 = 999, 936$ workers launch honge aur 1,024 output elements kabhi compute hi nahi honge — silent data corruption. Tum hamesha round up karo taaki har real element cover ho.
Padded NDRange $= 3907 \times 256 = 1, 000, 192$ . Yeh real 1,000,000 se $192$ extra workers hain. Yeh step kyun? Hardware sirf full squads run kar sakta hai; tum partial squad launch nahi kar sakte. Toh launched army data se badi hogi.
Guard clause if (i < n) C[i] = A[i] + B[i]; jahan i = get_global_id(0), n = G_data = 1000000, aur C output array hai. Global id $1, 000, 000 \dots 1, 000, 191$ wale workers kuch nahi karenge. Yeh step kyun? Woh 192 padding workers ke paas i >= n hai, toh woh array slots par point karte hain jo exist hi nahi karte. Guard ke bina woh out of bounds read/write karenge — crash ya garbage. if unhe harmlessly idle bana deta hai.

Answer: 3907 groups, padded NDRange 1,000,192, 192 idle guarded workers ke saath.

Verify: $3907 \times 256 = 1, 000, 192 \geq 1, 000, 000$ ✓ aur $3906 \times 256 = 999, 936 < 1, 000, 000$ (proves 3906 bahut kam hai) ✓. Padding $= 1, 000, 192 - 1, 000, 000 = 192$ ✓.

Worked example Ex 3 — Cell C: degenerate inputs

Statement. $L_{x} = 64$ ke saath do edge cases: (a) $G_{x} = 1$ (ek single scalar op). (b) $G_{x} = 0$ (empty input, jaise ek filtered array empty nikli). Kya launch hoga?

Forecast: ek element ke liye, kya tum ek worker launch karte ho ya 64 ka poora squad? Zero elements ke liye, kya 0 groups launch karna legal hai?

Case $G_{x} = 1$ . Ceiling groups $= ⌈ 1/64 ⌉ = 1$ . Tum ek full squad of 64 launch karte ho, lekin sirf worker i = get_global_id(0) = 0 guard if (i < n) pass karta hai with n = 1; baaki 63 (jahan i >= 1) idle hain. Yeh step kyun? Hardware apne wavefront se chhota squad run nahi kar sakta. Ek real element phir bhi ek poore squad ki cost lagata hai — yeh reminder hai ki tiny launches inefficient hoti hain.
Case $G_{x} = 0$ . Ceiling groups $= ⌈ 0/64 ⌉ = 0$ . Tumhe kuch nahi launch karna chahiye aur early return karna chahiye. Yeh step kyun? $0/64 = 0$ exactly, toh koi squad zaroorat nahi. 0-sized NDRange ke saath kernel launch karna ya toh no-op hai ya error, runtime par depend karta hai — safe pattern yeh hai ki host-side if (n == 0) return; likho jahan n = G_data hai.

Answer: $G_{x} = 1 \Rightarrow$ 1 group (63 idle workers); $G_{x} = 0 \Rightarrow$ 0 groups (launch skip karo).

Verify: $⌈ 1/64 ⌉ = 1$ ✓, active workers $= 1$ ✓, idle $= 63$ ✓. $⌈ 0/64 ⌉ = 0$ ✓.

Worked example Ex 4 — Cell D: 2D NDRange

Statement. Ek $2048 \times 1024$ image process karo (width $G_{x} = 2048$ , height $G_{y} = 1024$ ) 2D work-group $L_{x} = 16, L_{y} = 16$ ke saath. Group grid dhundho, aur width 2048 ke row-major array mein global coords $(x, y) = (100, 50)$ wale pixel ka linear buffer index dhundho.

Forecast: total work-groups kitne honge — hundreds ya thousands? Aur $(100, 50)$ ka flat index guess karo.

Figure 2 neeche ek chhota $8 \times 4$ NDRange dikhata hai jo $2 \times 2$ work-groups (red borders) mein tiled hai; green square woh pixel hai jise flatten kiya ja raha hai. Step 2 padhte waqt ise dekhna.

Groups per axis. $G_{x} / L_{x} = 2048/16 = 128$ , $G_{y} / L_{y} = 1024/16 = 64$ . Dono cleanly divide hote hain (Cell A style, 2D mein). Yeh step kyun? Har axis independently chop hoti hai — Figure 2 grid ko $16 \times 16$ squares mein tiled dikhata hai (clarity ke liye $2 \times 2$ drawn). Total groups $= 128 \times 64 = 8192$ .
Flat buffer index. Row-major image ke liye, index $= y \times width + x = 50 \times 2048 + 100 = 102, 500$ . Yeh step kyun? Memory ek 1D line hai, 2D sheet nahi. Pixel $(x, y)$ dhundne ke liye tum $y$ full rows skip karte ho (har row width 2048 ki hai), phir current row mein $x$ step karte ho. Same "shelf phir slot" logic jo 1D global-id formula mein tha, ab rows shelves ki tarah hain.

Answer: $128 \times 64 = 8192$ groups; pixel $(100, 50)$ flat index 102,500 par hai.

Verify: $128 \times 64 = 8192$ ✓; total work-items $8192 \times (16 \times 16) = 2, 097, 152 = 2048 \times 1024$ ✓; index $50 \times 2048 + 100 = 102, 500$ ✓.

Worked example Ex 5 — Cell E: 3D NDRange (full formula)

Statement. $64 \times 64 \times 64$ voxel cube ke upar ek physics sim, work-group $4 \times 4 \times 4$ . Total work-items, total groups, aur $x$ fastest stored cube mein voxel $(x, y, z) = (3, 2, 1)$ ka flat index dhundho.

Forecast: total work-items — kya yeh roughly quarter-million hai? Kisi early voxel jaise $(3, 2, 1)$ ka flat index guess karo.

Groups per axis $= 64/4 = 16$ har $x, y, z$ axis pe. Total groups $= 1 6^{3} = 4096$ . Yeh step kyun? Same clean divide, teesri axis tak extend hoti hai. Kuch naya nahi — bas ek aur factor.
Total work-items $= 6 4^{3} = 262, 144$ . Yeh step kyun? Har voxel per ek worker, toh army size poore cube ka volume hai: teen axis lengths ko multiply karo $64 \times 64 \times 64$ . Hum ise ab compute karte hain kyunki hum ise group count ko sanity-check karne ke liye use karenge — agar squads cube ko perfectly tile karein, toh (groups) × (workers per squad) is volume ke barabar hona chahiye, jo neeche Verify line confirm karta hai.
Flat index, $x$ fastest $= z \cdot (64 \cdot 64) + y \cdot 64 + x = 1 \cdot 4096 + 2 \cdot 64 + 3 = 4227$ . Yeh step kyun? "Shelf phir slot" ko teen levels tak generalize karo: $z$ whole slices skip karo (har ek $64 \times 64$ ka), phir $y$ rows, phir $x$ columns. Multipliers un sab cheezon ki sizes hain jo current axis se faster hain.

Answer: 262,144 work-items; 4096 groups; voxel $(3, 2, 1)$ flat index 4227 par.

Verify: $1 6^{3} = 4096$ ✓; $6 4^{3} = 262144$ ✓; group tiling check: $4096 \times 4^{3} = 4096 \times 64 = 262, 144$ = volume ✓; $1 \cdot 4096 + 2 \cdot 64 + 3 = 4227$ ✓.

Worked example Ex 6 — Cell F: same code, do vendors

Statement. Tum identical kernel $G_{x} = 1, 048, 576$ elements par run karte ho. Tum $L_{x}$ = ek wavefront/warp choose karte ho taaki koi partial batch lanes waste na kare. AMD (wavefront 64) vs NVIDIA (warp 32) par, kitne groups each, aur total hardware batches kitne?

Forecast: kya NVIDIA zyada groups launch karta hai ya kam? Kya "more groups" ka matlab "slower" hai?

AMD: $L_{x} = 64 \Rightarrow$ groups $= 1, 048, 576/64 = 16, 384$ ; har group = exactly 1 wavefront, toh 16,384 hardware batches. Yeh step kyun? Humne $L_{x}$ ko AMD ki physical batch width (64) ke barabar set kiya taaki har work-group exactly ek wavefront se map ho — koi lane idle na rahe. Kyunki 64 cleanly $G_{x}$ ko divide karta hai (Cell A), plain division directly group count deta hai.
NVIDIA: $L_{x} = 32 \Rightarrow$ groups $= 1, 048, 576/32 = 32, 768$ ; har group = 1 warp, toh 32,768 hardware batches. Yeh step kyun? Portable source byte-for-byte identical hai; sirf launch config change hoti hai physical batch width match karne ke liye. HIP/OpenCL ISA difference hide karta hai, lekin tum phir bhi $L_{x}$ ko vendor ki batch size se tune karte ho idle lanes avoid karne ke liye.
Note: NVIDIA par double batches ka matlab double slow nahi hai — har NVIDIA warp 32 lanes wide hai, AMD ka 64. Total lanes processed = $32, 768 \times 32 = 16, 384 \times 64 = 1, 048, 576$ dono taraf. Yeh step kyun? Throughput lanes mein hai, batch count mein nahi. Same total work, alag tarah se slice kiya.

Answer: AMD 16,384 groups; NVIDIA 32,768 groups; equal total lanes = 1,048,576.

Verify: $16384 \times 64 = 32768 \times 32 = 1, 048, 576$ ✓.

Worked example Ex 7 — Cell G: real-world word problem

Statement. Tumhe ek AMD GPU par $1920 \times 1080$ (Full HD) photo blur karni hai, ek work-item per pixel. Ek sensible 2D work-group choose karo aur padded NDRange aur idle-worker count report karo. (Full HD height axis par 16 ka clean multiple nahi hai.)

Forecast: kaunsi axis ko padding chahiye hogi — width 1920 ya height 1080? Wasted workers ki sankhya guess karo.

Figure 3 neeche Full HD frame dikhata hai: blue block real pixels hain, bottom ke saath red strip padded, guarded, idle workers hai. Step 4 ke dauran ise dekhna.

$L_{x} = L_{y} = 16$ choose karo (256 items/group = 4 AMD wavefronts, ek common sweet spot). Yeh step kyun? 256, 64-wide wavefront ka multiple hai, toh ek group ke andar koi lane waste nahi hogi.
Width: $1920/16 = 120$ exactly — clean, koi padding nahi. Yeh step kyun? Hum har axis ko raggedness ke liye alag check karte hain. Kyunki 16, 1920 ko zero remainder ke saath divide karta hai ( $1920 mod 16 = 0$ ), width 120 groups mein perfectly tile ho jaati hai aur extra column of workers ki zaroorat nahi.
Height: $1080/16 = 67.5$ — ragged. Round up: $⌈ 1080/16 ⌉ = 68$ groups tall, padded height $= 68 \times 16 = 1088$ . Yeh step kyun? Sirf ek axis par Cell B again: 16, 1080 ko divide nahi karta ( $1080 mod 16 = 8$ ), toh hum leftover 8 rows cover karne ke liye round up karte hain. Figure 3 neeche guarded, idle workers ki extra 8-pixel-tall red strip dikhata hai.
Padded NDRange $= 1920 \times 1088 = 2, 088, 960$ workers. Real pixels $= 1920 \times 1080 = 2, 073, 600$ . Idle workers $= 1920 \times (1088 - 1080) = 1920 \times 8 = 15, 360$ . Yeh step kyun? Sirf bottom strip padded hai; 1920 columns mein se har ek ke 8 extra rows hain (Figure 3 mein red band). Woh extra workers guard if (i < n) fail karte hain aur kuch nahi likhte.

Answer: work-group $16 \times 16$ ; padded NDRange $1920 \times 1088$ ; 15,360 guarded idle workers.

Verify: $1920/16 = 120$ exact ✓; $⌈ 1080/16 ⌉ = 68$ ✓; padded $= 1920 \times 1088 = 2, 088, 960$ ✓; idle $= 2, 088, 960 - 2, 073, 600 = 15, 360$ ✓.

Worked example Ex 8 — Cell H: exam twist (naive answer galat hai)

Statement. Exam question: "Ek kernel mein $G_{x} = 1, 500, 000$ aur $L_{x} = 512$ hai. Ek student likhta hai groups = G/L = 1500000/512 = 2929 (integer division). Kya 2929 correct hai? Agar nahi, toh sahi count do aur bug explain karo."

Forecast: abhi ek answer commit karo — kya 2929 sahi hai, ek kam hai, ya ek zyada?

Exact quotient compute karo. $1, 500, 000/512 = 2929.6875$ . Integer division truncate karta hai $⌊ 2929.6875 ⌋ = 2929$ par (floor = round down). Yeh step kyun? C++ mein integer / floor karta hai. Student ne silently fractional part drop kar diya.
Coverage check. $2929 \times 512 = 1, 499, 648$ . Yeh $1, 500, 000 - 1, 499, 648 = 352$ real elements chodta hai jo kabhi process nahi honge — classic off-by-one-group bug. Yeh step kyun? Fractional $0.6875$ group actually 352 workers ka ek real partial squad hai; ise drop karne se last 352 outputs corrupt ho jaate hain.
Correct formula: ceiling division. Integer code mein tum ise $groups = (G + L - 1) / L$ likho. Plug in: $⌊(1, 500, 000 + 511) /512 ⌋ = ⌊ 1, 500, 511/512 ⌋ = 2930$ . Yeh step kyun? Hum floats par ceil() safely use nahi kar sakte (exact multiples par rounding error aa sakti hai), toh hum integer trick $(G + L - 1) / L$ use karte hain. $L - 1$ add karna kyun kaam karta hai? Agar $G$ already $L$ ka clean multiple hai, toh $G + L - 1$ next multiple se just short hai, toh flooring exact quotient par land karta hai — koi spurious extra group nahi. Agar $G$ mein koi bhi remainder hai, toh $L - 1$ add karna total ko next multiple boundary se push karta hai, toh flooring exactly ek se upar round karta hai — leftover partial squad capture karte hue. Ek line mein yeh hai "floor, lekin jab bhi remainder ho toh bump up."
Coverage confirm karo. $2930 \times 512 = 1, 500, 160 \geq 1, 500, 000$ , toh ab har real element ke paas ek worker hai; last $160$ workers padding hain, guard if (i < n) C[i] = A[i] + B[i]; se silence kiye gaye hain with n = 1500000. Yeh step kyun? Round up karna tabhi safe hai jab hum padding ko bhi guard karein — warna woh 160 extra workers out of bounds run karenge. Ceiling division aur if (i < n) guard ek matched pair hain.

Answer: 2929 galat hai (ek group kam, 352 elements drop ho gaye). Correct = 2930 using (G + L - 1) / L, plus if (i < n) guard.

Verify: $2929 \times 512 = 1, 499, 648 < 1, 500, 000$ (bahut kam) ✓; $(1500000 + 511) //512 = 2930$ ✓; $2930 \times 512 = 1, 500, 160 \geq 1, 500, 000$ ✓; dropped $= 352$ ✓; padding $= 160$ ✓.

Recall drill

Recall Kaunsa formula (group_id, local_id) ko wapas global_id mein convert karta hai?

global_id = group_id × local_size + local_id ::: saare full squads skip karo, phir apne local seat se andar jaao.

Recall Work-group count compute karte waqt tum UP kyun round karte ho?

Kyunki neeche round karne se real elements ka leftover partial squad uncomputed reh jaata — silent data loss. ::: (G + L - 1) / L use karo aur if (i < n) se guard karo.

Recall AMD vs NVIDIA par, same source compile kyun hota hai lekin tum local_size phir bhi change karte ho?

HIP/OpenCL ISA hide karta hai, lekin tum $L_{x}$ ko physical batch width (wavefront 64 / warp 32) se match karte ho taaki koi lane idle na baithe. ::: Dono taraf same total lanes.

Ceiling-division bug in a launch config

G/L (floor) launch karna instead of ceil(G/L) final partial group drop kar deta hai — hamesha round up karo.

Padding workers ko silence kiya jaana chahiye

if (get_global_id(0) < n) guard se taaki out-of-bounds lanes kuch na karein.