Warps and warp scheduling: Exercises — Hardware deep dive

This page is a self-test ladder. Each rung is harder than the last. Cover the solution, try the problem, then reveal. Every number you compute here is machine-checked.

Before we begin, here is the only vocabulary you need — all inherited from Warps and warp scheduling:

The two arithmetic tools we lean on:

Level 1 — Recognition

Exercise 1.1 (L1)

A block is launched with blockDim = (32, 1, 1). How many warps does it use, and are they full?

Recall Solution

What we do: apply the warp-count formula. Threads $= 32 \times 1 \times 1 = 32$ . $warps = ⌈ \frac{32}{32} ⌉ = 1$ What it looks like: one full "row of 32 seats", every seat taken. Answer: 1 warp, fully occupied (32/32 active).

Exercise 1.2 (L1)

A block is launched with blockDim = (48, 1, 1). How many warps? How many thread slots are wasted?

Recall Solution

Threads $= 48$ . $warps = ⌈ \frac{48}{48 \to 48/32 = 1.5} ⌉ = ⌈ 1.5 ⌉ = 2$ Two warps provide $2 \times 32 = 64$ slots but only 48 are real threads. Wasted slots $= 64 - 48 = 16$ . Answer: 2 warps, 16 wasted slots (warp 1 runs 32/32, warp... wait — recount): warp 0 = threads 0–31 (full), warp 1 = threads 32–47 (16 active, 16 idle). So the waste all sits in warp 1.

Exercise 1.3 (L1)

True or false: an SM schedules individual threads onto its execution units.

Recall Solution

False. The scheduling quantum is the warp. All 32 threads of a warp are issued together in lockstep; the scheduler never picks a single thread. (See SIMT-vs-SIMD for why SIMT keeps per-thread state but still issues per-warp.)

Level 2 — Application

Exercise 2.1 (L2)

blockDim = (16, 4, 1). (a) How many warps? (b) Which warp does thread $(t_{x}, t_{y}, t_{z}) = (3, 2, 0)$ land in?

Recall Solution

(a) Threads $= 16 \times 4 \times 1 = 64 \Rightarrow ⌈ 64/32 ⌉ = 2$ warps, both full. (b) Linear index with $D_{x} = 16, D_{y} = 4$ : $i = 0 \cdot (16 \cdot 4) + 2 \cdot 16 + 3 = 32 + 3 = 35$ Warp $= ⌊ 35/32 ⌋ = 1$ . Answer: 2 warps; thread (3,2,0) is in warp 1.

Exercise 2.2 (L2)

An SM can hold at most 64 warps. You currently run 48 active warps. What is the occupancy (as a percentage)?

Recall Solution

$Occupancy = \frac{48}{64} = 0.75 = 75%$ Answer: 75% occupancy. See Occupancy-vs-Performance — 75% is often plenty; chasing 100% can backfire if it forces register spills.

Exercise 2.3 (L2)

Global memory latency is $L = 400$ cycles. Each warp does $W = 8$ independent instructions between memory operations. Roughly how many warps must be resident to fully hide this latency (single scheduler)?

Recall Solution

Why this formula: while one warp waits $L$ cycles, other warps must keep the scheduler fed, one instruction each per cycle. Each warp offers $W$ useful instructions before it too stalls. $N_{warps} = ⌈ \frac{L}{W} ⌉ = ⌈ \frac{400}{8} ⌉ = 50$ Answer: 50 warps to hide the 400-cycle latency.

Level 3 — Analysis

Exercise 3.1 (L3)

blockDim = (17, 8, 1) (136 threads). How many warps, and exactly how many active threads does the last warp carry?

Recall Solution

Threads $= 17 \times 8 = 136 \Rightarrow ⌈ 136/32 ⌉ = ⌈ 4.25 ⌉ = 5$ warps. First 4 warps use indices 0–127 (128 threads, all full). Last warp = indices 128–135. Active in last warp $= 136 - 128 = 8$ . Answer: 5 warps; the last warp runs only 8/32 threads (24 idle lanes → 75% of that warp wasted). Fix the launch to (16,8,1) = 128 threads = exactly 4 full warps.

Exercise 3.2 (L3)

A warp hits if (threadIdx.x % 2 == 0) A(); else B();. Branch A costs 30 cycles, B costs 20 cycles. What is the divergent execution time, and the slowdown versus the no-divergence cost of just A?

Recall Solution

Because odd/even lanes split, both branches are present in the warp → they run sequentially with masks: $T_{diverged} = T_{A} + T_{B} = 30 + 20 = 50 cycles$ No-divergence baseline (all lanes take A) $= 30$ cycles. $Slowdown = \frac{50}{30} \approx 1.67 \times$ Answer: 50 cycles, ≈1.67× slower. See Branch-Divergence-Patterns: a % 2 predicate is the worst kind — it guarantees every warp diverges.

Exercise 3.3 (L3)

Same code as 3.2 but the predicate is threadIdx.x < 32 inside a 32-thread warp where every thread's threadIdx.x is 0–31. Does this warp diverge? What is its time?

Recall Solution

Every one of the 32 threads satisfies < 32, so all take branch A, none take B. No divergence → $T = T_{A} = 30$ cycles. Answer: No divergence; 30 cycles. The lesson: divergence depends on whether lanes within one warp disagree, not on the branch existing in the source.

Level 4 — Synthesis

Exercise 4.1 (L4)

An SM has 4 warp schedulers, each issues 1 instruction/cycle ( $T_{issue} = 1$ ). Average stall latency $L = 300$ cycles. How many warps must be resident to keep all four schedulers busy, assuming each warp yields 1 instruction before stalling (worst case)?

Recall Solution

Each scheduler needs a fresh ready warp every cycle for $L$ cycles: $warps \geq N \times \frac{L}{T _{issue}} = 4 \times \frac{300}{1} = 1200$ Answer: 1200 warps' worth of independent work. Since a real SM caps at ~64 warps, this is impossible with 1-instruction warps — you also need instruction-level parallelism inside each warp (more independent instructions per warp before it stalls).

Exercise 4.2 (L4)

Continue 4.1, but now each warp has $W = 20$ independent instructions before stalling. How many warps (not instructions) are needed? Is it achievable with a 64-warp cap?

Recall Solution

$warps = \frac{N \cdot L / T _{issue}}{W} = \frac{1200}{20} = 60$ $60 \leq 64$ , so yes, achievable. Answer: 60 warps, achievable within the 64-warp limit (occupancy $= 60/64 \approx 93.8%$ ).

Exercise 4.3 (L4)

You have a kernel needing 32 registers/thread. The SM has 65536 registers total and a max of 2048 resident threads. Which limit — registers or the thread cap — decides occupancy, and what is the max resident warp count?

Recall Solution

Register-limited thread count: $65536/32 = 2048$ threads. Hardware thread cap: 2048 threads. They tie exactly → limit is either (both give 2048). Warps $= 2048/32 = 64$ → full occupancy. Answer: ==Both limits give 2048 threads = 64 warps = 100% occupancy==. If registers rose to 40/thread, then $65536/40 = 1638$ threads $\Rightarrow ⌊ 1638/32 ⌋ = 51$ warps → register-limited. (See Register-Pressure.)

Level 5 — Mastery

Exercise 5.1 (L5)

A grid processes a $100 \times 100$ image with blockDim = (16, 16, 1). (a) Threads per block? (b) Warps per block? (c) Wasted slots per block? (d) If the grid is $⌈ 100/16 ⌉^{2}$ blocks, total blocks?

Recall Solution

(a) $16 \times 16 = 256$ threads. (b) $⌈ 256/32 ⌉ = 8$ warps, all full (256 divides evenly by 32). (c) $8 \times 32 - 256 = 0$ wasted slots — perfect, since 16×16 is a multiple of 32. (d) Blocks per dimension $= ⌈ 100/16 ⌉ = ⌈ 6.25 ⌉ = 7$ . Total $= 7 \times 7 = 49$ blocks. Answer: 256 threads, 8 full warps, 0 waste, 49 blocks. (The grid wastes edge threads outside 100×100, but each block is internally warp-clean — handled by a bounds if, which see 5.2.)

Exercise 5.2 (L5)

The edge blocks in 5.1 guard with if (x < 100 && y < 100). For a warp that straddles the boundary — say 20 lanes inside the image and 12 outside — what is the divergence cost if the "inside" work costs $T = 40$ cycles and the "outside" lanes do nothing (0 cycles)?

Recall Solution

Two groups: 20 lanes run the body ( $T = 40$ ), 12 lanes take the empty else (0). They still serialize, but the empty branch costs 0: $T_{diverged} = 40 + 0 = 40 cycles$ Baseline (no divergence, all 32 do the body) would also be 40. Slowdown $= 1 \times$ — the divergence is free here because the disabled lanes had nothing to do. Answer: 40 cycles, no real penalty. The masked-off lanes idle "for free" during the single 40-cycle body pass. Divergence only costs when both paths have real work.

Exercise 5.3 (L5, Synthesis + Analysis)

Kernel A: each warp diverges 50% (half the warp costs $T_{A} = 100$ , half costs $T_{B} = 100$ ). Kernel B: no divergence, but only 32 resident warps vs Kernel A's 64. Memory latency $L = 400$ , warps need $\geq ⌈ L / W ⌉$ with $W = 25$ to hide it. Which kernel more likely hides latency, and what's kernel A's per-warp divergence slowdown?

Recall Solution

Latency-hiding need: $⌈ 400/25 ⌉ = 16$ warps.

Kernel A: 64 warps $\geq 16$ ✓ hides latency.
Kernel B: 32 warps $\geq 16$ ✓ also hides latency. Both clear the bar, so latency isn't the deciding factor here. Kernel A divergence slowdown per warp: $(T_{A} + T_{B}) / T_{A} = (100 + 100) /100 = 2 \times$ . Answer: Both hide latency; but Kernel A pays a 2× divergence tax per warp. With latency covered either way, Kernel B wins — its warps are 2× faster in the divergent region. Lesson: fixing divergence often beats piling on more warps once latency is already hidden.