GPU 工作原理：和 CPU 有何不同？

发表于 更新于 分类于 ， 阅读次数： 本文字数： 6.4k 阅读时长 ≈ 12 分钟

如果把 CPU 比作“多才多艺的总管”，GPU 更像“高效的流水线工厂”：CPU 擅长复杂分支与少量任务的低延迟处理，GPU 擅长大量相同/相似任务的高吞吐处理。本文先给出关键概念，再用 CPU 对比串起 GPU 的架构、执行模型、内存层次与性能要点。

📚 核心概念

  • 指令级并行 ILP（Instruction-Level Parallelism）
    • 定义：单线程内部，多条彼此独立的指令可乱序/并行执行以提高吞吐。
    • 省流：一个人能同时“烧水+等电梯”，就更快。
  • 数据级并行 DLP（Data-Level Parallelism）
    • 定义：对大量数据元素执行相同算子，实现并行化。
    • 省流：100 份相同盒饭，开 100 条流水线一起装。
  • SIMD 与 SIMT
    • 定义：SIMD（Single Instruction Multiple Data）单指令多数据；SIMT（Single Instruction Multiple Threads）以线程为抽象、底层按 SIMD 成组执行。
    • 省流：SIMT 就是把一群线程捆成小队，表面各自执行，实际“一声令下齐步走”。
  • Warp/Wavefront 与 SM/CU（以 NVIDIA/AMD 为例）
    • 定义：NVIDIA 将 32 个线程编成一个 Warp，AMD 将 64 个线程编成一个 Wavefront；它们在 SM（Streaming Multiprocessor）/CU（Compute Unit）上被调度执行。
    • 省流：先分“班级”（Warp/Wavefront），再分“年级”（Block），最后“全校”（Grid）一起上课，老师（调度器）轮流点名。
  • 分歧 Divergence
    • 定义：同一 Warp 内线程走不同控制分支，会被拆分顺序执行，导致利用率下降。
    • 省流：本来排队一起过闸，结果有人走旁门，只能分两拨依次过。
  • 访存合并 Coalescing
    • 定义：相邻线程访问相邻地址能够合并为更少、更宽的内存事务，提高带宽利用率。
    • 省流：把零碎快递打包成一箱寄，省钱省时间。
  • 占用率 Occupancy
    • 定义：某 SM 上活跃 Warps 数量与硬件上限之比，影响隐藏延迟能力。
    • 省流：教室能坐 100 人，只来了 30 人，空位越多越“冷清”。
  • 内存层次
    • 定义：寄存器 → 共享内存/L1 → L2 → 全局显存（GDDR/HBM） → 主机内存（经 PCIe/NVLink）。容量越大通常延迟越高。
    • 省流：越近越快、越远越慢；把常用东西放在手边。

🆚 CPU vs GPU：设计取舍一图看懂

维度 CPU GPU
目标 低延迟、强控制、通用 高吞吐、大规模并行
核心 少量大核，复杂控制与大缓存 大量简化计算单元，轻控制、多并行
缓存 多级大缓存（L1/L2/L3） 更小的 L1/共享 + 中等 L2 + 超高显存带宽
并行 强 ILP/少量线程 强 DLP/海量线程（Warp/Wavefront）
分支 复杂分支友好 分歧代价高，需要规避
典型场景 系统服务、事务处理、复杂逻辑 图形渲染、矩阵/向量计算、深度学习

🧩 GPU 架构速写（抽象）

  • 指挥与前端：命令处理、硬件队列、前端调度。
  • 计算集群：多个 SM（NVIDIA）/CU（AMD），含标量/向量 ALU、特殊函数单元（SFU）、张量单元（新架构）。
  • 调度：Warp/Wavefront 由调度器轮换发射，遇访存延迟切换到其他 Warp 以隐藏延迟。
  • 存储：每 SM 有寄存器与共享内存/L1，全局共享 L2，高带宽显存承载大数据；主机经 PCIe/NVLink 互联。


graph LR
    Host[主机/驱动] --> Queue[命令队列]
    Queue --> GPU[GPU]
    GPU --> SM1[计算单元 SM/CU]
    GPU --> SM2[计算单元 SM/CU]
    SM1 --> W1[Warp/Wavefront 调度]
    SM2 --> W2[Warp/Wavefront 调度]
    GPU --> L2[L2 缓存]
    L2 --> VRAM[显存 GDDR/HBM]

SM/CU 详解（做什么 + 怎么调度）

  • 组成：一个 SM/CU 内通常包含多条算术逻辑管线（标量/向量 ALU）、特殊函数单元（SFU，计算 sin/cos/exp 等）、加载/存储单元（LD/ST）、有的架构还包含张量/矩阵加速单元；配套有寄存器文件与可配置的共享内存/L1。
  • 调度：每个 SM 有多个 Warp 调度器，会从“就绪”的 Warps 中挑选并在每个周期向执行管线发射指令；当某 Warp 因访存/依赖而“阻塞”时，切换到其他 Warp 以隐藏延迟。
  • 资源约束：每个 SM 的寄存器总量与共享内存容量是固定的。单线程寄存器用得越多、单 Block 共享内存占得越多，可并发常驻的 Warps/Blocks 就越少（占用率下降）。
  • 省流：
    • SM 就像一间“装了多套设备”的大教室：算术机、函数机、搬运机（LD/ST）。
    • 老师（调度器）盯着好几班（Warps），谁准备好了就先让谁上机做题；有人卡住等资料（访存），就先让别的班做题。
    • 班里每个学生（线程）要用的课本（寄存器）越多，教室能同时容纳的班级就越少。

⚙️ 执行模型：SIMT 与 Warp 如何配合

  • 编程视角：写“每线程”的内核函数（kernel），通过 Grid/Block 组织。这里的“每线程”指逻辑线程（thread），非 warp；硬件会把若干逻辑线程打包为 warp 锁步执行，但不改变以单线程为编程抽象的模型。
  • 硬件视角：线程被分成 Warp/Wavefront 成组执行；遇分歧需串行化；遇访存等待就“切 Warp”。
  • Warp 尺寸：NVIDIA 常见 32，AMD 常见 64（概念相同）。

为什么分歧会降低利用率

  • SIMT 锁步：一个 warp（如 32 线程）必须“同一步执行同一条指令”。if 分叉后，硬件用掩码把不相关线程关掉，先跑一支，再跑另一支。
  • 串行两段各自的并行：例如 20 线程走 if、12 线程走 else → 先以 20/32 利用率跑 if，再以 12/32 利用率跑 else；总时间≈两段时间相加，平均利用率下降。
  • 不能靠预测/乱序补救：GPU 的面积和功耗主要花在“海量算力+寄存器”上，没有为单条控制流配备很强的分支预测/乱序执行硬件；它擅长通过“切到别的 warp”隐藏内存延迟，但同一 warp 的分支分歧仍要分段跑完，无法被这种切换抵消。

📦 内存层次与数据路径（越近越快）

  • 寄存器：每线程私有，最快；数量限制影响占用率。
  • 共享内存（近似 L1）：Block 内共享、低延迟；注意 bank 冲突与容量。
  • L2：片上缓存，所有 SM 共享。
  • 全局显存：带宽高、延迟大；优先合并访存、减少随机访问。
  • 主机内存：经 PCIe/NVLink；数据迁移应成批、异步化。


flowchart TD
    Reg["寄存器 / 每线程"] --> SMem["共享内存 / 每SM"]
    SMem --> L2["L2 缓存"]
    L2 --> VRAM["显存"]
    VRAM --> Host["主机内存（PCIe/NVLink）"]

内存层次补充（共享内存 vs 缓存、事务与 bank）

  • 共享内存 vs L1：共享内存是软件可控的片上 SRAM，线程块内可见；L1/统一缓存由硬件自动管理。很多架构将“共享内存与 L1”设计为可配比的同一物理池。
  • 合并访存：Warp 内线程按连续地址访问时，可合并为较少的 32/64/128B 事务，降低总事务数；随机/跨步访问会放大事务数与延迟。
  • Bank 冲突：共享内存被划分为多个 bank，同时访问落在同一 bank 的多个地址会产生串行化（冲突），需要通过数据排布/步长调整避免。
  • 省流：
    • 共享内存像你们班的“黑板贴”，大家商量好怎么贴，就能少跑回后仓（显存）。
    • 合并访存像“集中下单一起发货”，比一个个零碎下单更省事。
    • bank 冲突像“同时挤同一扇门”，需要分流。
  • PCIe：主机与 GPU 常用的通用互联总线，按代际与通道数提供不同带宽（例如 PCIe 4.0 x16 单向理论带宽约 ~32 GB/s，PCIe 5.0 x16 约 ~64 GB/s），延迟相对更高。
  • NVLink：面向 GPU-GPU/GPU-CPU 的高带宽、低延迟互联，支持点对点直连与多链路聚合，典型带宽远高于同代 PCIe，并支持内存一致性/地址空间统一的特性（随架构而异）。
  • 省流：
    • PCIe 像“城市主干道”，什么车都能走，但高峰会慢；
    • NVLink 像“GPU 之间修的高速专用通道”，路更宽、收费站更少，GPU 彼此传东西更快更省时间。

延迟 vs 带宽

  • 延迟 Latency
    • 定义：从发出请求到收到第一个结果的等待时间，单位常用 ns/µs/ms，越小越好。
    • 省流：起步反应快不快。
  • 带宽 Bandwidth
    • 定义：单位时间可传输的数据量，单位常用 GB/s，越大越好。
    • 省流：路有多宽、一次能过多少车。
  • 典型量级（直觉参考）
    • 寄存器/共享内存：延迟极低（纳秒级），带宽很高。
    • 显存（GDDR/HBM）：延迟更高（数百纳秒），带宽极高。
    • NVLink：延迟低于 PCIe、带宽高于 PCIe。
    • PCIe：延迟微秒级，带宽取决于代际与通道数。
  • 为什么重要（GPU 如何对抗高延迟）
    • 省流：通过“切 Warp/提占用率”隐藏等待；用共享内存/合并访存把热数据放近身；用异步拷贝与计算重叠，减少干等。

🚀 性能要点

  • 访存合并：相邻线程访问相邻地址。
  • 共享内存分块（tiling）：把重复数据留近身，减显存往返。
  • 占用率权衡：寄存器/共享内存用得越多，并发度可能越低。
  • 控制流优化：减少同 Warp 内分歧，可用重排/掩码等方式。
  • 传输重叠：异步拷贝、多 Stream 藏 PCIe 延迟。

🧪 最小工作流

  1. 主机准备/加载数据；
  2. 复制到显存（或使用统一内存/页迁移）；
  3. 启动 kernel（grid, block）；
  4. 同步/异步收集结果，可与下批次流水线并行。

🧾 示例：最小 CUDA Kernel（直觉建立）

1
2
3
4
5
6
7
8
9
10
11
12
13
14
// vadd.cu
__global__ void vadd(const float* a, const float* b, float* y, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x; // 逻辑“每线程”索引
  if (i < n) y[i] = a[i] + b[i];
}

__global__ void thresh_abs(const float* x, float* y, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n) y[i] = (x[i] > 0.f) ? x[i] : -x[i]; // 分支：若 warp 内正负混杂会产生分歧
}

// 启动（host）
// int threads = 256; int blocks = (n + threads - 1) / threads;
// vadd<<<blocks, threads>>>(d_a, d_b, d_y, n);

省流：每个逻辑线程处理一个元素；Grid/Block 决定并行规模；分支一致更快。

🤔 何时用 GPU，何时用 CPU？

  • 适合 GPU：高算术密度、强 DLP、可分块重用数据（如矩阵乘、卷积、光线追踪）。
  • 适合 CPU：复杂分支、状态机密集、数据量小且延迟敏感（如事务、编译器前端）。

小清单： - 单请求首响应（µs 级）/强交互/大量系统调用 → 倾向 CPU - 高算术密度/规则批处理/大规模并行 → 倾向 GPU

🧨 常见误区

  • “GPU 频率高就一定更快？”——吞吐取决于并行度与访存模式。
  • “线程越多越好？”——资源配额有限，盲堆线程会掉占用率或引入开销。
  • “分支无所谓？”——同 Warp 分歧会串行化，是吞吐杀手。

👀 分歧可视化（无分歧 vs 有分歧）



flowchart TD
  A["Warp（32） 无分歧"] --> A1["一步完成 同步执行"]
  B["Warp（32） 有分歧：20 if ／ 12 else"] --> B1["先跑 if：20 活动 ／ 12 掩码"]
  B1 --> B2["再跑 else：12 活动 ／ 20 掩码"]

✅ 小结

  • 一句话记忆：GPU 用“成群线程”换吞吐，用“切 Warp”藏延迟；真正的性能来自“对齐访问 + 合理占用 + 少分歧 + 数据重用”。

🧭 访存模式对比：AoS vs SoA（合并访存更友好）

  • AoS（Array of Structs） 
    1
    2
    struct Pt { float x,y,z; };
    Pt* pts; // pts[i].x 连续性差，warp 取 x 会跨步

  • SoA（Struct of Arrays） 

    1
    2
    struct Pts { float* x; float* y; float* z; };
    // 连续访问 x[i] 更易合并为少量内存事务

  • 假设一个 warp 有 32 个线程，每个线程要读一个 float x（4 字节）。
  • 如果是 AoS：内存像这样排 
    1
    x0 y0 z0 | x1 y1 z1 | x2 y2 z2 | ...   // 每个点的 x、y、z 挨在一起
    第 i 个线程要读 xi，地址间隔是“一个点的大小”（约 12 字节），即 0, 12, 24, 36, ... 32 个地址跨越范围≈ 32 × 12 = 384 字节，比较分散，需要多次内存事务才能取完。

  • 如果是 SoA：内存像这样排 

    1
    x0 x1 x2 x3 ... x31 | y0 y1 ... | z0 z1 ...   // 所有 x 连在一起
    32 个线程读取 x0..x31 对应地址正好是连续的 128 字节（32 × 4B），硬件可以“一次/极少数几次”就打包取齐（称为合并访存）。

经验法则：让“相邻线程 → 访问相邻地址”。读 x 用 SoA 更容易一次性打包；读整点（x,y,z 都要）时 AoS 也未必差，但要注意对齐与访问顺序。

省流：相邻线程访问相邻地址（SoA）→ 合并事务更容易，带宽利用更高。

🧱 张量核心与混合精度（DL 常见）

  • 张量核心/Tensor Core：面向矩阵乘累加（MMA）的专用单元，吞吐远高于标量/向量 ALU。
  • 混合精度：TF32/FP16/BF16 以轻微精度损失换显著吞吐提升（常见于训练/推理）。 省流：矩阵算子走“专用快车道”，吞吐更猛；前提是算子形状/对齐满足要求。

🧰 性能指标与工具

  • 指标：
    • warp_execution_efficiency（分歧/掩码影响）
    • achieved_occupancy（占用率）
    • dram__throughput / l2__hit_rate（显存带宽与 L2 命中）
    • sm__throughput（核心忙闲）
  • 工具：
    • Nsight Compute（内核级剖析）
    • Nsight Systems（全局时间线与重叠）
    • compute-sanitizer（memcheck/racecheck）

⚠️ 常见坑清单（避坑速览）

  • 寄存器溢出到“本地内存”（栈）→ 延迟飙升
  • 共享内存 bank 冲突/容量不足 → 吞吐打折
  • 未对齐/跨步访存 → 合并失败，事务增多
  • kernel 太小 → 启动/调度开销占比过大
  • 统一内存频繁页迁移 → 带宽/延迟抖动

🔗 主机-设备数据通道优化

  • 固定页内存（pinned）+ cudaMemcpyAsync → 提升带宽、支持拷贝与计算重叠
  • 双缓冲/多 stream → pipeline 化生产与消费
  • P2P/NVLink（多 GPU）→ 绕开主机内存，直连更快

🧬 多 GPU 与通信（入门向）

  • NCCL：多 GPU 通信库，提供 all-reduce/all-gather/broadcast 等原语
  • 拓扑：同节点内优先 NVLink/同交换域，跨节点看 IB/RoCE；拓扑差异决定带宽与延迟

📒 术语速查

术语 定义 省流
ILP 单线程内的指令级并行 一人多活
DLP 多数据上的同算子并行 开多条线
SIMT 线程抽象下的 SIMD 执行 小队齐步走
Warp 成组锁步的线程集合 班级
Occupancy 活跃 warps/上限 比例 教室坐满没
Coalescing 相邻访存合并为少事务 集中发货
Latency 首字节等待时间 起步反应
Bandwidth 单位时间传输量 路有多宽