Warp Scheduling

The warp scheduler operates at the hardware level with minimal overhead. When a warp stalls—typically due to a long-latency memory access or a synchronization barrier—the scheduler does not save its architectural state to memory. Instead, it instantly switches to another warp that is ready to execute, as all warps' contexts (registers, program counters) are already resident in the SM's on-chip resources. This allows latency to be hidden without the costly context-switch penalties seen in CPU operating systems.

This is the core purpose of warp scheduling. By maintaining a large pool of active warps per SM, the scheduler can interleave their execution.

• When Warp A issues a load from global memory (hundreds of clock cycles latency), the scheduler immediately issues instructions from Warp B, Warp C, and others.
• The goal is to always have at least one arithmetic-logic unit (ALU)-ready warp to execute while others wait for data. This turns memory latency into a throughput problem, maximizing the utilization of expensive execution hardware.

Warps execute according to the Single Instruction, Multiple Threads (SIMT) model. The scheduler issues one instruction for all 32 threads (typical warp size) in a lockstep manner. A critical challenge is control flow divergence (e.g., an if/else statement). When threads within a warp take different paths, the scheduler must serialize execution: it first executes the threads in the 'if' branch (masking off the 'else' threads), then executes the 'else' branch. This divergence reduces effective parallelism within the warp and is a key performance consideration for kernel developers.

Different GPU architectures implement specific scheduling policies to select the next warp to issue.

• Greedy-Then-Oldest (NVIDIA): Prioritizes warps that are ready to issue their next instruction. Among ready warps, it selects the one that has been waiting the longest. This policy maximizes instruction throughput.
• Round-Robin: A simpler, fair-share policy that cycles through all active warps in order, issuing an instruction if the warp is ready. This can prevent starvation but may be less efficient at hiding latency if many warps are stalled. The policy is fixed in hardware and directly impacts kernel performance.

Modern GPU SMs often contain multiple, independent warp schedulers (e.g., four schedulers per SM). Each scheduler manages a subset of the warps assigned to the SM and can issue instructions to a dedicated set of execution units (e.g., INT32, FP32, tensor cores). This allows a single SM to issue multiple independent instructions per clock cycle, significantly increasing instruction-level parallelism (ILP) and overall throughput. The effectiveness depends on the kernel having sufficient independent warps and instruction mix to keep all schedulers busy.

Occupancy is the ratio of active warps on an SM to the maximum number supported. It is a key resource constraint for the warp scheduler. High occupancy (many active warps) provides more candidates for the scheduler to interleave, improving latency hiding. However, occupancy is limited by:

• Register file size per thread.
• Shared memory allocation per thread block.
• Thread block and warp limits per SM. Optimizing a kernel often involves trading maximum occupancy for other optimizations (e.g., increased register usage for loop unrolling) to find the peak performance point.

The fundamental execution model that warp scheduling serves. In SIMT, a single instruction is issued to a warp (or wavefront) of threads, each executing it on its own data. The scheduler's job is to manage these warps, hiding latency by keeping Stream Multiprocessors (SMs) busy despite individual threads stalling on memory accesses or divergent control flow.

The programmable computing core where warp scheduling physically occurs. Each SM contains:

• Execution units (CUDA cores, Tensor Cores)
• A fixed pool of registers and shared memory
• Multiple warp schedulers The scheduler on an SM selects which resident warp is ready to execute and dispatches its instruction to the execution units. SM resources directly constrain occupancy.

The logical group of threads scheduled together on a single SM. A block is divided into warps (typically 32 threads). The warp scheduler manages all warps from all active blocks on the SM. Key constraints:

• Block size determines the number of warps.
• Block resources (shared memory, registers) affect how many blocks/warps can be co-scheduled.
• Threads within a block can synchronize via barriers, causing warp stalls the scheduler must work around.

A critical performance metric measuring how effectively the warp scheduler utilizes the SM's hardware. It's the ratio of active warps on an SM to the maximum possible active warps. High occupancy provides more warps for the scheduler to choose from, improving its ability to hide instruction and memory latency. Occupancy is limited by:

• Register usage per thread
• Shared memory usage per block
• Thread block size and grid configuration

Defines the rules for when memory writes from one thread become visible to others. GPU memory models (like NVIDIA's weak consistency) allow for high performance but require explicit synchronization (e.g., __syncthreads()). The warp scheduler must respect these ordering constraints. Improper synchronization can lead to data races and non-deterministic results, as warps execute in a non-fixed order.

A coordination primitive (e.g., __syncthreads()) that forces all threads in a block to reach the barrier before any proceed. This directly impacts the warp scheduler:

• Warps reaching the barrier stall, freeing scheduler slots.
• The scheduler switches to other, non-stalled warps.
• Once the last warp hits the barrier, all warps in the block become eligible for scheduling again. Excessive barriers can limit the scheduler's flexibility to hide latency.