Chapter 8 GPU Architecture for LLM Engineers

8.1 The Compute and Memory Hierarchy

8.1.1 SMs, Warps, Threads, and Blocks

An interviewer designing a custom CUDA kernel or probing your understanding of FlashAttention will not ask “what is an SM?” - they will ask you to walk through what happens when a warp executes a matrix multiply, why warp divergence tanks throughput, and why tile counts that are not multiples of the SM count cause a cliff drop in utilization. Getting this right requires internalizing the three-level execution hierarchy cold. This question appears at the ML infrastructure engineer and systems engineer level, and confidently deriving performance predictions from first principles is what separates a strong answer from a mediocre one.

A GPU is a throughput machine, not a latency machine. A CPU is designed to finish any single task as fast as possible - its die area goes to branch predictors, out-of-order execution engines, and large caches that minimize single-thread latency. A GPU sacrifices per-thread latency to maximize the total number of operations completed per second across thousands of concurrent threads. If you give a CPU four tasks, it finishes the first quickly, then the second, and so on. A GPU finishes all four simultaneously - later in wall-clock time than the CPU’s first task, but far sooner in aggregate.

The atomic unit of GPU computation is the Streaming Multiprocessor (SM). An SM is a self-contained engine: it has its own control logic, register file, shared memory, and specialized arithmetic units including tensor cores. An A100 SXM4 has 108 SMs; an H100 SXM5 has 132. Each SM can hold up to 2048 threads in flight simultaneously, which it uses to hide memory latency by switching between warps while some are stalled.

The execution hierarchy has three levels:

Block. A block is a group of threads assigned to a single SM for the duration of its execution. All threads in a block share the SM’s shared memory and can synchronize with barriers. A block cannot span SMs. The programmer chooses how many threads are in each block (must be a multiple of 32) and how many blocks to launch (the grid).

Warp. A warp is 32 consecutively numbered threads within a block. This is the unit at which the SM dispatches instructions. All 32 threads in a warp execute the same instruction simultaneously, each on different data. This execution model is called SIMT: Single Instruction, Multiple Threads. The SM maintains many warps concurrently; when one warp stalls on a memory read, the SM instantly switches to another ready warp with zero overhead (each warp’s live state is stored in registers, not on a call stack).

Thread. A thread is the individual execution lane. Each thread has its own registers and runs its own instance of the kernel code. Threads within a warp are always in lockstep; threads in different warps are independent.

The performance consequence of SIMT is warp divergence: if threads within a warp take different branches through an if-else, the SM must execute both branches serially, masking off the inactive threads for each. A 50/50 branch split halves effective throughput. Kernels should be written to avoid conditional logic that splits warps.

The SM is the atom of GPU computation: blocks are assigned to SMs, warps are the 32-thread instruction-dispatch units within a block, and any branch that splits a warp forces serial execution - design kernels to maintain warp uniformity.

The SIMT model wins because matrix multiplies and tensor operations are perfectly amenable to it: every thread does exactly the same arithmetic on a different slice of the matrix. The GPU dedicates the vast majority of its die area to arithmetic units rather than control logic, which is why an H100 delivers $\sim 2 , 000$ TFLOP/s of bfloat16 tensor-core throughput while a CPU core delivers $\sim 1$ TFLOP/s at most - not because the GPU runs at a higher clock, but because it runs the same instruction across thousands of parallel threads at once.

The lightweight thread model is also central to the design. A CPU thread requires a full stack, OS scheduling, and context switch overhead measured in microseconds. A GPU thread’s entire state is a fixed-size register file entry; switching between warps takes zero cycles because the SM simply picks up where the other warp left off. This is how an SM with 2048 in-flight threads achieves high utilization despite the 300-cycle latency of a global memory access: 300 cycles is enough time to cycle through many other warps and keep the arithmetic units fed.

The alternative - a CPU-style design with a small number of heavy threads and large caches - would give low latency on a single problem but would not scale to the thousands of simultaneous accumulators that a large matrix multiply requires.

A Worked Example

Consider a bfloat16 matrix multiply $C = A \times B$ where $A , B \in \mathbb{R}^{4096 \times 4096}$ on an A100 (108 SMs).

Tiling. The output matrix $C$ has $4096 \times 4096 = 16 . 8$M elements. A typical output tile size for an A100 GEMM kernel is $128 \times 128$ elements (chosen to fit the tensor core’s preferred shape and the SM’s shared memory). The number of output tiles is:

Block assignment. Each tile is one block, so 1024 blocks are launched. With 108 SMs, the scheduler fills all 108 SMs in the first wave of 108 blocks. The second wave takes the next 108, and so on:

The 10th wave dispatches only $1024 - 9 \times 108 = 52$ blocks, leaving 56 SMs idle for the duration of that wave. GPU time is approximately $10 \times$ the single-wave time rather than $( 1024 / 108 ) \approx 9 . 5 \times$. This is wave quantization: adding even one tile beyond a wave boundary doubles the cost of that marginal tile.

Warp execution within a block. Each block contains 256 threads = 8 warps. The SM schedules these 8 warps in rotation. When warp 0 issues a shared-memory load and stalls for $\sim$20 cycles, warps 1-7 are immediately dispatched. By the time the SM cycles back to warp 0, the load has completed. The SM stays busy.

Fix for wave quantization. Round $4096$ to $4224 = 33 \times 128$, giving $33^{2} = 1089$ tiles. That is still 11 waves with 108 SMs, but the last wave has $1089 - 10 \times 108 = 9$ tiles - even more waste. A better fix: choose a dimension that makes the tile count a multiple of 108, e.g., $4096$ with tile size $\approx 86$ (non-round) is impractical. In practice, run the problem with a tile count of $1080 = 10 \times 108$ by padding the dimension to $\lceil 4096 /$tile$\rceil \times$tile and choosing the tile dimension to make the product a multiple of 108.

What You Actually Decide in Practice

• Always make thread block size a multiple of 32. Any block size that is not a multiple of the warp size (32) wastes the last partial warp’s lanes. Standard choices: 128, 256, 512, or 1024 threads per block depending on register pressure.

• Avoid conditional branches that split warps. An if (thread_idx < 16) construct where both branches do real work halves the throughput of that warp. Restructure using masking (compute both branches, select the result with a conditional move) or ensure all threads in a warp always take the same branch.

• Target tile counts that are multiples of the SM count. For an A100 (108 SMs), target grid sizes that are multiples of 108. For an H100 (132 SMs), multiples of 132. If you cannot choose the grid size, at minimum ensure the final wave has high occupancy (close to 108 or 132 blocks).

• Tensor core input must be a multiple of 16. A100 tensor cores operate on $16 \times 16$ matrix fragments in bfloat16. Dimensions that are not multiples of 16 require padding; dimensions that are multiples of 128 are ideal because they align with the preferred GEMM tile size and the 128-byte L2 cache line. This is the engineering reason behind the well-known speedup from rounding vocabulary sizes to multiples of 64.

• Maximize occupancy by controlling register pressure. Each SM on an A100 has 65,536 registers. With 256 threads per block and 128 registers per thread, the block uses $256 \times 128 = 32 , 768$ registers, leaving room for exactly two blocks per SM (occupancy $= 2 \times 256 / 2048 = 25 \%$). Reduce registers per thread (by spilling to local memory or restructuring the kernel) to fit more blocks per SM and increase the warp pool for latency hiding.

How This Shows Up in Interviews

• “What is a warp and why does its size matter?” A warp is 32 threads that execute the same instruction simultaneously on different data. Its size matters because (1) it is the unit of instruction dispatch, so thread block sizes that are not multiples of 32 waste compute, and (2) branches within a warp cause serial execution (warp divergence), reducing throughput by the fraction of threads on the minority branch.

• “How does an SM hide memory latency?” By maintaining many warps in flight simultaneously. When one warp issues a memory read and stalls for 300+ cycles, the SM instantly switches to another ready warp. No context switch overhead because each warp’s register state is always live in the register file. This is why occupancy (ratio of active warps to the SM’s maximum) matters: higher occupancy means more warps to switch to, and the SM stays busy even with long-latency memory operations.

• Senior follow-up: “An A100 has 108 SMs. Your kernel launches 109 thread blocks. Quantify the performance impact and describe the fix.” 108 blocks execute in wave 1, using all 108 SMs in parallel. Block 109 executes alone in wave 2, using 1/108 of the SM capacity. Total GPU time $\approx 2 \times$ the single-wave time rather than $109 / 108 \approx 1 . 009 \times$. The fix: either reduce the problem to 108 blocks (coarsen the tile size), or pad to 216 blocks (the next multiple of 108) and ensure the added work is trivial (zero-padding the output matrix). In practice this means choosing tile dimensions that make the grid size $\equiv 0 ( \mathrm{mod} 108 )$, which often requires rounding the problem dimension up by a small amount.

Key Takeaways

• GPU execution hierarchy: Block (assigned to one SM) $\rightarrow$ Warp (32 threads, same instruction, different data) $\rightarrow$ Thread (individual register state and computation lane).

• A100 has 108 SMs, H100 has 132. Each SM can hold up to 64 warps (2048 threads) in flight simultaneously, switching between them with zero overhead to hide memory latency.

• SIMT means all 32 threads in a warp execute the same instruction; any branch that splits a warp serializes execution for the two groups, reducing throughput proportionally.

• Grid sizes (number of blocks) should be multiples of the SM count; a remainder causes the last wave to severely underutilize the GPU - this is wave quantization.

• Tensor core operations require input dimensions to be multiples of 16 (ideally 128) in bfloat16; misaligned dimensions require padding and reduce effective throughput.

8.1.2 HBM, L2, Shared Memory, and Registers

When an interviewer asks why FlashAttention is roughly 3$\times$ faster than a naive PyTorch attention implementation, the correct answer is not “it avoids materializing the full attention matrix” - the correct answer is “it moves roughly 7$\times$ less data through HBM.” GPU compute throughput has scaled $\sim$10,000$\times$ from K20 to H100; HBM bandwidth has scaled $\sim$100$\times$. Every generation, the gap widens, and every generation, the dominant bottleneck in LLM kernels is memory movement, not arithmetic. Understanding the memory hierarchy in numbers - latency ratios, bandwidth, capacity at each level - is foundational for any performance conversation at the research engineer or systems engineer level.

A GPU has five distinct levels of memory, organized from fastest and smallest near the compute units to slowest and largest farther away:

Registers (fastest, $\sim$1 cycle): per-thread, $\sim$255 32-bit registers available per thread on an A100. Holds live scalar values during execution. No programmer overhead to use them - the compiler allocates registers automatically. Spilling to local memory (HBM) when register pressure is too high is a serious performance hazard.

Shared memory / L1 cache ($\sim$20 cycles, per-SM): configurable up to 164 KB total per SM on an A100, split between shared memory and L1 cache. The programmer explicitly controls shared memory by declaring __shared__ arrays. Shared by all threads in a block; not visible to other blocks. This is the primary staging area for data that multiple threads will reuse.

L2 cache ($\sim$200-300 cycles): 40 MB on the A100, shared across all 108 SMs. Hardware-managed; the programmer cannot directly address it. Acts as a buffer between repeated HBM reads and the SMs. Effective for weight matrices in batched inference where the same weights are reused across many requests.

HBM (High Bandwidth Memory) / global memory ($\sim$300-600 cycles latency, 2 TB/s bandwidth on A100, 3.35 TB/s on H100 SXM): 40-80 GB capacity. This is what CUDA programmers call “global memory.” All inter-block communication must go through HBM. This is also where model parameters, activations, and KV caches live.

Host DRAM (CPU side, via PCIe 4.0): $\sim$64 GB/s in each direction on a PCIe 4.0 $\times$16 link - roughly 30$\times$ slower than HBM. This level matters for data loading, model transfers, and CPU-GPU coordination, but is too slow to be in the critical path of any compute-intensive kernel.

The critical ratio is between compute throughput and HBM bandwidth. An A100 delivers 312 TFLOP/s of bfloat16 tensor-core throughput and 2 TB/s of HBM bandwidth. A bfloat16 element is 2 bytes, so HBM delivers $10^{12}$ elements/second. The compute unit can consume $312 \times 10^{12} / 10^{12} = 312$ FLOPs per element delivered. A kernel is compute-bound only if its arithmetic intensity (FLOPs per byte loaded from HBM) exceeds $312 / 2 = 156$ FLOP/byte. Most element-wise operations (softmax, layer norm, ReLU) have arithmetic intensity below 5 FLOP/byte - they are deeply memory-bound.

HBM uses burst-mode access. When a read is issued to any address, the DRAM controller moves a full 128-byte cache line (one burst section) to the L2 or SM. If the thread only needed 2 bytes, the remaining 126 bytes are either cached or discarded. Memory coalescing exploits burst mode: when 32 threads in a warp issue reads to 32 consecutive 2-byte addresses that together span one 64-byte or 128-byte aligned burst section, the hardware services all 32 reads with a single HBM transaction. Non-coalesced access - 32 threads reading 32 scattered addresses - requires up to 32 separate transactions, each fetching 128 bytes and using only 2. Effective bandwidth drops by up to 32$\times$.

The performance gap between compute throughput and HBM bandwidth grows with every hardware generation; every optimization strategy for GPU kernels ultimately reduces to minimizing data movement through HBM.

The memory hierarchy is not a design choice - it is physics. Faster memory requires transistors physically closer to the compute units and costs exponentially more die area per bit. Registers and shared memory live on-SM because the SM needs sub-nanosecond access ($\sim$1 ns at 1 GHz). HBM is a stack of DRAM dies bonded adjacent to the GPU package; it trades latency for large capacity. Host DRAM sits on the CPU board and is an order of magnitude further away.

The strategic consequence is that compute has become so abundant that it is nearly free compared to data movement. Moving a bfloat16 value from HBM to an SM costs roughly 30$\times$ more energy per operation than performing a multiply-add on it. Kernels that move data unnecessarily waste almost all of their energy and time waiting, not computing. This is why kernel fusion - combining multiple operations into a single kernel that keeps data in registers and shared memory across all operations - is one of the highest-leverage optimizations available. A fused operation that applies 5 element-wise transformations reads and writes HBM once; the unfused equivalent reads and writes HBM 5 times for the same arithmetic result.

The contrast with the memory-bound profile of CPU workloads is instructive. CPU caches are hardwaremanaged and opaque to the programmer. On a GPU, shared memory is programmer-managed and explicit - you load a tile, compute over it, and write back. This explicitness is a burden during development but enables the precise control needed to achieve 80-90% of peak HBM bandwidth on well-structured kernels.

A Worked Example

A100 SXM4. Layer normalization over a tensor of shape $( T , D ) = ( 2048 , 4096 )$ in bfloat16.

Data volume. $2048 \times 4096 \times 2$ bytes $= 16 . 8$ MB.

Naive three-pass implementation. Pass 1: read $x$, compute mean per row (1 HBM read). Pass 2: read $x$ again, compute variance per row (1 HBM read, 1 HBM write of variance). Pass 3: read $x$, read mean, read variance, compute normalized output, write output (3 HBM reads, 1 HBM write). Total HBM traffic: $5 \times 16 . 8$ MB $+ 2048 \times 4$ bytes for mean/variance $\approx 84$ MB.

Fused one-pass implementation. Load each row of $x$ once into shared memory. Compute mean and variance in registers using an online Welford accumulator. Normalize and write the output in the same pass. Total HBM traffic: 1 read $+$ 1 write $= 2 \times 16 . 8$ MB $= 33 . 6$ MB. Speedup: $84 / 33 . 6 \approx 2 . 5 \times$ from eliminating HBM traffic alone, at zero change in arithmetic.

Time at peak HBM bandwidth.

Coalescing example. Same tensor in row-major layout ($T$ rows, $D$ columns). Kernel assigns one thread per row; a warp of 32 threads processes 32 consecutive rows. At the start of pass 1, warp reads element $[ 0 , 0 ] , [ 1 , 0 ] , . . . , [ 31 , 0 ]$ (one element per thread from column 0). These 32 addresses are stride-$D$ apart ($D = 4096$ elements $= 8192$ bytes). Each address falls in a different burst section: 32 separate HBM transactions required. Effective bandwidth: $32 \times 2$ bytes $/ ( 32 \times 128$ bytes fetched) $= 1 / 64$ of peak. Fix: tile the reduction so that within each warp, threads access consecutive elements within a row (stride 1), enabling full coalescing.

What You Actually Decide in Practice

• Configure shared memory vs. L1 split for your kernel’s access pattern. On A100, you can allocate up to 164 KB of shared memory per SM (the rest goes to L1). For kernels that explicitly manage tiling (FlashAttention, fused GEMMs, layer norm), maximize shared memory. Use cudaFuncSetAttribute with cudaFuncAttributeMaxDynamicSharedMemorySize. For kernels that rely on hardware caching of irregular accesses, bias toward L1.

• Keep tile sizes within 48-96 KB of shared memory per block. An A100 SM has 192 KB of on-chip memory. With two concurrent blocks per SM (the common case), each block can use up to 96 KB. With three blocks, the limit is 64 KB. Choose tile dimensions $T_{r} \times T_{c}$ such that the shared memory footprint stays below your budget; exceeding it reduces SM occupancy.

• Align all HBM accesses to 128-byte boundaries. The L2 cache line is 128 bytes. Accesses that straddle cache line boundaries require two transactions. Pad matrix dimensions to multiples of 64 for bfloat16 (128 bytes = 64 $\times$ 2-byte elements) or multiples of 32 for float32. Padding a vocabulary size from 50,257 to 50,304 (the nearest multiple of 64) produces a measurable throughput improvement for the token embedding lookup.

• Profile with memory throughput as the primary metric, not FLOP utilization. Use ncu --metrics l1tex__t_bytes.sum,lts__t_bytes.sum,gpu__dram_throughput.avg to measure L1, L2, and HBM traffic. A kernel at 90% FLOP utilization that is memory-bound is not a compute-bound kernel - it means the memory traffic is the bottleneck and compute is hiding behind it. Reduce HBM traffic first; FLOP utilization is secondary.

• Exploit L2 reuse for weights in batched inference. In multi-query or grouped-query inference, the key/value weight matrices are shared across the batch. At batch size 64, a 4096$\times$4096 bfloat16 weight matrix fits in the 40 MB A100 L2 ($4096 \times 4096 \times 2 = 33 . 6$ MB). Structure the kernel to process all batch elements against one weight slice before moving to the next, so the L2 cache holds the weight during the full batch processing and HBM is only read once per weight matrix.

How This Shows Up in Interviews

• “What is the difference between shared memory and global memory, and why does it matter?” Shared memory is on-SM, $\sim$20 cycle latency, 96 KB per SM on A100. Global memory (HBM) is off-chip, $\sim$300-600 cycle latency, 80 GB, 2 TB/s bandwidth. The $15$-$30 \times$ latency ratio means any kernel that reads the same data from HBM multiple times instead of caching it in shared memory wastes most of its time waiting. The key design pattern: load from HBM once into shared memory, compute over the tile in registers/shared memory, write back once.

• “What is memory coalescing and when does it fail?” When 32 threads in a warp read 32 consecutive addresses that together fit in one 128-byte HBM burst section, the hardware issues a single transaction - one read brings all 32 values. It fails when threads access strided or random addresses: each thread’s address falls in a different burst section, requiring up to 32 transactions for the same amount of useful data. Column-major access of a row-major matrix is the canonical failure case: threads in a warp read elements with stride $D$ (the row width), touching 32 different cache lines.

• Senior follow-up: “Compute throughput scales 10,000$\times$ from K20 to H100 but HBM bandwidth scales only $\sim$100$\times$. What does this imply for algorithm design going forward?” The arithmetic intensity threshold for being compute-bound (FLOPs per byte) increases by 100$\times$ over this span. Operations that were memory-bound on K20 are even more memory-bound on H100. Three concrete implications: (1) kernel fusion has compounding value - each fused operation eliminates a full HBM round-trip at a penalty that grows 100$\times$ over this hardware range; (2) quantization (INT8, FP8) reduces bytes moved by 2-4$\times$ and its benefit grows proportionally as the memory wall tightens; (3) any new architecture that requires sequential operations with no fusion opportunity (e.g., state-space models with per-element recurrence) will be increasingly penalized relative to block-structured operations like attention. Designing for memory efficiency is not a current concern - it is an evergreen constraint that only becomes more important.

Key Takeaways

• Memory hierarchy latency ladder: registers ($\sim$1 cycle) $\rightarrow$ shared memory/L1 ($\sim$20 cycles, 96 KB per SM) $\rightarrow$ L2 ($\sim$200-300 cycles, 40 MB) $\rightarrow$ HBM ($\sim$300-600 cycles, 2 TB/s, 40-80 GB) $\rightarrow$ host DRAM (PCIe, $\sim$64 GB/s).

• An A100 delivers 312 TFLOP/s bfloat16 compute and 2 TB/s HBM bandwidth; a kernel is compute-bound only when its arithmetic intensity exceeds $156$ FLOP/byte. Most LLM non-GEMM operations are at $< 5$ FLOP/byte - deeply memory-bound.

• Memory coalescing: 32-thread warps reading 128-byte-aligned consecutive addresses get one HBM transaction; strided or random access requires up to 32 transactions for the same data volume.

• Compute scales $\sim$10,000$\times$ from K20 to H100; HBM bandwidth scales $\sim$100$\times$; the gap grows every generation, making kernel fusion and quantization increasingly high-leverage.

• Shared memory is the programmer-controlled fast cache: load from HBM once, compute in shared memory, write back once. This is the design principle behind FlashAttention and every high-performance transformer operator.

8.2 Arithmetic Intensity and the Roofline Model

8.2.1 Compute-Bound vs. Memory-Bound Regimes

The most common GPU performance question in an LLM systems interview is some variant of “why is autoregressive decoding slow, and what can you do about it?” Candidates who answer purely in terms of algorithmic complexity-“it is $O ( n^{2} )$ attention”-miss the real bottleneck. During generation the constraint is almost never FLOPs; it is memory bandwidth. Internalizing this requires two concepts: arithmetic intensity and the roofline model. This pair of ideas shows up at every seniority level; senior answers add numbers instead of stopping at the qualitative claim.

Arithmetic intensity is the ratio of floating-point operations performed to bytes transferred between the accelerator’s compute units and its main memory (HBM on a GPU):

AI $=$FLOPs Bytes transferred from HBM The roofline model says a kernel’s achievable throughput in FLOPs/sec is bounded by two separate ceilings. The memory ceiling is bandwidth$\times$AI: regardless of how many tensor cores you have, you cannot process data faster than you can deliver it from HBM. The compute ceiling is the chip’s peak FLOPs. The actual achievable throughput is:

A kernel is memory-bound when AI $<$ peak compute$/$bandwidth and compute-bound when the reverse holds. The crossover is the ridge point:

ridge point $=$peak compute (FLOPs/s) memory bandwidth (bytes/s) On an H100 SXM5 running BF16 tensor-core operations:

• Peak tensor-core throughput: $\approx 989$ TFLOPs/s

• HBM3 bandwidth: $\approx 3 . 35$ TB/s

• Ridge point: $989 \times 10^{12} / 3 . 35 \times 10^{12} \approx 295$ FLOPs/byte

To saturate the H100’s compute, a kernel must execute at least 295 floating-point operations per byte fetched from HBM-and most LLM operations during inference do not come close.

GPU generations scale compute throughput far faster than memory bandwidth. From the V100 to the H100, BF16 tensor-core peak grew roughly $10 \times$ while HBM bandwidth grew roughly $3 \times$. The ridge point climbs with every generation: operations that were barely compute-bound in 2018 are memory-bound today, and the trend continues. Designing an efficient kernel means understanding which side of this ridge your operation sits on, because the optimization lever is completely different in each regime.

Below the ridge point the GPU is starved for data. An SM loads a tile from HBM, executes a handful of arithmetic instructions, and stalls waiting for the next load. Tensor cores sit dark. Adding more FLOPs to the kernel does not help; only raising AI does. The productive levers are: fuse kernels to eliminate intermediate HBM round-trips, lower numerical precision to cut bytes per element, and restructure data access to maximise reuse within SRAM.

Above the ridge point the opposite holds. Instruction throughput is the constraint. The levers are tile-size tuning, register allocation, and ensuring the tensor cores have a continuous stream of work with no pipeline bubbles. Reducing bytes transferred has no effect if compute is the bottleneck.

A Worked Example

Consider two kernels operating on tensors stored in BF16 (2 bytes per element).

ReLU on a vector of $n$ elements:

• Read $n$ elements: $2 n$ bytes

• One comparison per element: $n$ FLOPs

• Write $n$ elements: $2 n$ bytes

• Total bytes: $4 n$, total FLOPs: $n$

• AI $= n / ( 4 n ) = 0 . 25$ FLOPs/byte This is $0 . 25 / 295 \approx 0 . 08 \%$ of the H100 ridge point-purely memory-bound at any tensor size.

Batched matrix multiply $\mathbf{X}_{B \times D} \cdot \mathbf{W}_{D \times F}$ in BF16:

• Reads: $2 B D + 2 D F$ bytes; writes: $2 B F$ bytes

• FLOPs: $2 B D F$ (multiply-accumulate)

• AI $\approx 2 B D F / ( 2 B D + 2 D F + 2 B F ) \approx B$ (when $B \ll D , F$) For a projection layer in a 7B model with $D = 4096$, $F = 11008$:

• $B = 1$ (single decode step): AI $\approx 1$ FLOPs/byte. The H100 is working at $1 / 295 < 0 . 5 \%$ of its compute capacity; throughput is entirely determined by the $3 . 35$ TB/s HBM bandwidth.

• $B = 295$ (medium batch): AI $\approx 295$ FLOPs/byte. At the ridge point; switching from memory-bound to compute-bound.

• $B = 512$: AI $\approx 440$ FLOPs/byte. Compute-bound; tensor cores are fully utilized. This single calculation explains why serving throughput scales almost linearly with batch size up to around $B = 300$, then saturates: below $B = 295$ you are paying bandwidth; above it you are paying compute.

What You Actually Decide in Practice

• Identify your regime before optimizing. Run a profiler and check whether the kernel is bandwidth-limited or compute-limited. Spending hours tuning tile sizes on a memory-bound softmax yields nothing; spending hours fusing element-wise ops on a compute-bound matmul is equally wasted.

• Rule of thumb for inference: if $B \times T < 300$ on an H100, assume your linear layers are memory-bound. Budget latency as bytes-transferred / bandwidth, not FLOPs / peak-compute.

• Halving precision doubles AI. Going from FP32 to BF16 cuts bytes per element from 4 to 2, doubling AI for any memory-bound kernel. This is one of the highest-leverage single changes for inference performance and costs almost nothing in quality for the weight and activation tensors that dominate bandwidth.

• Every HBM round-trip is expensive. An elementwise activation written to HBM and read back for the next layer pays the full 295 FLOPs/byte tax twice for zero additional compute. Fuse it. PyTorch’s torch.compile and custom Triton kernels both eliminate this by keeping intermediate results in registers or shared memory.

• Attention during decode is the worst-case scenario. For a single decode step, the attention kernel must load the entire KV cache ($2 \times L \times T_{\text{past}} \times N_{\text{kv}} \times H$ bytes) from HBM and do $O ( T_{\text{past}} )$ dot products per head. The AI for this is essentially 1 regardless of batch size-batch size does not help because each sequence has its own KV cache. This is the structural reason GQA, MLA, and local attention reduce decode latency: they shrink the bytes, not the FLOPs.

How This Shows Up in Interviews

• “Why is autoregressive decoding slow compared to prefill?” Strong answer: prefill processes all tokens in parallel; AI $= O ( T )$ for attention and $\approx B$ for projections, so it is compute-bound at typical sequence lengths. Decode generates one token at a time; AI $\approx 1$ for projections and AI $\approx 1$ for attention regardless of sequence length. Throughput is capped by HBM bandwidth, not tensor cores.

• “If the next GPU generation doubles peak FLOPs but keeps memory bandwidth the same, does token generation speed improve?” Strong answer: no. Decode is already memory-bound; the ridge point doubles (from 295 to 590 FLOPs/byte), making the regime gap even worse. The only improvement comes from raising AI-larger batch sizes, fewer KV-cache bytes per token (GQA, quantization), or architectures with constant-size state.

• Senior follow-up: “How would you use the roofline model to evaluate a proposed architectural changesay, replacing the FFN with a depth-wise convolution?” Sketch the answer: estimate AI for the proposed op (depth-wise conv has AI $\approx$ kernel_size$/ 2$ at best, typically $\ll 295$), compare to the matmul it replaces (AI $\approx B$, compute-bound at $B > 300$). If the replacement is memory-bound where the original was compute-bound, you likely made the architecture slower despite reducing FLOPs-a classic efficiency trap that roofline analysis catches immediately.

Key Takeaways

• AI $=$ FLOPs$/$bytes; ridge point $=$ peak compute$/$peak bandwidth $\approx 295$ FLOPs/byte on H100 BF16.

• Memory-bound kernels benefit from fusion, lower precision, and reduced data movement; compute-bound kernels benefit from better tiling and occupancy.

• Matrix multiply AI $\approx B$ (batch size); single-stream decode ($B = 1$) has AI $\approx 1$, using less than $0 . 5 \%$ of the GPU’s compute capacity.

• The ridge point rises with every GPU generation because compute scales faster than bandwidth-memory-bound problems stay memory-bound and architectural changes are the only durable fix.

• Attention during decode has AI $\approx 1$ independent of batch size, making KV-cache size reduction (GQA, MLA, quantization) the primary lever for decode throughput.

8.2.2 Why Matrix Multiplies Are “Blessed” Operations

When an interviewer asks why every production LLM architecture concentrates its computation in matrix multiplications, the answer is not “GPUs are fast at matmuls.” The precise answer is that matrix multiplies are the only common deep-learning operation whose arithmetic intensity grows with problem size. Every other operation you might add-elementwise activations, layer norms, softmax, attention masking-has AI that is $O ( 1 )$ or $O ( 1 / N )$: it is memory-bound at any scale and stays memory-bound no matter how large you make the tensor. Matmuls have AI $= O ( N )$. This geometric fact, combined with dedicated hardware that delivers 10-14$\times$ the FLOPs/s of general ALUs, means non-matmul FLOPs are not just slower-they are categorically slower. This topic tests at the senior level whenever architecture or kernel design is on the table.

For a square $N \times N$ matrix multiply $\mathbf{C} = \mathbf{A} \mathbf{B}$ in BF16:

• FLOPs: $2 N^{3}$ (one multiply plus one add per output element, $N^{2}$ outputs each needing a dot product of length $N$)

• Bytes transferred: $3 \times 2 N^{2} = 6 N^{2}$ (read $\mathbf{A}$, $\mathbf{B}$, write $\mathbf{C}$)

• AI $= 2 N^{3} / ( 6 N^{2} ) = N / 3$ At $N = 885$ the AI crosses the H100 ridge point of $\approx 295$ FLOPs/byte. At $N = 4096$ (a typical model width), AI $\approx 1365$-more than $4 \times$ above the ridge, solidly compute-bound.

The contrast with elementwise operations is stark. A ReLU on an $N \times N$ tensor has $N^{2}$ FLOPs and $4 N^{2}$ bytes (AI $= 0 . 25$) regardless of $N$. Softmax on an $N$-element vector has $\approx 3 N$ FLOPs and $4 N$ bytes (AI $\approx 0 . 75$). These operations will never be compute-bound on any conceivable GPU because their FLOPs and bytes scale at the same rate.

Modern GPUs reinforce this asymmetry with hardware specialization. On an H100, general-purpose FP32 ALUs deliver $\approx 67$ TFLOPs/s. BF16 tensor-core throughput is $\approx 989$ TFLOPs/s-a $\mathbf{1} \mathbf{4} . \mathbf{8} \times$ gap. A kernel that cannot use tensor cores (because it performs something other than a matrix multiply-accumulate) runs at $\leq 1 / 14$ the FLOPs/s of one that can.

Matrix multiplies are “blessed” because they are the only common DL operation where making the problem larger raises the GPU’s efficiency, and the only one with dedicated hardware that delivers $14 \times$ higher FLOPs/s than every other operation class.

The geometric intuition is data reuse. An $N \times N$ matmul has $O ( N^{2} )$ input data but $O ( N^{3} )$ work. If you load a tile of $\mathbf{A}$ and a tile of $\mathbf{B}$ into fast shared memory, you can perform all $O ( N )$ dot products that involve that tile before evicting it-reusing each loaded byte $O ( N )$ times. The number of HBM loads grows as $O ( N^{2} )$ while the FLOPs grow as $O ( N^{3} )$, so the ratio AI $= O ( N )$.

No other operation has this structure. Layer normalization computes a mean and variance over each row- $O ( N )$ work on $O ( N )$ data, AI $= O ( 1 )$. Attention score accumulation (the softmax and weighted sum) has $O ( T^{2} )$ FLOPs and $O ( T^{2} )$ bytes for sequence length $T$, so AI $= O ( 1 )$. Depth-wise convolution with kernel size $k$ achieves AI $\approx k / 2$-bounded by a small constant.

The implication for architecture design is direct: if you add a non-matmul component that processes the same activations as a matmul, you are adding a memory-bound bottleneck alongside a compute-bound one. The new component runs at $1 / 14$ the throughput per FLOPs, and if it touches the same bytes as the matmul it “replaces,” you have paid the bandwidth cost twice.

The performance gap between matmul-dominant and non-matmul-dominant architectures is empirically large. The same compute budget spent on a matmul-based FFN versus a recurrent or convolutional alternative generally yields worse wall-clock performance for the alternative, even when the alternative has fewer theoretical FLOPs-because those FLOPs are running at $1 / 14$ the hardware speed.

A Worked Example

Consider the gate projection in a SwiGLU FFN layer of a 7 B model: $D = 4096$, $F = 11008$, BF16 storage.

Weight matrix $\mathbf{W} \in \mathbb{R}^{4096 \times 11008}$: $4096 \times 11008 \times 2$ bytes $= 90 . 2$ MB.

$B = 1$ (single decode token):

• FLOPs: $2 \times 1 \times 4096 \times 11008 = 90 . 2$ GFLOPs

• Bytes: $90 . 2$ MB (weights) $+ 2 \times 4096$ (input) $+ 2 \times 11008$ (output) $\approx 90 . 2$ MB

• AI $\approx 90 . 2$ GFLOPs$/ 90 . 2$ MB $= 1$ FLOPs/byte

• Throughput: $3 . 35$ TB/s $\times 1 = 3 . 35$ TFLOPs/s attainable vs. $989$ TFLOPs/s peak-$0 . 34 \%$ utilization

$B = 512$ (batched inference):

• FLOPs: $512 \times 90 . 2$ GFLOPs $= 46 . 2$ TFLOPs

• Bytes: $90 . 2$ MB (weights, loaded once and reused for all 512 tokens) $+ 2 \times 512 \times 4096 \approx 4 . 2$ MB (inputs) $+ \approx 11 . 3$ MB (outputs) $\approx 105 . 7$ MB

• AI $= 46200$ GFLOPs$/ 105 . 7$ MB $\approx 437$ FLOPs/byte

• Attainable: $\min ( 3 . 35 \times 437 , 989 ) = \min ( 1464 , 989 ) = 989$ TFLOPs/s-fully compute-bound

The $512 \times$ increase in batch size did not increase the weight bytes at all-the same 90.2 MB is broadcast to all 512 tokens. That data reuse is precisely why batching improves hardware utilization for linear layers and why serving systems target high request concurrency.

Now compare to a standalone SwiGLU activation applied to $B = 512$, $F = 11008$ elements:

• FLOPs: $\approx 5 \times 512 \times 11008 \approx 28 . 2$ GFLOPs (multiply, sigmoid, gate product, etc.)

• Bytes: $3 \times 2 \times 512 \times 11008 \approx 33 . 9$ MB (read two inputs, write one output)

• AI $\approx 28200 / 33900 \approx 0 . 83$ FLOPs/byte-memory-bound at any batch size This activation runs in $33 . 9$ MB$/ 3 . 35$ TB/s $\approx 10 \mu$s when fused into the projection kernel. As a standalone kernel it incurs an additional HBM round-trip and takes twice as long for the same bytes. Fusing it is not optional at production scale.

What You Actually Decide in Practice

• Target $> 80 \%$ of FLOPs in matmuls. If your architecture allocates more than $\sim 20 \%$ of FLOPs to non-matmul operations, profile immediately. Those operations run at $\leq 1 / 14$ the hardware speed and will dominate latency despite their minority share of theoretical FLOPs.

• Use dimensions that are multiples of 128 (minimum 64). Matmul tile sizes on H100 are typically $128 \times 128$ or $256 \times 128$. A matrix dimension of 4097 requires the same number of tiles as 4224 but leaves the edge tile almost empty. Padding the vocabulary from 50,257 to 50,304 (nearest multiple of 64) is documented to yield a 25% throughput gain-the tile alignment effect is not academic.

• $B = 295$ is the practical break-even for linear layers on H100 BF16. Below that, optimize for bandwidth (GQA, quantization, fewer HBM reads). Above that, optimize for compute (better tile sizes, higher occupancy). Do not conflate the two regimes.

• Fuse non-matmul work into adjacent matmul kernels. Activations, scales, biases, and residual adds should never be standalone kernels in a throughput-sensitive path. torch.compile handles this automatically for simple cases; Triton is the right tool when you need custom fusion boundaries.

• FlashAttention improves attention throughput by reducing bytes, not FLOPs. The attention score matrix ($T \times T$ in FP32) is never materialized in HBM; instead the softmax is computed tile-by-tile in SRAM. Attention FLOPs are $O ( T^{2} )$ in both cases-FlashAttention moves attention from bandwidth-limited to closer to compute-limited by eliminating most HBM traffic.

How This Shows Up in Interviews

• “Why should most FLOPs in an LLM be in matrix multiplications?” Strong answer: matmul AI $= N / 3$ for square $N \times N$, growing linearly with size and crossing the H100 ridge point at $N \approx 885$. All other common operations have constant AI regardless of size, so they are permanently memory-bound. Additionally, tensor cores deliver $14 \times$ more FLOPs/s than general ALUs-non-matmul operations cannot access this hardware.

• “An engineer proposes replacing one FFN layer with a bank of depth-wise convolutions to reduce FLOPs. What’s your concern?” Strong answer: depth-wise conv has AI $\approx k / 2$ where $k$ is the kernel size (typically $3$ or $5$), which is well below the ridge point. The replacement is memory-bound while the matmul it replaces was compute-bound at batch size $> 300$. Even if it has half the FLOPs, it may well be slower, because those FLOPs run at $1 / 14$ the speed. Profile first; the roofline model predicts the answer before running any code.

• Senior follow-up: “Why does nanoGPT’s throughput increase by 25% after increasing vocab size from 50,257 to 50,304?” Sketch the answer: the vocabulary embedding lookup and the final unembedding projection are both matmuls with one dimension equal to the vocab size. When that dimension is not a multiple of the tile size (64 or 128), the last tile is sparsely populated, meaning the assigned SM executes a full tile’s worth of instructions for a fraction of useful work. Padding to the nearest multiple of 64 eliminates this waste. At the scale of an unembedding matmul ($D_{\text{model}} \times V$ ), even a $0 . 1 \%$ alignment improvement propagates into the entire forward and backward pass because this is the largest single matmul in the model.

Key Takeaways

• Matmul AI $= N / 3$ for $N \times N$ matrices, scaling linearly with $N$; all other common LLM operations have AI $= O ( 1 )$ and remain memory-bound at any size.

• Tensor cores deliver $\approx 14 \times$ more FLOPs/s than general ALUs on H100; operations that cannot use tensor cores pay this penalty unconditionally.

• $B \approx 295$ is the batch-size break-even for single linear layers on H100 BF16; below it, optimize memory movement; above it, optimize compute utilization.

• Keep matmul dimensions at multiples of 128 (minimum 64) to avoid tile-boundary underutilization; misalignment alone can cost 25%+ of throughput.

• FlashAttention is a memory-movement optimization, not a FLOPs reduction; it moves attention toward the ridge point by eliminating HBM traffic, not by changing the operation count.