Question bank — GPU architecture — SIMT, warps, CUDA model

Thread ::: one running copy of the kernel with its own index and registers. Warp ::: the hardware bundle of 32 consecutive threads that fetch/decode ONE instruction together and advance in lockstep. Block ::: a group of threads that live on one Streaming Multiprocessor (SM), share on-chip shared memory, and can call __syncthreads(). Occupancy ::: active warps on an SM divided by the max warps that SM supports — a measure of how many "spare" warps exist to hide stalls. Divergence ::: when threads in the same warp take different branches, forcing the hardware to run both paths one after the other.

SIMT means every thread in a warp is forced to execute the identical operation, exactly like SIMD lanes. ::: False. Each thread has its own program-counter state and can branch independently; lockstep is the default, not a hard rule. Divergence costs speed but never breaks correctness. A block of 32 threads and a block of 33 threads use the same number of warps. ::: False. 32 threads = 1 warp; 33 threads spills into a second warp that runs 31 idle (masked) lanes, so it uses 2 warps. Increasing occupancy always increases performance. ::: False. Once you have enough warps to hide memory latency, more occupancy gives diminishing returns and can hurt — more threads means fewer registers per thread, causing spills to slow local memory. __syncthreads() synchronizes the whole grid. ::: False. It only barriers threads within one block (which lives on one SM). Cross-block synchronization requires a new kernel launch or cooperative groups. If all 32 threads in a warp take the then branch of an if, there is zero divergence penalty. ::: True. Divergence is measured per warp: a warp-uniform branch runs one path only, so the cost is $\max$, not the sum of both paths. A GPU hides memory latency mainly with large caches, like a CPU does. ::: False. A GPU hides latency by oversubscribing threads and switching to a ready warp when one stalls; a CPU avoids latency with big caches and out-of-order engines. Two threads in the same warp always share the same registers. ::: False. Threads share a fetched instruction, but each has its own private registers and its own PC state — that is exactly what makes SIMT programmable as scalar code. Making block size a multiple of 32 never wastes warp lanes. ::: True. Warps are exactly 32 wide; a block size divisible by 32 fills every warp completely, leaving no masked-off idle lanes. A stride-1 access by a warp and a stride-32 access move the same amount of data. ::: False. Stride-1 packs 32 reads into ~1 transaction (efficiency 1); stride-32 may touch 32 separate cache lines for the same 128 useful bytes (efficiency $1/32$), moving ~32× more data.

A student writes the bounds guard as if (i < N) return;. ::: Wrong condition. return is the early exit, so you must exit the invalid threads: if (i >= N) return;. As written it kills every valid thread and keeps only the out-of-range ones. A student computes the global index as i = threadIdx.x * blockDim.x + blockIdx.x. ::: Terms swapped. The block offset is blockIdx.x * blockDim.x (skip whole earlier blocks), then add threadIdx.x (position inside your block): i = blockIdx.x*blockDim.x + threadIdx.x. Someone says "blockDim is the number of blocks in the grid." ::: Name confusion. blockDim = threads per block; gridDim = number of blocks. Check units: blockIdx*blockDim must produce a thread offset. To avoid divergence, a student branches on if (threadIdx.x % 2 == 0). ::: Maximal divergence. Even/odd interleaves the two paths within every warp, so each warp runs both paths. Branch on threadIdx.x / 32 (or block-level conditions) so each warp is uniform. Code launches exactly N/B blocks (integer division) for N elements. ::: Drops the tail. Integer division truncates, so a partial last block is lost. Use $\lceil N/B\rceil =(N+B-1)/B$ blocks plus an if (i >= N) return; guard. A student claims divergence makes the result incorrect. ::: Overstated. Divergence only serializes the two paths and costs time; the masking guarantees each thread still executes its own correct path. It is purely a performance issue. "Use __syncthreads() to make block 0 finish before block 1 starts." ::: Impossible primitive. __syncthreads() cannot order two different blocks; blocks are independent and may run in any order. Enforce the dependency by splitting into two kernel launches.

Why is the warp size a fixed 32 rather than something you choose? ::: Because it is a hardware constant — one instruction is physically broadcast to 32 lanes. Your block size can be anything, but it is always chopped into 32-wide warps. Why does a GPU spend transistors on ALUs instead of caches and branch predictors? ::: Its target workloads (graphics, matmul, ML) have huge amounts of identical independent work, so it bets that parallelism is cheaper than locality — keep many threads in flight rather than making one thread fast. Why must divergent branches be serialized within a warp? ::: There is only one instruction fetch/decode unit per warp, so both paths cannot be issued at the same cycle; the hardware runs one path (masking the others), then the other. Why do we usually launch more threads than there are data elements? ::: Because we round the block count up to cover the last partial block, some threads land past the data; the if (i >= N) return; guard retires exactly those extras. Why does higher occupancy help hide latency? ::: More active warps means more "spare tires" ready to run; when one warp stalls on a ~400-cycle load, the scheduler swaps in a ready warp so the ALUs never sit idle. Why is SIMT described as "SIMD execution with a scalar programming model"? ::: You write code for one scalar thread (c[i]=a[i]+b[i]); the hardware gathers 32 such threads into a vector warp — so you get vector throughput without hand-writing vector instructions. Why does memory coalescing depend on the access pattern, not just the number of bytes? ::: The hardware fuses a warp's 32 addresses into as few fixed-size transactions as possible; consecutive addresses fit one line, scattered ones spill across many lines, so pattern decides how much data actually moves.

What happens to a block of exactly 1 thread? ::: It still occupies a full warp — 1 active lane, 31 masked idle lanes — wasting 31/32 of that warp's throughput. What is occupancy when only 1 warp is active on an SM that supports 64? ::: $1/64\approx 1.6\%$ — almost no spare warps, so a single stall exposes the full memory latency with nothing to hide it. A warp where all 32 threads take different branches of a many-way switch — cost? ::: Worst case: the warp serializes through every distinct path taken, so time ≈ sum of all those paths (this is the divergence ceiling). N is an exact multiple of block size B — is the if (i >= N) return; guard still needed? ::: Harmless but not strictly needed there, since no thread exceeds N; keep it anyway because it costs nothing and protects against non-multiple cases. What if a load's latency is longer than any warp's arithmetic can cover, even at max occupancy? ::: Latency is only partially hidden — the SM runs out of ready warps and stalls, so throughput drops. This is when memory-bound kernels need coalescing or fewer/cheaper accesses, not more threads. A warp reads a[i] but half the threads are masked off by a bounds guard — does coalescing still work? ::: Yes; the active addresses are still consecutive, so they coalesce into ~1 transaction. Masked lanes simply do not contribute addresses. Does zero divergence guarantee zero performance loss for a warp? ::: No — the warp still pays for memory latency, uncoalesced accesses, or shared-memory bank conflicts. Divergence-free only removes the branch-serialization penalty, not every stall.