Clock Derating and the Roofline That Moves

The roofline is usually drawn with two fixed ceilings. On a device that manages power, neither is fixed. The arithmetic ceiling scales with the core clock, the memory ceiling scales with the memory clock and with the share of interface capacity available, and both depend on how long the kernel has been running and on how much of the device is active. A kernel measured for a few milliseconds and the same kernel measured for a second can sit at different points relative to the ridge, and the difference is not noise.

Where the power goes

Dynamic power in CMOS scales as

$$ P \approx \alpha C V^2 f $$

where $\alpha$ is the activity factor, $C$ the switched capacitance, $V$ the supply voltage, and $f$ the clock. The quadratic voltage term is why clock scaling is not linear in power: raising the clock requires raising the voltage, so power rises faster than frequency. The practical consequence is that a device has a power budget, and staying within it requires either lowering the clock, lowering the voltage, or leaving parts of the chip unclocked.

The budget is set by the package’s thermal design and by the board’s power delivery, and it is enforced by a controller that observes temperature, current, and the requested clock. When the observed value approaches a limit, the controller reduces the clock. When it falls below, the controller may raise it again. The result is a feedback loop whose behaviour depends on the time constants of the thermal path, which are much slower than the electrical ones: a die heats over milliseconds to seconds, and the controller’s response lags accordingly.

There is also a boost behaviour. Many parts are rated at a base clock that they can sustain indefinitely within the power budget, and can run above it for a limited period while the thermal reserve allows. The duration of that period is a firmware parameter, and it is not exposed as a stable number. A kernel that runs for a few milliseconds may execute entirely within the boost window; a kernel that runs for a second may spend most of its time at the sustained clock.

What this does to the two ceilings

The arithmetic ceiling is proportional to the core clock, so a derate from a boost clock to a sustained clock lowers it by the clock ratio. That ratio is not small: a part that boosts well above its base clock can lose a large fraction of its arithmetic throughput once the boost window closes.

The memory ceiling is more complicated. DRAM clock may be managed separately, and the interface’s available bandwidth depends on the mix of read and write traffic and on contention with other traffic. On a device with a shared power budget, a kernel that saturates the arithmetic units leaves less budget for memory and vice versa, so the two ceilings are not independent. A compute-bound kernel can depress the memory ceiling by consuming the budget, and a bandwidth-bound kernel can run at a higher core clock because it is not heating the arithmetic units.

The ridge point, being the ratio of the two, therefore moves during a run. A kernel whose intensity is near the ridge can cross from compute bound to memory bound as the clock derates, and its optimal tuning changes with it. This is the sense in which the roofline “moves”: it is not a static chart but a family of charts parameterized by power state.

Interval	Core clock	Arithmetic ceiling	Likely regime
First few milliseconds	Boost	High	Compute bound, high ridge
Sustained load	Base or thermally limited	Lower	Memory bound, low ridge
Memory-saturated kernel	Near boost	High	Memory bound regardless
Idle-then-burst	Boost	High	Measurement not representative

Measurement artifacts that follow

The most common error in GPU performance measurement is to time a kernel in a loop and report the average, when the first iterations ran at a different clock from the rest. The average is then a number that describes neither state.

A second error is to measure a short kernel once, immediately after an idle period, and treat the result as the kernel’s cost. The die starts cool, the boost window is open, and the kernel runs at a clock it cannot sustain. The same kernel in a longer sequence runs slower. Neither measurement is wrong; they answer different questions, and a report that does not say which was measured is not interpretable.

A third error is to compare two kernels measured at different times. If kernel A was measured during a boost window and kernel B after a sustained load, the comparison reflects the thermal history rather than the code. The remedy is to measure both in the same thermal state, which in practice means either measuring both after a warm-up period long enough to reach steady state, or measuring both in short bursts from the same starting state, and stating which.

The counters help here. A profiler that reports the achieved clock alongside the kernel duration makes the state visible. A measurement without the clock cannot be compared to any ceiling, because the ceiling is a function of that clock.

// Sampling the clock around a launch makes derating visible in the log.
cudaEvent_t start, stop;
cudaEventCreate(&start); cudaEventCreate(&stop);
cudaEventRecord(start);
kernel<<<grid, block>>>(args);
cudaEventRecord(stop);
cudaEventSynchronize(stop);
float ms; cudaEventElapsedTime(&ms, start, stop);
// Compare ms across iterations; a rising trend indicates a falling clock.

Derating interacts with everything else

The interaction with occupancy is indirect but real. A lower core clock does not change the number of resident warps, but it changes the ratio of memory latency in cycles to memory latency in nanoseconds. If the memory clock is unchanged while the core clock falls, a fixed nanosecond latency corresponds to fewer core cycles, so the latency is easier to cover in cycle terms. The same physical memory latency becomes less costly relative to the arithmetic work, which shifts the balance toward memory-bound behaviour in time terms even as it looks better in cycle terms. Cycle-based reasoning and time-based reasoning diverge whenever the two clocks move independently, and the profiler’s cycle counts are only convertible to time with the clock that was in effect.

The interaction with the tail effect is more direct. A kernel whose blocks take a fixed number of cycles will see the per-block duration rise as the clock falls, so a wave-quantization penalty measured in milliseconds grows. The quantization ratio from wave quantization is unchanged, but its absolute cost is larger, and a kernel that was comfortably above the launch-overhead floor may no longer be.

The interaction with fusion is that a fused kernel has a longer duration and a higher power density than either of its parts, because it does both stages’ arithmetic. A fusion that halves traffic but pushes the kernel into a sustained derate can lose more than it gains, particularly if the unfused version was running in the boost window. The traffic argument in kernel fusion is computed in bytes and the derate is computed in watts, and the two have to be reconciled before the decision is made.

Designing an experiment that survives the derate

The measurement protocol that produces comparable numbers has four properties.

It warms the device to a steady state first, with a load that is representative in power density rather than merely in duration. A memory-bound warm-up does not heat the arithmetic units and does not close the boost window that a compute-bound kernel will encounter.

It reports the achieved clock for each measured interval, and discards intervals where the clock was still changing. A duration measured while the clock is ramping is not a steady-state duration.

It interleaves the variants being compared, rather than measuring all of one and then all of the other, so that any residual thermal drift affects both equally. Alternating A and B and comparing paired differences removes the drift to first order.

It separates the cold-start result from the sustained result and reports both when the difference matters. For a kernel invoked once in a program, the cold result is the relevant one. For a kernel invoked in a loop, the sustained result is.

The last point is the one most often skipped, and it is the one that determines whether a reported number describes the application or the benchmark.

What remains when the clock is accounted for

With the clock pinned to the sustained value and the ceilings recomputed for it, the roofline becomes usable again, and the diagnostic procedure is the same one that applies at a fixed clock. The kernel’s intensity places it relative to a ridge computed at the sustained clock; the achieved bandwidth is compared against the sustained memory ceiling; the stall histogram is read against the resulting expectation.

Two things remain that the clock correction does not fix. The first is that the power budget is shared, so a kernel that saturates arithmetic may be unable to simultaneously saturate memory even at a fixed clock. The ceiling is then a surface rather than two lines, and the achievable point is constrained by a power contour that the roofline does not draw. The second is that the derate is a function of the kernel itself: a more efficient kernel at the same clock may run at a higher clock because it generates less heat per unit of work. This is why two kernels with the same roofline position can have different measured ceilings, and why a comparison across kernels should always report the clock each achieved rather than assuming a common one.

The practical summary is that a performance number is a triple: the rate, the ceiling it is compared against, and the clock at which both were established. Any two of the three without the third produces a conclusion that will not reproduce.