ROCm - OpenCL alternatives: Worked examples — Hardware deep dive

This page is a drill. The parent topic told you what the work-item model and HIP porting are. Here we run them through every case class — clean divisions, ragged leftovers, degenerate 1-element launches, 2D and 3D grids, and the exam-twist where the naive formula gives the wrong answer.

Before we start, one promise: every symbol is earned. If you have never seen get_global_id, a "work-group", or the index formula, read the first two callouts — they build the whole vocabulary from a single picture.

Definition The three IDs, from a single row of boxes

Imagine you must add two arrays of numbers. You have a huge army of tiny workers, one worker per output number. Each worker needs to know: "which output slot is mine?"

The global ID is the worker's seat number counting from the very start: $0, 1, 2, \dots$ across the whole army. Written get_global_id(0).
The army is split into equal work-groups (squads) because the hardware physically runs threads in fixed-size batches (NVIDIA calls them warps of 32, AMD calls them wavefronts of 64). The local ID is the worker's seat inside its own squad: it resets to $0$ at the start of each squad. Written get_local_id(0).
The group ID is the squad's number: $0, 1, 2, \dots$ . Written get_group_id(0).

Figure 1 below shows all three IDs on a single row of workers (squad size 4). The green box is worker global_id 9 = group 2 × 4 + local 1.

The scenario matrix

Every launch you will ever configure falls into one of these cells. The examples below hit each one at least once.

#	Case class	What is special	Example
A	Clean 1D divide	$G_{x}$ is an exact multiple of $L_{x}$	Ex 1
B	Ragged 1D divide	$L_{x}$ does not divide $G_{x}$ — need padding + guard	Ex 2
C	Degenerate input	$G_{x} = 1$ (single element) or empty $G_{x} = 0$	Ex 3
D	2D NDRange	index math on rows × columns	Ex 4
E	3D NDRange	volume, full $(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, choose the launch config	Ex 7
H	Exam twist	the naive `G/L` answer is wrong; ceiling division	Ex 8

Worked examples

Throughout, we use two short names inside kernel code: i is that worker's global id (i.e. int i = get_global_id(0);) and n is the number of real data elements (i.e. n = G_data, the true array length, before any padding). The output array is C (defined above).

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

Statement. You launch a vector add over $G_{x} = 1, 048, 576$ elements with local size $L_{x} = 64$ (AMD wavefront). Find (a) the number of work-groups, (b) which group and local seat holds worker get_global_id(0) = 4160.

Forecast: guess the group count — is it near sixteen thousand or near a million? And guess: is worker 4160 near the start of a group or the end?

Number of groups $= G_{x} / L_{x} = 1, 048, 576/64 = 16, 384$ . Why this step? Because $L_{x} = 64$ divides $G_{x}$ exactly (Cell A), the whole army splits into equal squads with nothing left over — plain division works.
Invert the tie-together formula to locate worker 4160. We have $global_id = group_id \times 64 + local_id$ . So $group_id = ⌊ 4160/64 ⌋ = 65$ (floor = round down) and $local_id = 4160 mod 64 = 0$ (mod = remainder). Why this step? Dividing by the squad size tells you how many full squads you jumped over (the quotient = group), and the remainder is where you landed inside the squad (local seat). This is exactly the "which shelf, which slot" idea.

Answer: 16,384 groups; worker 4160 = group 65, local seat 0 (the very first worker of squad 65).

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

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

Statement. Now $G_{data} = 1, 000, 000$ real elements, local size $L_{x} = 256$ . But $256$ does not divide $1, 000, 000$ . How many work-groups must you launch, and how big is the padded NDRange? Why does the kernel need the guard if (i < n) (recall i = get_global_id(0), n = G_data)?

Forecast: will you launch fewer groups (throw away leftovers) or one extra group (waste some workers)? Guess before reading.

You must round UP, not down. Groups $= ⌈ 1, 000, 000/256 ⌉$ . Plain division gives $3906.25$ ; ceiling gives $3907$ . Why this step? If you rounded down to 3906 groups you would launch $3906 \times 256 = 999, 936$ workers and 1,024 output elements never get computed — silent data corruption. You always round up so every real element is covered.
Padded NDRange $= 3907 \times 256 = 1, 000, 192$ . That is $192$ extra workers beyond the real 1,000,000. Why this step? The hardware only runs full squads; you cannot launch a partial squad. So the launched army is bigger than the data.
Guard clause if (i < n) C[i] = A[i] + B[i]; where i = get_global_id(0), n = G_data = 1000000, and C is the output array. Workers with global id $1, 000, 000 \dots 1, 000, 191$ must do nothing. Why this step? Those 192 padding workers have i >= n, so they point at array slots that do not exist. Without the guard they read/write out of bounds — a crash or garbage. The if makes them idle harmlessly.

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

Verify: $3907 \times 256 = 1, 000, 192 \geq 1, 000, 000$ ✓ and $3906 \times 256 = 999, 936 < 1, 000, 000$ (proves 3906 is too few) ✓. Padding $= 1, 000, 192 - 1, 000, 000 = 192$ ✓.

Worked example Ex 3 — Cell C: degenerate inputs

Statement. Two edge cases with $L_{x} = 64$ : (a) $G_{x} = 1$ (a single scalar op). (b) $G_{x} = 0$ (empty input, e.g. a filtered array came out empty). What launches?

Forecast: for one element, do you launch one worker or a whole squad of 64? For zero elements, is it legal to launch 0 groups?

Case $G_{x} = 1$ . Ceiling groups $= ⌈ 1/64 ⌉ = 1$ . You launch one full squad of 64, but only worker i = get_global_id(0) = 0 passes the guard if (i < n) with n = 1; the other 63 (where i >= 1) are idle. Why this step? Hardware cannot run a squad smaller than its wavefront. One real element still costs one whole squad — a reminder that tiny launches are inefficient.
Case $G_{x} = 0$ . Ceiling groups $= ⌈ 0/64 ⌉ = 0$ . You should launch nothing and return early. Why this step? $0/64 = 0$ exactly, so no squads are needed. Launching a kernel with a 0-sized NDRange is either a no-op or an error depending on runtime — the safe pattern is a host-side if (n == 0) return; where n = G_data.

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

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

Worked example Ex 4 — Cell D: 2D NDRange

Statement. Process a $2048 \times 1024$ image (width $G_{x} = 2048$ , height $G_{y} = 1024$ ) with a 2D work-group $L_{x} = 16, L_{y} = 16$ . Find the group grid, and the linear buffer index for the pixel at global coords $(x, y) = (100, 50)$ in a row-major array of width 2048.

Forecast: how many total work-groups — hundreds or thousands? And guess the flat index of $(100, 50)$ .

Figure 2 below shows a small $8 \times 4$ NDRange tiled into $2 \times 2$ work-groups (red borders); the green square is the pixel being flattened. Consult it while reading step 2.

Groups per axis. $G_{x} / L_{x} = 2048/16 = 128$ , $G_{y} / L_{y} = 1024/16 = 64$ . Both divide cleanly (Cell A style, in 2D). Why this step? Each axis is chopped independently — Figure 2 shows the grid tiled into $16 \times 16$ squares (drawn as $2 \times 2$ for clarity). Total groups $= 128 \times 64 = 8192$ .
Flat buffer index. For a row-major image, index $= y \times width + x = 50 \times 2048 + 100 = 102, 500$ . Why this step? Memory is a 1D line, not a 2D sheet. To find pixel $(x, y)$ you skip $y$ full rows (each of width 2048), then step $x$ into the current row. Same "shelf then slot" logic as the 1D global-id formula, now with rows as shelves.

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

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. A physics sim over a $64 \times 64 \times 64$ voxel cube, work-group $4 \times 4 \times 4$ . Find total work-items, total groups, and the flat index of voxel $(x, y, z) = (3, 2, 1)$ in a cube stored with $x$ fastest.

Forecast: total work-items — is it about a quarter-million? Guess the flat index of an early voxel like $(3, 2, 1)$ .

Groups per axis $= 64/4 = 16$ on each of $x, y, z$ . Total groups $= 1 6^{3} = 4096$ . Why this step? Same clean divide, extended to a third axis. Nothing new — just one more factor.
Total work-items $= 6 4^{3} = 262, 144$ . Why this step? One worker per voxel, so the army size is the whole cube's volume: multiply the three axis lengths $64 \times 64 \times 64$ . We compute it now because we will use it to sanity-check the group count — if squads tile the cube perfectly, then (groups) × (workers per squad) must equal this volume, which the Verify line below confirms.
Flat index, $x$ fastest $= z \cdot (64 \cdot 64) + y \cdot 64 + x = 1 \cdot 4096 + 2 \cdot 64 + 3 = 4227$ . Why this step? Generalise "shelf then slot" to three levels: skip $z$ whole slices (each $64 \times 64$ ), then $y$ rows, then $x$ columns. The multipliers are the sizes of everything faster than the current axis.

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

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, two vendors

Statement. You run the identical kernel on $G_{x} = 1, 048, 576$ elements. You pick $L_{x}$ = one wavefront/warp so no partial batches waste lanes. On AMD (wavefront 64) vs NVIDIA (warp 32), how many groups each, and how many hardware batches total?

Forecast: does NVIDIA launch more groups or fewer? Does "more groups" mean "slower"?

AMD: $L_{x} = 64 \Rightarrow$ groups $= 1, 048, 576/64 = 16, 384$ ; each group = exactly 1 wavefront, so 16,384 hardware batches. Why this step? We set $L_{x}$ equal to AMD's physical batch width (64) so each work-group maps to exactly one wavefront — no lane sits idle. Since 64 divides $G_{x}$ cleanly (Cell A), plain division gives the group count directly.
NVIDIA: $L_{x} = 32 \Rightarrow$ groups $= 1, 048, 576/32 = 32, 768$ ; each group = 1 warp, so 32,768 hardware batches. Why this step? The portable source is byte-for-byte identical; only the launch config differs to match the physical batch width. HIP/OpenCL hides the ISA difference, but you still tune $L_{x}$ to the vendor's batch size to avoid idle lanes.
Note: twice the batches on NVIDIA does not mean twice slower — each NVIDIA warp is 32 lanes wide, AMD's 64. Total lanes processed = $32, 768 \times 32 = 16, 384 \times 64 = 1, 048, 576$ either way. Why this step? Throughput is lanes, not batch count. Same total work, sliced differently.

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. You must blur a $1920 \times 1080$ (Full HD) photo, one work-item per pixel, on an AMD GPU. Choose a sensible 2D work-group and report the padded NDRange and idle-worker count. (Full HD is not a clean multiple of 16 on the height axis.)

Forecast: which axis will need padding — width 1920 or height 1080? Guess the number of wasted workers.

Figure 3 below shows the Full HD frame: the blue block is the real pixels, the red strip along the bottom is the padded, guarded, idle workers. Consult it during step 4.

Choose $L_{x} = L_{y} = 16$ (256 items/group = 4 AMD wavefronts, a common sweet spot). Why this step? 256 is a multiple of the 64-wide wavefront, so no lanes are wasted inside a group.
Width: $1920/16 = 120$ exactly — clean, no padding. Why this step? We check each axis separately for raggedness. Since 16 divides 1920 with zero remainder ( $1920 mod 16 = 0$ ), the width tiles perfectly into 120 groups and needs no extra column of workers.
Height: $1080/16 = 67.5$ — ragged. Round up: $⌈ 1080/16 ⌉ = 68$ groups tall, padded height $= 68 \times 16 = 1088$ . Why this step? Cell B again but on one axis only: 16 does not divide 1080 ( $1080 mod 16 = 8$ ), so we round up to cover the leftover 8 rows. Figure 3 shows the extra 8-pixel-tall red strip of guarded, idle workers along the bottom.
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$ . Why this step? Only the bottom strip is padded; each of the 1920 columns has 8 extra rows (the red band in Figure 3). Those extra workers fail the guard if (i < n) and write nothing.

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: the exam twist (naive answer is wrong)

Statement. Exam question: "A kernel has $G_{x} = 1, 500, 000$ and $L_{x} = 512$ . A student writes groups = G/L = 1500000/512 = 2929 (integer division). Is 2929 correct? If not, give the right count and explain the bug."

Forecast: commit to an answer now — is 2929 right, one too few, or one too many?

Compute exact quotient. $1, 500, 000/512 = 2929.6875$ . Integer division truncates to $⌊ 2929.6875 ⌋ = 2929$ (floor = round down). Why this step? Integer / in C++ floors. The student silently dropped the fractional part.
Coverage check. $2929 \times 512 = 1, 499, 648$ . That leaves $1, 500, 000 - 1, 499, 648 = 352$ real elements never processed — the classic off-by-one-group bug. Why this step? The fractional $0.6875$ of a group is a real partial squad of 352 workers; dropping it corrupts the last 352 outputs.
Correct formula: ceiling division. In integer code you write it as $groups = (G + L - 1) / L$ . Plug in: $⌊(1, 500, 000 + 511) /512 ⌋ = ⌊ 1, 500, 511/512 ⌋ = 2930$ . Why this step? We cannot use a ceil() on floats safely (rounding error can bite at exact multiples), so we use the integer trick $(G + L - 1) / L$ . Why does adding $L - 1$ work? If $G$ is already a clean multiple of $L$ , then $G + L - 1$ is just short of the next multiple, so flooring lands back on the exact quotient — no spurious extra group. If $G$ has any remainder, adding $L - 1$ pushes the total past the next multiple boundary, so flooring rounds up by exactly one — capturing the leftover partial squad. In one line it is "floor, but bumped up whenever there is a remainder."
Confirm coverage. $2930 \times 512 = 1, 500, 160 \geq 1, 500, 000$ , so every real element now has a worker; the last $160$ workers are padding, silenced by the guard if (i < n) C[i] = A[i] + B[i]; with n = 1500000. Why this step? Rounding up is only safe if we also guard the padding — otherwise those 160 extra workers run out of bounds. Ceiling division and the if (i < n) guard are a matched pair.

Answer: 2929 is wrong (one group too few, 352 elements dropped). Correct = 2930 using (G + L - 1) / L, plus an if (i < n) guard.

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

Recall drill

Recall Which formula converts (group_id, local_id) back to global_id?

global_id = group_id × local_size + local_id ::: skip past all full squads, then step in by your local seat.

Recall Why do you round UP when computing work-group count?

Because rounding down would leave the leftover partial squad of real elements uncomputed — silent data loss. ::: Use (G + L - 1) / L and guard with if (i < n).

Recall On AMD vs NVIDIA, why does the same source compile but you still change local_size?

HIP/OpenCL hides the ISA, but you match $L_{x}$ to the physical batch width (wavefront 64 / warp 32) so no lanes sit idle. ::: Same total lanes either way.

Ceiling-division bug in a launch config

launching G/L (floor) instead of ceil(G/L) drops the final partial group — always round up.

Padding workers must be silenced by

an if (get_global_id(0) < n) guard so out-of-bounds lanes do nothing.