TL;DR: A kernel launch is a grid of thread blocks; each block is placed whole on one streaming multiprocessor and stays there until it finishes; inside the SM the hardware schedules in warps of 32 threads, one instruction per warp per issue slot. Latency is hidden because an SM holds up to 64 resident warps and switches between them every cycle at zero cost, so while one warp waits 600 ns on HBM, dozens of others issue. The programmer's job is to give the SM enough independent warps, and enough independent work per warp, that the memory pipe never drains.
How to approach it
Say the three programmer-side levels first (grid, block, thread), then immediately translate each into what the hardware does with it, because the interviewer is scoring the translation, not the vocabulary. Ask whether they want the Hopper numbers (132 SMs, 64 warps per SM) or the general model. Work one example kernel through placement: how many blocks land per SM and how many warps that gives the scheduler. Close with the latency-hiding sentence, because that is the idea everything else in GPU performance is built on.
A strong answer
A typical situation: a kernel is launched with 1,024 threads per block because that was the largest number the docs mentioned, it runs at two thirds of the speed of the same kernel at 256, and nobody in the review can say why. The answer is in the placement arithmetic below, and it needs both vocabularies.
The GPU Execution Model has two vocabularies that map onto each other.
On the software side, a launch is kernel<<<gridDim, blockDim>>>. The grid is the whole problem, the block is a group of threads that can share memory and synchronize with __syncthreads(), and the thread is the unit the code is written for. Threads inside a block have an index; blocks inside a grid have an index; nothing in the language says how many run at once.
On the hardware side, an H100 has 132 SMs. When the grid is launched, a hardware work distributor hands blocks to SMs. A block is placed whole on one SM: all of its threads live there, its shared memory is carved from that SM's 228 KB, its registers come from that SM's 64K-entry register file, and it does not migrate. When every thread in the block exits, the resources are freed and the distributor places the next block. Blocks that have not been placed yet are simply waiting; that is why launching 10,000 blocks on 132 SMs is fine and why blocks must not depend on each other running concurrently.
Inside the SM, the unit that gets scheduled is the warp: 32 consecutive threads of a block. The SM has four warp schedulers (four partitions), and each cycle each scheduler picks one eligible resident warp and issues one instruction for all 32 lanes. A thread does not exist to the scheduler; only warps do. This is why a block of 33 threads costs two warps, why divergence inside a warp serializes, and why memory access patterns are judged per warp.
The placement arithmetic is a chain a candidate should be able to do out loud:
SM limits (H100 SXM): 2,048 threads, 64 warps, 32 blocks, 65,536 registers, 228 KB shared per SM
kernel: 256 threads per block, 64 registers per thread, 16 KB shared per block
threads limit: 2,048 ÷ 256 = 8 blocks
register limit: 65,536 ÷ (256 × 64) = 65,536 ÷ 16,384 = 4 blocks
shared limit: 228 KB ÷ 16 KB = 14 blocks
resident blocks = min(8, 4, 14) = 4 blocks
resident warps = 4 × (256 ÷ 32) = 32 warps of a possible 64 (50% occupancy)
sanity: registers are the binding limit here, which is the usual case for anything
beyond a trivial kernel; halving registers per thread would double resident warps
Now the latency question. An HBM access on H100 is on the order of 600 ns, which at 1.8 GHz is about 1,000 cycles. A CPU core spends most of its transistors making sure one thread rarely waits that long: big private caches, out-of-order windows, prefetchers, branch predictors. An SM does none of that per thread. Its L1 is 256 KB shared by up to 2,048 threads, about 128 bytes each, and there is no out-of-order execution. Instead, when a warp issues a load, the scheduler marks it not-eligible and, on the very next cycle, issues from a different warp. Switching warps costs nothing because every resident warp's registers are already in the register file; there is no context to save. With 32 to 64 resident warps each holding several loads in flight, the SM keeps enough bytes moving to cover the latency.
How many bytes must be in flight is another short chain:
per-SM share of bandwidth = 3.35 TB/s ÷ 132 SMs ≈ 25 GB/s per SM
bytes in flight to cover latency = bandwidth × latency = 25 GB/s × 600 ns ≈ 15 KB per SM
with one 16-byte load per thread outstanding: 15 KB ÷ 16 B ≈ 950 threads ≈ 30 warps
sanity: 30 warps is about half the SM's 64, so a kernel at 50% occupancy with one load in
flight per thread is roughly at the edge; more loads per thread (unrolling) buys
the same coverage with fewer warps
That last line is the practical lesson: latency hiding comes from parallelism, whether across warps (occupancy) or within a warp (independent loads per thread), and the two are interchangeable. Caches on a GPU exist to save bandwidth on reuse, not to make a single thread fast.
The decision this drives: choose block size and register budget so that the SM has enough resident warps, or enough loads per thread, to cover latency, and no more. The reversal condition: a kernel whose per-thread state is large (attention, GEMM tiles), where forcing higher occupancy spills registers to local memory and the extra traffic costs more than the extra warps save. Occupancy and Register Pressure is the whole of that trade, and ptxas -v (via nvcc -Xptxas -v) prints the register count and the spill bytes that decide it, which is faster than reasoning about the source.
What interviewers probe next
- "Why is a warp 32 threads?" It is the width of the SIMT datapath the scheduler issues to; the software has no say. A block of 100 threads is 4 warps, the last one 4/32 utilized, so pick multiples of 32.
- "Can two blocks on the same SM talk to each other?" Not through shared memory, which is per block. They can through global memory and atomics, but nothing guarantees they are resident at the same time unless you use cooperative launch.
- "What happens when a block calls
__syncthreads()and one warp has exited?" On current hardware exited threads are treated as arrived; the classic hang is a barrier inside a branch that only some threads take. - "Where does a thread block cluster fit?" Hopper adds a level between block and grid: up to 16 blocks guaranteed co-resident on one GPC with distributed shared memory. It is the first time the model has let blocks depend on each other's residency.
Common mistakes
- Saying the GPU "runs all the threads in parallel." It runs up to 2,048 per SM at a time and queues the rest; the model only works because blocks are independent.
- Describing warps as a software construct or a "group of 32 threads you pick." Warps are formed by the hardware from consecutive thread indices.
- Claiming the L1 or L2 cache hides memory latency. It reduces traffic on reuse; latency is hidden by having other warps to issue.
- Computing occupancy from threads only and forgetting registers, which are the binding limit for almost every real kernel.
Key takeaways
- Grid is the request, block is the residency unit (whole block on one SM, never moves), warp of 32 is the scheduling unit.
- H100 SM: 132 SMs, up to 64 warps and 2,048 threads resident, 65,536 registers, 228 KB shared; registers usually decide how many blocks fit.
- Latency is hidden by switching between resident warps at zero cost; you need about 15 KB in flight per SM to cover 600 ns at H100 bandwidth.
- Occupancy and per-thread independent loads are interchangeable ways to get those bytes in flight.