Below the API

Execution Model: Threads, Warps, Blocks, Grids

Intermediate Intermediate 1h 15m Difficulty 3/5

Prerequisites 01, II.04

1. What is it?#

The hierarchy CUDA uses to organize parallel work.

GRID                    all the work for one kernel launch
 ├─ BLOCK 0             runs entirely on ONE SM; threads can cooperate
 │   ├─ WARP 0          32 threads executing in LOCKSTEP  ← the real unit
 │   │   ├─ thread 0
 │   │   ├─ thread 1
 │   │   └─ ... thread 31
 │   ├─ WARP 1
 │   └─ ...
 ├─ BLOCK 1
 └─ ...

The warp is the unit that matters. Threads are a programming convenience; the hardware schedules and executes warps.

Diagram — How a launch decomposes#

flowchart LR
  K["Kernel launch"] --> G["Grid<br/>all the blocks"]
  G --> B["Block<br/>runs on one SM, shares memory"]
  B --> W["Warp<br/>32 threads in lockstep"]
  W --> T["Thread<br/>one lane, own registers"]

  class K neutral
  class G io
  class B queue
  class W compute
  class T memory

2. Why does it exist?#

Because you need a model that maps naturally onto both the problem (millions of independent elements) and the hardware (SMs with SIMD lanes and limited shared resources).

The three levels correspond to three hardware realities:

  • Thread ↔ a lane in a SIMD unit.
  • Block ↔ resources on one SM (shared memory, synchronization).
  • Grid ↔ work distributed across all SMs.

3. Simple analogy#

An army.

A warp is a squad of 32 soldiers who march in perfect lockstep — they all take the same step at the same time. If the order is “go left,” and some soldiers should go right, the squad goes left with the right-goers standing still, then goes right with the left-goers standing still. That’s warp divergence, and it costs you exactly what it sounds like.

A block is a platoon assigned to one building (SM). They share a supply room (shared memory) and can coordinate (“everyone wait here” = __syncthreads()).

A grid is the whole operation. Platoons are assigned to buildings as buildings become available; they cannot coordinate with each other.

4. Tiny example#

__global__ void add(const float* a, const float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;   // the canonical index
    if (i < n) c[i] = a[i] + b[i];                    // the canonical guard
}

int main() {
    int n = 1'000'000;
    int threads = 256;
    int blocks  = (n + threads - 1) / threads;        // ceil division
    add<<<blocks, threads>>>(d_a, d_b, d_c, n);
}

Every CUDA kernel you write starts with those two lines. The index computation flattens the 2-level hierarchy into a global element id; the guard handles the case where n isn’t a multiple of blockDim.x.

n = 1,000,000, threads = 256
blocks = 3907
total threads = 3907 × 256 = 1,000,192   (192 idle, handled by the guard)
warps = 1,000,192 / 32 = 31,256

5. Technical explanation#

SIMT: threads are a fiction, warps are real#

Programming model:  32 independent threads
Hardware reality:   1 instruction stream, 32 data lanes, per-lane predication

This is why:

  • Threads in a warp cannot truly diverge — they mask.
  • There’s no per-thread program counter (pre-Volta; Volta+ has independent thread scheduling, but the execution is still warp-wide with masking).
  • Warp-level primitives (__shfl_sync, __ballot_sync, __reduce_add_sync) exist and are very fast, because the data is already in the same physical register file.

Choosing block size#

Constraints:
  - Must be a multiple of 32 (warp size). Non-multiples waste lanes.
  - Max 1024 threads per block.
  - Shared memory per block ≤ 48 KB (or up to 228 KB with opt-in on Hopper)
  - Registers per thread × threads per block ≤ 65,536 per SM partition

Practical guidance:
  128 or 256 threads is right for most kernels.
  Larger blocks: better for shared-memory-heavy kernels (more reuse).
  Smaller blocks: better granularity for load balancing, but more scheduling overhead.

Occupancy — the resource arithmetic#

An SM can host up to 64 warps (2048 threads). What limits you:

  Registers:  65,536 registers per SM partition ÷ (registers/thread × 32)
              = max warps from registers
  Shared mem: 228 KB ÷ shared_per_block = max blocks
  Blocks:     max 32 blocks per SM
  Threads:    max 2048 threads per SM

Occupancy = active_warps / 64

Example: a kernel using 64 registers/thread and 16 KB shared per 256-thread block:

Registers: 65536 / (64 × 32) = 32 warps per partition... 
           → with 4 partitions and 8 warps per block, ~fine
Shared:    228 KB / 16 KB = 14 blocks, but max 32 → not limiting
Threads:   2048 / 256 = 8 blocks
→ 8 blocks × 8 warps = 64 warps = 100% occupancy

Now raise registers to 128/thread:

Registers: 65536 / (128 × 32) = 16 warps per partition
→ occupancy drops to ~50%

Register pressure is the usual occupancy limiter. nvcc --ptxas-options=-v reports register usage; __launch_bounds__ lets you cap it (at the cost of spilling to local memory).

Grid sizing and wave quantization#

GPU has 132 SMs, each hosting (say) 4 blocks concurrently → 528 block slots.

Grid of 500 blocks:  one wave, 95% of slots used. Good.
Grid of 530 blocks:  two waves, second has 2 blocks. ~50% effective utilization.
Grid of 5000 blocks: ~9.5 waves, the fractional last wave costs ~5%. Fine.

Rule: either make the grid much larger than the SM count (so the tail is negligible), or size it to exactly fill one or a few waves. The bad case is a grid slightly larger than a whole number of waves.

For a decode kernel with batch 8 and 32 heads, the natural grid is 256 blocks — on a 132-SM GPU that’s ~2 waves with poor packing. This is why decode kernels use split-K parallelization (Section IV.06): to manufacture more blocks.

Thread block clusters (Hopper)#

New in Hopper: blocks can be grouped into clusters that are co-scheduled on the same GPC and can access each other’s shared memory (distributed shared memory, DSMEM).

__global__ void __cluster_dims__(2, 1, 1) my_kernel(...) { ... }

Useful for larger cooperative tiles than a single block’s shared memory allows. FlashAttention-3 and modern GEMMs use it.

6. Under the hood#

How a warp scheduler works, per cycle:

1. Look at all resident warps (up to 16 per scheduler)
2. Find those with their next instruction's operands ready
   (not waiting on a memory load, not waiting on a barrier)
3. Pick one (greedy-then-oldest policy)
4. Issue its instruction to the appropriate functional unit

If no warp is ready, the scheduler issues nothing — a stall cycle. The purpose of high occupancy is to make stalls rare, by ensuring some warp always has ready operands.

You can see this directly in Nsight Compute’s “Warp State Statistics”: it shows why warps were stalled (long scoreboard = waiting on memory, barrier = __syncthreads, etc.).

7. Performance implications#

Issue                        Cost
Block size not multiple of 32  up to 31/32 lanes idle in the last warp
Low occupancy from registers   memory latency exposed; 20-50% slower
Wave quantization              up to 50% on small grids
Warp divergence                up to 32x on the divergent region
Too few blocks (< #SMs)        SMs idle
Excessive __syncthreads()      serialization

8. Production implications#

You will rarely write kernels in production, but you will read profiles that report these metrics. Being able to interpret “achieved occupancy 23%” or “warp execution efficiency 40%” turns a profile from noise into a diagnosis.

When you do write kernels (custom fused ops, novel attention variants), these are the constraints you’re designing against.

9. Common mistakes#

Block sizes not a multiple of 32.

Chasing 100% occupancy. Occupancy is a means to hide latency. A kernel at 50% occupancy that uses more registers per thread (more instruction-level parallelism) often beats a 100%-occupancy version. Measure, don’t assume.

Forgetting the bounds guard. Out-of-bounds writes corrupt memory silently.

Assuming threads in a block run simultaneously. They run as warps, scheduled independently. Only __syncthreads() gives you an ordering guarantee.

Using __syncthreads() inside divergent code. Undefined behavior — all threads in the block must reach it.

10. Hands-on exercise#

A. Write the vector add. Implement, compile, and run the kernel in section 4. Verify correctness. Sweep block size over {32, 64, 128, 256, 512, 1024} and plot bandwidth. Explain the shape.

B. Occupancy calculator. For a kernel of your choosing, get register and shared memory usage from nvcc --ptxas-options=-v. Compute the theoretical occupancy by hand. Compare to what cudaOccupancyMaxActiveBlocksPerMultiprocessor reports and to Nsight’s measured occupancy.

C. Wave quantization. Write a kernel and launch it with grid sizes from 1 to 4×SM_count. Plot time vs grid size. Find the wave boundaries.

D. Warp primitives. Implement a block-wide sum reduction using __shfl_down_sync for the warp level and shared memory for the block level. Compare to a naive shared-memory-only version.

E. Occupancy vs performance. Take a kernel and artificially inflate its register usage (add unused local variables the compiler can’t eliminate). Measure occupancy and time. Does lower occupancy always mean slower?

11. Interview questions#

  1. Explain the thread/warp/block/grid hierarchy and what each maps to in hardware.
  2. What is SIMT and how does it differ from the thread abstraction it presents?
  3. What limits occupancy? How would you compute it for a given kernel?
  4. Is 100% occupancy always the goal? Explain.
  5. What is wave quantization and how do you avoid it?
  6. Why do decode attention kernels parallelize over the key dimension?
  7. What are warp-level primitives and why are they fast?

12. Further reading#

  • [REFERENCE] CUDA C++ Programming Guide, ch. 5 (Hardware Implementation) and the occupancy sections
  • [REFERENCE] NVIDIA CUDA Occupancy Calculator (now integrated into Nsight Compute)
  • [ESTABLISHED] Volkov, “Better Performance at Lower Occupancy” (GTC 2010) — the classic counter-argument to occupancy maximization
  • Next: 03 — Memory hierarchy
esc

↑↓ navigate ↵ open