6.2.4 · D5GPU Architecture
Question bank — SIMT (single instruction multiple thread)
This is the misconception-hunting deep dive for SIMT (single instruction multiple thread). Every item below is a place where smart people quietly get SIMT wrong. Read the question, say your answer out loud, then reveal.
Before we start, the words and symbols we lean on constantly. Read these once — the whole page assumes them:
- warp — the fixed group of 32 threads the hardware schedules together, one instruction at a time.
- lockstep — all active threads in a warp execute the same instruction in the same cycle.
- predication mask — a 32-bit flag, one bit per thread (1 = active, 0 = masked off / paused).
- — the number of distinct control-flow paths taken inside one warp at a branch. If every thread agrees, ; if all 32 threads pick different destinations, . Divergence cost scales with .
- — the latency of an operation, measured in clock cycles (e.g. a memory load might be cycles before the data comes back).
- — the throughput target, measured in instructions issued per cycle that you want the SM to sustain (often : keep issuing one instruction every cycle).
- — the number of resident warps on an SM available to switch between.
With , and named, the latency-hiding rule below reads plainly: to keep issuing at rate while each in-flight op takes cycles, you need
Three pictures to hold in your head
Almost every trap on this page is really a confusion between three separate pictures. Study these once and you'll answer half the items on sight.
(1) Divergence — a lane/mask picture. Which lanes are active on a given cycle?
(2) Coalescing — an address picture. Where in memory the 32 lanes point, and how many hardware transactions (memory segments) that costs.
(3) Occupancy — a scheduler picture. How many warps are resident so the scheduler always has a ready one while others wait on latency .
True or false — justify
A warp always contains exactly 32 active threads.
False. A warp is allocated 32 lanes, but some may be permanently inactive (a block of 100 threads leaves 28 idle lanes in its last warp) or temporarily masked off during divergence — see the greyed lanes in figure 1.
All 32 threads in a warp share one program counter, so they can never be on different instructions.
False in the way people mean it. They share one scheduled instruction stream, but during divergence the hardware tracks which threads are at which path and issues the paths one at a time — so at a given cycle only some lanes are active (figure 1).
SIMD and SIMT produce identical machine behaviour, just with different marketing names.
False. SIMD lanes must all execute (no per-lane branching); SIMT lanes can be individually masked, letting threads take different control-flow paths. See SIMD vs SIMT Comparison — the crux is that SIMT adds a per-lane mask, which SIMD lacks.
If a warp diverges into two branches, total time is the average of the two branch times.
False. Divergent branches run serially, so time is roughly (a sum), not an average or a max. Figure 1 shows the two phases stacked in time.
Fully converged code (all 32 threads take the same branch) has zero divergence penalty.
True. With one path taken, the mask is all-ones the whole time and path executes — no serialization, full lane utilization.
Higher occupancy always means faster kernels.
False. Occupancy only helps up to the point where latency is hidden (once ); beyond that, more warps compete for the same registers/cache and can hurt. The core idea from GPU Occupancy is that occupancy is a means to hide latency, not a goal in itself — see figure 3.
Warp size (32) is defined by the CUDA C++ standard and you should hard-code it.
False. 32 is an NVIDIA hardware architectural constant, not a language guarantee; portable code queries
warpSize rather than assuming it.Memory coalescing is a divergence problem in disguise.
False. They're independent. Coalescing is about the addresses 32 threads touch in one instruction (figure 2); divergence is about which threads execute an instruction at all (figure 1). A perfectly converged warp can still be badly uncoalesced. See GPU Memory Coalescing.
Because GPUs use SIMT, they must use out-of-order execution to hide the 400-cycle memory latency.
False. GPUs hide latency by switching to other ready warps (massive multithreading, figure 3), not by CPU-style speculation or out-of-order issue. This is the headline contrast in GPU vs CPU Architecture: CPUs hide latency with big caches and reordering, GPUs hide it with many warps.
A block of 1024 threads runs as one giant lockstep unit.
False. 1024 threads = 32 warps; each warp is independently scheduled. Only the 32 threads within a warp are in lockstep. See CUDA Thread Hierarchy, where thread → warp → block → grid is the key ladder.
Spot the error
"To use all lanes, launch 100 threads so every one is busy." — what's wrong?
100 isn't a multiple of 32. The 4th warp holds only 4 threads (12.5% utilization); 28 lanes are wasted. Launch a multiple of 32.
"if (tid % 2 == 0) alternates work so both branches run in parallel, doubling speed." — what's wrong?
Divergence runs branches serially, not in parallel ( paths). Half the lanes are masked off in each phase, so you get slower (sum of both paths), not a 2× speedup — figure 1.
"Increasing registers per thread makes threads faster, so always maximize register use." — what's wrong?
More registers per thread reduces how many warps fit on an SM (registers are a shared budget), lowering occupancy and hurting latency hiding. It's a trade-off, not a free win — the register limit is one of the three occupancy ceilings in GPU Occupancy.
"data[threadIdx.x * 32] is a nice contiguous read since we multiply by the warp size." — what's wrong?
That's a strided pattern. Each thread lands on a different memory segment → many separate transactions instead of one packed load (figure 2, bottom row). Contiguous means
data[threadIdx.x]."Divergence penalty is fixed at 32× because a warp has 32 threads." — what's wrong?
32× is the worst case (every thread a unique path, ). Actual penalty is the number of distinct paths taken , from (none) to (maximum).
"Threads in different warps that hit the same if branch help each other converge." — what's wrong?
Convergence is a per-warp property. Divergence only matters within one warp's 32 lanes; other warps are scheduled independently and don't interact with this warp's mask. See Warp Divergence.
"A warp scheduler issues to whichever warp finished first, so ordering is deterministic." — what's wrong?
The scheduler picks any ready warp; the choice depends on runtime latencies and is not guaranteed deterministic across runs, which is why race conditions must be avoided explicitly.
Why questions
Why 32 and not, say, 30 threads per warp?
32 is a power of two, so
tid & 31 cheaply extracts the lane index, memory addresses fall neatly into fixed-size segments for coalescing (on current NVIDIA GPUs the segment is typically 32 or 128 bytes — the exact size is architecture-dependent), and one fetch/decode unit maps cleanly onto 32 parallel ALU lanes.Why does SIMT let the programmer write plain scalar code instead of explicit vector code?
Each thread has its own logical registers and program counter, so you write "what one thread does" and the hardware broadcasts that instruction across 32 lanes — unlike SIMD, where you must vectorize by hand.
Why does memory latency ( cycles) not stall the whole GPU?
While one warp waits on memory, the scheduler issues from other ready warps (figure 3). Enough resident warps () keeps the ALUs busy — latency is hidden, not eliminated. See GPU Occupancy.
Why does the divergence penalty come from serialization and not from wasted silicon?
Both branch phases must run because some threads need each one; the hardware can't run two different instructions on the 32 lanes simultaneously, so it runs them one after another — the cost is time, in cycles (figure 1).
Why is high occupancy necessary but not sufficient for good performance?
Occupancy provides enough warps to hide latency (), but if those warps have divergence, uncoalesced memory, or are bandwidth-bound, they'll still run slowly. The takeaway from GPU Occupancy and Instruction-Level Parallelism: hiding latency is one lever among several, and a single warp with independent instructions can also hide some latency on its own.
Why do GPUs use many SMs rather than one very deep pipeline?
A single SM caps at ~32–64 resident warps, far short of the hundreds needed to hide deep latency; spreading work across many SMs delivers the total warp count for real throughput.
Why can warp-level primitives like __shfl_sync replace some branches?
They exchange data directly between lanes without control-flow divergence, so a computation that would have branched (and serialized into phases) becomes a single converged instruction across all 32 lanes.
Edge cases
What happens to warp utilization when a block has 1 thread?
That single thread occupies a full warp: 1/32 ≈ 3.1% utilization. The other 31 lanes are masked off the entire time — extremely wasteful (imagine only one lane lit in figure 1).
What if every thread in a warp diverges to a unique path (all 32 different)?
The warp serializes into phases, each with exactly one active lane — the worst-case 32× slowdown and 1/32 utilization per phase.
What happens if a branch condition is thread-independent (same for all 32 threads at runtime)?
No divergence occurs even though there's an
if: all lanes evaluate the condition identically, take one path together (), and the mask stays all-ones. The compiler/hardware sees a uniform branch.What is the occupancy when a kernel uses so many registers that only one warp fits per SM?
Occupancy = 1 / (max warps per SM), e.g. 1/64 ≈ 1.6%. With almost no warps to switch to ( is tiny), memory latency can no longer be hidden and the SM stalls.
What happens at the boundary where blockDim is not a multiple of 32 — say 48 threads?
You get 2 warps: one full (32) and one half-full (16 active, 16 idle). The idle lanes in the second warp waste execution slots every cycle that warp runs.
What if all warps on an SM are simultaneously waiting on memory?
There is no ready warp to switch to, so the SM genuinely stalls until a memory response returns — this is the failure mode high occupancy () exists to prevent. See Warp Divergence and GPU Occupancy.
What does coalescing efficiency approach as stride grows very large (random access)?
It approaches its floor: each thread pulls a separate memory segment, so bytes-transferred balloons while bytes-used stays fixed at one warp's worth of data. The exact floor depends on the architecture's segment size (figure 2, bottom row).
Recall Quick self-check before you leave
Divergence penalty scales with what? ::: The number of distinct control-flow paths taken within one warp, i.e. , from 1 (none) up to 32 (worst case). Latency hiding scales with what? ::: The number of resident, ready warps per SM — captured by , where is latency in cycles and is target instructions per cycle. Coalescing depends on what? ::: The address pattern of the 32 threads in a single memory instruction (figure 2), independent of divergence.