第 3 章 GPU 与计算基础

3.1 GPU vs CPU

为什么 LLM 需要 GPU

一句话：LLM 推理的核心操作是矩阵乘法，GPU 做矩阵乘法比 CPU 快 100-1000 倍。

CPU 的设计哲学是”少量核心、高频率、复杂控制逻辑”——适合分支多、逻辑复杂的串行任务。GPU 的设计哲学是”大量核心、较低频率、简单控制逻辑”——适合同一操作并行应用到大量数据。

具体数字：

Llama 2 7B 生成一个 token 需要约 140 亿次浮点运算（14 GFLOPS）。CPU 用 FP16 需要约 3ms（理论最优，实际更慢），A100 用 Tensor Core 只需约 0.045ms。batch size 越大，差距越明显。

CUDA Core vs Tensor Core

CUDA Core：通用浮点运算单元，一个 CUDA Core 每个时钟周期做 1 次 FP32 乘加（FMA）。A100 有 6912 个 CUDA Core。

Tensor Core：专门为矩阵乘法设计的硬件单元。一个 Tensor Core 一个时钟周期可以做一个 4x4 的矩阵乘加（也就是 128 次 FMA），效率远超 CUDA Core。A100 有 432 个 Tensor Core。

CUDA Core（1个周期做 1 次运算）:
  a × b + c → result

Tensor Core（1个周期做 4×4 矩阵乘加）:
  D = A × B + C
  其中 A, B, C, D 都是 4×4 矩阵
  = 64 次乘法 + 64 次加法 = 128 次 FMA

Transformer 的核心操作（Q@K^T、Attention@V、FFN 的矩阵乘法）天然适合 Tensor Core。实际上，现代推理引擎（vLLM、TensorRT-LLM）都会尽量把计算放到 Tensor Core 上。

A100 的 Tensor Core 支持的精度：FP16、BF16、TF32、INT8、INT4。H100 新增了 FP8 支持，B200 新增了 FP4——精度越低，吞吐越高：

FP16 到 INT8 吞吐翻倍——这就是量化（Quantization）的硬件基础。

3.2 显存（VRAM）

模型参数占多少显存

公式很简单：显存 = 参数量 x 每个参数的字节数

这只是模型权重的显存。推理时还有三块额外开销：

推理时的显存组成

┌────────────────────────────────────────────┐
│              总显存占用                      │
├────────────────────────────────────────────┤
│                                            │
│  模型权重        (固定)                     │
│  ████████████████████  14 GB (7B FP16)     │
│                                            │
│  KV Cache         (随 seq_len × batch 增长) │
│  █████████████         8 GB (示例)          │
│                                            │
│  激活值/中间结果  (临时，前向传播时占用)       │
│  ████                  2 GB (示例)          │
│                                            │
│  CUDA Context + 碎片                        │
│  ██                    1-2 GB               │
│                                            │
├────────────────────────────────────────────┤
│  总计:                 ~25 GB               │
└────────────────────────────────────────────┘

以 Llama 2 7B FP16 在 A100 80GB 上为例：

• 模型权重：14 GB（固定）
• KV Cache：取决于 batch size 和 seq_len，可能 1-50 GB
• 激活值：约 2 GB（取决于 batch size）
• CUDA overhead：1-2 GB
• 剩余显存可用于更大的 batch 或更长的 context

显存不够时的表现

1. OOM（Out of Memory）：最常见。CUDA 直接报错 RuntimeError: CUDA out of memory。通常发生在请求突然变长或并发变大时。
2. 无法加载模型：模型权重都放不下。解决：用量化（INT8/INT4）或切到更大卡。
3. 吞吐下降：显存紧张时 batch size 被迫缩小，GPU 利用率下降。

解决方案（按优先级）：

3.3 计算瓶颈 vs 显存瓶颈

这是理解 LLM 推理性能的关键概念。

Compute-bound vs Memory-bound

每个计算操作都涉及两步：(1) 从显存读数据到计算单元，(2) 计算。瓶颈在哪一步，决定了优化方向。

用一个简单的模型来思考——Arithmetic Intensity（计算强度）：

Arithmetic Intensity = FLOPs / Bytes accessed
单位: FLOP/Byte

GPU 有一个临界值：计算吞吐 / 显存带宽。A100 的临界值：

A100: 312 TFLOPS / 2.0 TB/s = 156 FLOP/Byte

如果一个操作的 Arithmetic Intensity > 156，它是 compute-bound（计算单元忙不过来）。如果 < 156，它是 memory-bound（计算单元在等数据）。

Prefill vs Decode

Prefill 阶段： 输入 n 个 token 一起做矩阵乘法。batch 维度（n）越大，每个权重参数被复用的次数越多，Arithmetic Intensity 越高。n=2048 时，每读一次权重矩阵可以做 2048 次运算——远超 156 的临界值，GPU 的计算单元成为瓶颈。

Decode 阶段： 每次只处理 1 个 token（batch=1 时）。[1, 4096] @ [4096, 4096] 的矩阵乘法中，权重矩阵（4096x4096x2 = 32MB）被完整读一遍，但只做了 1x4096x4096 ≈ 33M 次运算。Arithmetic Intensity ≈ 33M / 32M ≈ 1——远低于 156 的临界值，显存带宽成为瓶颈。

这意味着：

• 优化 Prefill → 提升算力（用更快的 GPU、更高效的 kernel）
• 优化 Decode → 提升显存带宽 或 减少要读取的数据量（量化、GQA 减少 KV Cache）

这也解释了为什么 batch 很重要：Decode 阶段如果同时处理 32 个请求（batch=32），权重矩阵读一次就能服务 32 个请求，Arithmetic Intensity 提升 32 倍，GPU 利用率大幅提高。这就是 continuous batching 的核心意义。

3.4 CUDA 编程的最小认知

你不需要写 CUDA，但需要理解三个概念，才能看懂推理引擎的优化思路。

Kernel、Thread、Block

CUDA 的执行模型：

GPU
 └── Grid（一次 kernel 调用）
      ├── Block 0
      │    ├── Thread 0
      │    ├── Thread 1
      │    ├── ...
      │    └── Thread 255
      ├── Block 1
      │    ├── Thread 0
      │    ├── ...
      │    └── Thread 255
      └── ...

Kernel：一个在 GPU 上运行的函数。比如”把两个矩阵相加”就是一个 kernel。 Thread：最小的执行单元，一个 thread 处理一个（或几个）数据元素。 Block：一组 thread，共享一块高速的 Shared Memory（SRAM）。 Grid：所有 block 的集合。

一个实际例子——向量加法：

// CUDA Kernel: C = A + B
// __global__ 表示这个函数在 GPU 上执行，由 CPU 端发起调用
__global__ void vector_add(float *A, float *B, float *C, int n) {
    // blockIdx.x 是当前 block 的编号，threadIdx.x 是当前 thread 在 block 内的编号
    // 两者组合计算出全局唯一的线程坐标
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        C[i] = A[i] + B[i];  // 每个 thread 处理 1 个元素
    }
}
 
// 启动: 256 个 thread/block, 需要多少 block 取决于数据大小
// <<<blocks, threads>>> 是 CUDA 的内核启动语法，指定 grid 和 block 的大小
vector_add<<<(n+255)/256, 256>>>(A, B, C, n);

GPU 内存层次

┌─────────────────────────────────────────────┐
│ HBM (显存, 80GB, ~2 TB/s)                    │ ← 模型权重、KV Cache 在这里
│                                             │
│   ┌─────────────────────────────────────┐   │
│   │ L2 Cache (40MB, ~5 TB/s)            │   │
│   │                                     │   │
│   │   ┌─────────────────────────────┐   │   │
│   │   │ Shared Memory / L1 Cache    │   │   │ ← 每个 Block 有 ~192KB
│   │   │ (~192KB per SM, ~19 TB/s)   │   │   │    FlashAttention 利用的就是这层
│   │   │                             │   │   │
│   │   │   ┌─────────────────────┐   │   │   │
│   │   │   │ Registers           │   │   │   │ ← 每个 Thread 私有
│   │   │   │ (最快, ~64KB/SM)    │   │   │   │
│   │   │   └─────────────────────┘   │   │   │
│   │   └─────────────────────────────┘   │   │
│   └─────────────────────────────────────┘   │
└─────────────────────────────────────────────┘

越上层容量越大但速度越慢，越下层容量越小但速度越快。

FlashAttention 为什么是 “IO 感知” 的

标准 Attention 的实现：

S = Q @ K.T          # [n, n] 写入 HBM
P = softmax(S)       # [n, n] 读 HBM，写 HBM
O = P @ V            # [n, d] 读 HBM，写 HBM

中间结果 S 和 P 是 [n, n] 的大矩阵（n=4K 时 = 64MB），每次都要在 HBM 里读写。这些 HBM 访问是 Attention 计算的主要瓶颈。

FlashAttention 的核心思路：把 Attention 分成小块（tile），在 Shared Memory（SRAM）里完成计算，避免把中间结果写回 HBM。

标准 Attention:  HBM → 计算 S → 写回 HBM → 读 S → 计算 P → 写回 HBM → 读 P → 计算 O
                 ^^^                ^^^       ^^^              ^^^       ^^^
                 5 次 HBM 读写（n×n 级别）

FlashAttention: HBM → SRAM → 分块计算 S,P,O → 写回 HBM
                ^^^                             ^^^
                只有 2 次 HBM 读写（n×d 级别）

FlashAttention 2 在 A100 上把 Attention 的速度提升了 2-4 倍，不是靠减少计算量（FLOPs 一样多），而是靠减少 HBM 读写次数。这就是 “IO-aware” 的含义——算法设计考虑了硬件的内存层次。

3.5 GPU 选型

主要 GPU 对比

选型建议：

个人开发/学习： RTX 4090（24GB）可以跑 7B FP16 或 13B INT4。性价比最高的学习卡。

生产推理（小规模）： L40S 或 A100 40GB。L40S 价格约为 A100 的一半，推理性能接近，但显存带宽低，Decode 速度会差一些。

生产推理（大规模）： H100 或 H200。H200 的 141GB 显存意味着单卡就能跑 70B INT4 模型。如果跑长 context（128K+），H200 的大显存优势明显。

旗舰方案： B200。FP16 算力是 H100 的 2.3 倍，显存带宽是 2.4 倍。FP4 推理可以在单卡上跑 70B 模型。

国内常见 GPU

由于出口管制，A100/H100 在国内受限。常见的替代方案：

A800/H800 和 A100/H100 的计算能力完全一样，只是卡间互联（NVLink）带宽被削减了。对单卡推理场景（7B-13B 模型）影响很小，对多卡并行训练影响较大。

云 GPU 实例参考价格

以下价格数据截至 2026 年初，仅供数量级参考，请以云厂商官网最新定价为准。

以按量计费为参考：

对于学习和开发，AutoDL 等 GPU 云平台的竞价实例性价比最高。生产环境建议用包月实例或阿里云/腾讯云的专业 GPU 实例。

延伸阅读：

• NVIDIA A100 Whitepaper
• FlashAttention: Fast and Memory-Efficient Exact Attention
• Making Deep Learning Go Brrrr From First Principles - Horace He 的经典文章
• GPU Puzzles - 用游戏学 CUDA 概念

代码示例