Almost every performance problem in deep learning is a memory problem because GPU can do math far faster than it can fetch the numbers to do math on.
You have to learn thinking in bytes, not FLOPs.
In this part, I will give the model that explains why and the vocabulary to debug memory problems on a typical GPU or a cluster.
This is the Part II in our GPU programming series, you can read the Part I here:
GPU Programming P1: Why is this kernel only hitting 30% of peak FLOPs?
A modern server CPU and a datacenter GPU are both made of transistors on similar process nodes, but they spend those…
agentnativedev.medium.com
The hierarchy, with real bandwidths
A GPU’s memory is a hierarchy of tiers that trade capacity for speed.
As you move outward from the arithmetic units, each tier is roughly an order of magnitude larger and an order of magnitude slower.
The whole game of GPU optimization is keeping the data you’re actively working on in the fast inner tiers and minimizing trips to the slow outer ones.
Let’s have a look at the most important tiers here:
- Shared memory is a slice of the on-chip SRAM that the programmer controls explicitly, this is visible to every thread in a block. Unlike a cache, you decide what goes in it and when. The standard optimization pattern is cooperatively loading a tile of data from slow HBM into fast shared memory once, then have all the threads in the block reuse it many times, amortizing the expensive HBM read. We’ll see this in the following series for tiled matrix multiply.
- L2 cache is shared across all SMs and is the last line of defense before HBM. On H100 it is a large 50 MB, which can hold meaningful working sets and matters for inference where the same weights are read repeatedly.
- HBM (High-Bandwidth Memory) is the GPU’s main memory, the 80GB or 141 GB you see on the datasheet. It is stacked DRAM connected by an extremely wide bus, which is how it reaches multi-terabyte-per-second bandwidth that ordinary DDR cannot. But multi-terabyte-per-second is still glacial next to the arithmetic units, and that gap is the whole story of this part.
Coalescing: the first rule of HBM
HBM is read in transactions, which are chunks of contiguous bytes (typically 32, 64 or 128 bytes).
When the 32 threads of a warp issue a load, the hardware coalesces their addresses into as few transactions as possible.
If those 32 threads access 32 consecutive 4-byte floats, that is one clean 128-byte transaction, which is fully coalesced and peak bandwidth.
But if those same 32 threads access scattered addresses, the hardware must issue many separate transactions, most of whose bytes are thrown away, and the effective bandwidth collapses to a fraction of peak.
Let’s have a look at an example:
// COALESCED: thread i reads element i. Consecutive threads -> consecutive // addresses -> the warp's 32 reads fuse into one 128-byte HBM transaction. global void copy_coalesced(const float* in, float* out, int n) { int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) out[i] = in[i]; // stride-1: perfect }
// STRIDED: thread i reads element istride. With stride=32, each thread lands // in a different 128-byte line -> up to 32 transactions for one warp. // Effective bandwidth can drop ~10-30x for large strides. global void copy_strided(const float in, float* out, int n, int stride) { int i = (blockIdx.x * blockDim.x + threadIdx.x) * stride; if (i < n) out[i] = in[i]; // scattered: pathological }
This is why data layout is a first-class performance concern.
Storing a batch of vectors so that the threads of a warp naturally walk contiguous memory (structure-of-arrays rather than array-of-structures) can be the difference between 80% and 8% of peak bandwidth without changing a single arithmetic operation.
Bank conflicts: the first rule of shared memory
Shared memory has its own performance trap.
It is divided into 32 banks, one per warp lane, so that 32 threads can each access a different bank simultaneously at full speed.
But if two threads in a warp access different addresses that fall in the same bank, those accesses serialize, a two-way bank conflict halves throughput, an N-way conflict cuts it to 1/N.
The classic cause is a stride that is a multiple of 32 (so column-major access of a 32-wide tile hits one bank).
The classic fix is padding: declare the tile one element wider than needed so the access pattern skews across banks.
// BAD: a 32x32 tile. Column access tile[threadIdx.y][k] for fixed k hits the // same bank for every row -> 32-way conflict. shared float tile_bad[32][32];
// GOOD: pad the inner dimension by 1. Now consecutive rows are offset by 33, // which is co-prime with 32, so column accesses spread across all 32 banks. shared float tile_good[32][33]; // <-- the "+1 padding" trick
Coalescing and bank conflicts are about different memories.
Coalescing is an HBM/global-memory concern (consecutive threads to consecutive addresses).
Bank conflicts are a shared-memory concern (avoid 32 threads hitting one of 32 banks).
The canonical optimization: tiled matrix multiply
Matrix multiply is the heart of deep learning, and the tiled-matmul kernel is the teaching example for why shared memory exists.
A naïve matmul reads each element of A and B from HBM many times, for an N×N×N multiply, O(N³) HBM reads.
The tiled version loads small tiles of A and B into shared memory once, then every thread in the block reuses those tiles for many multiply-accumulates before fetching the next tile.
The HBM traffic drops by a factor equal to the tile width, converting a memory-bound kernel into a compute-bound one.
In production you use cuBLAS or CUTLASS, which do this and far more but you must understand the idea.
#define TILE 16
// C = A * B, all NxN, row-major. Each block computes one TILExTILE tile of C.
__global__ void matmul_tiled(const float* A, const float* B, float* C, int N) {
__shared__ float As[TILE][TILE];
__shared__ float Bs[TILE][TILE];
int row = blockIdx.y * TILE + threadIdx.y;
int col = blockIdx.x * TILE + threadIdx.x;
float acc = 0.0f;
// March across the K dimension one tile at a time.
for (int t = 0; t < N / TILE; ++t) {
// Cooperative load: each thread brings ONE element of A and B from HBM
// into shared memory. Coalesced because threadIdx.x is the fast index.
As[threadIdx.y][threadIdx.x] = A[row * N + (t * TILE + threadIdx.x)];
Bs[threadIdx.y][threadIdx.x] = B[(t * TILE + threadIdx.y) * N + col];
__syncthreads(); // all loads done before compute
// Now do TILE multiply-accumulates entirely from FAST shared memory.
// Each HBM value loaded above is reused TILE times here.
#pragma unroll
for (int k = 0; k < TILE; ++k)
acc += As[threadIdx.y][k] * Bs[k][threadIdx.x];
__syncthreads(); // before overwriting the tiles
}
if (row < N && col < N) C[row * N + col] = acc;
}The arithmetic is identical to the naïve version, the only change is where the data lives during reuse.
Arithmetic intensity & the roofline model
Now we make compute-bound vs memory-bound quantitative.
We define the arithmetic intensity (AI) of a computation as the number of floating-point operations it performs per byte it moves from memory:
arithmetic intensity = FLOPs performed / bytes moved [FLOP/byte]
A computation with high AI does a lot of math per byte, i.e., it can keep the ALUs fed. A computation with low AI is starved, it spends its time waiting on memory.
Every GPU has a peak compute rate (FLOP/s) and a peak memory bandwidth (bytes/s).
Their ratio defines a threshold AI called the ridge point.
Below it, you are memory-bound, and above it, compute-bound.
Plotting achievable performance against AI gives the roofline: a diagonal memory roof (bandwidth × AI) that rises until it hits the flat compute roof (peak FLOP/s).
An H100’s ridge point is ~295 FLOP/byte (peak compute / peak HBM bandwidth). Anything to the left of the ridge is limited by memory bandwidth no matter how fast the math units are. LLM decoding at batch=1 sits at AI≈6, which is ~40x to the left of the ridge, hopelessly memory-bound. Kernel fusion (FlashAttention) and larger batches push workloads rightward toward the compute roof. A100’s ridge is lower, ~156.
The roofline is the most useful single diagram in this series because it dictates strategy:
- If a kernel is memory-bound (left of the ridge), buying more FLOPs does nothing, you must move fewer bytes (fusion, lower precision, better reuse) or move them faster (more bandwidth, i.e. H200).
- If it is compute-bound (right of the ridge), more or faster math units (Blackwell, FP8/FP4) is exactly the lever.
When your GPUs sit at low utilization, the very first question is which side of the ridge your workload is on.
When you ask “Inference is slow, should I buy faster GPUs?”, the roofline answer is “Token generation is memory-bandwidth-bound, at batch size 1, you read the entire model from HBM for every single token, doing almost no math per byte. A higher-FLOPS GPU barely helps and what helps is higher bandwidth (H200 over H100), bigger batches to amortize the weight reads across more sequences, quantization to move fewer bytes, or a serving stack like vLLM that batches continuously. You should measure your tokens/sec and batch size before you talk hardware.”
FlashAttention: fusion as a bandwidth play
FlashAttention is the most cited example of raising arithmetic intensity by not touching HBM.
Standard attention computes softmax(QKᵀ)V by materializing the full N×N score matrix in HBM, where N is the sequence length.
For long sequences that matrix is enormous, and writing it out and reading it back dominates the runtime, the operation is memory-bound, with traffic scaling as O(N²).
FlashAttention restructures the computation to tile Q, K and V into blocks that fit in shared memory and computes the softmax incrementally (an online-softmax trick that keeps running max/sum statistics) so the full score matrix never has to exist in HBM at all.
The result is roughly an order of magnitude fewer HBM accesses, turning a memory-bound kernel into a near-compute-bound one, the same tiling idea we see before applied to attention.
A huge fraction of real-world speedups come not from faster silicon but from not moving data you don’t have to move.
Sizing model memory
How much GPU memory does it take to train an N-parameter model?
There is a clean back-of-envelope answer because training memory has four components:
So a 7B model needs ~112 GB of state, already more than one 80 GB H100, before a single activation.
A 70B model needs ~1.1–1.2 TB; a 405B model ~6.5 TB.This is the arithmetic that forces the sharding and parallelism, and it is the reason distributed training exists.
On top of this base sits activation memory, which scales with
batch size x sequence length x hidden size x number of layers
and for large batches and long context, can rival or exceed the weight memory, which is why activation checkpointing (recomputing activations in the backward pass instead of storing them) is a standard lever.
def training_memory_gb(num_params_billions, bytes_per_param=16):
"""Base state for mixed-precision Adam training, excluding activations."""
return num_params_billions * 1e9 * bytes_per_param / 1e9 # -> GB
for n in [7, 13, 70, 405]:
gb = training_memory_gb(n)
h100s = -(-gb // 80) # ceil divide by 80GB
print(f"{n:>4}B model: ~{gb:6.0f} GB state -> >= {h100s:.0f}x H100 just to hold it")# 7B model: ~ 112 GB state -> >= 2x H100 just to hold it
# 13B model: ~ 208 GB state -> >= 3x H100 just to hold it
# 70B model: ~ 1120 GB state -> >= 14x H100 just to hold it
# 405B model: ~ 6480 GB state -> >= 81x H100 just to hold itInference memory is much smaller, roughly 2 bytes/param for BF16 weights (or 1 for INT8/FP8) plus the KV cache, which grows with (batch size x sequence length) and often becomes the binding constraint for long-context serving, which we will cover in following parts.