Compute-bound and memory-bound
Every kernel eventually hits one of two limits: the time it spends waiting on arithmetic units, or the time it spends waiting on memory traffic. NVIDIA calls the arithmetic side of that pair math bandwidth. A kernel is flop bound when there are stretches where nothing is moving through memory; it is memory bound when there are stretches where no floating-point operation is in flight. The ratio between a chip’s peak arithmetic throughput and its peak memory bandwidth sets exactly where that boundary falls, in operations performed per byte moved.
On an A100, that works out to 208 floating-point operations per byte moved. Below that arithmetic intensity, a kernel is memory bound: it finishes exactly as fast as its bytes can be streamed in, no matter how many arithmetic units sit idle. Above it, the kernel is flop bound, and streaming data faster would not help. This is the same reasoning the roofline model formalizes across an entire chip, not just one kernel.
That picture is Williams, Waterman, and Patterson’s, and reading it as a diagram leaves most of it unused. Their bound is one formula, attainable performance as the minimum of peak floating-point performance and peak memory bandwidth multiplied by operational intensity, and the term is chosen with care. Operational intensity means operations per byte of DRAM traffic, counting the bytes that reach main memory after the cache hierarchy has filtered them. The paper uses it rather than arithmetic intensity or machine balance because those measure traffic between the processor and the cache, while the roofline wants traffic between the cache and DRAM. The distinction is invisible when a working set fits on chip and decisive when it does not, which is the case the model was built for. Both limits are also created once per computer, not once per kernel: given a roofline, you can use it repeatedly on different kernels, since the roofline does not vary.
The knee in that figure has a name and a reading. Williams et al. call it the ridge point, where the diagonal and horizontal roofs meet, and its x-coordinate is the minimum operational intensity required to achieve maximum performance. That makes its position a statement about the machine rather than about any kernel. If the ridge point is far to the right, only kernels with very high operational intensity can achieve the maximum performance of that computer; if it is far to the left, almost any kernel can potentially hit maximum performance. The ridge point, in the paper’s words, suggests the level of difficulty for programmers and compiler writers to achieve peak performance on a given machine. Their own example is generational: moving from the two-core Opteron X2 to the four-core X4, which doubles peak floating-point performance on the same DRAM channels, shifts the ridge point right from 1.0 to 4.4, so a kernel needs an operational intensity higher than 1 to see any gain at all from the newer chip. An A100 demanding 208 flops per byte before it will run at peak is that same reading, applied to an accelerator.
What turns the roofline from a picture into a procedure is the set of ceilings underneath it. Suppose a program is performing far below its roofline: which optimizations should it apply, and in what order? The paper answers by drawing each optimization as a ceiling below the appropriate roof, under a rule that gives the stack its teeth. You cannot break through a ceiling without performing the associated optimization, and to break through one you must have already broken through all the ones below it. On their Opteron X2 the computational ceilings are, from the bottom up, improving instruction-level parallelism and applying SIMD, then balancing the floating-point operation mix, since peak typically requires an equal number of simultaneous additions and multiplications. The memory ceilings are restructuring loops for unit-stride accesses, ensuring memory affinity so that threads touch the DRAM local to their own chip, and using software prefetching.
Two rules make that stack actionable. The order of the ceilings suggests the optimization order, because they are ranked from bottom to top: those most likely to be realized by a compiler or with little effort by a programmer at the bottom, those difficult for a programmer to implement or inherently lacking in a kernel at the top. And the height of the gap between a ceiling and the next higher one is the potential reward for trying that optimization, so a tall gap is worth the work and a short one is not. Operational intensity then picks the region: a kernel far to the left should be reading the memory ceilings, one far to the right the computational ceilings, one in the middle both. Like the roofs above them, the ceilings need be measured only once per computer. That is the difference between owning a roofline chart and owning an ordered list of things to try.
One caveat travels with every roof drawn this way. The compute roof is itself a computed number, maximum clock frequency multiplied by the number of tensor cores multiplied by the flops each tensor core performs per clock, a product often called the speed of light. It holds only at the maximum clock, and in practice the peak throughput depends on the actual clock frequency, which can vary under power or thermal throttling. When the clock drops, so does the effective speed of light. Every percentage-of-peak figure in this book is therefore a percentage of a ceiling that moves.
The bandwidth in that ratio is a peak, and whether a kernel gets it depends on the shape of its accesses. Global memory loads and stores by the threads of a warp are coalesced by the device into as few transactions as possible, and on devices of compute capability 6.0 or higher the rule is simple: a warp’s concurrent accesses coalesce into however many 32-byte transactions are needed to service them all. Thirty-two threads reading adjacent 4-byte floats need four 32-byte transactions, and every byte fetched is a byte used. NVIDIA’s best practices guide marks keeping accesses coalesced as a high-priority recommendation, and the penalty for breaking the pattern is arithmetic: when adjacent threads access memory with a stride of two elements, half of every fetched segment goes unused, a 50 percent load and store efficiency, and as the stride grows the decay continues until a warp of 32 threads loads 32 separate 32-byte segments. A strided kernel is not paying the streaming rate the roofline promises; it is paying for bytes it never touches.
Coalescing decides how many bytes of each transaction get used. A second lever decides how many instructions it takes to ask for them. The default load and store instructions move 32 bits per thread, and NVIDIA’s pro tip on vectorized memory access opens by noting that many CUDA kernels are bandwidth bound and that the rising ratio of flops to bandwidth in new hardware is producing more of them. Its remedy is the vectorized load and store instructions, which move 64 or 128 bits at a time, and the way to reach them from CUDA C++ is not an intrinsic but a cast: the vector types in the standard headers, int2, int4, float2, float4, pack several values into one unit, and dereferencing a recast pointer makes the compiler emit the wide instruction.
__global__ void device_copy_vector2_kernel(int* d_in, int* d_out, int N) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
for (int i = idx; i < N/2; i += blockDim.x * gridDim.x) {
reinterpret_cast<int2*>(d_out)[i] = reinterpret_cast<int2*>(d_in)[i];
}
// in only one thread, process final element (if there is one)
if (idx==N/2 && N%2==1)
d_out[N-1] = d_in[N-1];
}Three things changed against the scalar version of the same copy. The loop runs N/2 times because each iteration moves two elements; a tail check handles the final element when N is odd; and the host launches half as many threads. In the disassembly, the scalar kernel’s LDG.E and STG.E become LDG.E.64 and STG.E.64, every other instruction unchanged, and the loop executes half as many times: a 2x reduction in instruction count that the post calls very important in instruction-bound or latency-bound kernels. The int4 version cuts it by a factor of four. The catch is alignment. These instructions require aligned data, and while device allocations are aligned to a multiple of the data type’s size, an offset pointer must be aligned too: reinterpret_cast<int2*>(d_in+1) is invalid where d_in+2 is fine. Structures work as long as they are a power of two bytes in size.
The post’s own verdict is that in almost all cases vectorized loads are preferable to scalar loads, with one reservation worth keeping: they increase register pressure and reduce overall parallelism, so a kernel that is already register limited or has very low parallelism may be better off with scalar loads. It is worth being precise about where the win comes from. The best practices guide notes there is no register-related reason to pack data into vector types like float4; the gain is in the memory instructions, not the register file.
The best practices guide prescribes a workflow for applying any of this: Assess, Parallelize, Optimize, Deploy, a cycle rather than a checklist. Assess profiles the application to find the code responsible for the bulk of the execution time and uses Amdahl’s and Gustafson’s laws to bound what accelerating it can possibly buy. Parallelize exposes the parallelism, sometimes as simply as calling an existing GPU library. Optimize is explicitly iterative (identify an opportunity, apply and test it, verify the speedup, repeat), so no one needs to memorize every strategy before seeing gains. Deploy ships each partial speedup to production before the next hotspot is tackled, so every pass around the loop pays for itself.
The critical batch size#
The boundary ratio has a second reading, and for inference it is the more useful one. Take the matrix multiply every MLP and every attention projection performs: a batch of B tokens against a weight matrix of shape D by F, in bf16. The arithmetic is 2 * B * D * F flops. The traffic is the batch, the weights, and the result, and when the model dimensions are large compared to the batch, which they almost always are, the weights dominate it. Being compute-bound then reduces to a condition on B alone: the batch must exceed the accelerator’s flops per second divided by its bytes per second. The arithmetic intensity of the hardware is not only a boundary in flops per byte. It is a batch size in tokens, the point at which a matmul stops being memory bound. How to Scale Your Model works it for a TPU v5e, 1.97e14 flops per second over 8.2e11 bytes per second, and gets 240 tokens. For an H100 it is about 280.
Precision moves that threshold in both directions. Writing beta for bits per parameter divided by bits per activation, and alpha_hbm for the accelerator’s flops divided by its memory bandwidth, the critical batch size is B_crit = beta * alpha_hbm. Quantizing the weights to int8 or fp8 halves it, because half as many bytes cross for the same arithmetic. Doing the flops in int8 or fp8 doubles it, because the arithmetic side gets faster. On a TPU v5e, int8 parameters with bf16 compute drop the threshold to 120, while int8 activations against int8 parameters put it back at 240, since the chip supplies 400 TOPs/s of int8 by int8 arithmetic.
One number then splits inference into two workloads with different hardware characters, which is the fact the rest of this book keeps leaning on. Prompts are hundreds to thousands of tokens long and all of those tokens go through the model at once, so during prefill all matrix multiplications are basically always compute-bound, and simply maximizing hardware utilization is enough to get both throughput per chip and time to first token. Generation has the opposite shape. Each request advances one token per step, and the steps are sequentially dependent, so the only route to B_crit is batching concurrent requests. A generate batch of 240 means 240 requests decoding at once and 240 separate KV caches, which is difficult to achieve in practice except in some bulk inference settings. Pushing more than 240 tokens through a prefill, by contrast, is routine.
Attention behaves differently again, and worse. Its arithmetic intensity works out to S * T / (S + T) for a cache of S past tokens and a query block of T. In prefill the two are equal, so the intensity is about T / 2 and grows with the sequence: attention is usually compute-bound for any reasonable sequence length, roughly above 480 tokens. In generation T is 1 while S is large, the ratio collapses to about 1, and nothing a kernel does raises it. The reason is structural. Parameters are reused across every item in the batch, so batching amortizes them, but every item carries its own KV cache, so batching buys attention nothing.
Those facts combine into the estimate worth carrying out of this chapter. For the small generate batches that are common, where both the attention and the MLP blocks are bandwidth bound, the theoretical minimum step time is (batch size * KV cache size + parameter size) / total memory bandwidth. As the batch grows the MLP crosses into the compute-bound regime and the general form separates the two terms: attention stays at batch size times KV cache size over bandwidth, never compute-bound and so never needing a flops roofline at all, while the MLP term becomes the larger of 2 * batch size * parameter count / total flops per second and parameter bytes over bandwidth.
LLaMA 2-13B shows what those terms weigh. One copy of the parameters in bf16 is 26 GB. A KV cache for a single 8,192-token sequence, at 40 KV heads by 128 dimensions by 40 layers, two bytes per number and two vectors per token, is 6.7 GB, so four concurrent sequences outweigh the entire model. On eight TPU v5e chips holding 128 GiB of HBM between them, generation runs out of memory beyond batch size 16, and theoretical throughput reaches only 964 tokens per second even at the batch of 240 that will not fit. Shrink the cache five times over by sharing 8 KV heads across the 40 query heads and the picture changes completely: a batch of 64 fits, latency improves at every batch size, and throughput keeps climbing to about 4,500 tokens per second at 240. Later LLaMA generations made exactly that change, and LLaMA 3 8B ships 32 query heads over 8 KV heads. The structure that produces those bytes is the next section’s subject.
Occupancy as a tuning knob#
The other lever is keeping the memory system busy at all. Thread instructions execute sequentially in CUDA, so executing other warps when one warp is paused or stalled is the only way to hide latency and keep the hardware busy, and occupancy names how much of that capacity is in play: the ratio of active warps per multiprocessor to the maximum number of possible active warps. The best practices guide is careful about the direction of the claim, and the care is the whole point. Higher occupancy does not always equate to higher performance, and there is a point above which additional occupancy does not improve performance. But low occupancy always interferes with the ability to hide memory latency. Occupancy is a floor to clear, not a score to maximize.
The direction of that claim makes more sense once you know what the ratio feeds. An SM is physically divided into four quadrants, each housing a subset of its compute units, and each quadrant has a warp scheduler that can issue one warp instruction per cycle. So an SM issues instructions from at most four warps in a given cycle, 128 threads in true parallel execution, while hosting up to 2,048 concurrent threads, 64 warps, that are resident and scheduled in and out over time. Parallelism and concurrency are separate quantities here: the first is capped at 128 threads per SM, the second runs to 2,048, and occupancy measures the second. It is the size of the pool those four schedulers draw from. Raising it does not make the SM retire more instructions per cycle; it raises the odds that when the warps currently issuing stall on memory, some other warp is eligible to take their slot. That is also where the floor under block size comes from. A block is resident on a single SM, and with four schedulers to keep fed, a block wants at least four warps, which is 128 threads.
What sets it is how much of an SM one block consumes. Every SM publishes its limits through cudaGetDeviceProperties, and the programming guide works an example on a compute capability 10.0 device with 2,048 resident threads per SM and a ceiling of 32 resident blocks. Launch that device with 768 threads per block and only two blocks fit, because a third would exceed the thread ceiling: occupancy is (768 x 2) / 2048, or 75 percent. Launch it with 32 threads per block and the thread ceiling never binds; the block ceiling does, so 32 blocks of 32 threads leave 1,024 of a possible 2,048 threads resident, or 50 percent. Same kernel, same SM, occupancy halved by a block size.
Registers are the subtler constraint, because the programmer does not set them directly. They are allocated to an entire block at once out of a file that every resident thread shares, so on a compute capability 7.0 device with 65,536 32-bit registers per SM and a maximum of 2,048 simultaneous threads, 100 percent occupancy leaves each thread at most 32 registers. Granularity then bends the arithmetic. On that same device, a kernel using 37 registers per thread reaches 75 percent occupancy with 12 resident 128-thread blocks, but the identical 37 registers in 320-thread blocks reach only 63 percent, because just four such blocks fit. Register allocations are also rounded up to the nearest 256 registers per warp. The exact relationship between register usage and occupancy, the guide concedes, can be difficult to determine.
The knobs are -maxrregcount, which caps registers per thread for a whole file, and __launch_bounds__, which does it per kernel. Turning either one down does not make the values disappear; it pushes them into local memory. The programming guide states the trade without softening it: a kernel that needs more registers than the cap is likely to spill, which changes its performance characteristics, and yet in some cases even though spilling occurs, limiting registers allows more thread blocks to be scheduled, which increases occupancy and may result in a net increase in performance. That is a measurement, not a rule, and it is why this book treats occupancy as a knob rather than a target.
One use of __launch_bounds__ is not optional. The best practices guide states that to maintain forward compatibility with future hardware and toolkits, and to ensure that at least one thread block can run on an SM, developers should include the single-argument form __launch_bounds__(maxThreadsPerBlock) naming the largest block size the kernel will ever be launched with. Failure to do so could lead to “too many resources requested for launch” errors. The two-argument form, which adds a desired minimum number of resident blocks per multiprocessor, can improve performance in some cases, and the right value should be determined by a detailed per-kernel analysis. Because the optimal bounds typically differ across major architecture revisions, the guide keys them off the architecture macro:
#define THREADS_PER_BLOCK 256
#if __CUDA_ARCH__ >= 900
#define MY_KERNEL_MAX_THREADS (2 * THREADS_PER_BLOCK)
#define MY_KERNEL_MIN_BLOCKS 3
#else
#define MY_KERNEL_MAX_THREADS THREADS_PER_BLOCK
#define MY_KERNEL_MIN_BLOCKS 2
#endif
__global__ void
__launch_bounds__(MY_KERNEL_MAX_THREADS, MY_KERNEL_MIN_BLOCKS)
MyKernel(...) {
...
}The qualifier is a promise the compiler plans against. From the bounds it derives an upper limit L on registers, the number that still lets the requested blocks of that size be resident. If the kernel’s initial register usage exceeds L, the compiler reduces it until it fits, which usually results in increased local memory usage, a higher instruction count, or both. If usage is already below L and both arguments were given, the compiler may go the other way and raise register usage up to L in order to reduce the number of instructions and better hide the latency of single-threaded instructions. The arguments survive into PTX as the .maxntid and .minnctapersm directives, and a kernel launched with more threads per block than its bound fails to launch at all.
None of this is worth guessing at when it can be read off. Compiled with --resource-usage (or --ptxas-options=-v), nvcc reports registers per thread and total local memory per kernel, and a variable that landed in local memory shows up in the PTX declared .local and accessed through ld.local and st.local. NVIDIA ships an occupancy calculator inside Nsight Compute for projecting the result before running anything, and the CUDA runtime exposes an occupancy API, cudaOccupancyMaxActiveBlocksPerMultiprocessor, so an application can select launch configurations from runtime parameters rather than from a constant baked in at build time.
What remains is a short list of defaults the guide is willing to state outright. Threads per block should be a multiple of the warp size, both to avoid wasting computation on under-populated warps and to facilitate coalescing. A minimum of 64 threads per block should be used, and only if there are multiple concurrent blocks per multiprocessor. Between 128 and 256 threads per block is a good initial range for experimentation. Prefer several smaller blocks over one large one when latency affects performance, which particularly helps kernels that call __syncthreads() often. The grid should have more blocks than the device has multiprocessors so every SM has work, and to scale to future devices the number of blocks per launch should be in the thousands. Past that the guide stops prescribing, because a lower-occupancy kernel has more registers available per thread and may spill less, and with a high degree of exposed instruction-level parallelism it is in some cases possible to fully cover latency at low occupancy. The cheapest way to find out which case a kernel is in costs no code change at all: raise the dynamic shared memory argument in the execution configuration, which lowers occupancy without modifying the kernel, and measure what happens.