gpu warp scheduling

**GPU Warp Scheduling** is the **hardware mechanism that selects which ready warp to execute each clock cycle on a Streaming Multiprocessor (SM) — where the warp scheduler's ability to find a ready warp among dozens of resident warps every cycle is what hides the 400+ cycle memory latency of global memory accesses, effectively converting memory latency into throughput by overlapping useful computation from one warp with memory stalls from another**. **Warp Scheduler Architecture** Each SM contains 2-4 warp schedulers (depending on GPU generation). Each scheduler: 1. Examines its pool of assigned warps (16-32 warps per scheduler). 2. Identifies ready warps — warps that have their next instruction ready to issue (no dependencies stalled). 3. Selects one ready warp and issues its next instruction. 4. The selected warp's instruction executes on the SM's functional units (INT, FP, SFU, Tensor Core, Load/Store). **Scheduling Policies** - **Greedy-Then-Oldest (GTO)**: Continue issuing from the same warp until it stalls, then switch to the oldest ready warp. Promotes temporal locality — the active warp benefits from L1 cache hits before switching. - **Round-Robin**: Cycle through warps in order, issuing one instruction per warp per turn. Fair but poor locality. - **Two-Level Scheduler (Volta+)**: Warps divided into pending (stalled) and active (ready) pools. Scheduler only considers the active pool, reducing selection latency. Stalled warps are moved to the pending pool and reactivated when their memory request completes. **Dual-Issue Capability** Some GPU generations can issue two independent instructions from the same warp in one cycle (dual-issue or instruction pairing): - Pair an integer instruction with a floating-point instruction. - Pair a load/store with a compute instruction. - Dual-issue increases IPC from 1.0 to up to 2.0 for instruction-parallel code. **Warp Stall Reasons** NVIDIA Nsight Compute reports why warps are stalled: - **Long Scoreboard**: Waiting for a long-latency operation (global memory load, texture fetch). Most common stall — indicates the kernel is memory-bound. - **Short Scoreboard**: Waiting for a short-latency operation (shared memory, L1 cache). Indicates shared memory bank conflicts or L1 misses. - **Not Selected**: Warp is ready but another warp was selected by the scheduler. Not a problem — indicates sufficient warp occupancy. - **Wait**: Barrier synchronization (__syncthreads()). Threads in the warp have reached the barrier but other warps in the block have not. - **Dispatch Stall**: Functional unit busy — too many warps requesting the same unit (e.g., SFU for transcendental math). **Occupancy and Scheduling Interaction** Warp scheduling effectiveness depends on having enough warps to hide latency: - **Memory-bound kernel**: Need enough warps so that while 75% are stalled on memory, 25% are executing. With ~30 cycle pipeline and ~400 cycle memory latency, need ~13 warps minimum per scheduler. - **Compute-bound kernel**: Fewer warps needed — functional unit throughput is the bottleneck, not memory latency. Even 2-4 warps per scheduler may suffice. GPU Warp Scheduling is **the zero-cost context switching mechanism that converts GPU memory latency into throughput** — the hardware scheduler that makes thousands of threads appear to execute simultaneously by rapidly switching between warps, hiding memory access delays behind useful computation from other warps.