LLM Basics 7 — The Hardware Fit

Series: 1 · What a model is on disk → 2 · Tokens and embeddings → 3 · A layer is matrix multiplication → 4 · The forward pass → 5 · Generating text → 6 · Attention and the KV cache → 7 · The hardware fit → Capstone · Why-LLM-Decode-Is-Memory-Bound

Three earlier lessons ended at the same place.

• Lesson 3: multiplying a matrix by one token’s vector was paced by how fast the matrix could be read.
• Lesson 4: generating text ran at 33 tokens/s, while reading a prompt ran at 2,400.
• Lesson 6: a longer context slowed generation, because the KV cache was more bytes to read.

This lesson opens the machine to see why. It covers what a GPU is, where its memory sits, and what it measures out to on the DGX Spark. Every figure marked observed comes from a small CUDA C benchmark run on the Spark.

CPU and GPU: two answers to waiting for memory

Main memory is slow compared with arithmetic. A processor can do hundreds of operations in the time one value arrives from memory. CPUs and GPUs handle that wait in opposite ways.

The GPU’s approach is the one an operating system uses for disk I/O, applied at every clock cycle. When a process blocks on a read, the OS runs another process instead of waiting. A GPU does this in hardware, between groups of threads, and the switch costs nothing: every resident thread’s registers stay loaded in the chip at all times, so there is no state to save and restore.

That only works when there are thousands of independent pieces of work, and the matvecs that make up 99.97% of an LLM (lesson 3) are exactly that. Each of the 11,008 outputs of ffn_gate is its own dot product.

Inside the GPU

The streaming multiprocessor

A streaming multiprocessor (SM) is the GPU’s unit of execution, roughly what a core is on a CPU. The Spark’s GB10 has 48. Observed, from the GPU’s own device query, each SM has:

• CUDA cores: simple arithmetic units that each do one multiply-add at a time.
• Tensor cores: units that multiply a small matrix tile — for example 16 × 16 — in a single operation. A tensor is an array of numbers with any number of dimensions; a matrix is a 2-D one.
• A register file of 256 KB (65,536 registers of 32 bits). This is where every resident thread keeps its working values, and the reason switching threads is free.
• 100 KB of shared memory: a small, very fast scratchpad that the threads of one group can share.
• Up to 1,536 resident threads.

Kernels, threads, and warps

A program that runs on the GPU is a kernel: a C function that thousands of threads all run at once, each working on a different part of the data. Here is the benchmark’s memory-reading kernel, complete:

__global__ void read_kernel(const float4 *in, size_t n, float *out) {
    float acc = 0.f;
    // each thread starts at its own index, then strides by the total number of threads
    for (size_t i = blockIdx.x * blockDim.x + threadIdx.x; i < n; i += gridDim.x * blockDim.x) {
        float4 v = in[i];                 // read 16 bytes
        acc += v.x + v.y + v.z + v.w;
    }
    if (acc == 123.456f) out[0] = acc;    // only so the compiler can't skip the reads
}

• It’s ordinary C with one difference: blockIdx, blockDim, threadIdx, and gridDim tell each thread which thread it is, so it can pick its own slice of the data.
• Threads run in groups of 32 called warps. All 32 threads of a warp execute the same instruction at the same time, on different data. When a warp stalls on memory, the SM issues instructions from another warp that is ready.

The memory hierarchy

Each level down is larger and slower than the one above it. Observed sizes on the Spark:

The problem in one comparison: the model is about 7 GB, and the largest cache is 24 MB — about 300 times smaller. Nothing from one token’s pass is still in cache for the next, so every generated token streams the whole model out of main memory again.

On the Spark, the CPU and GPU share that main memory. The device query reports the GPU as integrated. A discrete GPU, like an H100, has its own memory on its own board: much faster (HBM, about 3.35 TB/s on an H100) but much smaller (80 GB), with the CPU’s memory across a link.

What the Spark measures

1. How fast the GPU can read memory

Observed: the kernel above read 2.15 GB in 8.29 ms — 259 GB/s, 95% of the Spark’s specified 273 GB/s.

For comparison, lesson 3’s matvec on the Spark’s CPU reached about 125 GB/s, from the same memory chips. The GPU gets twice as much out of them, which fits its design: it keeps far more reads in flight at once. That is an explanation, not something this benchmark isolates.

2. Tensor cores versus plain arithmetic

Observed: an 8192 × 8192 × 8192 matrix multiply — big enough that memory is not the limit:

• A TFLOP/s is a trillion floating-point operations per second. A FLOP is one multiply or one add.
• BF16 is a 16-bit (2-byte) floating-point format with the same range as 32-bit FP32 but less precision. It is the usual format for LLM weights.
• Tensor cores did the same work 4.6× faster. They also handle half-size numbers, which halves the bytes to move.

3. One weight matrix times N tokens

This is lesson 3’s matvec-versus-matmul experiment, on the GPU, at real scale. The matrix has ffn_gate’s shape, 11008 × 4096 in BF16 (90 MB — far larger than the L2 cache). It is multiplied by N token vectors at once, for N from 1 to 4,096. Observed:

Read it in two halves.

• From 1 to 128 tokens, the time barely moves — 0.37 ms to 0.42 ms for 128 times the work.
  • The GPU is reading the matrix at nearly full memory speed, and the arithmetic fits in the gaps.
  • This is memory-bound: the step time is set by bytes moved.
  • At N = 1 the GPU uses 0.24 of its roughly 80 TFLOP/s: 0.3%.
• From about 512 tokens on, the time grows in step with N, and throughput levels off near 76 TFLOP/s.
  • The tensor cores are now the bottleneck.
  • This is compute-bound: the step time is set by arithmetic.
  • From here on, per-token cost stops falling.

Where the two halves meet

The crossover can be predicted from two numbers.

• Arithmetic intensity is a job’s FLOPs per byte moved from memory.
  • One token through a BF16 matrix does 2 FLOPs (a multiply and an add) per 2-byte weight: 1 FLOP per byte.
  • N tokens sharing each weight do N FLOPs per byte.
• The machine’s balance point is its peak compute divided by its memory bandwidth: about 80 TFLOP/s ÷ 259 GB/s ≈ 300 FLOPs per byte.
  • A job below that is memory-bound; a job above it is compute-bound.
  • This is called the ridge point, from the roofline chart engineers use to plot it.

So the crossover should sit near N ≈ 300. The measurement puts it between 256 and 512.

Putting the course’s measurements together

• Decode (lessons 4–5) is N = 1 through every matrix of the model, every token.
  • About 7 GB is read per token. At the measured 259 GB/s, that sets a ceiling near 37 tokens/s.
  • llama.cpp measured 33.2–33.7 tokens/s, about 90% of that ceiling.
  • Meanwhile, the tensor cores sit almost entirely idle.
• Prefill (lesson 4) is N = 1,921: the whole prompt shares each weight read, which puts it past the ridge point.
  • The 2,400 tokens/s measured there works out to about 32 TFLOP/s of useful work: 2,400 × 13.5 billion FLOPs per token.
  • That is compute-bound territory, though below this benchmark’s peak. llama.cpp’s kernels for its 8-bit format are not cuBLAS’s BF16 routine, and attention adds work of its own.
• The KV cache (lesson 6) adds to the bytes read per token, without adding enough arithmetic to change which side of the ridge decode sits on.

Decode leaves almost all of the GPU’s arithmetic unused. The capstone follows that conclusion through to its consequences: why batching, quantization, mixture-of-experts models, and speculative decoding are all ways of getting more tokens out of each byte read.

On a DGX Spark

Reproduce it: every number in this lesson is printed by lessons/07_hardware/bench.cu in the course’s companion code repository, llm-bytes-up-code. Build it with make and run ./bench; it needs the CUDA toolkit, which provides nvcc and cuBLAS. The run this lesson quotes is recorded in that repository’s outputs/2026-10-12/, produced from a fresh clone on the Spark.

• The machine: NVIDIA GB10, compute capability 12.1, CUDA 13.0, driver 580.
  • The GPU shares 128 GB of LPDDR5X with a 20-core Arm CPU.
  • The device query reports 130.7 GB as visible to the GPU, of which 34.4 GB was free at the time. The Spark’s resident model servers held the rest of the GPU-visible memory, but were idle.
• The benchmark is about 150 lines of CUDA C.
  • The bandwidth test is the hand-written kernel shown above, reading a 2 GiB buffer 20 times.
  • The matrix multiplies use cuBLAS (cublasGemmEx with BF16 inputs and FP32 accumulation), timed with CUDA events after warm-up runs.
  • The FP32 comparison turns off cuBLAS’s TF32 shortcut (CUBLAS_PEDANTIC_MATH), so it measures true FP32.
• “Peak” here means what this benchmark reached, not NVIDIA’s specification. The specified figures for GB10 are quoted for lower-precision formats, such as FP4, and are not comparable with BF16.
• The CPU’s details (core types, caches) came from lscpu.