Knowledge 00 — Hardware & Accelerator Foundations

The one idea that governs everything below: a modern AI accelerator can do arithmetic far faster than it can move data. Almost every performance problem in inference is, at bottom, a data-movement problem. If you internalize the roofline model and the memory hierarchy, you can reason about any chip — NVIDIA, AMD, Qualcomm, TPU, Groq — without memorizing spec sheets.

This module assumes zero prior hardware knowledge and ends at the level where you can derive whether a given operation is compute-bound or memory-bound, estimate its runtime, and explain why on a whiteboard.


Table of Contents


1. Why a CPU is the wrong tool, and what a GPU changed

A CPU is a latency-optimized machine. It has a handful of very powerful cores, each with deep pipelines, branch prediction, out-of-order execution, and large caches — all designed to finish one thread's work as fast as possible. That is ideal for code full of branches and dependencies (a web server, a compiler) and terrible for the one thing neural networks are made of: enormous, regular, independent arithmetic (matrix multiplies).

A GPU is a throughput-optimized machine. It trades single-thread speed for thousands of simple arithmetic units running in lockstep. It does not care about finishing one operation quickly; it cares about finishing millions per second in aggregate. A matrix multiply of two 4096×4096 matrices is ~137 billion floating-point operations with no branches and total independence between output elements — the perfect GPU workload.

The deep reason GPUs won AI: neural network inference is ~90% matrix multiplication and the rest is elementwise/reduction ops, all of which are massively parallel. The hardware shape (many weak lanes) matches the problem shape (many independent multiply-adds).

Principal framing: "CPU optimizes time-to-finish-one-thing; GPU optimizes things-finished-per-second. Inference is a things-per-second problem, so we use throughput machines and our enemy becomes keeping them fed with data."


2. The anatomy of a GPU (NVIDIA as the reference)

We use NVIDIA's vocabulary because it's the lingua franca; §10 maps it to other vendors.

A GPU is a hierarchy:

GPU (e.g. H100)
 ├── ~132 Streaming Multiprocessors (SMs)         ← the "cores" that actually compute
 │     ├── CUDA cores (FP32/INT ALUs)             ← scalar/vector math
 │     ├── Tensor Cores (matrix-multiply units)   ← the AI workhorse, does D = A·B + C in one op
 │     ├── Register file (huge, ~256 KB/SM)       ← fastest storage, per-thread
 │     ├── Shared memory / L1 (~228 KB/SM)        ← software-managed scratchpad, per-block
 │     └── Warp schedulers                         ← issue instructions for 32-thread "warps"
 ├── L2 cache (~50 MB)                             ← shared across all SMs
 └── HBM (High Bandwidth Memory, 80 GB @ ~3.35 TB/s on H100) ← the big, "slow" main memory

Key terms:

  • SM (Streaming Multiprocessor): the unit of parallel execution. An H100 has ~132. Each runs many threads concurrently.
  • Warp: a group of 32 threads that execute the same instruction in lockstep (SIMT — Single Instruction, Multiple Threads). If threads in a warp take different branches ("warp divergence"), the warp serializes the branches — a real performance pitfall.
  • CUDA core: a scalar ALU. Does ordinary FP32/INT math.
  • Tensor Core: a dedicated unit that computes a small matrix multiply-accumulate (e.g., a 16×16 tile) in a single hardware operation. This is where ~95% of an LLM's FLOPs go. Tensor cores are the reason FP16/BF16/FP8 matter — they only reach peak throughput in those formats.

You do not program SMs directly in this role most of the time — but you must know the structure because every profiler, every "occupancy" warning, and every vendor's architecture talks in these terms.


3. The memory hierarchy — the thing that actually matters

This is the single most important section in the module. The whole game of inference optimization is moving data as little as possible and reusing it as much as possible.

LevelSize (H100-class)BandwidthLatencyScope
Registers~256 KB / SM (~33 MB total)~hundreds of TB/s~1 cycleper-thread
Shared mem / L1~228 KB / SM~tens of TB/s (SM-local SRAM)~tens of cyclesper-block
L2 cache~50 MB~tens of TB/s~hundreds of cycleswhole GPU
HBM (VRAM)80 GB~3.35 TB/s~hundreds of nswhole GPU
Host RAM (over PCIe)TBs~64 GB/s (PCIe5 x16)~µssystem
NVMe / diskTBs~GB/s~ms+system

Read the bandwidth column carefully. On-chip SRAM is ~10–30× faster than HBM, but HBM is ~1000× larger. And HBM is ~50× faster than PCIe to host. Every order of magnitude down this list is a cliff.

The consequence: an operation that reads a value from HBM, does one multiply, and writes it back is bottlenecked by HBM, not by the multiply. The tensor cores sit idle waiting for data. This is memory-bound execution, and it dominates LLM token generation.

This is why the two most important software techniques in modern inference exist:

  • FlashAttention keeps the attention computation in SRAM instead of materializing the giant N×N attention matrix in HBM — it's a memory-movement optimization, not a math one (the math is identical). (See Knowledge 01.)
  • Operator fusion combines many small ops (e.g., bias-add + GELU + dropout) into one kernel so intermediate results stay in registers/SRAM instead of round-tripping through HBM. (See Knowledge 03.)

Whiteboard line: "HBM bandwidth is the budget. Every byte you read from HBM that you didn't have to is latency you're paying for nothing. FlashAttention and fusion are both just 'touch HBM less.'"


4. FLOPs, bandwidth, and arithmetic intensity from scratch

Three quantities. Master them and the roofline falls out for free.

4.1 FLOPs of a matmul

A matrix multiply C = A·B where A is M×K and B is K×N produces an M×N result. Each output element is a dot product of length K: K multiplies and K-1 adds ≈ 2K floating-point operations. There are M·N output elements. So:

FLOPs(matmul) = 2 · M · N · K

The factor of 2 is "multiply + add." Burn this in — you will use it constantly. (Example: 4096×4096 times 4096×4096 = 2·4096³ ≈ 137 GFLOP.)

For a linear layer y = xW with input batch B tokens, input dim d_in, output dim d_out: that's a (B×d_in)·(d_in×d_out) matmul = 2·B·d_in·d_out FLOPs.

For a whole transformer forward pass, a famous and useful approximation: processing one token through a model with P parameters costs ≈ 2P FLOPs (each parameter is used in roughly one multiply-add). Training is ~6P (forward + backward + weight update is ~3× forward). Memorize: inference ≈ 2P FLOPs/token.

4.2 Bytes moved

To do that matmul, the hardware must read A, B, and write C from/to memory. In FP16 (2 bytes/element):

Bytes(matmul) ≈ 2 · (M·K + K·N + M·N)

For the LLM decode case the dominant cost is reading the weights: a P-parameter model in FP16 must stream 2P bytes of weights from HBM for every token generated (because each weight is used once per token and the matrices are too big to cache).

4.3 Arithmetic intensity (AI)

The ratio that decides everything:

Arithmetic Intensity = FLOPs performed / Bytes moved   (units: FLOP/byte)

It measures how much compute you extract per byte of memory traffic. High AI = you reuse data a lot = compute-bound = good (you're using the expensive tensor cores). Low AI = you barely reuse data = memory-bound = the tensor cores starve.

  • A big square matmul (M=N=K large) has AI ≈ N/3 → grows with size → compute-bound. Good.
  • LLM decode (batch=1, multiplying a 1×d vector by d×d weights) reuses each weight exactly once → AI ≈ 1 (in FP16, ~0.5 FLOP/byte after the factor of 2) → deeply memory-bound. This is the central fact of LLM serving.

Principal one-liner: "Decoding tokens one at a time is the lowest-arithmetic-intensity thing you can do on a GPU. That's why we batch — batching is how you raise arithmetic intensity and stop wasting the tensor cores." (Full batching math in Knowledge 01.)


5. The Roofline Model — the principal's core tool

The roofline (Williams et al., 2009) puts the two hardware limits on one chart and lets you read off the bound for any kernel.

Two hardware numbers define a chip:

  • Peak compute π (FLOP/s) — e.g., H100 ≈ ~990 TFLOP/s dense FP16 via tensor cores (numbers vary by SKU and sparsity; treat as ~1 PFLOP/s).
  • Peak memory bandwidth β (byte/s) — e.g., H100 ≈ 3.35 TB/s HBM3.

Their ratio is the ridge point (a.k.a. machine balance):

ridge point = π / β   (FLOP/byte)

For H100: ~990e12 / 3.35e12 ≈ ~295 FLOP/byte.

Now plot attainable FLOP/s (y) against arithmetic intensity (x), both log scale:

attainable FLOP/s
   ▲
 π ┤            ________________  ← compute roof (flat): you're limited by tensor cores
   │           /
   │          /  ← memory roof (slope = β): attainable = AI · β
   │         /
   │________/______________________► arithmetic intensity (FLOP/byte)
          ridge
          point (~295 for H100)

The rule:

  • AI < ridge point → you land on the sloped part → memory-bound. Your speed = AI · β. Buying more FLOPs (a faster chip's tensor cores) does nothing; you need more bandwidth or higher AI.
  • AI > ridge point → you land on the flat part → compute-bound. You're using the tensor cores well; only more FLOPs (or lower-precision math) helps.

This single chart is how you diagnose a customer's workload in 30 seconds:

  • "Their decode AI is ~1, ridge point is ~295 → they're ~300× off the compute roof → wildly memory-bound → batching, KV-cache quantization, or a higher-bandwidth chip helps; a higher-FLOPs chip won't."
  • "Their prefill on long contexts is a big matmul, AI in the hundreds → near compute-bound → FP8/INT8 and tensor-core utilization is where the wins are."

Memorize the shape and the ridge-point arithmetic, not the exact TFLOP numbers (those change yearly). The reasoning is timeless; the constants aren't.


6. Tensor cores and the numeric formats

Tensor cores only hit peak throughput in reduced precision. Knowing the formats is mandatory because quantization (Knowledge 02) is half your job.

FormatBitsExponent/MantissaRangeUse
FP32328 / 23hugereference accuracy, rarely for inference math
TF3219*8 / 10FP32 range, less precisionNVIDIA tensor-core FP32 path
FP16165 / 10±65504 (narrow!)classic inference; can overflow
BF16168 / 7FP32-like range, less precisionpreferred for stability (same exponent as FP32)
FP8 (E4M3 / E5M2)84/3 or 5/2narrowHopper/Blackwell; ~2× FP16 throughput
INT88integerneeds scale factorclassic quantization, ~2× FP16
INT4 / FP44integer / tiny floatneeds scale + careweight-only quant, ~4×

Key intuitions:

  • More exponent bits = more dynamic range (can represent very large/small values). More mantissa bits = more precision (finer steps). FP16 has narrow range (overflows on large activations); BF16 trades precision for FP32-like range, which is why training and many inference stacks prefer it.
  • Every halving of bit width roughly doubles tensor-core throughput and halves memory traffic. That's why FP8/INT8/INT4 are the throughput levers. The cost is accuracy, which §[Knowledge 02] teaches you to recover.
  • FP8 is the current frontier for inference on Hopper/Blackwell: near-FP16 quality with ~2× speed if you handle scaling (per-tensor/per-channel scales) correctly.

Trade-off framing for customers: "Going FP16→FP8 roughly doubles your throughput and halves your memory, for typically <1% task-quality loss if we calibrate. That can be the difference between needing 8 GPUs and 4."


7. Occupancy, latency hiding, and why batch size matters

A single memory load from HBM takes hundreds of cycles. A CPU would stall. A GPU instead hides latency by having many warps resident per SM: when one warp stalls waiting on memory, the scheduler instantly switches to another warp that's ready. With enough independent work in flight, the memory latency is hidden behind useful compute.

Occupancy = (active warps) / (max possible warps per SM). It's limited by register and shared-memory usage per thread (use too many registers → fewer warps fit → can't hide latency). High occupancy is necessary but not sufficient for performance.

The serving consequence: small batches → not enough parallel work → poor latency hiding → idle SMs. This is another angle on why decode wants batching. A batch of 1 token can't keep an H100's 132 SMs busy; a batch of 256 sequences decoding in parallel can. Continuous batching (see Knowledge 04) exists precisely to keep occupancy high.

Don't over-index on "maximize occupancy" — a well-tuned FlashAttention kernel may run at modest occupancy because it deliberately uses lots of shared memory. Occupancy is a means (latency hiding), not the goal (throughput).


8. The execution model: warps, kernels, launch overhead, CUDA Graphs

  • A kernel is a function that runs on the GPU, launched from the CPU ("host"). You launch a grid of thread-blocks; each block runs on one SM; each block's threads run as warps.
  • Kernel launch overhead is real: each launch costs a few microseconds of CPU↔GPU coordination. In decode, where each token may trigger dozens of tiny kernels and each kernel runs for only microseconds, launch overhead can dominate — the GPU spends more time being told what to do than doing it.
  • CUDA Graphs fix this: you record the whole sequence of kernel launches once into a "graph," then replay it with a single submission. This eliminates per-launch CPU overhead and is a major decode-latency win — vLLM, TensorRT-LLM, and others use it heavily. (When you hear "we enabled CUDA graphs and decode got 20% faster," this is why.)

Profiling tell: if nsys shows lots of small gaps between tiny kernels on the GPU timeline and the CPU is busy launching, you're launch-bound → fuse kernels and/or use CUDA Graphs. (See Knowledge 06.)


When a model doesn't fit on one chip, or you want more aggregate throughput, you split it across GPUs (see Knowledge 05). Then the interconnect bandwidth between chips becomes a roofline of its own.

LinkBandwidth (per GPU, approx)Scope
NVLink / NVSwitch~900 GB/s (H100, bidirectional aggregate)GPU↔GPU inside a node
PCIe Gen5 x16~64 GB/sGPU↔CPU, GPU↔GPU without NVLink
InfiniBand / RoCE~400 Gb/s (~50 GB/s) per portnode↔node (scale-out)

The principle: tensor parallelism inserts an all-reduce every layer, and that collective's time is bounded by the interconnect. On NVLink it's cheap; over PCIe it can erase the speedup. This is why "TP within a node over NVLink, pipeline/data parallel across nodes over InfiniBand" is the canonical topology. A principal sizes parallelism to the interconnect, not in the abstract. (Collective cost math is in Knowledge 05.)

Customer reality: a customer who bought 8 PCIe-only GPUs (no NVLink) and expects 70B-model tensor-parallel performance will be disappointed — the all-reduces choke on PCIe. Catching this in the sizing conversation is exactly the JD2 value-add.


10. The accelerator zoo: how non-NVIDIA chips differ

JD2 (Qualcomm-archetype) explicitly spans "accelerator-based" platforms. The roofline framework transfers to all of them; only the constants and software change.

Vendor / chipArchitecture ideaMemory storySoftwareThe pitch
NVIDIA H100/H200/B200SMs + tensor cores + HBMbig HBM, NVLinkCUDA/TensorRT-LLM/Tritonecosystem & peak FLOPs
AMD MI300XCDNA CUs + matrix cores192 GB HBM3 (huge)ROCm/hipBLASLt/vLLM-ROCmfit big models on fewer chips
Qualcomm Cloud AI 100/80many AI cores + on-die SRAM, low powerlarge on-chip SRAM, modest HBMqaic SDK, ONNX/qaic-exec, AIMET quantperf/watt, INT8/edge+DC
AWS Inferentia2/TrainiumNeuronCores, systolic-array matmulHBMNeuron SDK (compile from PyTorch/XLA)captive-cloud $/token
Google TPU v5e/v5psystolic MXU arrayHBM, ICI interconnectXLA / JAXdense matmul throughput, pod-scale
Groq LPUdeterministic, SRAM-only, no HBMall weights in on-chip SRAM (small per chip → many chips)GroqFlow / ONNXlowest single-stream token latency
Cerebras WSEwafer-scale, weight streaminghuge on-wafer SRAMtheir stackgiant models, weight streaming
Etched Sohu / otherstransformer-as-ASICextreme specialization bet

Two recurring architectural patterns to recognize:

  1. Systolic arrays (TPU, Inferentia, parts of others): a 2D grid of multiply-accumulate cells where data flows ("systolically") through, computing matmuls with minimal memory traffic. Extremely efficient for dense matmul, less flexible for irregular ops.
  2. SRAM-vs-HBM bets: Groq/Cerebras put everything in SRAM to crush the memory wall for latency, paying with tiny per-chip capacity (so they need many chips). NVIDIA/AMD lean on big HBM for capacity. This is the central architectural fork in the industry, and it's exactly the roofline trade-off (bandwidth vs capacity) made physical.

Why this matters for the role: when a customer says "we benchmarked on an A100, will it be faster on chip X?", the answer is never the spec sheet — it's "what's the arithmetic intensity of your workload, and does chip X's roofline favor it?" A decode-heavy chat workload favors bandwidth and big SRAM (Groq, MI300X); a prefill/embedding/vision workload favors raw FLOPs (TPU, H100). That diagnosis is the job.

Qualcomm Cloud AI 100 specifics (since it's the JD2 archetype)

  • Design point: high TOPS/watt via many AI cores backed by large on-die SRAM, targeting INT8/INT4 inference. Strong for power-constrained data center and edge where perf/watt and density beat raw peak.
  • Toolchain: export to ONNX, compile/optimize with the qaic/Apps SDK (qaic-exec, the AIC compiler), quantize with AIMET (AI Model Efficiency Toolkit — PTQ/QAT). Contributing to the "Cloud AI GitHub repo and developer docs" (a JD bullet) means living in this toolchain. (Compilation flow generalized in Knowledge 03.)
  • The customer story you'll tell: "Same throughput at a fraction of the power and rack space" — which only lands if you can prove it with honest benchmarks (Knowledge 06) and recover INT8 accuracy (Knowledge 02).

11. The driver and software stack (bring-up reality)

JD2 demands "Linux-based systems, low-level software, drivers, and system bring-up." You won't write a driver, but you must understand the stack so you can debug it and not look lost in a war room.

Your Python (PyTorch / vLLM)
   │
Framework runtime (libtorch, etc.)
   │
Vendor runtime + libraries (CUDA runtime, cuBLAS/cuDNN, NCCL  |  qaic runtime  |  Neuron RT)
   │
User-mode driver (CUDA driver API, libcuda)
   │
Kernel-mode driver (nvidia.ko / amdgpu / qaic kernel module)
   │
Hardware (GPU/NPU over PCIe)

Things that actually go wrong (and that you'll be the one to triage):

  • Driver/runtime version mismatch — the #1 deployment failure. CUDA 12.x runtime vs an older driver; nvidia-smi works but torch.cuda.is_available() is False; container's CUDA ≠ host driver. Know nvidia-smi, nvcc --version, ldd, and the CUDA compatibility matrix.
  • ECC / Xid errors — hardware faults surface as Xid codes in dmesg/nvidia-smi -q. Learn to read them; a flaky GPU throttling under ECC errors looks like a "random latency spike."
  • MIG (Multi-Instance GPU) — an H100 can be partitioned into isolated slices; a customer's "GPU" might be 1/7th of one. Affects sizing.
  • Thermal throttling / power capsnvidia-smi -q -d POWER,TEMPERATURE. A "performance regression" is sometimes just a hot rack or a lowered power limit.
  • NUMA / PCIe topologynvidia-smi topo -m shows which GPUs share NVLink vs cross a CPU socket. Wrong process pinning tanks multi-GPU performance.

Bring-up checklist instinct: driver↔runtime↔container versions aligned → GPUs visible and healthy (nvidia-smi, no Xid) → topology understood (topo -m) → power/thermal headroom confirmed → only then benchmark. Skipping this is how teams chase phantom model bugs that are really infrastructure bugs.


12. Worked examples

Example A — Is Llama-3-8B decode compute- or memory-bound on an H100?

  • Decode generates 1 token at a time (batch 1). FLOPs/token ≈ 2P = 2·8e9 = 16 GFLOP.
  • Bytes moved ≈ read all weights once in FP16 = 2P = 16 GB.
  • Arithmetic intensity ≈ 16e9 FLOP / 16e9 B = ~1 FLOP/byte.
  • H100 ridge point ≈ ~295 FLOP/byte. Since 1 ≪ 295 → deeply memory-bound.
  • Predicted time per token ≈ bytes / bandwidth = 16e9 / 3.35e12 ≈ ~4.8 ms → ~210 tokens/s ceiling for batch 1, bandwidth-limited. (Real systems get less due to KV-cache reads, attention, and overhead.)
  • Lever: batch. At batch B, you still read the weights once but do B× the FLOPs → AI ≈ B → at B≈295 you approach compute-bound and ~B× the throughput for the same weight reads. This is the entire economic argument for batched serving.

Example B — Does FP8 help this workload?

Decode is memory-bound, and FP8 weights are half the bytes of FP16 → halves the dominant weight-read traffic → ~2× decode throughput, and halves memory so you fit 2× the KV cache / batch. Here FP8 helps via bandwidth + capacity, not via its 2× FLOPs (which a memory-bound kernel can't use). Correctly attributing which benefit applies is principal-level reasoning.

Example C — A customer's prefill of 8k-token prompts is slow.

Prefill processes all 8k tokens at once → it's a big batched matmul → AI in the hundreds → compute-bound. Here the levers are the opposite of decode: FP8/INT8 for more FLOPs, better tensor-core utilization (tiling, FlashAttention), and chunked prefill to overlap with decode (Knowledge 04). Same customer, two stages, two completely different bottlenecks — this is why "make inference faster" is never one answer.


13. References

  • Williams, Waterman, Patterson, "Roofline: An Insightful Visual Performance Model for Multicore Architectures", CACM 2009 — the source of §5.
  • NVIDIA Hopper (H100) and Blackwell (B200/GB200) Architecture Whitepapers — SM counts, tensor-core formats, NVLink bandwidth.
  • NVIDIA CUDA C++ Programming Guide — execution model, occupancy, CUDA Graphs (§7–8).
  • Jouppi et al., "In-Datacenter Performance Analysis of a Tensor Processing Unit" (ISCA 2017) — systolic arrays (§10).
  • Qualcomm Cloud AI 100 product brief and AIMET docs; AWS Neuron, Google TPU/XLA, Groq architecture overviews (§10).
  • Horace He, "Making Deep Learning Go Brrrr From First Principles" (2022) — the best plain-language treatment of compute-bound vs memory-bound vs overhead-bound.
  • Pope et al., "Efficiently Scaling Transformer Inference" (2022) — applies all of this to transformers (bridge to Knowledge 01).

Next: Knowledge 01 — LLM Inference Internals, where the roofline meets the actual transformer and we derive KV-cache size, prefill/decode costs, and the batching math that pays the bills.