How Scoreboards, Dependencies, and Pipeline Availability Stall Warps

Hello! Last lesson showed that a warp can spend issue opportunities on different branch paths or masked instructions. Even when control flow is uniform, though, a resident warp may be unable to issue its next instruction. Today you will distinguish three causes: an unfinished operand dependency, a busy execution pipeline or instruction queue, and a scheduler that has no ready warp to choose. This distinction turns a stall report into a useful performance hypothesis rather than a list of alarming counters.

A resident warp is not necessarily a ready warp

A warp remains active while it occupies resources on an SM. To issue an instruction, it must also be eligible: its next instruction must be available, its required inputs must be ready, and the relevant execution resource must be able to accept it. Of the eligible warps assigned to a scheduler, one may be selected to issue; another may be ready but not selected. That second warp is not blocked by a dependency.

The SM architecture schematic shows why the scheduler is the right level at which to reason about this. Warps are assigned to an SM sub-partition with a warp scheduler and register file; different instructions use different execution paths, while some resources are shared more widely across the SM. Its depiction of four schedulers is a model of the architecture shown, not a specification to assume for every GPU generation.

For a precise definition of the states, read the short architectural description in the profiler documentation.

At a given instant for one scheduler, active warps are either eligible or stalled. Thus, with the same scope and sampling basis,

That identity says nothing by itself about whether an issue slot was wasted. If several warps are eligible and one issues, the others were not selected; the scheduler may be working at its available issue rate. Conversely, eight resident warps offer little latency hiding if all eight repeatedly wait for operands.

A scoreboard protects dependencies, not an entire warp after every load

Consider the instruction-level shape of a simple kernel operation:

float a = x[i];
float b = y[i];
float product = a * b;
out[i] = product;

Think of the two loads as producing register values and the multiplication as consuming them. A load can be issued before its result has returned. The multiplication cannot issue until both inputs are ready. A scoreboard tracks the outstanding dependency so that the scheduler does not send an instruction to an execution unit with unavailable operands.

Crucially, an outstanding load does not mean the warp is forbidden from issuing anything else. If the compiler has placed an independent instruction next—perhaps arithmetic using registers already available—the warp may issue it while the load remains in flight, provided the required pipeline can accept it. If the next instruction needs the load result, that warp must wait; the scheduler can instead choose another eligible warp. This is the operational reason both instruction-level parallelism within a warp and ready warps alongside it can hide latency.

The dependency has a producer and a consumer. In the example, the load produces a; the multiplication consumes it. A sampled stall at the multiplication does not establish that multiplication is slow. It may show that the multiplication cannot begin because the load has not delivered its value.

Profiler stall names add useful detail, but they name the kind of wait, not a complete diagnosis:

• Long scoreboard denotes a dependency on an L1TEX-path operation, such as a global, local, surface, or texture access. “Long” does not prove that a request went to DRAM; determine where the access was served separately.
• Short scoreboard denotes a dependency on a memory-input/output (MIO) operation rather than the L1TEX path. Shared-memory loads are a common source, but special math and some dynamic branches can also be involved. It is not simply a label for “a short wait.”
• Wait covers a fixed-latency execution dependency, such as one arithmetic instruction consuming the result of a preceding arithmetic instruction.

These are distinct from a block barrier. A warp waiting for sibling warps at __syncthreads() is blocked by synchronization, even if its own arithmetic inputs are ready. As with divergence last lesson, the useful question is what work and dependencies occurred before the waiting point.

Ready operands do not guarantee an available pipeline

Now imagine that the multiplication’s inputs have arrived, but many warps want the same arithmetic pipeline. The warp can still fail to issue because that pipeline cannot accept another instruction at that moment. This is a resource-availability problem, not a missing-value problem.

The distinction also applies to instruction queues. A warp may be unable to submit another local/global memory instruction because the relevant queue is full. That differs from a warp whose consumer waits for the result of a memory instruction already submitted. In profiler terminology, examples include math_pipe_throttle for an oversubscribed math pipeline, lg_throttle for a full local/global instruction queue, and mio_throttle for a full MIO queue.

Read the stall definitions alongside the scheduler-level diagnostic guidance. The goal is to associate each state with a mechanism; there is no need to memorize every category on the page.

There is a subtle consequence. A pipeline can be fully occupied while many warps report a throttle. In that case, eliminating some throttle samples without reducing the work demanded of the limiting pipeline may not shorten the kernel. By contrast, if issue activity is low because most resident warps await load results, finding independent work or more useful latency-hiding concurrency may raise throughput. Neither “more warps” nor “fewer stalls” is a universal objective; keeping the limiting resource productively supplied is.

This also explains why occupancy alone is insufficient. More resident warps provide additional candidates for the scheduler, but they cannot increase a pipeline’s maximum throughput. Later kernel-optimization lessons will examine how changing instruction mix, independent work, and resource use affects this balance.

Read a stall report as evidence, then locate the cause

The following brief profiler walkthrough shows a long-scoreboard finding and how a source view helps localize it.

A sampled timeline can display several kinds of waiting at once. In the Warp States (Aggregated) image, bands include long scoreboard, math-pipe throttle, MIO throttle, and wait states. Their presence shows what was observed during sampling; the picture alone does not establish which change would reduce kernel time.

A defensible interpretation proceeds in this order:

1. Establish the limiting situation. Inspect issue activity, active and eligible warps per scheduler, and whether a compute or memory pipeline is already near its throughput limit. A low eligible count together with low issue activity suggests insufficient issuable work. Many not_selected warps instead indicate ready work was available.
2. Classify why warps could not issue. A scoreboard points to a producer–consumer dependency; a throttle points to pipeline or queue pressure; a barrier points to other threads or warps that have not met a synchronization condition.
3. Find the producer or pressured resource. For a long-scoreboard sample on a multiply, locate the load supplying its operand. For math-pipe throttle, examine the instruction mix and utilization of that pipeline. Source lines help navigate, but generated instructions and memory behavior determine what actually ran.

Suppose a kernel has many active warps, few eligible warps, low issue activity, and dominant long-scoreboard samples on arithmetic consumers. That combination supports a load-to-use latency hypothesis. The next evidence concerns the producing loads: their access pattern, cache behavior, and how much independent work can run before their results are needed. It does not yet establish a DRAM-bandwidth bottleneck.

Suppose instead issue activity is high and the relevant math pipeline is saturated, with math-pipe-throttle samples elsewhere. The scheduler is already feeding a busy resource. Adding warps to “hide the stalls” is unlikely to increase that pipeline’s capacity. These two profiles can both contain large stall counts yet call for different explanations.

Takeaways

An active warp occupies SM resources; an eligible warp can issue its next instruction; a selected warp actually issues. Scoreboards prevent consumers from using unfinished results, while pipeline and queue throttles prevent otherwise ready work from entering a busy resource. A sampled stall is therefore a clue about why this warp did not issue, not a standalone measure of wasted kernel time.

When interpreting a profile, pair stall reasons with scheduler issue activity and resource throughput, then trace a dependent consumer back to its producer. Next lesson, you will estimate a kernel’s data movement from its access pattern—a useful way to test whether the memory behavior implicated by a stall hypothesis is plausible.