Warps and warp scheduling

Overview

Warps are the fundamental unit of execution in NVIDIA GPUs, representing a group of threads that execute the same instruction in lockstep. Understanding warp scheduling is critical for optimizing GPU performance because the GPU scheduler operates at warp granularity, not individual threads.

Core Concepts

How Warps Map to Thread Blocks

Derivation from thread indices:

Given a 3D thread block with dimensions (blockDim.x, blockDim.y, blockDim.z), threads are assigned to warps in row-major order:

\text{Linear Thread Index} = \text{threadIdx.z} \times (\text{blockDim.x} \times \text{blockDim.y}) + \text{threadIdx.y} \times \text{blockDim.x} + \text{threadIdx.x}

$\text{Warp ID within Block} = \left\lfloor \frac{\text{Linear Thread Index}}{32} \right\rfloor$

Why this matters: If blockDim.x = 64, threads0-31 form warp 0, threads 32-63 form warp 1. If blockDim.x = 17, threads 0-31 form warp 0, but only 17 threads are active—the other 15 slots are wasted. This is why block dimensions should be multiples of 32.

Warp Scheduling Mechanics

The Scheduling Problem

An SM can hold multiple thread blocks simultaneously (limited by registers, shared memory, and max blocks/SM). Each block contributes multiple warps. If an SM can hold 2048 threads and has 64 warps resident, the scheduler must decide which warp executes next.

Derivation of why zero-overhead switching works:

When warp W₁ stalls on memory, the scheduler immediately switches to warp W₂ in the same cycle. This is possible because:

Register files are partitioned: Each warp has dedicated register space (no context save/restore needed)
Program counters are per-warp: Switching warps = updating one pointer (1 cycle)
Instruction buffers are separate: Each warp's next instruction is pre-fetched

Latency hiding capacity:

If a memory operation takes $L$ cycles and each warp performs one memory operation per $W$ instructions, you need:

N_{\text{warps needed}} = \left\lceil \frac{L}{W} \right\rceil

Why? If $L = 400$ cycles and $W = 10$ instructions between memory ops, you need 400/10 = 40 warps to keep the SM busy while one warp waits.

Scheduling Policies

Modern GPUs use variations of these policies:

Round-robin with priority: Cycle through eligible warps, prioritizing long-stalled warps
Gredy-then-oldest (GTO): Execute one warp until it stalls, then pick the oldest stalled warp
Two-level scheduling: Prioritize warps likely to hit in cache

Formula Occupancy and Scheduling Efficiency

Occupancy is the ratio of active warps to maximum warps per SM:

\text{Occupancy} = \frac{\text{Active Warps per SM}}{\text{Max Warps per SM}}

Why occupancy matters: Higher occupancy → more warps available → better latency hiding. But occupancy alone doesn't guarantee performance; you also need warp diversity (warps at different instruction stages).

Deriving optimal occupancy:

Let $T_{\text{issue}}$ = cycles between instruction issues per scheduler, $L$ = average stall latency, $N$ = number of schedulers.

To keep all schedulers busy:

\text{Warps needed} \geq N \times \frac{L}{T_{\text{issue}}}

For $N=4$ schedulers, $L=400$ cycle memory latency, $T_{\text{issue}}=1$ (one instruction per cycle scheduler):

\text{Warps needed} \geq 4 \times \frac{400}{1} = 1600 \text{ instructions' worth}

If each warp has 20 independent instructions before stalling: 1600/20 = 80 warps. But max is typically 64, so you need instruction-level parallelism within warps too.

Warp Divergence

Derivation of divergence penalty:

Consider an if-else where $n_{\text{if}}$ threads take the if branch and $n_{\text{else}} = 32 - n_{\text{if}}$ take the else branch.

Without divergence (all threads take same path):

T_{\text{normal}} = T_{\text{if}}

With divergence (assuming $n_{\text{if}} > 0$ and $n_{\text{else}} > 0$ ): $T_{\text{diverged}} = T_{\text{if}} + T_{\text{else}}$

Worst case: Both branches have equal cost $T$ and are equally populated:

\text{Slowdown} = \frac{T + T}{T} = 2\times

Why this happens: The warp executes the if branch with a mask (disabling threads that took else), then executes the else branch with the opposite mask. Both paths run sequentially even if only1 thread needs each path.

Worked example Divergence Example: Boundary Handling

__global__ void kernel(float* data, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {  // Boundary check
        data[idx] = data[idx] * 2.0f;
    }

Case 1: N = 128, blockDim.x = 128, 4 warps

Warps 0-3: All 32 threads pass idx < N → no divergence, full speed Case 2: N = 100, blockDim.x = 128, 4 warps
Warp 0 (threads 0-31): All pass → no divergence
Warp 1 (threads 32-63): All pass → no divergence
Warp 2 (threads 64-95): All pass → no divergence
Warp 3 (threads 96-127): Only threads 96-99 pass (4/32) → divergence, 28 threads disabled

Performance: Warp 3 executes the if body (4 active threads) but still occupies the execution pipeline for a full instruction. If the body takes 10 cycles, warp 3 takes 10 cycles for 4 threads instead of 32→ 87.5% waste

Why this step? The hardware cannot skip inactive threads' instruction slots; it executes the instruction with most threads masked off.

Common mistake Common Divergence Mistake: Thread-ID-Based Branching

Wrong approach:

if (threadIdx.x % 2 == 0) {
    // Even threads
} else {
    // Odd threads
}

Why it feels right: "I'm just splitting threads evenly, that's balanced!"

Why it's wrong: Everyarp has 16 even and 16 odd threads (threadIdx.x goes 0-31). Every warp diverges and runs both branches serially → 2× slowdown across the entire kernel.

The fix: Organize work so entire warps take the same path:

int warpId = threadIdx.x / 32;
if (warpId % 2 == 0) {
    // Even warps
} else {
    // Odd warps
}

Now warp 0 executes one branch, warp 1 executes the other—no intra-warp divergence.

Steel-man the mistake: The even/odd split seems fine because you're distributing work evenly. The hidden cost is that warps execute in SIMT—the unit of execution is the warp, not the thread. Thinking in threads instead of warps causes this error.

Optimizing for Warp Scheduling

1. Maximize Occupancy (Within Reason)

Constraints per SM:

Max threads: 2048 (varies by architecture)
Max blocks: 32 (varies)
Max warps: 64 (compute capability 7.0+)
Register file size: e.g., 65,536 registers
Shared memory: e.g., 64 KB (configurable split with L1 cache)

Occupancy formula:

\frac{\text{Threads per Block} \times \text{Blocks per SM}}{\text{Max Threads per SM}}, \frac{\text{Blocks per SM}}{\text{Max Blocks per SM}}, \frac{\text{Warps per SM}}{\text{Max Warps per SM}} \right)$$ Subject to:

\text{Registers per Block} \times \text{Blocks per SM} \leq \text{Total Registers}

\text{Shared Memory per Block} \times \text{Blocks per SM} \leq \text{Total Shared Memory}

2. Minimize Divergence

Use warp-wide operations: __ballot_sync(), __shfl_)
Structure conditionals so whole warps take the same path
Pad data structures to warp boundaries

3. Balance Instruction Mix

Instruction throughput varies by type:

FP32: typically 1 op/cycle/SM
FP64: 1/2 rate (or 1/32 on consumer GPUs)
Special functions (sin, cos): 1/4 rate

If all warps need FP64, schedulers compete for limited FP64 units → underutilization of FP32 units.

Recall Explain to a 12-Year-Old

Imagine your teacher wants the class to solve math problems. Instead of calling on each student individually (slow!), she groups you into teams of 32. Each team gets the same problem, and everyone works on their copy at the same time. That's a warp!

The teacher (warp scheduler) walks around and checks which teams are ready. If Team A is stuck waiting for a book from the library (memory access), she immediately helps Team B instead. By the time she cycles through several teams, Team A's book has arrived, so no one wastes time waiting.

But here's the tricky part: if Team A's problem says "students with even IDs do addition, odd IDs do subtraction," the team has to split up. First, the even-ID students solve while odd-ID students sit idle. Then they switch. The team takes twice as long because they can't work together anymore. That's divergence!

The lesson: keep your whole warp team doing the same thing at the same time for maximum speed.

Connections

SIMT-vs-SIMD - How GPU SIMT differs from CPU SIMD
Thread-Blocks-and-Grids - How warps fit into the larger hierarchy
GPU-Memory-Hierarchy - Why warp scheduling matters for hiding memory latency
Branch-Divergence-Patterns - Common divergence scenarios and solutions
Occupancy-vs-Performance - Why 100% occupancy isn't always optimal
Cooperative-Groups - Modern warp-synchronous programming primitives
Register-Pressure - How register usage limits occupancy and warp count

#flashcards/hardware

What is a warp in GPU architecture? :: A group of 32 consecutive threads that execute the same instruction simultaneously in SIMT fashion. Warps are the fundamental scheduling unit in NVIDIA GPUs.

How is warp ID calculated from thread indices?

Linear thread index divided by 32 (floor division). Linear index = threadIdx.z * blockDim.x * blockDim.y + threadIdx.y * blockDim.x + threadIdx.x, then warpId = linearIdx / 32.

Why should thread block dimensions be multiples of 32?

To avoid wasted warp slots. If blockDim is not a multiple of 32, the last warp will have inactive threads that still consume resources but do no work.

What is warp divergence?

When threads within the same warp take different execution paths (different branches). The GPU must serialize the divergent paths, executing each with some threads masked off, reducing performance.

What is the worst-case performance penalty of warp divergence?

2× slowdown when both branches are equally costly and equally populated. The warp executes both branches sequentially instead of one path.

How does warp scheduling hide memory latency?

When one warp stalls on memory (100s of cycles), the scheduler immediately switches to another ready warp with zero overhead. With enough resident warps, the SM stays busy while stalled warps wait.

What is GPU occupancy?

The ratio of active warps to maximum warps per SM. Formula: Active Warps / Max Warps per SM. Higher occupancy provides more warps for latency hiding.

What factors limit occupancy?

Registers per thread, shared memory per block, threads per block, max blocks per SM, and max warps per SM. The most restrictive limit determines actual occupancy.

Why does warp context switching have zero overhead?

Register files are partitioned per warp, program counters are per-warp, and instructions are pre-fetched. Switching warps only updates a pointer, no context save/restore needed.

What is the difference between warp schedulers and cores?

Warp schedulers select which warp executes next. Cores (CUDA cores) actually execute instructions. Modern SMs have multiple schedulers (e.g., 4) to issue instructions from multiple warps per cycle.

How many warps are in a 256-thread block?

256 / 32 = 8 warps. Each block's threads are divided into groups of 32 consecutive threads.

What makes a warp eligible for scheduling?

The warp must be not stalled (no memory wait, dependency wait, or synchronization barrier) AND resources must be available (execution units, memory ports).

Why is threadIdx.x % 2 branching bad for performance?

It causes divergence in every warp because each warp contains both even (16) and odd (16) threadIdx values. Every warp must execute both branches serially.

What is the formula for warps needed to hide latency L?

Warps needed ≥ (Number of schedulers × L) / Instructions between stalls. For 4 schedulers and 400-cycle latency with 10 instructions between memory ops: 4 × 400 / 10 = 160 warps worth of work.

What is a warp-synchronous operation?

Operations that act across all threads in a warp simultaneously, like __shfl_sync() for data exchange or __ballot_sync() for voting, without requiring explicit synchronization.

Concept Map

Hinglish (regional understanding)

Intuition Hinglish mein samjho

Hinglish (regional understanding)

Intuition Hinglish mein samjho

Dekho, GPU mein jab hazaaron threads chalte hain, toh unhe individually manage karna bahut mehnga pad jaata. Isliye NVIDIA ne ek smart trick lagayi hai - 32 threads ko ek group mein daal diya jise warp kehte hain. Ye saare 32 threads ek hi instruction ko ek saath, lockstep mein execute karte hain. Socho jaise class mein 32 students ko ek hi math problem samjhaani ho - alag-alag 32 baar explain karna bekaar hai, toh teacher ek baar samjhaata hai aur sab saath follow karte hain. Bas yahi warp ka concept hai: ek baar instruction fetch aur decode hoti hai, lekin 32 parallel executions milti hain. Yeh control overhead ko saare threads mein baant deta hai, jisse GPU massive parallelism handle kar paata hai.

Ab yahan sabse important baat samajhne wali ye hai ki GPU scheduler warp ke level pe kaam karta hai, individual thread pe nahi. Matlab agar aap block dimensions ko 32 ke multiple mein nahi rakhte, toh waste ho jaata hai. Jaise agar aapka blockDim.x = 17 hai, toh warp toh phir bhi 32 threads ka banega, par sirf 17 active honge aur baaki 15 slots khaali barbaad ho jaayenge. Isiliye humesha block size 32 ka multiple rakhna chahiye - jaise (128,1,1) perfect hai kyunki 128/32 = exactly 4 full warps bante hain bina kisi waste ke.

Yeh cheez practically kyun matter karti hai? Kyunki CPU aur GPU ka philosophy alag hai - CPU ek single thread ko fast chalane pe focus karta hai (latency), jabki GPU bahut saare threads ko ek saath throughput ke liye optimize karta hai. Jab aap GPU code likhoge (CUDA vagera), toh agar aapne warp-friendly dimensions choose kiye, toh aapki performance kaafi better hogi. Warp waste kam karna, aur scheduler ko efficiently warps dena - yahi GPU optimization ka core funda hai. Isliye ye concept seekhna zaroori hai agar aapko serious GPU programming karni hai.

Overview

Core Concepts

How Warps Map to Thread Blocks

Warp Scheduling Mechanics

The Scheduling Problem

Scheduling Policies

Warp Divergence

Optimizing for Warp Scheduling

1. Maximize Occupancy (Within Reason)

2. Minimize Divergence

3. Balance Instruction Mix

Connections

Concept Map

Hinglish (regional understanding)

Hinglish (regional understanding)

Go deeper — visual, from zero

Test yourself — GPU Architecture

Connections