Warps and warp scheduling: Question bank — Hardware deep dive

This is a rapid-fire misconception buster. Each line is a Question ::: Answer reveal — read the question, commit to an answer out loud, then reveal. If your reasoning didn't match the justification, that's a trap you needed to hit.

Everything here builds on Warps and warp scheduling. Before starting, make sure you can say in one breath what these three words mean:

warp — a fixed group of 32 consecutive threads that share one instruction fetch and step through the program together, one instruction at a time.
lockstep — all 32 threads are always pointing at the same instruction; if some threads shouldn't run it, they are masked off (present but doing nothing), not skipped.
stall — a warp that cannot make progress this cycle because it is waiting for something (memory data, a previous result, a barrier).

Related deep dives you'll want open: SIMT-vs-SIMD, Branch-Divergence-Patterns, Occupancy-vs-Performance, Register-Pressure, Thread-Blocks-and-Grids, GPU-Memory-Hierarchy.

True or false — justify

A warp always contains exactly 32 threads on every GPU ever made.

False. On NVIDIA GPUs a warp is 32 threads, but AMD "wavefronts" are 64. Always tie the number to the vendor; never hard-code 32 in portable reasoning.

If a thread block has 40 threads, it uses exactly 2 warps.

True in count but with waste. 40 threads round up to 2 warps (64 slots); warp 1 has only 8 active threads and 24 masked-off — the SM still reserves a full warp's resources. See Occupancy-vs-Performance.

Two threads in the same warp can execute two different instructions in the same cycle.

False. That is the defining property of lockstep — one instruction is fetched per warp per issue. Divergent threads are handled by running paths sequentially with masks, never simultaneously.

Threads in the same warp always finish their kernel at the same time.

False. If they diverge, some threads sit masked-off while others run a branch; the warp only "reconverges" when paths meet again, so effective per-thread finish times differ even though they share a program counter.

Higher occupancy always means higher performance.

False. Occupancy only enables latency hiding by giving the scheduler more warps to switch to. A kernel can be memory-bandwidth-bound or already hiding latency at 50% occupancy, so pushing occupancy higher does nothing. See Occupancy-vs-Performance.

Switching from a stalled warp to a ready warp costs the GPU a context-save like a CPU thread switch.

False. Each warp keeps its own registers and program counter resident in the SM, so switching is just choosing a different pointer — zero-overhead, same cycle. That is the whole reason GPUs tolerate huge memory latency.

A warp scheduler can only run one warp per SM at a time.

False. Modern SMs have multiple schedulers (often 4), each issuing from its own pool of warps, so several warps make progress in the same cycle.

__syncthreads() synchronises threads within a warp.

Misleading. Threads in a warp are already in lockstep; __syncthreads() is a block-wide barrier synchronising all warps of the block. For within-warp coordination you use warp-level primitives, see Cooperative-Groups.

SIMT and SIMD are the same thing with different names.

False. SIMD exposes one wide vector register the programmer packs explicitly; SIMT lets each thread have its own registers and appear independent, with the hardware ganging 32 of them per instruction. Divergence is invisible in SIMD but a real cost in SIMT — see SIMT-vs-SIMD.

Using more registers per thread can lower the number of resident warps.

True. Register file per SM is fixed; more registers per thread means fewer threads (hence warps) fit at once, reducing occupancy. This is register pressure — see Register-Pressure.

Spot the error

"To avoid divergence, I made blockDim = (30, 1, 1) so each warp is nearly full."

The error: 30 is not a multiple of 32, so warp 0 has 30 active and 2 masked threads — you added permanent waste, and this has nothing to do with divergence. Divergence is about branches, not block size; pick 32.

"My kernel branches on threadIdx.x < 16, so half the warp diverges — unavoidable."

The error: divergence depends on how the branch splits within each 32-thread warp. Since threads 0–15 and 16–31 are in the same warp, yes it diverges. Branch on something aligned to warp boundaries (e.g. warpId % 2) and there is no divergence. See Branch-Divergence-Patterns.

"A memory load takes 400 cycles, so my program is 400× slower."

The error: latency is hidden, not paid per warp serially. While warp 0 waits 400 cycles, the scheduler runs dozens of other ready warps, so throughput barely drops if occupancy is sufficient.

"Divergence with all 32 threads taking the if branch still costs 2× because there's an if-else."

The error: if no thread takes the else branch, the warp skips the masked-out else entirely — no divergence, no penalty. Cost only appears when both branches have at least one active thread.

"Warp ID equals threadIdx.x / 32."

The error: that only holds for 1D blocks. In general you must first flatten to a linear index using all three dimensions, then divide by 32. The x-only formula silently breaks for 2D/3D blocks — see Thread-Blocks-and-Grids.

"Two warps from different blocks can never run on the same SM."

The error: an SM commonly holds several blocks at once (limited by registers and shared memory), and mixes their warps in one scheduling pool — that variety is exactly what feeds latency hiding.

Why questions

Why does the scheduler operate at warp granularity instead of per thread?

Because instruction fetch/decode is amortised across 32 threads: one fetch drives 32 executions, so control-hardware cost is paid once per 32 lanes instead of 32 times. Scheduling individual threads would throw that saving away.

Why do block dimensions being multiples of 32 matter?

Any partial final warp still reserves a full warp's slots and resources but leaves lanes masked-off, wasting execution slots and lowering effective occupancy for zero benefit.

Why does memory latency hiding need many warps rather than a few?

You need enough independent work to fill the entire stall window: if a load costs

L

cycles and a warp does useful work only

W

cycles between loads, you need roughly

⌈ L / W ⌉

warps per scheduler to always have someone ready. Too few and the SM idles.

Why can high occupancy still leave the SM idle?

If every resident warp stalls on the same dependency at the same time (poor "warp diversity"), the scheduler has many warps but none eligible. Occupancy counts resident warps; only ready warps hide latency.

Why does branch divergence hurt even when both branch bodies are short?

Because the two paths run sequentially under complementary masks; the warp pays T_if + T_else in time even though at most 32 threads' worth of work happened. Idle lanes during each masked pass are pure waste. See Branch-Divergence-Patterns.

Why is per-warp register partitioning what makes zero-overhead switching possible?

Each warp's registers stay physically resident in the SM's register file, so there is nothing to save or reload on a switch — the scheduler just points its issue logic at a different warp. See Register-Pressure.

Edge cases

A block of exactly 32 threads — how many warps and any waste?

Exactly 1 warp, fully packed, zero waste. This is the smallest "clean" block size and a safe default.

A block of 1 thread — what happens to the warp?

It still occupies a full warp; 1 lane active, 31 masked-off. Enormously wasteful, so single-thread blocks are almost never sensible.

What if only ONE thread in a warp takes the else branch?

The whole warp still serialises both paths — 31 threads sit idle during the else pass for that single thread. Divergence cost is set by whether both paths are taken, not by how many threads take each. See Branch-Divergence-Patterns.

What if every resident warp is stalled on global memory at once?

The scheduler has zero eligible warps, so the SM issues nothing and idles until the first memory reply arrives — this is the failure mode of insufficient warp diversity, not low occupancy.

A block with 1024 threads and a warp size of 32 — how many warps, and is that automatically good?

32 warps. That satisfies the count, but if each thread hogs registers, not all 1024 may be resident simultaneously; large blocks can reduce occupancy via register/shared-memory limits. See Register-Pressure and Occupancy-vs-Performance.

A warp where all 32 threads read consecutive global addresses vs. scattered addresses — same latency?

No. Consecutive (coalesced) accesses collapse into few memory transactions; scattered accesses issue many, multiplying effective latency. Warp scheduling hides latency; access pattern determines how much latency there is. See GPU-Memory-Hierarchy.

Recall Fast self-test before you close this page

Name the one situation where an if-else inside a kernel costs nothing extra. ::: When every active thread in the warp takes the same branch, so no path is serialised. State the difference between "resident" and "eligible" warps in one sentence. ::: Resident = loaded on the SM with registers reserved; eligible = resident and not stalled, i.e. actually schedulable this cycle.