Warp Stall Reasons and Cycle Attribution

A kernel that occupies the device fully and still runs at a third of its ceiling is not short of warps. It is short of issuing warps. The scheduler picks a warp to issue from each cycle, and every cycle in which it finds nothing issuable is a wasted issue slot. Attributing those wasted slots to reasons is the most direct diagnostic available, and also the most misread one.

The scheduler’s view of a cycle

A warp scheduler holds a small set of resident warps, selects one per cycle, and issues one instruction from it. A warp is not eligible if it is waiting for something. The hardware tracks the reason for each ineligible warp, and the profiler aggregates the reasons into a histogram, usually expressed as cycles per issued instruction or as a fraction of total cycles.

The categories are mutually exclusive per cycle and jointly exhaustive over the kernel, which is what makes the histogram add up to 100%. That property is also the trap: a kernel can have a dominant category for reasons that are consequences of a different problem, and the largest bar is not always the root cause.

The standard categories, in roughly the order they are usually encountered:

Category	Meaning at the scheduler	Typical cause
Long scoreboard	Waiting on a global memory dependency	Cache miss, insufficient memory-level parallelism
Short scoreboard	Waiting on a shared memory or local dependency	Bank conflict, spill reload
Wait	Waiting on a fixed-latency dependency	Arithmetic latency not covered by other warps
Barrier	Waiting at a block-wide synchronization point	Uneven work between warps before the barrier
Not selected	Eligible but not picked this cycle	More eligible warps than issue slots
Dispatch	Waiting for an execution unit that is busy	Pipeline throughput limit, often special-function units
MIO throttle	Queue to the memory-input/output pipeline is full	Too many shared-memory or special instructions in flight
No instruction	Waiting for an instruction fetch	Instruction cache miss, very large unrolled loops

The first three are latency. The next four are throughput and scheduling. The distinction matters because latency stalls are fixed by adding concurrency or reducing dependencies, while throughput stalls are fixed by reducing the number of instructions or changing their mix. Applying one remedy to the other problem produces no improvement and often a regression.

Why “long scoreboard” is not the same as “memory bound”

The long scoreboard category counts cycles a warp spends waiting for a global load whose result is needed. A kernel with a high fraction here is waiting on memory. That is not the same as being limited by memory bandwidth. A kernel can spend most cycles waiting on memory while using a small fraction of the interface, because it has too few requests outstanding. The interface is idle and the warps are stalled, and both facts are true simultaneously.

The discriminator is the achieved bandwidth from achieved versus peak throughput. If long scoreboard is high and bandwidth is near the practical ceiling, the kernel is genuinely bandwidth limited and the stall is the correct, unavoidable consequence. If long scoreboard is high and bandwidth is well below the ceiling, the kernel is latency limited and the fix is concurrency or independence, which is the subject of memory-level parallelism.

Reporting the stall reason without the bandwidth figure is the single most common way to reach a wrong conclusion from a profiler.

Separating latency from throughput numerically

Two quantities make the distinction concrete. The issue rate is the fraction of cycles on which the scheduler issues an instruction. The eligible rate is the fraction of cycles on which at least one warp was eligible. If eligibility is high but issue is low, the scheduler or an execution pipe is the limit. If eligibility is low, the warps are stalled, and the stall histogram says why.

For a scheduler that can issue one warp-instruction per cycle, the issue rate over an interval is the number of warp-instructions it issued divided by the number of cycles in the interval, and its ceiling is one. The practical reading is therefore simple: IPC below one per scheduler means the scheduler is idle, and the histogram explains the idle cycles. IPC at one per scheduler means the issue path is saturated, and no amount of latency hiding will help; the kernel needs fewer instructions or a different instruction mix.

This is the boundary between the latency regime and the issue regime, and it corresponds exactly to the vertical part of the roofline. A kernel whose IPC is at the maximum is compute bound in the model’s sense, and its remaining inefficiency is in the instruction stream rather than in the memory system.

Barriers as an amplifier

A barrier stall is a warp waiting at a synchronization point for the slowest warp in its block. The wait itself is not the problem; the imbalance that produces it is. A block where all warps do identical work reaches the barrier together and the barrier costs almost nothing. A block where one warp has twice the work of the others has the other warps waiting for half the block’s duration.

The histogram reports this as barrier cycles, and the natural but wrong conclusion is that barriers are expensive and should be removed. The right conclusion is that the work distribution before the barrier is uneven, and the fix is to change the distribution, or to eliminate the barrier by restructuring so warps do not need to reconverge.

There is a second source of barrier cost that is genuinely the barrier’s fault: a barrier imposes a full block-wide dependency, so all the latency hiding that was happening across warps collapses at that point. A kernel with frequent barriers has effectively serialized its warps at each one, and the average concurrency over the kernel is lower than the resident warp count suggests. The mitigation is fewer and larger barriers, with more work between them.

// Warp-level reduction defers the block barrier to the last step,
// so only one warp per block waits at it.
__device__ float warp_reduce(float val) {
    for (int offset = 16; offset > 0; offset >>= 1)
        val += __shfl_down_sync(0xffffffffu, val, offset);
    return val;                       // lane 0 holds the warp total
}

__global__ void block_sum(const float *in, float *out, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    float v = (i < n) ? in[i] : 0.0f;
    v = warp_reduce(v);
    __shared__ float partial[32];
    if ((threadIdx.x & 31) == 0) partial[threadIdx.x >> 5] = v;
    __syncthreads();                  // eight warps arrive here, not 256 threads
    if (threadIdx.x < 32) {
        v = (threadIdx.x < (blockDim.x >> 5)) ? partial[threadIdx.x] : 0.0f;
        v = warp_reduce(v);
        if (threadIdx.x == 0) out[blockIdx.x] = v;
    }
}

The “not selected” bar is not a defect

A high fraction of not-selected cycles means warps were ready to issue and the scheduler chose another. That is the healthy state: it is the scheduler doing its job, and it indicates the kernel has more eligible warps than issue slots, which is exactly what latency hiding is supposed to produce.

A kernel with a large not-selected fraction and a good achieved rate should be left alone. A kernel with a large not-selected fraction and a poor rate is issuing instructions that do not retire useful work, which points at the instruction mix rather than at concurrency. The distinction is invisible in the histogram alone and requires the instruction counts alongside it.

The corollary is that occupancy tuning aimed at reducing not-selected cycles is aimed at the wrong target. Reducing residency until warps are never not-selected would simply move the kernel to a state where it has no latency cover.

Sampling bias and what the histogram hides

Profiler stall attribution is normally collected by sampling: the hardware records the state of a set of warps at intervals, and the histogram is the aggregate of those samples. Two biases follow.

Duration bias. Long stalls are sampled more often than short ones, in proportion to their duration. That is correct for attributing cycles and wrong for attributing events. A kernel that issues a million loads each stalling briefly and a thousand loads each stalling for a long time will show the long stalls dominating the histogram, even if the short ones are the ones that matter for the total instruction count.

Warp-position bias. Sampling picks warps, and warps that are stalled are more likely to be sampled as stalled than warps that are issuing, simply because issuing warps spend less time in a sampleable stalled state. The resulting histogram is a time-weighted average over warps, which is the quantity wanted, but it is not a count of stall events.

The practical consequence is that the histogram should be read as “where the cycles went”, never as “what happened most often”. A kernel with 60% long scoreboard spent 60% of its warp-cycles waiting on global memory, which may correspond to a small number of very slow accesses.

There is also a category the histogram omits: cycles when no warp was resident at all, during the ramp and drain of the grid. Those cycles are not attributed to any warp because there was no warp to sample. A kernel whose measured duration exceeds the sum of its attributed warp cycles has a launch or tail problem, and wave quantization is where to look.

Working from the histogram to a fix

The sequence that produces a defensible conclusion has four steps, in order.

First, establish whether the kernel is near its roofline ceiling at all. A kernel at 90% of its practical bandwidth ceiling has a stall histogram that is a description of the ceiling, not a list of defects.

Second, check the issue rate. If schedulers are issuing on most cycles, the kernel is issue limited and the histogram’s latency categories are being successfully hidden; the remaining work is in the instruction stream.

Third, for kernels that are not issue limited, compare the dominant latency category against the achieved bandwidth and occupancy. Long scoreboard plus low bandwidth means insufficient concurrency; long scoreboard plus high bandwidth means the kernel is at its ceiling; short scoreboard means shared-memory or spill traffic, and the relevant counters are the bank-conflict and local-memory ones; barrier means imbalance.

Fourth, confirm by changing one thing. A stall histogram is a hypothesis generator. The confirmation is that altering the hypothesised cause moves the measured rate in the predicted direction, and that the histogram changes accordingly. A kernel whose long-scoreboard fraction falls while its rate stays flat has not been fixed; the concurrency was never the binding constraint.