Achieved Bandwidth Versus the Quoted Peak

A datasheet bandwidth figure is a property of the memory interface, computed from bus width, signalling rate, and the number of channels. It bounds the interface rather than promising what a kernel can extract: reaching it needs continuous full-width transfers with no gaps, and real access patterns leave gaps. The distance between that figure and a stream rate measured on the same part decomposes into parts with different causes and different fixes, and separating the parts is the whole of the diagnostic.

How the peak number is derived

A memory interface with a bus width of $W$ bits per channel, $N$ channels, and a signalling rate of $R$ transfers per second per wire carries

$$ B_{theoretical} = \frac{W \cdot N \cdot R}{8} $$

bytes per second. Every term is a physical constant of the board layout and the memory device’s rated speed. Nothing in the formula involves the processor, the kernel, or the access pattern.

Two corrections apply before comparing against it. The first is that the rate is a transfer rate; on interfaces that move a burst per access, the usable payload is the burst size, and small or scattered accesses pay for full bursts. The second is that some interfaces carry error-correcting bits alongside data, and the fraction of the bus devoted to them is unavailable to the kernel. Where ECC is inline rather than out of band, the practical ceiling is below the datasheet number by that fraction.

A third correction is thermal and operational: the memory clock the part sustains under a compute load is not always the clock at which the interface was rated, and power management can lower it. That effect belongs with clocks and the moving roofline, because it changes the ceiling over time rather than acting as a fixed derate.

Counting bytes rather than assuming them

Any comparison requires a measured byte count. The honest denominator comes from memory-system counters, not from the source code. A counter that reports DRAM read bytes and write bytes at the partition level answers the question directly, and it disagrees with a source-level count for several reasons worth enumerating.

Write allocation. A partial write to a line the memory system does not hold may cause the line to be fetched before it is modified, so a kernel that writes 4 bytes per element can generate a read of 32 or 128 bytes per element as well. Full-line writes avoid this. The counter sees the fetch; the source does not.

Sector granularity. Transfers are accounted in sectors, not in bytes requested. A warp that requests 32 consecutive 4-byte words from 32 threads touches 128 bytes; a warp whose 32 threads each touch a different 128-byte line touches 32 lines. Both are 32 loads in the instruction stream. The first touches four 32-byte sectors; the second touches 32, so the counter reports eight times the traffic for the same instruction count, and that difference is the whole subject of sector counting.

Redundant traffic. A line evicted before all its consumers read it is fetched again. The counter reports both fetches; the source-level count reports one read.

Metadata and bookkeeping. Page-table entries fetched for address translation consume interface capacity without appearing in any kernel’s byte count. On some designs this overhead scales with the number of distinct pages touched, which makes it sensitive to allocation layout rather than to the kernel.

The rule that follows: compute intensity for the roofline from counter bytes, and use source-level bytes only as a sanity check on whether the counters are reporting something unexpected.

What a stream loop can and cannot show

The standard way to establish a practical ceiling is a triad-like loop that reads two arrays and writes one, with no reuse. Such a loop is pure bandwidth: its arithmetic intensity is around 0.17 FLOP/byte, so on every current device it is far below the ridge and the memory system is unambiguously the binding resource. Whatever rate it achieves is a lower bound on the achievable ceiling, and it is usually taken as the reference.

Two caveats make it a lower bound rather than an equality. The first is the read/write mix. A loop that is two-thirds reads and one-third writes does not exercise the interface the way a read-only reduction does, because write turnaround on the bus costs cycles and because writes and reads contend differently for the same channels. A kernel whose mix is more write-heavy will see a different ceiling, and a kernel that is almost entirely reads may exceed the triad rate.

The second caveat is concurrency. A stream loop with one element per thread and a low arithmetic count per element may not have enough outstanding requests to cover memory latency, in which case it measures latency tolerance rather than bandwidth. The fix is a grid-stride loop with multiple independent elements in flight per thread, which raises memory-level parallelism without changing the byte count. That mechanism is treated separately in memory-level parallelism.

// Grid-stride triad: the loop keeps blocks resident long enough to
// sustain requests, and the stride guarantees coalesced accesses.
__global__ void triad(int n, float a, const float *x, const float *y, float *z) {
    int stride = gridDim.x * blockDim.x;
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < n; i += stride) {
        z[i] = a * x[i] + y[i];
    }
}

Decomposing the gap

A measured rate below the datasheet figure is not a single defect. It is a sum of contributions, and each has a different remedy.

Contribution	Signature	Remedy
Insufficient requests in flight	Rate rises with more blocks or more elements per thread, then plateaus	Increase parallelism or per-thread independent accesses
Read/write turnaround	Write-heavy kernels underperform read-heavy ones at the same byte count	Restructure to batch writes; accept the cost
Partial-line writes	Write traffic far exceeds the bytes written	Pad or rearrange so each warp writes full lines
Sector over-fetch	Read traffic exceeds the bytes requested	Improve coalescing or change the data layout
Refresh and maintenance	Fixed percentage independent of the kernel	None; it is part of the practical ceiling
Clock derate under load	Rate falls as the kernel runs longer	Treat the sustained clock as the real ceiling

The first row is the one most often mistaken for a bandwidth limit. A kernel that reaches a plateau as parallelism increases has found its ceiling. A kernel whose rate keeps rising with parallelism has not yet reached the memory system’s limit at all, and comparing it to the datasheet at that point measures the kernel’s concurrency, not the hardware.

The last two rows are the reason a practical ceiling should be established on the same device, under the same power state, with a known-good kernel. Transferring a ceiling from another part or another configuration is the most common source of a diagnostic that points at the wrong resource.

The refresh row deserves a note because it is irreducible. DRAM cells leak charge and must be rewritten periodically, and the fraction of interface time devoted to that maintenance is fixed by the device technology rather than by the access pattern. On some interfaces the refresh is scheduled around traffic and costs little; on others it consumes a few percent of the available cycles regardless of load. Whatever the figure, it belongs inside the practical ceiling and should never be attributed to the kernel.

The clock row is different in kind: it is not a fixed derate but a moving one. A kernel that runs for milliseconds may execute entirely at a boost clock, while the same kernel in a longer sequence settles to a lower sustained clock. Comparing the two durations yields a ratio that looks like a kernel inefficiency and is actually a power-state difference. That mechanism, and the measurement protocol that neutralizes it, is treated in clocks and the moving roofline.

Reporting a fraction of peak honestly

A statement of the form “the kernel achieves 62% of peak bandwidth” carries three implicit claims: which peak, which bytes, and which clock. Each needs to be fixed before the number means anything.

The peak should be the practical ceiling established by a reference loop on the same device, and the datasheet figure should be quoted alongside it as context rather than as the denominator. The bytes should come from counters and should be reported as read and write separately, since the mix is what makes the comparison meaningful. The clock should be the sustained clock during the measured interval.

With those fixed, a fraction below the practical ceiling has a definite interpretation: the kernel is not saturating the interface, and the remaining time is being spent waiting rather than transferring. A fraction at or near the practical ceiling closes the bandwidth question and moves the diagnosis to the arithmetic side or to the launch geometry. The fraction is a diagnostic instrument, not a score, and the useful output of measuring it is the decision it forces about where to look next.