6.2.2 · D5GPU Architecture
Question bank — Streaming multiprocessors (SM)
This is a misconception hunt. Each line below is a claim, question, or edge case about the Streaming Multiprocessor. Read the left side, commit to an answer before looking right, then reveal. Every answer gives you the reasoning, not just a verdict — that reasoning is what you are meant to keep.
Before we start, two words we will lean on repeatedly, in plain language:
If either of those feels fuzzy, revisit Warp-scheduling and Occupancy-optimization first — the traps below assume you can picture both.
True or false — justify
More cores per SM always means proportionally faster execution.
False — cores are useless if warps stall on memory; a kernel with too few resident warps leaves cores idle regardless of how many exist. Throughput is limited by whichever runs out first: work, warps, or bandwidth.
A block, once assigned to an SM, can migrate to a less busy SM mid-flight.
False — a block runs on exactly one SM from start to finish; its registers and shared memory live physically on that SM, so migration would mean copying private state, which the hardware never does.
100% occupancy guarantees peak performance.
False — occupancy only measures how many warps could run, not whether they need to. A compute-bound kernel with 50% occupancy can beat a 100% one; extra warps only help by hiding latency you actually have.
Two threads in the same warp can execute different instructions in the same cycle.
False — that is the core SIMT rule: one instruction is issued to all 32 lanes per cycle. Different branches are handled by masking lanes off and running the paths one after another, not simultaneously.
Increasing threads per block always increases occupancy.
False — bigger blocks consume more registers and shared memory per block, which can lower the number of resident blocks and thus reduce total resident warps. Occupancy is a resource tug-of-war, not a monotonic dial.
Shared memory and L1 cache are just two names for the same thing.
False — they often sit on the same physical SRAM, but shared memory is explicitly indexed by you with no tag lookup, while L1 is hardware-managed with tags and eviction. Same bricks, different rulebook and different latency.
Registers can suffer bank conflicts just like shared memory.
False — registers are private per thread, so no two threads ever contend for the same register slot. Bank conflicts are a shared-memory phenomenon, where a warp's 32 lanes hit 32 banks and collisions serialize.
A stalled warp wastes the SM's cycles while it waits.
False (that's the whole point of GPUs) — when a warp stalls, the scheduler instantly issues from another eligible warp, so the SM keeps working. The stall is only wasted if there is no other eligible warp to switch to.
If a kernel never branches, warp divergence cannot happen.
True — divergence is caused strictly by lanes in one warp taking different control-flow paths. No data-dependent branch means all 32 lanes always agree, so no serialized paths.
Doubling the register count per thread halves the maximum resident warps.
Usually true — the register file is fixed size, so warps-per-SM scales roughly inversely with registers-per-thread until some other cap (max warps, max blocks) becomes the binding limit first.
Spot the error
"My block has 40 threads, so it forms exactly one warp with 8 threads unused."
The error is the count: 40 threads need two warps (⌈40/32⌉ = 2). The first warp is full (32), the second holds 8 active lanes and 24 permanently masked-off lanes wasting slots — a reason to pick block sizes that are multiples of 32.
"I have 8 warps in my block, so I can fully hide a 300-cycle memory stall."
8 warps cover only about 8 cycles of latency per scheduler pass, not 300. Hiding a 300-cycle stall needs dozens of resident warps across many blocks — one block's 8 warps is nowhere near enough.
"I used __shared__ so threads automatically see each other's writes immediately."
Missing the barrier: without
__syncthreads() after the writes, a thread may read a neighbor's slot before that neighbor has written it. Shared visibility is possible, but ordering is your job."To avoid bank conflicts I made thread i access address i * 32."
That is the worst case: multiplying by 32 (the bank count) maps every lane to the same bank, forcing a 32-way serialized access. You want
i % 32 spread across banks, i.e. consecutive addresses, not strided-by-32 ones."My kernel is compute-bound, so I should still crank occupancy to the max."
If you are compute-bound, extra warps do not help — the cores are already busy every cycle. Pushing occupancy may even hurt by shrinking registers per thread and forcing spills to slower memory.
"An SM runs one block at a time, then loads the next."
An SM holds multiple resident blocks simultaneously (as many as registers, shared memory, and block-count caps allow). Concurrency across blocks is exactly how it accumulates enough warps to hide latency.
"Both branches of an if-else run, so a if...else doubles cost even when all threads go the same way."
If all 32 lanes in a warp take the same branch, only that one path executes — no divergence, no doubling. The 2× cost appears only when lanes within one warp split across the branches.
Why questions
Why is the warp size 32 and not, say, 40?
It is a hardware constant tied to the SIMT datapath width and scheduler design; the whole scheduling machinery, masking, and shuffle instructions assume 32-lane warps. Block sizes should therefore be multiples of 32 to avoid dead lanes.
Why do GPUs need thousands of threads when an SM only issues a few instructions per cycle?
Because their job is latency hiding: with a 400–800 cycle memory stall, you need a deep pool of ready warps to switch to while data is in flight. Excess threads are fuel for keeping the pipeline busy, not for issuing them all at once.
Why keep variables in registers instead of memory?
Registers are ~1-cycle access, versus hundreds of cycles for global memory. Keeping a thread's working set in registers removes the latency entirely — the trade-off is that heavy register use lowers occupancy.
Why does shared memory enable algorithms like tiled matrix multiply?
It lets threads in a block cooperatively load a tile once and reuse it many times at ~100× global-memory speed, turning many slow global reads into one shared read plus fast on-chip reuse — the reason Tensor-cores pipelines feed from shared memory.
Why can raising block size lower performance even below the register limit?
Larger blocks reduce granularity for the scheduler and can hit the shared-memory-per-block or max-blocks cap, so fewer blocks fit and the SM ends up with fewer resident warps overall. Bigger is not automatically fuller.
Why does the SM have Special Function Units separate from CUDA cores?
Transcendentals (sin, exp, sqrt) would cost dozens of cycles on a general ALU; dedicating a small SFU pipeline does them in 1–2 cycles, freeing the main cores for ordinary arithmetic instead of clogging them.
Edge cases
What is the occupancy of an SM running zero warps?
Occupancy is — a valid but useless state, meaning the SM is idle. It happens between kernel launches or when no block has been assigned yet.
A block requests more shared memory than the SM physically has — what happens?
The kernel fails to launch (a launch-time error), because the SM cannot host even one such block. Resource requests are validated against per-SM limits before any thread runs.
A warp where all 32 lanes take the if and none take the else — is there divergence?
No — divergence requires disagreement within the warp. A unanimous branch executes a single path at full efficiency, exactly as if there were no branch.
Only 1 thread out of 32 is active in a warp (the rest masked) — how efficient is that lane utilization?
About — the warp still occupies a full 32-lane issue slot but does one lane's worth of work. This is the extreme cost of divergence and of block sizes badly misaligned to 32.
Launching more blocks than there are SMs — does anything break?
Nothing breaks; the Grid Distribution Unit queues surplus blocks and dispatches them as SMs free up. Blocks are designed to be independent precisely so they can run in any order or wave.
A memory-bound kernel with high occupancy but tiny arithmetic intensity — will occupancy save it?
Only up to the point where memory bandwidth saturates; beyond that, more warps just wait in longer queues. Occupancy hides latency, but it cannot manufacture bandwidth that the DRAM cannot deliver.
Recall Fast self-test before you close this
If a warp stalls and the SM still runs at full throughput, what condition made that possible? ::: There was at least one other eligible warp for the scheduler to switch to — latency hiding requires a spare, ready warp. Name the one resource whose over-use most directly crushes occupancy in a typical kernel. ::: Registers per thread — the register file is fixed, so heavy per-thread register use starves the SM of resident warps.