NVIDIA GPU Hardware for Deep Learning

Table of Contents

• 0. GPU Quick Reference
• 1. GPU Architecture Overview
  • 1.1 The Streaming Multiprocessor
  • 1.2 Warps and SIMT Execution
  • 1.3 Memory Hierarchy
• 2. Memory Types: GDDR6X vs HBM
  • 2.1 GDDR6X
  • 2.2 HBM2e and HBM3
  • 2.3 Why Memory Bandwidth Matters
• 3. Tensor Cores
  • 3.1 Matrix Multiply-Accumulate
  • 3.2 Supported Precision Modes
  • 3.3 Evolution Across Generations
• 4. GPU Generations
  • 4.1 Turing (2018): RTX 20xx
  • 4.2 Ampere (2020): RTX 30xx / A100
  • 4.3 Ada Lovelace (2022): RTX 40xx / L40S
  • 4.4 Hopper (2022): H100
  • 4.5 Blackwell (2024): B100 / B200
• 5. Consumer vs. Data Center GPUs
  • 5.1 Feature Comparison Table
  • 5.2 Key Differentiators
• 6. CUDA Specializations
  • 6.1 NVLink and NVSwitch
  • 6.2 Multi-Instance GPU (MIG)
  • 6.3 Structured Sparsity (2:4)
  • 6.4 FP8 and the Transformer Engine
• 7. The Roofline Model
  • 7.1 Arithmetic Intensity
  • 7.2 Roofline Analysis
  • 7.3 Deep Learning Operations on the Roofline
• 8. Practical GPU Selection
• References

0. GPU Quick Reference 📋

A bird’s-eye view of every major NVIDIA GPU line relevant to deep learning, organized by generation and use case.

1. GPU Architecture Overview 🏗️

A GPU is a massively parallel processor organized around a hierarchy of execution units and memory. Understanding this hierarchy is prerequisite for reasoning about performance.

1.1 The Streaming Multiprocessor

Definition (Streaming Multiprocessor). A Streaming Multiprocessor (SM) is the fundamental compute unit of an NVIDIA GPU. It contains its own register file, instruction schedulers, functional units, and a configurable on-chip memory pool shared between L1 cache and shared memory.

The total number of SMs varies by chip and generation:

Each SM is organized into processing blocks (also called sub-cores). In Ampere and later, each SM has four processing blocks. Every processing block contains one Tensor Core unit, one warp scheduler, and its share of FP32 and INT32 cores.

1.2 Warps and SIMT Execution

Definition (Warp). A warp is a group of exactly 32 threads that execute in SIMT (Single Instruction, Multiple Threads) fashion: all 32 threads execute the same instruction each cycle, applied to distinct data. A warp is the atomic unit of scheduling on an SM.

Each SM issues instructions at the warp granularity. When a warp stalls (e.g., waiting on a memory load), the SM’s warp scheduler selects a different resident warp to issue — this latency hiding is the primary mechanism by which GPUs tolerate long-latency memory accesses.

Warp divergence occurs when threads within a warp take different branches. The SM must serialize the divergent paths, reducing effective throughput. Avoiding divergence in hot loops is a practical concern for CUDA kernel writers.

1.3 Memory Hierarchy

The GPU memory hierarchy, from fastest to slowest:

flowchart TB
    R["Register File
(per-thread, ~256KB/SM,
~1 cycle latency)"]
    S["Shared Memory / L1
(per-SM, 128–256 KB configurable,
~20 cycle latency)"]
    L2["L2 Cache
(on-die, 4–80 MB depending on GPU,
~200 cycle latency)"]
    HBM["Global Memory
(off-chip: HBM or GDDR,
~600–800 cycle latency)"]

    R --> S --> L2 --> HBM

Figure: GPU memory hierarchy from register file to global memory, with representative latencies.

Definition (Shared Memory). Shared memory is on-chip SRAM local to each SM. It is software-managed (programmer explicitly loads and stores to it), and can be used as a scratchpad to stage data for reuse across threads in the same block, eliminating repeated global memory traffic.

In Ampere and later, the L1 cache and shared memory share a single unified 128–256 KB SRAM bank per SM. The split between L1 and shared memory is configurable at kernel launch time. Using cudaFuncSetAttribute to maximize shared memory is a common optimization for attention and matmul kernels.

Key capacity numbers:

2. Memory Types: GDDR6X vs HBM 💾

The off-chip memory technology is one of the most consequential architectural choices separating consumer from data center GPUs.

2.1 GDDR6X

GDDR6X (Graphics Double Data Rate 6X) is the memory standard used in NVIDIA’s consumer Ada Lovelace (RTX 40xx) cards and some Turing/Ampere consumer cards. It uses PAM4 (Pulse Amplitude Modulation with 4 levels) signaling over a wide parallel bus, achieving high bandwidth over a physically narrow interface.

RTX 4090 GDDR6X numbers: - Bus width: 384-bit - Memory clock: ~21 Gbps effective per pin - Peak bandwidth: \(384 \times 21 / 8 \approx 1{,}008 \text{ GB/s}\)

Limitation: The GDDR6X die sits adjacent to the GPU on the package substrate. Scaling to much higher bandwidth requires either a wider bus (expensive, large PCB area) or faster signaling (power-limited). These constraints cap practical consumer GPU bandwidth around 1–1.2 TB/s.

2.2 HBM2e and HBM3

High Bandwidth Memory (HBM) stacks multiple DRAM dies vertically and connects them to the GPU through a silicon interposer via thousands of microscopic through-silicon vias (TSVs). This dramatically increases pin count and bandwidth density.

HBM2e (Ampere A100): - Up to 80 GB per GPU - Peak bandwidth: ~2.0 TB/s (A100 SXM)

HBM3 (Hopper H100 SXM): - 80 GB per GPU - Peak bandwidth: ~3.35 TB/s

HBM3e (H200, B200): - 141 GB (H200) / 192 GB (B200) per GPU - Peak bandwidth: ~4.8 TB/s (H200), ~8.0 TB/s (B200)

2.3 Why Memory Bandwidth Matters

For many deep learning operations, performance is limited not by how many FLOPs the chip can execute but by how fast weights and activations can be read from and written to memory. This distinction is formalized in the roofline model (Section 7).

Intuition: A transformer layer inference at batch size 1 reads the full weight matrix once per token. At BF16 precision, a 7B parameter model has \(7 \times 10^9 \times 2 \approx 14 \text{ GB}\) of weights. On an RTX 4090 (1 TB/s), loading these weights takes at minimum \(14 \times 10^{-3} \text{ s} = 14 \text{ ms}\), corresponding to ~71 tokens/second as a bandwidth ceiling regardless of compute speed.

3. Tensor Cores ⚡

3.1 Matrix Multiply-Accumulate

Definition (Tensor Core MMA). A Tensor Core is a specialized execution unit that performs a small matrix multiply-accumulate (MMA) operation in a single warp-level instruction. The abstract operation is:

\[D = A \cdot B + C\]

where \(A\), \(B\), \(C\), \(D\) are tiles of a larger matrix multiplication. The precise tile dimensions vary by generation and precision mode (see below), but conceptually each Tensor Core computes a fused outer product and accumulates into \(D\).

Why tiles? A single CUDA FMA computes one scalar fused multiply-add per cycle. A Tensor Core computes many such FMAs simultaneously using dedicated datapath wiring, amortizing instruction decode overhead and enabling dramatically higher throughput per area.

The CUDA programmer accesses Tensor Cores via: 1. WMMA API (nvcuda::wmma) — warp-level intrinsics, available since CUDA 9 2. PTX MMA instructions (mma.sync.aligned.*) — lower-level, more control 3. cuBLAS / cuDNN / CUTLASS — library calls that use Tensor Cores automatically

3.2 Supported Precision Modes

NVIDIA Tensor Cores operate in mixed precision: inputs at lower precision, accumulator at higher precision. This allows smaller data types for throughput and memory savings while preserving numerical stability in the accumulation.

Definition (TF32). TF32 (TensorFloat-32) is a non-standard NVIDIA format: 1 sign bit, 8 exponent bits (same range as FP32), and 10 mantissa bits (same precision as FP16). Its purpose is to drop mantissa bits from FP32 inputs so they fit into the Tensor Core pipeline without explicit user-side format conversion, giving automatic 10x speedup over FP32 CUDA cores for matmul.

Definition (FP8 E4M3 / E5M2). The Hopper architecture introduced two FP8 formats: E4M3 (4 exponent bits, 3 mantissa bits, used for forward pass where precision matters) and E5M2 (5 exponent bits, 2 mantissa bits, used for gradients where range matters more). The Transformer Engine automatically selects between them per-tensor.

3.3 Evolution Across Generations

4. GPU Generations 📅

4.1 Turing (2018): RTX 20xx

Turing (microarchitecture codename TU10x) was introduced in 2018 and is NVIDIA’s first consumer architecture to include dedicated hardware for both ray tracing and deep learning inference.

Key specs (RTX 2080 Ti, TU102): - Process: TSMC 12nm - Transistors: 18.6B - SMs: 68, CUDA cores: 4,352 - Tensor Cores: 2nd-gen, 8 per SM (544 total) - Memory: 11 GB GDDR6, 616 GB/s - FP16 Tensor TFLOPS: ~113.8 - Compute capability: 7.5

What Turing added for DL: - 2nd-gen Tensor Cores with INT8 and INT4 modes (for inference quantization) - Concurrent execution of FP32 and INT32 operations on independent datapaths

Turing is now largely obsolete for training but remains viable for inference on models that fit in 11 GB and can exploit INT8 quantization via TensorRT.

4.2 Ampere (2020): RTX 30xx / A100

Ampere (GA10x) is a generational leap: NVIDIA moved to Samsung 8nm (consumer) and TSMC 7nm (A100), nearly doubling transistor count and introducing a redesigned SM.

Key additions: - 3rd-gen Tensor Cores: BF16, TF32 (automatic FP32 → TF32 conversion in matmul), FP64 - Structured sparsity: 2:4 sparsity doubles Tensor Core throughput (see Section 6.3) - MIG (Multi-Instance GPU): Hardware partition of A100 into up to 7 isolated instances - 3rd-gen NVLink: 600 GB/s bidirectional GPU-to-GPU bandwidth

Flagship models: - A100-40GB SXM: 108 SMs, 312 FP16 TFLOPS, 1.555 TB/s HBM2e, CC 8.0 - A100-80GB SXM: Same compute, 2.0 TB/s HBM2e, 80 GB - RTX 3090 (GA102): 82 SMs, 35.6 TFLOPS FP32, 936 GB/s GDDR6X, 24 GB, CC 8.6 - RTX A6000 (GA102): 84 SMs, 77.4 FP16 TFLOPS, 768 GB/s GDDR6, 48 GB, CC 8.6

4.3 Ada Lovelace (2022): RTX 40xx / L40S

Ada Lovelace (AD10x) is NVIDIA’s 2022 consumer and professional GPU architecture, fabricated on TSMC 4N (effectively 5nm class). It targets ray tracing and AI inference alongside training.

Key additions: - 4th-gen Tensor Cores with FP8 — 5× throughput vs. Ampere for supported ops - Massive L2 cache: 96 MB on AD102 (16× the GA102’s 6 MB), reducing GDDR bandwidth pressure - 3rd-gen RT Cores, DLSS 3 Frame Generation - Shader Execution Reordering (SER) for graphics

Flagship models: - RTX 4090 (AD102): 128 SMs, 16,384 CUDA cores, 165 FP16 TFLOPS (no sparsity), 1,008 GB/s GDDR6X, 24 GB, CC 8.9 - L40S (AD102 professional): 142 SMs, 362 FP16 TFLOPS, 864 GB/s GDDR6, 48 GB, CC 8.9

The RTX 4090’s 24 GB VRAM remains a hard constraint for large-model fine-tuning, despite its strong per-FLOP performance. The L40S doubles VRAM at 48 GB.

4.4 Hopper (2022): H100

Hopper (GH100) is NVIDIA’s dedicated data center architecture for the LLM era. Announced alongside Ada but targeting an entirely different market, it introduces the first native FP8 support and the Transformer Engine.

Key additions: - 4th-gen Tensor Cores with native FP8 (E4M3/E5M2 formats) - Transformer Engine (TE): Hardware + software layer that automatically manages per-tensor FP8/FP16 scaling, delivering up to 9× speedup over A100 on transformer training - 2nd-gen MIG: Up to 7 confidential compute instances with dedicated video decode per slice - 4th-gen NVLink: 900 GB/s bidirectional per GPU - HBM3 (SXM5): 3.35 TB/s — 67% more than A100 SXM - DPX instructions: 40× acceleration for dynamic programming algorithms (bioinformatics, path finding) - Compute capability: 9.0

H100 SXM vs H100 PCIe:

The SXM form factor requires a dedicated baseboard (HGX) and liquid or high-airflow cooling — it is a rack-level purchase. PCIe H100 fits in a conventional GPU slot but sacrifices ~24% of compute and 40% of bandwidth.

4.5 Blackwell (2024): B100 / B200

Blackwell (GB100/GB200) is NVIDIA’s 2024 data center architecture. The flagship B200 uses a dual-die design connected by NV-HBI (NVLink chip-to-chip interconnect at 10 TB/s), appearing as a single logical GPU.

Key additions: - 5th-gen Tensor Cores: Up to 2.5× training improvement over Hopper per GPU - FP4 precision: 2nd-gen Transformer Engine adds FP4 for inference (18 PFLOPS sparse FP4 on B200) - FP6 precision: Intermediate option between FP8 and FP4 - 5th-gen NVLink: 1.8 TB/s per GPU, scalable to 576 GPUs via NVLink Switch - HBM3e: Up to 192 GB (B200), 8 TB/s bandwidth - RAS Engine: Hardware reliability, availability, and serviceability monitoring - Compute capability: 10.0

B100 vs B200:

5. Consumer vs. Data Center GPUs 📊

5.1 Feature Comparison Table

*RTX A6000 supports NVLink bridge for 2-GPU configurations (112.5 GB/s bidirectional), but not full NVSwitch fabric. †RTX 4090 FP8 Tensor Core via Ada’s 4th-gen Tensor Cores; no Transformer Engine.

5.2 Key Differentiators

1. ECC (Error-Correcting Code) memory. Data center GPUs include ECC on HBM and SRAM, allowing single-bit error correction. For multi-week pretraining runs, uncorrected DRAM errors accumulate and can corrupt model state silently. Consumer GPUs omit ECC for cost and minor bandwidth overhead reasons.

2. NVLink for multi-GPU bandwidth. All-reduce operations in data-parallel training dominate communication cost. Consumer GPUs communicate over PCIe Gen 4/5 (64–128 GB/s bidirectional), which is 5–15× narrower than NVLink. A PCIe-limited 8×RTX 4090 cluster often spends more time communicating gradients than computing them for large models.

3. MIG (Multi-Instance GPU). Allows a single A100 or H100 to be partitioned into up to 7 isolated GPU instances, each with guaranteed memory bandwidth and compute. Critical for multi-tenant inference serving (e.g., running seven independent models on one A100).

4. Transformer Engine. Available only on Hopper (H100) and Ada professional cards (L40S), not consumer Ada (RTX 40xx). It manages per-tensor FP8 scaling automatically, unlocking the full FP8 throughput numbers on the spec sheet.

5. Thermal and power design. Consumer GPUs assume air-cooled tower cases. SXM data center GPUs assume server-grade cooling (forced air or direct liquid cooling). The H100 SXM at 700W in a 1U server requires purpose-built infrastructure.

6. CUDA Specializations 🔧

6.1 NVLink and NVSwitch

Definition (NVLink). NVLink is NVIDIA’s proprietary high-speed GPU-to-GPU interconnect, implemented as a set of parallel differential signal pairs on the GPU substrate. Each NVLink “link” provides 50 GB/s bidirectional bandwidth (4th gen), and GPUs expose multiple links simultaneously.

NVSwitch is a dedicated switching chip that allows all-to-all connectivity among up to 8 (or 72, with newer generations) GPUs at full NVLink bandwidth simultaneously, eliminating the bandwidth contraction of a ring topology.

DGX H100 contains 8 H100 SXMs connected via NVSwitch, providing each GPU with 900 GB/s of all-to-all bandwidth. This enables tensor parallelism within a node without PCIe bottleneck — critical for large models where a single layer does not fit on one GPU.

6.2 Multi-Instance GPU (MIG)

Definition (MIG). Multi-Instance GPU is a hardware partitioning feature introduced in Ampere A100 that allows a single physical GPU to be divided into up to 7 isolated GPU instances, each with: - Guaranteed, non-interfering memory bandwidth slices - Isolated L2 cache partitions - Dedicated SM groups (no cross-instance contention) - Independent hardware fault isolation

MIG is designed for cloud multi-tenancy and inference serving. A cloud provider can carve one A100 into seven 10 GB instances and sell each to a different customer with SLA-level performance guarantees — impossible with software partitioning, which allows one tenant to thrash the shared L2.

MIG is not available on any consumer NVIDIA GPU.

6.3 Structured Sparsity (2:4)

Definition (2:4 Structured Sparsity). NVIDIA’s 2:4 sparsity format requires that at most 2 of every 4 consecutive weights are non-zero. Hardware stores only the non-zero values plus a 2-bit index array, halving the storage and bandwidth cost of loading weight tiles. Tensor Cores execute on the compressed representation.

The effective speedup is exactly 2× for compute and ~2× for memory traffic to load weight tiles, only if the weight matrix satisfies the 2:4 constraint. In practice:

• Training: weights are pruned to 2:4 format using a magnitude-based selection at each step, then fine-tuned. The NVIDIA ASP (Automatic SParsity) library automates this.
• Inference: a pretrained dense model can be post-hoc pruned to 2:4 with minimal accuracy loss for many architectures.

2:4 sparsity doubles Tensor Core throughput on Ampere and later, but requires explicit opt-in. PyTorch supports 2:4 sparse training via torch.ao.pruning.

6.4 FP8 and the Transformer Engine

Definition (Transformer Engine). NVIDIA’s Transformer Engine (TE) is a software library and hardware mechanism that automates per-tensor dynamic scaling for FP8 matrix multiplications. It: 1. Maintains per-tensor scale factors that map FP8 range to the tensor’s actual value distribution 2. Automatically chooses between E4M3 (forward pass, higher precision) and E5M2 (backward pass, higher range) 3. Updates scale factors based on a delayed scaling scheme (using statistics from prior iterations)

Without TE, using FP8 naively causes catastrophic loss of accuracy because the narrow 8-bit dynamic range cannot represent both near-zero gradients and large weight values simultaneously. TE’s per-tensor scaling solves this.

Performance impact: - H100 SXM FP8 Tensor Core: ~3,958 TFLOPS (vs. 1,979 in FP16) — a 2× throughput gain - In practice, end-to-end training speedup is ~1.5–2×, as not all operations benefit from FP8

TE is available only on Hopper (H100), Ada professional (L40S), and Blackwell hardware. Consumer Ada cards (RTX 40xx) have the FP8 Tensor Core hardware but lack the Transformer Engine software integration in the same form.

7. The Roofline Model 📐

7.1 Arithmetic Intensity

Definition (Arithmetic Intensity). For a compute kernel, let \(F\) denote the total floating-point operations (FLOPs) performed and \(B\) denote the total bytes transferred between the processor and main memory. The arithmetic intensity is:

\[I = \frac{F}{B} \quad \text{(FLOP/byte)}\]

Higher arithmetic intensity means more computation is extracted per byte loaded, reducing sensitivity to memory bandwidth.

Example: Dense matrix multiply. For \(C = AB\) with \(A \in \mathbb{R}^{M \times K}\), \(B \in \mathbb{R}^{K \times N}\): - FLOPs: \(F = 2MKN\) (one multiply, one add per element of the inner product) - Bytes (assuming inputs/output are loaded/stored once, FP16): \(B = 2(MK + KN + MN)\)

For \(M = N = K = d\) (square):

\[I = \frac{2d^3}{6d^2} = \frac{d}{3} \quad \text{FLOP/byte (FP16)}\]

This grows linearly with \(d\): a \(d=4096\) matmul has \(I \approx 1365\) FLOP/byte, comfortably compute-bound on any modern GPU.

Example: Softmax. For a vector of length \(N\): - FLOPs: \(\approx 5N\) (exp + sum + divide + log, roughly) - Bytes: \(2 \times 2N = 4N\) (read input, write output, FP16) - \(I \approx 5N / 4N = 1.25\) FLOP/byte — extremely memory-bound

7.2 Roofline Analysis

Definition (Roofline Model). Given a GPU with peak compute throughput \(\Pi\) (FLOP/s) and peak memory bandwidth \(\beta\) (byte/s), the roofline model bounds attainable performance \(P\) for a kernel with arithmetic intensity \(I\) by:

\[P \leq \min(\Pi,\ \beta \cdot I)\]

The bound \(\beta \cdot I\) is the memory roof (diagonal line in log-log space): even if compute were infinite, performance is capped by how fast data can be streamed from memory. The bound \(\Pi\) is the compute roof (horizontal line).

Definition (Ridge Point). The ridge point \(I^*\) is the arithmetic intensity at which the two bounds are equal:

\[I^* = \frac{\Pi}{\beta}\]

Kernels with \(I < I^*\) are memory-bound; kernels with \(I > I^*\) are compute-bound.

Ridge points for key GPUs (FP16 Tensor Cores):

The H100’s ridge point is 3.6× higher than the A100’s: a kernel must achieve 591 FLOP/byte to be compute-bound on the H100, versus 156 on the A100. This means more workloads are memory-bound on the H100 than on the A100, and achieving peak Tensor Core utilization requires larger batch sizes or longer sequences.

7.3 Deep Learning Operations on the Roofline

Compute-bound (high arithmetic intensity): - Large-batch GEMM: \(I = d/3\) grows with dimension and batch size - Attention QKV projection (large batch): same as GEMM - Prefill / prompt processing in LLM inference (long sequences, \(I \propto \text{seq len}\))

Memory-bound (low arithmetic intensity): - Element-wise ops: ReLU, GELU, dropout, residual add (\(I \approx 1–3\) FLOP/byte) - LayerNorm, RMSNorm: reduction + broadcast (\(I \approx 5–10\) FLOP/byte) - Decoding (autoregressive generation) at batch size 1: all weight matrices loaded once per token, \(I \approx 1\) FLOP/byte - KV-cache attention with short sequences

8. Practical GPU Selection 🎯

Choosing a GPU requires matching workload characteristics to hardware capabilities. The dominant variables are: model size (determines VRAM floor), batch size (determines whether memory- or compute-bound), training vs. inference (determines whether precision/Tensor Cores or bandwidth dominate), and multi-GPU need (determines NVLink necessity).

flowchart TD
    A{"Workload type?"}
    B["Inference: single model,
low latency"]
    C["Fine-tuning:
< 13B params"]
    D["Fine-tuning:
13B–70B params"]
    E["Pretraining at scale
or > 70B params"]

    A -->|inference| B
    A -->|fine-tune small| C
    A -->|fine-tune large| D
    A -->|pretraining| E

    B --> F{"VRAM fits?
(24 GB = ~13B 4-bit
or ~7B 16-bit)"}
    F -->|yes| G["RTX 4090
Cost-optimal,
high perf/dollar"]
    F -->|no| H["A100-80GB or
H100 PCIe
(for 70B 4-bit)"]

    C --> I["RTX 4090 (LoRA / QLoRA)
or RTX A6000
(48 GB, no NVLink limit)"]

    D --> J["A100-80GB
or 2× RTX A6000
(with NVLink bridge)"]

    E --> K["H100 SXM cluster
(NVLink + MIG +
FP8 TE required)"]

Figure: Decision flowchart for GPU selection by workload type.

Summary of use cases:

References