Warps and warp scheduling: Worked examples — Hardware deep dive

This is the "roll up your sleeves" page for the parent topic. We will not re-teach the theory — instead we hunt down every kind of situation the parent note can throw at you and work each one to a number. If a term feels unfamiliar, the parent note built it; here we use it.

Before anything else, one plain-language recap of the three numbers we keep reusing, because every example below is built from them:

Recall The three quantities every example uses

Warp size ::: 32 threads (fixed on NVIDIA hardware). A warp is one bundle of 32 threads that share one instruction fetch.
Linear thread index ::: flatten a 3D (x,y,z) thread coordinate into a single counting number, x fastest.
Warp ID ::: that linear index divided by 32, rounded down. So threads 0–31 → warp 0, 32–63 → warp 1, and so on.

The scenario matrix

Think of this table as a checklist. Every cell is a class of problem. Below it, each worked example is stamped with the cell(s) it clears. By the end, no cell is left uncovered — you will never meet a warp scenario on an exam that we skipped.

#	Case class	What makes it tricky	Covered by
A	Clean 1D block, multiple of 32	The "no waste" baseline	Ex 1
B	Block not a multiple of 32	Partial (wasted) warp at the tail	Ex 2
C	2D block, row-major flatten	Warps cut across rows	Ex 3
C2	True 3D block, non-unit depth	The $t_{z}$ term actually fires	Ex 3b
D	Degenerate: block smaller than 32	One warp, mostly idle lanes	Ex 4
E	Latency hiding — enough warps?	Need $⌈ L / W ⌉$ warps	Ex 5
F	Limiting case: zero compute between loads	No hiding possible — memory bound	Ex 6
G	Warp divergence, both branches taken	Serialized paths, slowdown factor	Ex 7
H	Divergence edge: all threads same branch	Zero penalty (no real divergence)	Ex 8
I	Real-world word problem	Choose block dims to avoid waste	Ex 9
J	Exam twist: occupancy vs. warps-needed	Hardware cap beats the ideal number	Ex 10

Group 1 — Where do threads land? (cells A, B, C, C2, D)

The single formula behind all four:

Recall Total threads vs. total warps — the rounding rule made precise

Let $P = D_{x} D_{y} D_{z}$ be the total thread count. Two facts must never be confused:

The largest warp ID is $⌊(P - 1) /32 ⌋$ (floor of the last linear index, which is $P - 1$ since indices count from 0). Warp IDs run $0, 1, \dots, ⌊(P - 1) /32 ⌋$ .
The number of warps is therefore $⌊(P - 1) /32 ⌋ + 1$ , and this equals $⌈ P /32 ⌉$ (ceiling of the total).

Why do these two agree? Write $P = 32 q + r$ , where $q$ is the whole number of complete groups of 32 and $r$ is the remainder ( $0 \leq r \leq 31$ ). Two cases:

If $r = 0$ : the block is an exact multiple of 32. The last index is $P - 1 = 32 q - 1$ , so $⌊(P - 1) /32 ⌋ = q - 1$ , giving $q - 1 + 1 = q$ warps. And $⌈ P /32 ⌉ = ⌈ q ⌉ = q$ . Same.
If $r > 0$ : the last index is $P - 1 = 32 q + (r - 1)$ with $0 \leq r - 1 \leq 30$ , so $⌊(P - 1) /32 ⌋ = q$ , giving $q + 1$ warps. And $⌈ P /32 ⌉ = q + 1$ (round the fraction up). Same.

So in both cases they land on the identical value — that is why "floor the last index, then add one" (which tells you the top warp's ID) equals "ceil the total" (which tells you how many warps exist).

One-line takeaway ::: $⌊(P - 1) /32 ⌋ + 1 = ⌈ P /32 ⌉$ for every positive integer $P$ , because splitting $P = 32 q + r$ makes both sides equal $q$ (if $r = 0$ ) or $q + 1$ (if $r > 0$ ).

Example 1 — Cell A: clean 1D block

Worked example The baseline:

blockDim = (256, 1, 1) How many warps, and are they all full?

Forecast: guess the warp count before reading on. (Hint: divide.)

Flatten. With $D_{y} = D_{z} = 1$ , the formula collapses to $linear = t_{x}$ . Threads are just $0, 1, \dots, 255$ . Why this step? A 1D block needs no cross-row bookkeeping — the linear index is the x-index.
Divide by 32. $256/32 = 8$ exactly. Why this step? Warp count is $⌈ 256/32 ⌉ = 8$ ; equivalently the top warp ID is $⌊ 255/32 ⌋ = 7$ , and $7 + 1 = 8$ warps.
Check fullness. Each warp covers exactly 32 consecutive indices with none left over, since 256 is a multiple of 32. Why this step? Waste only appears when the count isn't divisible by 32.

Answer: 8 warps, all 32/32 active, 0 wasted lanes.

Verify: $8 \times 32 = 256$ = the block size. No thread is unaccounted for, no lane is idle. ✓

Example 2 — Cell B: not a multiple of 32

Worked example The tail warp:

blockDim = (100, 1, 1) How many warps, and how many lanes are wasted?

Forecast: is it 3 warps or 4? And what's the waste?

Two ways, one answer. Total threads $P = 100$ . Warp count $= ⌈ 100/32 ⌉ = ⌈ 3.125 ⌉ = 4$ . Cross-check via the last index: $⌊ 99/32 ⌋ = 3$ is the top warp ID, so warps are 0,1,2,3 → $3 + 1 = 4$ warps. Why this step? The recall box above proves ceil-the-total and floor-the-last-index-plus-one give the same 4. Either route rounds up, because even a single leftover thread needs a whole warp.
Find the tail. Warps 0,1,2 hold indices 0–95 (that's $3 \times 32 = 96$ threads, all full). Warp 3 holds indices 96–99 → only 4 active threads. Why this step? Threads fill warps in order; whatever is left after the last full group forms a partial warp.
Count waste. Warp 3 has $32 - 4 = 28$ idle lanes. Why this step? The SM reserves a full warp's scheduling slot for warp 3 even though 28 lanes do nothing — that's the cost of not being a multiple of 32.

Answer: 4 warps; last one 4/32 full; 28 wasted lanes.

Verify: $3 \times 32 + 4 = 100$ ✓. Waste fraction of the tail warp $= 28/32 = 87.5%$ .

Example 3 — Cell C: 2D block, warps cut across rows

Worked example Row-major reality:

blockDim = (48, 4, 1) Which threads fall into warp 1?

Forecast: guess whether warp 1 stays inside a single row.

Total & flatten rule. $48 \times 4 = 192$ threads. With $D_{z} = 1$ the $t_{z}$ term is zero, so linear index $= t_{y} \cdot 48 + t_{x}$ . Why this step? Row $t_{y} = 0$ occupies linear 0–47, row $t_{y} = 1$ occupies 48–95, and so on. Rows are 48 wide, not 32 — so warp boundaries won't line up with row boundaries.
Warp 1 = linear 32–63. Linear 32–47 come from row 0 ( $t_{x} = 32..47$ ). Linear 48–63 come from row 1 ( $t_{x} = 0..15$ ). Why this step? Warp 1 straddles the seam between row 0 and row 1 — a classic "warp cuts across rows" situation.
Count warps. $192/32 = 6$ exactly → 6 full warps, no waste (192 is a multiple of 32 even though the rows aren't). Why this step? Waste depends on the total thread count vs. 32, not on the row width.

Figure walkthrough: each colored cell is one thread, arranged as the block really is — 48 columns (x) across, 4 rows (y) down, x running fastest. Cells are tinted by warp ID, so warp 0 is one color, warp 1 another, etc. The single warp outlined in white is warp 1: notice how that white group starts partway along the top row (x=32–47, the yellow arrow), then jumps down and continues at the start of the second row (x=0–15). That visual "wrap-around" is the whole point — because a row is 48 wide but a warp is only 32, warp boundaries slice through rows instead of aligning with them.

Answer: warp 1 = the tail of row 0 (x=32–47) plus the head of row 1 (x=0–15); 6 warps total, 0 waste.

Verify: $6 \times 32 = 192 = 48 \times 4$ ✓. Warp 1 spans linear 32–63 = 16 threads from row 0 + 16 from row 1 = 32 ✓.

Example 3b — Cell C2: a true 3D block (depth actually matters)

Worked example The

$t_{z}$ term fires: blockDim = (8, 4, 3) Which slice and row does the thread with linear index 70 live in, and how many warps are there?

Forecast: guess whether index 70 is in the first, second, or third z-slice.

Total threads. $P = D_{x} D_{y} D_{z} = 8 \times 4 \times 3 = 96$ . This is our first block where $D_{z} \neq = 1$ , so the $t_{z} (D_{x} D_{y})$ term finally does work. Why this step? One z-slice holds $D_{x} D_{y} = 8 \times 4 = 32$ threads. So each new z-slice adds a full plane of 32 — three slices stacked.
Locate linear 70. A slice is 32 threads, so slice $t_{z} = ⌊ 70/32 ⌋ = 2$ (the third slice, since z counts from 0). Within that slice the offset is $70 - 2 \times 32 = 6$ . Then row $t_{y} = ⌊ 6/8 ⌋ = 0$ and column $t_{x} = 6 - 0 \times 8 = 6$ . Why this step? We invert the flatten formula: peel off the biggest place-value ( $t_{z}$ ) first, then $t_{y}$ , then $t_{x}$ — exactly like reading digits off a mixed-radix number.
Count warps. $⌈ 96/32 ⌉ = 3$ warps, and beautifully each z-slice is exactly one warp here (since a slice = 32 threads). Why this step? When $D_{x} D_{y}$ is itself a multiple of 32, warp boundaries land cleanly on slice boundaries — no cross-slice straddling.

Answer: linear 70 → thread coordinate $(t_{x}, t_{y}, t_{z}) = (6, 0, 2)$ , in the third z-slice; block has 3 warps, one per slice.

Verify: rebuild the index: $t_{z} (D_{x} D_{y}) + t_{y} D_{x} + t_{x} = 2 \times 32 + 0 \times 8 + 6 = 70$ ✓. Warp of index 70 $= ⌊ 70/32 ⌋ = 2$ , matching slice $t_{z} = 2$ ✓.

Example 4 — Cell D (degenerate): block smaller than a warp

Group 2 — Will the scheduler stay busy? (cells E, F, J)

Recall the parent's latency-hiding count. In plain words: while one warp waits $L$ cycles for memory, the scheduler runs other warps. If a warp does $W$ useful instructions between memory requests, you need enough warps so that "everyone else's work" fills the wait.

Example 5 — Cell E: enough warps to hide latency

Example 6 — Cell F (limiting case): zero compute between loads

Worked example The unhidable stall:

$L = 400$ , $W = 1$ A kernel that does nothing but stream loads back-to-back. How many warps to hide it?

Forecast: is this even possible?

Plug in. $N = ⌈ 400/1 ⌉ = 400$ warps. Why this step? With only 1 instruction between memory ops, each warp barely covers 1 cycle of the 400-cycle wait — so you'd need 400 warps.
Reality check. Max is 64 warps/SM ⇒ 400 ≫ 64 ⇒ cannot hide. The kernel is memory-bound. Why this step? This is the degenerate limit of the formula: as $W \to 1$ , $N \to L$ , which blows past the hardware ceiling. No occupancy can save a kernel with no arithmetic to overlap.

Answer: would need 400 warps; impossible → memory-bound, throughput limited by bandwidth, not scheduling.

Verify: $⌈ 400/1 ⌉ = 400 > 64$ ✓. See Occupancy-vs-Performance: past a point, more occupancy stops helping — this is why.

Example 10 — Cell J (exam twist): ideal warps vs. hardware cap

First, the scheduler-level formula this example leans on. The parent note runs the count per scheduler, so we must name its symbols before using them:

Formula Scheduler-level warp demand

$Warps needed \geq N_{sched} \times \frac{L}{T _{issue}} / I_{warp}$ where (all cycle-counts are in clock cycles, all counts are dimensionless whole numbers):

$N_{sched}$ = number of warp schedulers per SM (independent hardware units that each pick a warp to issue every cycle; modern SMs have 4). Unit: schedulers (a plain count).
$L$ = memory stall latency. Unit: cycles.
$T_{issue}$ = cycles between successive issues from one scheduler ( $T_{issue} = 1$ means a scheduler can fire one instruction every cycle). Unit: cycles per instruction.
$I_{warp}$ = independent instructions a warp can issue before it stalls. Unit: instructions per warp.

Reading it with units: $N_{sched} \cdot L / T_{issue}$ has units $schedulers \times cycles / (cycles/instruction) = instruction-slots$ — the total number of instruction slots that must stay filled during one stall. Dividing by $I_{warp}$ (instructions per warp) cancels the "instructions", leaving warps.

Worked example The trap:

$N_{sched} = 4$ , $L = 400$ cyc, $T_{issue} = 1$ cyc/inst, $I_{warp} = 20$ inst/warp; hardware max = 64 warps/SM Can the SM be fully fed?

Forecast: does 64 warps suffice? (The exam wants you to say yes.)

Instruction budget. $N_{sched} \cdot L / T_{issue} = 4 \times 400/1 = 1600$ instruction slots must be "in flight" per SM. Why this step? Four schedulers each firing every cycle for 400 cycles = $4 \times 400 = 1600$ slots to fill while the stall lasts.
Convert to warps. Divide by $I_{warp} = 20$ : $1600/20 = 80$ warps. Why this step? Warps are the unit you can actually schedule; this converts the instruction budget into a warp count via the unit cancellation shown above.
Hit the cap. The hardware allows at most 64 warps/SM, but we need 80. Since $80 > 64$ , warp-level parallelism alone falls short by 16 warps. Why this step? The twist: the ideal answer (80) exceeds what the SM can physically hold (64). The hardware cap, not the formula, is the binding limit — so you must also exploit instruction-level parallelism inside each warp (issuing independent instructions back-to-back within one warp) to make up the deficit. More ILP costs registers — see Register-Pressure.

Answer: need 80 warps ideally, but the cap is 64 → short by 16; close the gap with ILP, not just occupancy.

Verify: $4 \cdot 400/1 = 1600$ ; $1600/20 = 80$ ; $80 > 64$ (deficit $= 16$ ) ✓.

Group 3 — Divergence cost (cells G, H)

When threads in one warp disagree about which branch to take, the warp runs both branches one after the other, masking off the lanes that shouldn't participate each time. See Branch-Divergence-Patterns for pattern-level tricks.

Example 7 — Cell G: both branches populated

Worked example Split warp: 20 lanes take

if ( $T_{if} = 30$ cyc), 12 take else ( $T_{else} = 30$ cyc) What's the slowdown vs. an undivided warp?

Forecast: does the 20-vs-12 split change the answer? Guess yes or no.

Divergent time. Both branches run: $T_{diverged} = 30 + 30 = 60$ cycles. Why this step? The warp executes the if with 12 lanes masked off, then the else with 20 lanes masked off — sequentially. Lane count per branch doesn't change the time; masked lanes still tick through the same cycles.
Baseline. No divergence would cost one branch: $30$ cycles. Why this step? If all 32 lanes agreed, only one path runs.
Slowdown. $60/30 = 2 \times$ . Why this step? This is the classic worst case: equal-cost branches ⇒ exactly double time regardless of the split ratio.

Figure walkthrough: the top blue bar is the baseline — if all 32 lanes agreed, only the if block runs, 30 cycles, done. The bottom row is what divergence actually does: first the green if bar runs (20 lanes live, 12 masked off and idle), and only after it finishes does the red else bar run (12 lanes live, 20 masked off). Because they run end-to-end rather than side-by-side, the yellow double-arrow spans the full 60 cycles — visibly twice the blue bar. The masked lanes in each colored bar are "paying rent" (occupying cycles) without doing useful work, which is exactly the divergence penalty.

Answer: 60 cycles, 2× slowdown — and it would be 2× even for a 31-vs-1 split.

Verify: $30 + 30 = 60$ ; $60/30 = 2.0$ ✓. Key insight the figure shows: the 12 idle if-lanes and 20 idle else-lanes waste cycles but don't save any.

Example 8 — Cell H (edge case): all threads agree

Worked example No real divergence: all 32 lanes take

if ( $T_{if} = 30$ ), 0 take else What's the penalty?

Forecast: does the else block cost anything if nobody enters it?

Populated branches. $n_{else} = 0$ , so the else path is skipped entirely — the warp mask for else is all-zero and the hardware jumps past it. Why this step? Divergence only serializes branches that at least one lane takes. An empty branch is free.
Time. $T = T_{if} = 30$ cycles. Slowdown $= 30/30 = 1 \times$ . Why this step? This is the boundary of the divergence formula: the moment $n_{else} \to 0$ , $T_{diverged} \to T_{if}$ and the penalty vanishes.

Answer: 30 cycles, 1× (no penalty) — a fully convergent warp.

Verify: with $n_{else} = 0$ , $T = 30 + 0 \cdot 30 = 30$ ; $30/30 = 1.0$ ✓. Moral: aligning branch decisions to warp boundaries (all 32 agree) kills the penalty.

Group 4 — The word problem (cell I)

Example 9 — Cell I: choosing block dims for an image kernel

Worked example Real world: process a

$1920 \times 1080$ image, one thread per pixel You must pick a 2D block shape. Compare a "balanced-looking" (30, 30) against an "aligned" (32, 8). Which wastes fewer lanes, and how many warps does each block hold?

Forecast: the square block looks efficient — is it actually better? Guess before computing.

Block (30,30). Threads $= 30 \times 30 = 900$ . Warp count $= ⌈ 900/32 ⌉ = ⌈ 28.125 ⌉ = 29$ . The last (partial) warp holds $900 - 28 \times 32 = 900 - 896 = 4$ threads → $32 - 4 = 28$ wasted lanes in that tail warp. Why this step? 900 isn't a multiple of 32, so we land in Cell B inside a 2D block — a partial tail warp appears, just like Example 2.
Block (32,8). Threads $= 32 \times 8 = 256$ . Warp count $= 256/32 = 8$ exactly, 0 wasted lanes. Why this step? Making blockDim.x equal to the warp width (32) guarantees clean warp boundaries — the row width equals the warp width, so warps never straddle rows either (unlike Example 3).
Decide. (32,8) wastes nothing and keeps warps aligned to rows, which also helps coalesced memory access (see GPU-Memory-Hierarchy). (30,30) wastes 28 lanes and misaligns rows. Why this step? Fewer wasted lanes + aligned memory ⇒ higher effective throughput. The "balanced-looking" square block is the trap the question set for you.

Answer: choose (32,8) — 8 warps, 0 wasted lanes — over (30,30), which needs 29 warps with 28 wasted lanes.

Verify: $⌈ 900/32 ⌉ = 29$ and $900 - 896 = 4$ active (28 idle) ✓; $256/32 = 8$ exact, 0 idle ✓. (The image is tiled by many such blocks; here we compare per-block efficiency, which is what determines whether each block wastes lanes.)

Recall Quick self-test

A (72,1,1) block: how many warps and how many wasted lanes? ::: $⌈ 72/32 ⌉ = 3$ warps; tail warp holds $72 - 64 = 8$ active → 24 wasted lanes. A (8,4,3) block: how many warps? ::: $⌈ 96/32 ⌉ = 3$ warps (one per z-slice, since each slice is exactly 32 threads). $L = 300$ , $W = 15$ : warps needed to hide? ::: $⌈ 300/15 ⌉ = 20$ warps. 16 lanes take if (cost 40), 16 take else (cost 40): slowdown? ::: $(40 + 40) /40 = 2 \times$ .