TL;DR: A kernel launch creates a grid of thread blocks. Each ordinary block executes on one streaming multiprocessor (SM), and its threads form 32-thread warps. Resource limits determine which blocks can reside; dependencies determine which warps can issue. More resident warps can hide waiting, but occupancy alone does not tell you how fast a kernel runs.
How to approach it
Separate the programmer's hierarchy from the hardware scheduler. Explain where threads can share data and synchronize, then count warps and per-block resources. End by asking whether the stalled kernel lacks eligible work, bandwidth or something else.
A strong answer
You launch thousands of blocks, but a profiler shows few eligible warps. Those statements are compatible. Most blocks may still be pending, while resident warps wait for data or a dependency. GPU Execution Model helps separate those states.
A grid is the collection of blocks for a launch. A block is a group of threads that can cooperate through block-scoped shared memory and synchronization. A warp is the 32-thread grouping used for SIMT (single instruction, multiple threads) execution: the hardware issues one instruction for all 32 threads at once. An SM holds resources for resident blocks and schedules their work. CUDA does not promise a round-robin block-to-SM mapping or a particular block order. See the programming model.
For 100 threads per block, the number of warps is ceil(100/32) = 4: three full warps and a fourth with four valid threads. The remaining lanes do not become threads available to a neighboring block. A conditional can further reduce the active lanes for a particular instruction. Independent thread scheduling does not make arbitrary warp communication safe without the appropriate synchronization.
The next question is how many blocks can reside together. Consider an illustrative block with 256 threads, 64 registers per thread and 16 KiB of programmer-requested shared memory. NVIDIA's Hopper guide documents 64 warps, 65,536 32-bit registers and up to 32 blocks per SM.
warps per block = 256 / 32 = 8
registers per block = 256 × 64 = 16,384
register-limited blocks = floor(65,536 / 16,384) = 4
warp-limited blocks = floor(64 / 8) = 8
four blocks contain 32 warps: 32 / 64 = 50% theoretical occupancy
sanity: those four blocks use all 65,536 registers in this simplified count.
The arithmetic points to registers as a possible residency limit, but it simplifies how the hardware allocates resources. The compiler's register count, allocation granularity, shared-memory configuration and other launch limits determine the actual result. Four blocks request 64 KiB of shared-memory payload in this example; that does not establish how much shared memory this launch can use. Check the compiled kernel and the occupancy API before treating four resident blocks as its real limit.
| Observation | What to investigate |
|---|---|
| Low resident warps | Registers, shared memory, block size and available grid work |
| Many residents, few eligible | Data dependencies, barriers and memory waits |
| Many eligible, low throughput | Issue limits, instruction mix and the actual critical path |
| High occupancy, slow kernel | Traffic, redundant work and synchronization; occupancy is not the objective |
Occupancy measures residency. A waiting warp can retain registers while another eligible warp issues; this is one latency-hiding mechanism. Independent instructions within a thread can also create useful outstanding work. Neither guarantees that a fixed HBM latency disappears, and adding warps cannot raise bandwidth beyond the memory system's capacity.
GPU Memory Hierarchy separates the data stores those instructions access. Use Nsight Compute's scheduler and occupancy sections to compare eligible and active warps. Change one launch parameter, verify results, and time the whole kernel. If reducing registers introduces spills and increases runtime, a higher occupancy percentage is a worse result. More arithmetic reuse may justify fewer resident blocks.
Block synchronization has a scope. __syncthreads() coordinates the participating block; it is not a grid-wide barrier. A block that spins waiting for another pending block can prevent progress. Cooperative launches and thread-block clusters add explicit capabilities with their own residency and synchronization requirements. Query supported cluster sizes; do not assume a 16-block cluster is portable across devices and configurations.
What interviewers probe next
- “Does 33 threads mean one warp plus one cheap thread?” It requires two warps, with only one valid lane in the second. Measure whether that layout makes sense for the workload.
- “Should every block have 1,024 threads?” No. A large block can restrict residency or waste lanes. Select it from resource use and measured work distribution.
- “Does a divergent branch always halve performance?” No. Its active masks, path lengths and surrounding work determine the effect; there is no universal multiplier.
- “Can blocks share ordinary shared memory?” Ordinary shared memory belongs to a block. Cluster distributed shared memory is a separate, explicitly synchronized feature.
Common mistakes
- Treating grid size as the number of simultaneously running blocks.
- Equating resident warps with eligible warps.
- Using block synchronization as a global barrier.
- Maximizing occupancy even when spills or redundant work increase.
Key takeaways
- A 100-thread block needs four warps, with four valid threads in the last.
- The illustrative register budget limits residency to four 256-thread blocks, or 50% occupancy.
- Correct synchronization and measured kernel time decide whether a configuration is useful.