Warp Execution & Scheduling

Understanding warp execution is the key to understanding GPU performance. The GPU doesn’t see individual threads — it sees warps. Every performance optimization, from avoiding divergence to maximizing occupancy, operates at the warp level.

1. Warp Formation and Scheduling

How Warps Are Formed

When a block is assigned to an SM, its threads are partitioned into warps of 32:

Block (256 threads)
├── Warp 0:  threads   0 –  31
├── Warp 1:  threads  32 –  63
├── Warp 2:  threads  64 –  95
├── Warp 3:  threads  96 – 127
├── Warp 4:  threads 128 – 159
├── Warp 5:  threads 160 – 191
├── Warp 6:  threads 192 – 223
└── Warp 7:  threads 224 – 255

Grouping is by consecutive threadIdx — you cannot choose which threads form a warp.

Warp Schedulers

Each SM has multiple warp schedulers (4 on H100). Every cycle, each scheduler:

Checks which warps are ready (instruction and data available)
Selects one ready warp
Issues its next instruction

If a warp is waiting for memory (stalled), the scheduler immediately switches to another ready warp — zero-cost context switching. This is how GPUs hide memory latency without caches.

Cycle 1: Scheduler picks Warp 0 → issues load instruction (takes 600 cycles)
Cycle 2: Warp 0 stalled → Scheduler picks Warp 1 → issues compute
Cycle 3: Warp 1 ready  → Scheduler picks Warp 1 → issues compute
...
Cycle 600: Warp 0's data arrives → Scheduler picks Warp 0 → resumes

Key insight: The GPU hides latency by having more warps than execution units. While some warps wait for data, others compute. This is why occupancy matters.

2. SIMT: Single Instruction, Multiple Threads

All 32 threads in a warp execute the same instruction each cycle, but each thread has its own:

Register file (private data)
Program counter (since Volta — independent thread scheduling)
Predicate registers (for masking)

SIMT vs SIMD

Property	SIMD (CPU AVX-512)	SIMT (GPU warp)
Width	8-16 elements	32 threads
Branching	Must process all elements same way	Can diverge (with penalty)
Programming	Explicit vector instructions	Write scalar code, hardware parallelizes
Memory access	Fixed stride	Each thread can access different address

SIMT is more flexible — you write code as if for a single thread, and the hardware runs 32 copies. But that flexibility comes with the divergence penalty.

3. Warp Divergence: The Performance Tax

What Happens

Warp Divergence

When threads in a warp take different branches, both paths execute sequentially:

if (threadIdx.x % 2 == 0) {
    a = computeA();  // Even threads: active, odd threads: masked
} else {
    b = computeB();  // Odd threads: active, even threads: masked
}
// Total time: time(A) + time(B), not max(time(A), time(B))

Measuring the Cost

Divergence pattern	Active threads	Throughput
No divergence (all same path)	32/32	100%
50% divergence (half each way)	16/32 per path	50%
Full divergence (32 different paths)	1/32 per path	3%

Mitigation Strategies

1. Align branches with warp boundaries:

// BAD: divergence within every warp
if (threadIdx.x % 2 == 0) { ... }

// GOOD: entire warps take the same path
if (threadIdx.x / 32 < 4) { ... }  // warps 0-3 take one path, warps 4-7 take another

2. Predication for short branches: The compiler may convert short if-else into predicated instructions (both paths execute, results selected by predicate). No warp serialization.

3. Data reorganization: Sort data so threads in a warp process similar items (e.g., sort by class before classification).

Volta+ Independent Thread Scheduling

Before Volta, all threads in a warp shared a single program counter — diverged threads were strictly serialized. Since Volta (V100, 2017), each thread has its own program counter, enabling:

Finer-grained reconvergence: threads can rejoin after divergence without waiting for the entire warp
Better interleaving: the scheduler can interleave instructions from diverged paths
Starvation-free execution: no thread is permanently masked

In practice, this means divergence is less catastrophic on modern GPUs — but it’s still a performance loss and should be minimized.

4. Occupancy: Keeping the GPU Busy

Occupancy = active warps per SM / maximum warps per SM.

Higher occupancy → more warps available for the scheduler → better latency hiding.

What Limits Occupancy

Resource	Limit (H100)	How it limits occupancy
Threads per SM	2048 (64 warps)	Hard ceiling
Threads per block	1024	Block size caps warps per block
Registers per SM	65,536	More registers/thread → fewer threads fit
Shared memory per SM	228 KB	More smem/block → fewer blocks fit

Example Calculation

A kernel uses 64 registers/thread and 16 KB shared memory/block, with 256 threads/block:

Warps per block: 256 / 32 = 8
Blocks limited by registers: 65,536 / (64 × 256) = 4 blocks
Blocks limited by shared mem: 228 KB / 16 KB = 14 blocks
Blocks limited by threads: 2048 / 256 = 8 blocks
Actual blocks per SM: min(4, 14, 8) = 4 blocks (register-limited)
Occupancy: 4 × 8 warps = 32 active warps / 64 max = 50%

Using the Occupancy Calculator

# NVIDIA provides an occupancy calculator spreadsheet and API
nvcc --ptxas-options=-v kernel.cu  # shows register/shared memory usage

// Programmatic occupancy query
int minGridSize, bestBlockSize;
cudaOccupancyMaxPotentialBlockSize(&minGridSize, &bestBlockSize, kernel, 0, 0);

Is Higher Always Better?

Not always. Some kernels achieve peak performance at 50% occupancy because they’re compute-bound (every thread is busy computing, not waiting for memory). The goal is to have enough occupancy to hide memory latency — beyond that, more occupancy doesn’t help.

Rule of thumb: Aim for ≥50% occupancy. If a kernel is compute-bound, even 25% may be sufficient. Profile to confirm.

5. Latency Hiding: The Full Picture

The GPU hides latency at multiple levels:

Level	Mechanism	Latency hidden
Warp-level	Scheduler switches between warps	Memory latency (~600 cycles)
Instruction-level	Independent instructions in a warp execute while earlier ones wait	Short dependency chains
Thread-level	More threads per SM = more warps to switch between	All stalls

The formula for how many warps are needed to hide memory latency:

\text{Warps needed} \geq \frac{\text{Memory latency (cycles)}}{\text{Cycles per instruction}} = \frac{600}{4} = 150 \text{ warps (theoretical)}

In practice, instruction-level parallelism reduces this — 20-40 active warps typically suffice.

Self-Check

Explain how warps are formed from block threads
Explain how warp schedulers hide memory latency
Calculate the performance cost of 50% warp divergence
Name three strategies to minimize divergence
Calculate occupancy given register and shared memory usage
Explain why 50% occupancy can be sufficient for compute-bound kernels
Explain the relationship between occupancy and latency hiding