GPU Streaming Multiprocessor Warps

In short: A GPU is an array of Streaming Multiprocessors (SMs); each SM splits into four sub-partitions with their own warp scheduler, ALU lanes and register file slice, and the hardware places whole thread blocks onto SMs, slicing them into 32-thread warps that the schedulers swap every cycle to hide memory latency.

"Draw a GPU and explain how a thread block runs on it" is a favourite whiteboard question at NVIDIA, AMD, Qualcomm and every AI-accelerator startup. Interviewers are not testing whether you memorised a marketing slide. They want to see that you can connect the software model (grid, block, thread) to the silicon (SM, sub-partition, lane) and reason about why performance behaves the way it does. This walkthrough uses NVIDIA's Hopper H100 numbers as a concrete reference, and the same reasoning transfers to AMD compute units and other SIMT designs. For the chip-level picture first, read our GPU architecture explained article.

1. The top-level picture: GPCs, SMs and memory

Start the whiteboard from the outside in. A full H100 die is organised as graphics processing clusters (GPCs), each holding several texture processing clusters, each holding two SMs. The SXM part enables 132 SMs. All SMs share a large L2 cache (50 MB on H100) and reach HBM3 through memory controllers on the periphery. The SM is the unit of replication: designers verify one SM thoroughly, then tile it. That is also why the SM is the unit that matters to a programmer, since a thread block never spans two SMs.

Key fact: A thread block is always resident on exactly one SM from launch to completion. Blocks cannot migrate, and blocks on different SMs can only communicate through global memory (or clusters and distributed shared memory on Hopper).

2. Inside a sub-partition: scheduler, lanes and registers

Each of the four sub-partitions contains one warp scheduler and one dispatch unit. Every cycle, that scheduler can issue one instruction from one ready warp. A warp is 32 threads executing the same instruction on different data (SIMT). With 32 FP32 lanes per sub-partition, a simple FP32 add for a whole warp completes its issue in one cycle. Multiply by four sub-partitions and you get 128 FP32 "CUDA cores" per SM. Note that "CUDA core" is a marketing term for one FP32 lane, not an independent processor with its own instruction fetch.

Alongside the FP32 lanes sit INT32 lanes, FP64 units (64 per SM on H100), special function units for sin, rcp and rsqrt (16 per SM), a tensor core per sub-partition for matrix multiply-accumulate, and load/store units (32 per SM, so a warp memory instruction takes about four issue slots across the SM when all sub-partitions are busy). The register file is 64K 32-bit entries per SM, 256 KB in total, which is physically larger than the L1. That is a deliberate GPU choice: keep thousands of threads' state on chip so a context switch costs zero cycles.

• Warp scheduler: picks one eligible warp per cycle using a scoreboard; eligible means operands ready, no barrier pending, and the target unit free.
• Register file: statically partitioned among resident warps at launch. No spilling to a stack unless the compiler forces local memory.
• Load/store unit: coalesces the 32 addresses of a warp into as few 32-byte sectors (128-byte lines) as possible.
• Shared memory: software-managed scratchpad, 32 banks wide; two threads hitting the same bank with different addresses serialise.

3. Warps and divergence

Threads in a block are numbered linearly (x fastest) and cut into groups of 32 to form warps. A 100-thread block therefore becomes four warps, the last one only 4/32 occupied, wasting 28 lanes. This is why block sizes should be multiples of 32; 128 or 256 are common defaults.

Because all 32 threads share one program counter in the classic model, an if/else whose condition differs inside a warp executes both paths one after the other with inactive lanes masked off. Since Volta each thread has its own PC and call stack for forward-progress guarantees, but throughput still drops under divergence. Here is a classic interview trap:

// divergent: odd and even lanes take different paths in EVERY warp
if (threadIdx.x % 2) a(); else b();

// warp-uniform: whole warp takes the same path, no penalty
if ((threadIdx.x / 32) % 2) a(); else b();

Key fact: Divergence cost is per warp, not per block. Branching on threadIdx.x / 32 is free; branching on threadIdx.x % 2 halves throughput of that region.

4. How a thread block maps to hardware

When you launch kernel<<>>, a hardware block scheduler (the GigaThread engine) hands blocks to SMs whenever an SM has enough free resources. Resources, not just a count, decide whether a block fits. An SM on H100 can hold at most 2048 threads (64 warps), 32 blocks, 64K registers and about 228 KB of shared memory. Whichever limit is hit first caps occupancy.

Worked example: a kernel uses 256-thread blocks, 64 registers per thread and 24 KB of shared memory per block.

Registers win, so 4 blocks = 1024 threads = 32 warps, or 50% theoretical occupancy. Dropping to 32 registers per thread doubles it to 100%, at the risk of spills. Nsight Compute reports exactly this table under "Occupancy", and cudaOccupancyMaxActiveBlocksPerMultiprocessor returns the same answer programmatically.

5. Latency hiding: why warps beat caches

A global memory load takes roughly 400 to 800 cycles on HBM. A CPU fights that with big caches and out-of-order execution. A GPU simply has other warps to run. If a sub-partition holds 16 resident warps and each issues a load then waits, the scheduler rotates through the others, and by the time warp 0 gets its data back, the ALUs have stayed busy with warps 1 to 15. This is why occupancy matters, but only up to the point where there are enough ready warps to cover latency. Many tuned kernels, such as large GEMMs, run at 25 to 50% occupancy because they use more registers per thread for instruction-level parallelism and tile reuse.

6. Whiteboard checklist and common follow-ups

When you draw this in an interview, hit these beats in order: grid to blocks, blocks to SMs, SM to four sub-partitions, warps to schedulers, then the three limiters (registers, shared memory, thread slots). Then volunteer one performance consequence, such as coalescing or divergence, before they ask. Expect follow-ups on bank conflicts, why shared memory and L1 share a pool, how tensor cores change the lane count story, and how AMD's 64-wide wavefronts or Apple's SIMD groups differ.

Chip designers verify and implement SM-class blocks with the usual flow: RTL simulation in Synopsys VCS or Siemens Questa, synthesis with Synopsys Design Compiler or Cadence Genus, and place and route in Cadence Innovus or Synopsys IC Compiler II. The register file and shared memory SRAMs dominate SM area and power, which is why compilers and architects obsess over register pressure. If you want to see how such blocks reach silicon, the team at Avecas VLSI design services works on exactly this kind of physical-design problem. For a deeper look at the memory side, see GPU memory hierarchy from registers to HBM.

Want to go from block diagrams to real RTL and implementation flows? Explore the AI hardware and VLSI tracks in our CourseTron courses and practise architecture interview questions with mentors.

Frequently asked questions

How many CUDA cores are in one SM?

It depends on the generation. Hopper H100 has 128 FP32 lanes per SM (32 per sub-partition), Ampere GA100 has 64, and consumer Ampere and Ada have 128. A CUDA core is only a single-precision lane, not a full core.

What is the difference between a warp and a thread block?

A thread block is the programmer-defined group (up to 1024 threads) that shares shared memory and can synchronise with __syncthreads(). A warp is the hardware's 32-thread execution unit; a block is split into warps automatically.

Can two thread blocks share one SM?

Yes. Several blocks are resident on an SM simultaneously as long as registers, shared memory, thread slots and block slots suffice. They do not share shared memory with each other.

Why does warp divergence hurt performance?

The lanes in a warp issue together, so when threads take different branches the hardware runs each path in turn with inactive lanes masked. Throughput falls roughly in proportion to the number of distinct paths taken within the warp.