Wave Quantization and the Tail Effect

A grid of 5,000 blocks on a device that holds 2,048 blocks at once does not run 2.44 times as fast as a grid of 2,048. It runs in three waves: two full ones and a third holding 904 blocks, a little under half the machine. The last wave costs almost as much wall-clock time as a full one, and the arithmetic of that rounding is one of the largest avoidable inefficiencies in a kernel that otherwise has no problems at all.

Counting waves

The number of blocks resident simultaneously is

$$ \text{blocks}_{resident} = \text{SMs} \times \text{blocks per SM} $$

where blocks per SM is the minimum imposed by the warp limit, the block limit, the register limit, and the shared-memory limit, as computed in occupancy as a latency-hiding budget. The number of waves is then

$$ W = \left\lceil \frac{\text{blocks}{total}}{\text{blocks}{resident}} \right\rceil $$

and the efficiency of the launch, in the crude sense of useful occupancy over the whole run, is the ratio of the block count to the product of waves and residency. A grid of 5,000 blocks with 2,048 resident runs at $5000 / (3 \times 2048) \approx 0.81$. A grid of 4,096 blocks runs at exactly 1.0 with the same residency.

The crude figure overstates the loss, because the tail wave is not idle for its whole duration: blocks in it run at full speed, they just do not fill the machine, and the memory system may still be saturated by a quarter of the machine if the kernel is bandwidth bound. A bandwidth-bound kernel with 25% of the blocks resident can still saturate DRAM, in which case the tail wave costs almost nothing. A compute-bound kernel with 25% of the blocks resident runs at roughly a quarter speed. Which case applies is determined by the kernel’s position on the roofline, and this is the reason tail analysis cannot be done from the grid shape alone.

Why the tail is disproportionate for short kernels

The cost of the tail is fixed in absolute terms by the duration of one block. If a block takes 2 milliseconds, the third wave of a 5,000-block launch costs about 2 milliseconds regardless of how empty it is. For a kernel whose full waves take 2 milliseconds each, that tail is a third of a 6-millisecond run.

The same fraction applied to a kernel whose blocks take 20 microseconds is 20 microseconds on a 60-microsecond run, which is the same third of the total but now comparable to launch overhead. At that scale the quantization and the launch cost become the same order of magnitude, and the useful comparison is against the fixed per-launch cost rather than against the steady-state rate.

There is a second-order effect that makes small blocks worse. The device does not start and finish all blocks in lockstep; block scheduling is dynamic, and a block starts when a slot frees. With many small blocks the ramp and drain are smoother, and the effective tail is shorter than the ceiling calculation suggests. With a small number of long blocks the quantization is sharp and the tail is exactly as bad as the arithmetic says. This is the trade hidden inside block-size choice: small blocks smooth the tail but add scheduling events and reduce per-block shared-memory reuse.

Grid-stride loops as a fix and as a hazard

The standard remedy is to launch exactly enough blocks to fill the device, or a small multiple of it, and have each block process several elements in a grid-stride loop. That converts a block-count problem into a per-thread loop-count problem, and the quantization moves from the block level to the iteration level, where it is much finer.

// One full wave, with the block count taken from the occupancy the
// kernel actually achieves rather than from a rule of thumb.
int blocks_per_sm;
cudaOccupancyMaxActiveBlocksPerMultiprocessor(&blocks_per_sm, kernel, 256, 0);
int device;
cudaGetDevice(&device);
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, device);
int blocks = prop.multiProcessorCount * blocks_per_sm;
kernel<<<blocks, 256>>>(n, x, y);

The loop count per thread is then $\lceil n / (\text{blocks} \times \text{threads}) \rceil$, and if $n$ is not divisible by the total thread count, some threads execute one more iteration than others. That imbalance is at most one iteration, so for a kernel with a long inner loop it is negligible, and for a kernel with a short inner loop it is bounded by the cost of a single element.

The hazard is that a grid-stride loop with a fixed block count can under-occupy the machine if the per-SM residency estimate is wrong. If the chosen block count assumes 8 blocks per SM and the kernel actually supports 6, the launch fills only three quarters of the machine and the kernel is slower than a naive over-subscribed launch would have been. The block count must come from the occupancy the kernel actually achieves, which is why the runtime query matters more than a rule of thumb.

A second hazard is a kernel that is not bandwidth bound. A grid-stride loop with exactly one wave removes the ability of the scheduler to overlap the tail of one block with the start of another, and for a latency-bound kernel the over-subscription was providing useful extra parallelism. The remedy is to launch a small multiple of a wave, two to four, so the scheduler has slack to fill gaps. The optimum is empirical and shifts with the kernel.

Decomposing work when the grid cannot change

Some kernels are launched with a grid fixed by the problem, for instance one block per output tile in a matrix multiply or one block per image in a batch. There the wave count is not a free variable, and the tail has to be attacked differently.

Persistent blocks. A persistent kernel launches exactly one wave and has each block pull work items from a global counter until the queue is empty. The tail becomes a work-stealing problem: blocks that finish early take the remaining items. The cost is an atomic counter and the loss of any static mapping between block index and data, which for kernels relying on locality can be a real loss. Persistent blocks also keep the tail’s duration bounded by one work item rather than one block’s full workload.

Work redistribution. If the total work can be split unevenly, giving the first blocks more items and the last fewer equalizes finish times. The schedule is computed from the block count and the item cost, and it only works when item costs are known in advance and uniform.

Smaller blocks for the tail. Launching a smaller block size raises the block count and reduces the per-block duration, which shrinks the absolute tail. The cost is more scheduling and, often, lower occupancy per block. This is a coarse lever but it is available without changing the algorithm.

Uneven item costs. When items genuinely differ in cost, static distribution leaves the slowest item determining the tail regardless of wave count. That is a load-balancing problem rather than a quantization one, and the diagnostic signature is a tail wave that is not merely under-filled but is dominated by a single long-running block.

Measuring the wave structure

The wave structure is visible in a timeline view of block scheduling: blocks are grouped into bands separated by the moments when a wave’s blocks retire together. For a well-quantized launch the bands are equal and the last is full. For a badly quantized one the last band is short.

Two counters help. The average number of active blocks over the kernel’s lifetime, divided by the maximum resident count, gives a utilization figure that includes the tail. Comparing that against the theoretical residency from the occupancy calculation separates the tail from a residency shortfall: if achieved active blocks average below theoretical during the middle of the kernel, occupancy is the problem; if it is at theoretical in the middle and drops at the end, the tail is the problem.

A third measurement is more direct: run the same total work with two different grid shapes that have the same residency but different wave counts, and compare. If the difference matches the wave arithmetic, the tail is confirmed. If it does not, the kernel is bandwidth bound and the tail was never costing what the arithmetic suggested.

When quantization is not the bottleneck

Three conditions make wave analysis the wrong place to look.

The kernel is bandwidth bound and even a partial wave saturates memory. Then the tail costs the drain time of the last few blocks, not a full wave.

The kernel’s blocks have highly variable duration. Then the wave boundaries blur and the last block to finish determines the runtime, which is a load-balancing question rather than a rounding one.

The kernel is so short that launch latency dominates. Then the fixed cost per launch, which is on the order of microseconds, is larger than any quantization loss, and the fix is to reduce the number of launches or to batch work, not to reshape the grid. The related question of whether two kernels should be one is treated in fusion and launch boundaries.

The general point is that quantization is a rounding effect, and rounding only matters when the thing being rounded is a significant fraction of the total. Establishing that fraction first, by measuring the kernel’s duration and the block’s duration, decides whether the rest of the analysis is worth doing.