High Bandwidth Memory (HBM)

The FLOPS are fine. The bus is not.

You bought the top SKU. Tensor Cores light up on paper. Token generation still crawls. The profiler is not lying about compute util — there simply are not enough bytes arriving from device memory.

HBM is DRAM moved into the package: stacks of thinned dies, vertical copper vias, a silicon interposer with a bus a thousand lanes wide. DDR is a highway to a distant warehouse. HBM builds the warehouse next to the factory.

This page is why that packaging exists, how the tower is wired, and when your kernel actually hits the wall.

Two numbers

1. ~106× vs ~11×
From P100 (2016) to B200 (2024), dense FP16 compute grew roughly two orders of magnitude. HBM bandwidth grew about one — and that is with stacked DRAM. The gap is the product problem.

2. ~1 FLOP/byte
LLM decode at batch 1: each token reads the weights once and does almost nothing with them. On an H100-class ridge near ~300 FLOP/B, that is well under 1% of peak compute. Inference speed is HBM bandwidth.

The wall keeps growing

Even with HBM, FLOPS outrun bytes. Scrub generations and watch the gap sticky climb.

Three levers — HBM is width

More channels. Faster pins. Wider on-package bus. Flip the lever; the toy bus changes shape.

A thousand slow, short wires beat sixty-four fast, long ones when the job is terabytes per second next to a 700 W die.

What’s in the package

Stacks beside the GPU on an interposer — or GDDR chips out on the board. Click bays for the role of each layer.

How the tower wires itself

Through-silicon vias: etch, insulate, copper-fill, thin to ~50 µm, micro-bump, stack. One bad via can scrap the tower — a large part of HBM’s price.

Generations (per stack, peaks)

Device bandwidth = stacks × bin rate. H100 SXM ~5× HBM3 → ~3.35 TB/s; B200 ~8× HBM3E → ~8 TB/s — both under raw JEDEC stack peaks. Read the datasheet, not only this table.

HBM is the bottom of the on-package ladder

Registers → shared/L1 → L2 → HBM. Everything above exists to avoid paying ~500 cycles and the interposer trip.

Programmer rule: every byte from HBM should be reused as many times as possible from registers and shared memory before the next global load. Coalesce warps; tile through shared. (See Streaming Multiprocessor for warp occupancy and coalescing instruments.)

Roofline: is HBM your limit?

Arithmetic intensity (FLOP per byte from memory) places you under a roof: slanted memory limit, flat compute limit. The corner is the ridge.

Try LLM decode vs big GEMM on H100. Decode sits left of the ridge — buying more TFLOPS without bandwidth does nothing. GEMM can sit right if shapes and Tensor Core paths cooperate.

Programming habits (not a different API)

cudaMalloc already gave you HBM. Using it well is still:

1. Coalesce — consecutive addresses per warp → full-width transactions.
2. Tile and reuse — stage in shared memory; raise intensity toward the ridge.
3. Profile the bound — Nsight Compute: memory throughput, achieved occupancy, not just “kernel finished.”

// Teaching sketch: coalesced tile load + reuse from shared (production: cuBLAS / CUTLASS)
__global__ void tiled_gemm(const float* A, const float* B, float* C, int M, int N, int K) {
    const int TILE = 32;
    __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 sum = 0.f;
    for (int t = 0; t < K; t += TILE) {
        As[threadIdx.y][threadIdx.x] =
            (row < M && t + threadIdx.x < K) ? A[row * K + t + threadIdx.x] : 0.f;
        Bs[threadIdx.y][threadIdx.x] =
            (col < N && t + threadIdx.y < K) ? B[(t + threadIdx.y) * N + col] : 0.f;
        __syncthreads();
        #pragma unroll
        for (int k = 0; k < TILE; k++)
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        __syncthreads();
    }
    if (row < M && col < N) C[row * N + col] = sum;
}

Heat and cost (the real stack limits)

Heat. Eight to sixteen dies share one escape path into the same cooler as a 700 W+ GPU. DRAM leakage rises with temperature → more refresh → less useful bandwidth. Stacks throttle in the ~95 °C class; thermal design bounds height as much as process does.

Money. TSVs, interposers, and whole-tower yield make HBM far more expensive per GB than DDR or GDDR. The economics only close when bandwidth is the product — training and serving large models.

Prices are approximate and move with AI demand.

What to do

1. Measure intensity — is the kernel left or right of the ridge on your GPU?
2. Raise reuse before buying a bigger SKU — tiling, fusion, FlashAttention-style fewer HBM trips.
3. Coalesce and vectorize global loads; fix strided layouts first.
4. Size for occupancy when memory-bound — spare warps hide HBM latency (SM page).
5. Believe the datasheet BW and the profiler, not marketing TFLOPS alone, for decode-heavy work.
6. Budget multi-GPU traffic separately — HBM is top of the cliff; NVLink and Ethernet are further drops (multi-GPU).

Further reading

• JEDEC HBM4 (JESD270-4) announcement
• What Every Programmer Should Know About Memory — Drepper; DRAM fundamentals under every stack
• NVIDIA Hopper Architecture In-Depth — HBM3 integration on a flagship GPU

Related concepts