TL;DR: Software is grid of blocks of threads; hardware is GPU of SMs running warps of 32 threads in SIMT lockstep. The scheduler pins each block to one SM and hides hundreds of cycles of memory latency by switching among up to 64 resident warps, which is why GPUs want tens of thousands of threads in flight, not a few fast ones.
How to approach it
Answer in two passes: first the software hierarchy (grid, block, thread), then how each level lands on hardware (GPU, SM, warp). The interviewer is checking whether you understand that the GPU hides latency with massive parallelism rather than with the caches and branch predictors a CPU uses. Say that explicitly, because it is the idea everything else hangs on.
A strong answer
When you launch a kernel you specify a grid of thread blocks, each block containing up to 1,024 threads. That is the software view. On hardware the GPU is a collection of streaming multiprocessors (an H100 has 132 SMs), and the hardware scheduler assigns each block to exactly one SM, where it stays until it finishes. Within the SM the block is carved into warps of 32 threads, and the warp is the real unit of execution: a warp scheduler issues one instruction for all 32 lanes in SIMT fashion every cycle it can.
Two design consequences matter. First, blocks must be independent: they cannot reliably synchronize with each other, only threads within a block can (via __syncthreads() and shared memory). That independence is what lets the same binary scale from a laptop GPU with 20 SMs to a datacenter part with 132. The scheduler just distributes blocks across whatever SMs exist.
Second, latency hiding. A global memory load costs hundreds of cycles. A CPU hides that with big caches and out-of-order execution; an SM hides it by oversubscription. It keeps up to 64 warps resident (2,048 threads per SM) and whenever one warp stalls on memory, the scheduler issues from another in a zero-cost context switch. This is why you want far more threads than cores: an H100 wants on the order of 100k+ threads in flight to stay busy. It is also why occupancy, register pressure, and divergence are performance topics at all. They all gate how many warps the scheduler can choose from.
Register pressure is the best place to make "occupancy" a number instead of a vibe, because it is one division. An SM has 65,536 32-bit registers and caps at 2,048 resident threads. A kernel compiled to 64 registers per thread fits 65,536 / 64 = 1,024 threads, which is 32 warps: the scheduler is working with half its maximum hand of latency-hiding cards before your code runs a single instruction. Trim the kernel to 32 registers per thread and 2,048 threads fit, the full 64 warps. Neither number appears anywhere in your source code; the compiler chooses, -Xptxas -v reports it, and this is why "my kernel got slower when I added three local variables" is a real bug report with a real mechanism. It is also the honest answer to why occupancy is not something you maximize blindly: sometimes 32 warps of a register-rich kernel that never spills beats 64 warps of one that does.
A concrete mapping for a 4096×4096 matrix op: launch a grid of 128×128 blocks of 32×32 threads; each block tiles a chunk of the output, gets pinned to an SM, runs as 32 warps, and the scheduler interleaves thousands of blocks across the chip.
What interviewers probe next
- "Why 32 threads in a warp?" It is the hardware SIMT width; block sizes should be multiples of 32 or you pay for phantom lanes in partial warps.
- "Can block 5 wait for block 3?" No, and designs that assume inter-block ordering deadlock; grid-wide sync means ending the kernel (or cooperative groups with strict limits).
- "When is a GPU slower than a CPU?" Low parallelism, branch-heavy or pointer-chasing code, latency-sensitive serial work, or anything dominated by the PCIe transfer. If you cannot fill tens of thousands of threads, the SMs idle and a CPU's higher clocks and caches win.
- "What happens if my block needs more registers than the SM has?" Fewer blocks become resident, occupancy drops, latency hiding degrades; the compiler may also spill to local memory.
Common mistakes
- Treating GPU threads like CPU threads (heavyweight, independently scheduled). The warp is the scheduling unit; per-thread reasoning leads you wrong on divergence and coalescing.
- Saying blocks map to "cores." Blocks map to SMs; the CUDA-core marketing count refers to ALU lanes, not independent processors.
- Not knowing that blocks cannot synchronize with each other. This single fact is what interviewers use to separate people who have debugged real kernels from people who have read a slide deck.
- Skipping latency hiding entirely. If your explanation never mentions warp switching to cover memory stalls, you have described the org chart but not why it is shaped that way.
Key takeaways
- Grid/block/thread is software; GPU/SM/warp is silicon, and the block-to-SM pinning plus warp-as-scheduling-unit are the facts that decide performance.
- Latency hiding through warp oversubscription, not caches, is why occupancy and register pressure matter.
- Blocks are independent by design; assuming inter-block sync is the classic deadlock and the classic interview tell.