The Roofline Bound and Arithmetic Intensity

A kernel that performs one multiply and one add for every four bytes it fetches cannot run faster than the memory system allows, no matter how many arithmetic units the device carries. That statement is obvious in isolation and surprisingly easy to forget when a kernel is tuned. The roofline model turns it into a quantitative bound by pairing the two resources a kernel competes for, arithmetic throughput and data movement, and reducing the question of “why is this slow” to the position of one point on a chart.

Two ceilings and the ratio between them

Let $P$ be attainable floating-point throughput in FLOP/s, $P_{peak}$ the device’s peak arithmetic rate, $B$ the sustained DRAM bandwidth in bytes/s, and $I$ the arithmetic intensity, defined as floating-point operations divided by bytes transferred to and from DRAM. The bound is

$$ P \le \min\left(P_{peak},; B \cdot I\right) $$

Plotting $P$ against $I$ on log-log axes gives two lines. The bandwidth ceiling $B \cdot I$ is a straight line of slope 1: doubling the work per byte doubles the attainable throughput, as long as the arithmetic units can keep up. The compute ceiling $P_{peak}$ is horizontal. They intersect at

$$ I_{ridge} = \frac{P_{peak}}{B} $$

Kernels with intensity below the ridge are memory bound, and their throughput is set by the slope line. Kernels above the ridge are compute bound, and their throughput is capped by the flat line. The ridge point is the only device-specific quantity in the model, and it is a ratio, so it moves whenever either the arithmetic rate or the memory rate changes between hardware generations.

Counting the bytes that actually cross the interface

The intensity denominator is the source of nearly every mistake made with the model. It counts DRAM traffic, not traffic served by caches. A shared-memory tiled matrix multiply reads each input tile dozens of times, but from shared memory rather than from DRAM, and the DRAM traffic is far below the naive count. Computing intensity from the number of array accesses in the source overstates the denominator by whatever factor the caches absorb, and understates the intensity accordingly.

Three kernels make the arithmetic concrete.

A saxpy, $y_i = a x_i + y_i$, reads 4 bytes of $x$, reads 4 bytes of $y$, and writes 4 bytes of $y$ per element. That is 12 bytes and 2 FLOPs, so $I \approx 0.17$ FLOP/byte. No amount of tuning changes the number; the loop body performs too little work per element for any cache to help.

A dot product of two length-$n$ vectors performs $2n$ FLOPs and reads $8n$ bytes, giving $I = 0.25$ FLOP/byte. It writes a single scalar, so the write traffic is negligible for large $n$, and the intensity is still far below any ridge point on any current device.

A dense $N \times N$ matrix multiply performs $2N^3$ FLOPs. Its DRAM traffic depends entirely on blocking. Without blocking, the traffic is at best the compulsory $3N^2$ elements, and usually much worse because each output element re-reads a full row and column. With square tiles of edge $T$, the number of tiles is $(N/T)^3$, each tile multiplication touches three $T \times T$ blocks, and the total element traffic is $3N^3/T$. In bytes that is $12N^3/T$, so

$$ I = \frac{2N^3}{12N^3/T} = \frac{T}{6} $$

A 32-wide tile yields roughly 5 FLOP/byte; a 64-wide tile roughly 11. The tile edge is therefore a direct control on intensity, and the largest tile that still fits in shared memory with acceptable occupancy is the one that matters.

// The intensity of this kernel is fixed by the data layout, not by the loop body.
__global__ void saxpy(int n, float a, const float *x, float *y) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        y[i] = a * x[i] + y[i];   // 8 bytes read, 4 bytes written, 2 FLOPs
    }
}

Because the ratio is $2/12$, the kernel is memory bound on any device whose ridge point exceeds about 1 FLOP/byte, which covers every general-purpose device of the last two decades; on a part with a ridge near 10, it sits more than an order of magnitude below it. Recognizing that from the source, before any profiling, is the practical value of the model.

Where the ridge sits, and why it keeps moving

The ridge point is a property of the part, not of the kernel. Take a device with roughly 1 TB/s of DRAM bandwidth and roughly 10 TFLOP/s of single-precision throughput. The ridge is near 10 FLOP/byte, so anything doing fewer than about ten operations per byte is bandwidth limited. Pair the same memory system with four times the single-precision throughput, as successive generations have done, and the ridge moves to about 40 FLOP/byte. Kernels that were comfortably compute bound become memory bound without a line of code changing.

Reduced precision shifts the ridge further and in the same direction. A matrix-multiply path that executes at an order of magnitude more operations per second than the single-precision path pushes the ridge an order of magnitude to the right. The consequence is counterintuitive: the faster the arithmetic, the more of the workload is limited by memory, and the more valuable bandwidth optimizations become relative to instruction-level tuning.

Region	Binding resource	Lever that helps
Below the ridge	DRAM bandwidth	Raise intensity by blocking or fusion; reduce bytes by narrowing element width or improving layout
At the ridge	Both, simultaneously	Either lever gives diminishing returns; the other ceiling becomes binding immediately
Above the ridge	Arithmetic issue rate	Instruction mix, dual-issue opportunities, avoiding redundant work

The table is the reason the model is worth drawing at all. Below the ridge, optimizing the arithmetic is wasted effort; the memory system will not deliver the operands any faster. Above the ridge, optimizing memory access patterns changes nothing.

Nested roofs for each level of the hierarchy

A single roof treats DRAM as the only memory. A more informative version adds one ceiling per level: L2, shared memory, and registers each have their own bandwidth and their own capacity. The kernel’s intensity is then computed once per level, using the traffic that level sees, and the applicable bound is the lowest ceiling above the point.

This matters because a kernel can be comfortably under the DRAM roof and over the shared-memory roof at the same time. The tiled matrix multiply above is the standard case: its DRAM intensity rises with tile size, but its shared-memory traffic scales with the number of times each tile is re-read by the threads that need it. Once the tile is large enough that DRAM is no longer binding, the shared-memory bank structure becomes the limit, and the fix is a layout change or padding rather than a larger tile. Treating the two as one roof hides that transition entirely.

A point far below both roofs

The roofline is an upper bound on steady-state throughput. It assumes full occupancy of the arithmetic pipeline, enough concurrent memory requests to saturate the interface, and no launch or synchronization overhead. When a measured point sits well below both ceilings, the model has nothing to say about why, and the useful question changes from “which resource is saturated” to “which resource is idle”.

Three families of cause dominate. Insufficient concurrency leaves the memory system with too few outstanding requests to cover its latency, which is the subject of occupancy as a latency-hiding budget. Warps stalled on dependencies or on barrier waits show up as idle cycles rather than as bandwidth, and are attributed through warp stall reasons. And the gap between the ceiling used in the model and what the part actually sustains at its operating clock is large enough that achieved versus peak throughput has to be measured rather than assumed.

The practical reading order is therefore: place the kernel on the roofline to learn which resource it should be competing for, then compare the measured rate against the ceiling to see how much of that resource is actually being used, then use counters to find what is consuming the rest of the time. The model supplies the first step and constrains the interpretation of the others. It never supplies a performance prediction on its own.