GPU 架构深入——上万个核心如何分工协作 / GPU Architecture and Massive Parallel Execution
📅 创建时间：2026-06-02 🏷️ 标签：#GPU #CUDA #StreamMultiprocessor #Warp #HBM #TensorCore 📚 前置知识：[[01-computer-architecture-basics]]（CPU 架构） [[02-parallel-computing-theory]]（并行理论） 📚 相关知识：[[04-cuda-programming-model]]（CUDA 编程） [[05-cuda-kernel-and-memory]]（内存管理）
先抓住直觉
不要把一个 CUDA Core 当成一个 CPU Core。GPU 的优势来自大量轻量执行单元按组完成相似工作，并在一组线程等待数据时切换到另一组。理解顺序应是 Thread → Warp → Block → SM → 整张 GPU，不必先背芯片框图。
- 必须理解：Warp 是调度单位；同一 Warp 走不同分支会串行化；GPU 用并发隐藏访存延迟。
- 用到再查：A100 的核心数、GPC/TPC 数量、各级缓存容量。
- 读完能回答：为什么 GPU 可以容忍较高延迟？为什么分支很多的任务可能不适合 GPU？
场景：GPU 的 6912 个核心，到底长什么样？
┌─────────────────────────────────────────────────────────────┐
│ │
│ A100 GPU：6912 个 CUDA 核心 │
│ │
│ 你可能听说过这个数字，但大脑无法直观理解： │
│ │
│ 6912 是什么概念？ │
│ - 普通 CPU：8-64 核 │
│ - A100：6912 核 │
│ - 差距：100 倍！ │
│ │
│ 但这 6912 个核心不是杂乱排列的。 │
│ 它们有严格的组织结构：SM → Warp → Thread。 │
│ 理解这个结构，是写好 CUDA 代码的第一步。 │
│ │
└─────────────────────────────────────────────────────────────┘第1节：GPU 架构全景图
A100 完整架构
┌─────────────────────────────────────────────────────────────┐
│ A100 GPU 完整架构 │
├─────────────────────────────────────────────────────────────┤
│ │
│ ┌─────────────────────┐ │
│ │ A100 GPU │ │
│ │ (GA100 核心) │ │
│ └──────────┬──────────┘ │
│ │ │
│ ┌──────────┬─────────┼─────────┬──────────┐ │
│ │ │ │ │ │ │
│ ┌────┴───┐ ┌────┴───┐ ┌────┴───┐ ┌────┴───┐ │
│ │ GPC 0 │ │ GPC 1 │ │ GPC 2 │ │ GPC 3-7 │ │
│ │(图形处理)│ │(图形处理)│ │(图形处理)│ │ (×5) │ │
│ │ ×4 个 │ │ ×4 个 │ │ ×4 个 │ │ │ │
│ └────┬────┘ └────┬────┘ └────┬────┘ └────────┘ │
│ │ │ │ │
│ ┌───┴───┐ ┌───┴───┐ ┌───┴───┐ │
│ │ TPC 0 │ │ TPC 0 │ │ TPC 0 │ │
│ │(纹理处理)│ │(纹理处理)│ │(纹理处理)│ │
│ └────┬────┘ └────┬────┘ └────┬────┘ │
│ │ │ │ │
│ ┌───┴───┐ ┌───┴───┐ ┌───┴───┐ │
│ │ SM 0 │ │ SM 0 │ │ SM 0 │ │
│ │ SM 1 │ │ SM 1 │ │ SM 1 │ │
│ │ ... │ │ ... │ │ ... │ │
│ │ SM 15 │ │ SM 15 │ │ SM 15 │ │
│ └────────┘ └────────┘ └────────┘ │
│ │
│ ┌───────────────────────────────────────────────────────┐ │
│ │ HBM2e 显存 │ │
│ │ 40GB / 80GB，带宽 2 TB/s │ │
│ └───────────────────────────────────────────────────────┘ │
│ │
└─────────────────────────────────────────────────────────────┘
GPC = Graphics Processing Cluster（图形处理簇） × 7 个
TPC = Texture Processing Cluster（纹理处理簇） × 16 个
SM = Streaming Multiprocessor（流式多处理器） × 108 个（实际 108 SM）
HBM = High Bandwidth Memory（高带宽显存）A100 规格一览
┌─────────────────────────────────────────────────────────────┐
│ A100 vs 其他 GPU 规格对比 │
├─────────────────────────────────────────────────────────────┤
│ │
│ 规格 │ A100 SXM │ H100 SXM │ RTX3090│
│ ────────────────┼───────────────┼──────────────┼────────┤
│ CUDA 核心 │ 6912 │ 16896 │ 10496 │
│ SM 数量 │ 108 │ 132 │ 82 │
│ Tensor Core │ 432 │ 528 │ 328 │
│ 显存容量 │ 40/80 GB │ 80 GB │ 24 GB│
│ 显存带宽 │ 2.0 TB/s │ 3.35 TB/s │ 936 GB/s│
│ FP32 算力 │ 19.5 TFLOPS │ 67 TFLOPS │ 36 TFLOPS│
│ Tensor FP16 │ 312 TFLOPS │ 989 TFLOPS │ 71 TFLOPS│
│ 功耗 │ 400W │ 700W │ 350W │
│ NVLink 带宽 │ 600 GB/s │ 900 GB/s │ N/A │
│ │
│ 价格（参考） │ ~10 万人民币 │ ~25 万人民币│ ~1.5万 │
│ │
└─────────────────────────────────────────────────────────────┘第2节：SM——GPU 的基本执行单元
SM 内部结构（以 A100 为例）
┌─────────────────────────────────────────────────────────────┐
│ 流式多处理器（SM）内部结构 │
├─────────────────────────────────────────────────────────────┤
│ │
│ ┌─────────────────┐ │
│ │ Warp Scheduler │ ← 调度器，每个 SM ×4 │
│ └────────┬────────┘ │
│ │ │
│ ┌─────────────────────────┼─────────────────────────────┐ │
│ │ 算术单元阵列 │ │
│ │ │ │
│ │ ┌──────┐ ┌──────┐ ┌──────┐ ┌──────┐ │ │
│ │ │INT32 │ │INT32 │ │ FP32 │ │ FP32 │ │ │
│ │ │ ALU │ │ ALU │ │ FMA │ │ FMA │ │ │
│ │ └──────┘ └──────┘ └──────┘ └──────┘ │ │
│ │ ×16 个 │ │
│ │ │ │
│ │ ┌──────────┐ ┌──────────┐ ┌──────────┐ │ │
│ │ │TensorCore│ │TensorCore│ │TensorCore│ │ │
│ │ │ (矩阵乘) │ │ (矩阵乘) │ │ (矩阵乘) │ │ │
│ │ └──────────┘ └──────────┘ └──────────┘ │ │
│ │ ×4 个（每个 Tensor Core 每时钟执行 256 FLOP） │ │
│ │ │ │
│ │ ┌──────────┐ ┌──────────┐ │ │
│ │ │ LD/ST │ │ LD/ST │ ← 加载/存储单元 │ │
│ │ │ Unit │ │ Unit │ ×16 个 │ │
│ │ └──────────┘ └──────────┘ │ │
│ │ │ │
│ └─────────────────────────┬───────────────────────────┘ │
│ │ │
│ ┌─────────────────────────┼───────────────────────────┐ │
│ │ 存储单元 │ │
│ │ │ │
│ │ ┌──────────────┐ ┌──────────────┐ │ │
│ │ │ L1/Shared │ │ Registers │ │ │
│ │ │ Memory │ │ (65536 × │ │ │
│ │ │ 128 KB │ │ 32-bit) │ │ │
│ │ └──────────────┘ │ = 256 KB │ │ │
│ │ └──────────────┘ │ │
│ │ │ │
│ │ ┌──────────────┐ ┌──────────────┐ │ │
│ │ │ L1 Cache │ │ L2 Cache │ │ │
│ │ │ (per SM) │ │ (per GPU) │ │ │
│ │ └──────────────┘ │ 40 MB │ │ │
│ │ └──────────────┘ │ │
│ │ │ │
│ └──────────────────────────────────────────────────────┘ │
│ │
└─────────────────────────────────────────────────────────────┘每个 SM 的关键数字
┌─────────────────────────────────────────────────────────────┐
│ A100 每个 SM 的资源 │
├─────────────────────────────────────────────────────────────┤
│ │
│ 寄存器文件：65536 × 32-bit = 256 KB（65536 个寄存器） │
│ 每个线程最多使用 255 个寄存器 │
│ 每个 SM 最多同时运行 2048 个线程（= 64 Warp × 32 线程） │
│ │
│ L1/Shared Memory：128 KB（可配置比例） │
│ - Shared Memory：用户显式管理的片上存储 │
│ - L1 Cache：自动缓存 │
│ - 比例可调：0/8/16/32/64/128 KB │
│ │
│ Warp Scheduler：4 个 │
│ 每个 Scheduler 每时钟周期可以发射一条指令 │
│ │
│ 计算单元： │
│ - INT32 ALU × 16（通用计算） │
│ - FP32 FMA × 16（浮点乘加） │
│ - LD/ST Unit × 16（内存访问） │
│ - Tensor Core × 4（矩阵乘法加速） │
│ │
└─────────────────────────────────────────────────────────────┘第3节：Warp——GPU 的基本调度单位
什么是 Warp？
┌─────────────────────────────────────────────────────────────┐
│ Warp：GPU 的"班级" │
├─────────────────────────────────────────────────────────────┤
│ │
│ Warp = 32 个线程为一组，一起执行同一条指令 │
│ │
│ 想象一个班级（32 人）： │
│ - 老师（Warp Scheduler）发一条指令："站起来" │
│ - 32 个学生同时站起来（32 个线程同时执行同一指令） │
│ - 完美并行！ │
│ │
│ 这就是 SIMT（Single Instruction Multiple Thread）： │
│ 单指令，多线程 │
│ │
│ 为什么是 32？ │
│ - 硬件设计权衡：太多线程调度开销大，太少并行度不够 │
│ - 32 是一个经验最优值 │
│ │
└─────────────────────────────────────────────────────────────┘Warp 的执行模型
┌─────────────────────────────────────────────────────────────┐
│ Warp 执行流程 │
├─────────────────────────────────────────────────────────────┤
│ │
│ 时间 → │
│ │
│ Clock 1: Warp Scheduler 选择一个 Warp │
│ ↓ │
│ Warp Scheduler 发射指令到执行单元 │
│ ↓ │
│ 32 个线程同时执行同一条 FP32 指令 │
│ ↓ │
│ 16 × 2 = 32 FLOP（一次发射完成） │
│ │
│ Clock 2: 下一个 Warp 被调度 │
│ ... │
│ │
│ 关键：每个时钟周期，Warp Scheduler 可以发射一条指令 │
│ 4 个 Scheduler × 每周期 1 条 = 每周期 4 条指令 │
│ │
└─────────────────────────────────────────────────────────────┘Warp 分支分化（Branch Divergence）——性能杀手
┌─────────────────────────────────────────────────────────────┐
│ Warp 分支分化 │
├─────────────────────────────────────────────────────────────┤
│ │
│ 场景：32 个线程执行 if-else │
│ │
│ if (threadIdx.x < 16) { ← 前16个线程走这里 │
│ a = b + c; │
│ } else { ← 后16个线程走这里 │
│ a = b - c; │
│ } │
│ │
│ 问题：Warp 中 32 个线程必须执行同一条指令 │
│ │
│ Clock 1: 执行 if 分支 → 只前16个线程干活，后16个等 │
│ Clock 2: 执行 else 分支 → 后16个线程干活，前16个等 │
│ │
│ 效果：2 个时钟周期完成 1 个 if-else │
│ = 效率降为 50%！ │
│ │
│ ┌─────────────────────────────────────────────────────┐ │
│ │ 理想：32 线程同时执行 = 100% 效率 │ │
│ │ 分支分化：只有 16 线程同时工作 = 50% 效率 │ │
│ │ │ │
│ │ 教训：GPU 代码中尽量避免 Warp 内的分支分化 │ │
│ │ 如果无法避免，让分支的线程数量是 32 的倍数 │ │
│ └─────────────────────────────────────────────────────┘ │
│ │
└─────────────────────────────────────────────────────────────┘数值例子
┌─────────────────────────────────────────────────────────────┐
│ 分支分化导致的性能损失 │
├─────────────────────────────────────────────────────────────┤
│ │
│ 场景：判断一个数是正数还是负数 │
│ │
│ // ❌ 差：每个线程都可能走不同分支 │
│ if (x > 0) { sum += sqrt(x); } │
│ else { sum += x * x; } │
│ │
│ // ✅ 好：强制所有线程走同一分支（通过 blockIdx 分组） │
│ if (blockIdx.x == 0) { │
│ // 所有 blockIdx==0 的线程走这里 │
│ sum += sqrt(x); │
│ } else { │
│ sum += x * x; │
│ } │
│ │
│ // ✅ 更好：完全没有分支 │
│ // 用 step 函数替代 if-else │
│ int branch = (x > 0); // 0 或 1 │
│ sum += branch * sqrt(x) + (1 - branch) * x * x; │
│ │
└─────────────────────────────────────────────────────────────┘第4节：GPU 内存层次——为什么 HBM 带宽是关键
GPU 内存与 CPU 内存的对比
┌─────────────────────────────────────────────────────────────┐
│ GPU vs CPU 内存对比 │
├─────────────────────────────────────────────────────────────┤
│ │
│ CPU 内存（DDR4/DDR5）： │
│ ┌──────────────────────────────────────────────────────┐ │
│ │ 容量：256 GB 带宽：~200 GB/s 延迟：~100ns │ │
│ │ │ │
│ │ 特点：容量大，延迟低，适合随机访问 │ │
│ │ 但带宽有限，多核同时访问容易成为瓶颈 │ │
│ └──────────────────────────────────────────────────────┘ │
│ │
│ GPU 显存（HBM2e）： │
│ ┌──────────────────────────────────────────────────────┐ │
│ │ 容量：40-80 GB 带宽：2.0 TB/s 延迟：~500ns │ │
│ │ │ │
│ │ 特点：带宽极高（CPU 的 10 倍！），但延迟高 │ │
│ │ 适合顺序的大块数据访问（矩阵乘法、卷积） │ │
│ └──────────────────────────────────────────────────────┘ │
│ │
└─────────────────────────────────────────────────────────────┘GPU 内存层次详解
┌─────────────────────────────────────────────────────────────┐
│ GPU 内存层次（由快到慢） │
├─────────────────────────────────────────────────────────────┤
│ │
│ L1 Cache (per SM) 128 KB ~1 ns ↑ 最快 │
│ Shared Memory (per SM) 128 KB ~1 ns （可编程） │
│ L2 Cache (per GPU) 40 MB ~30 ns │
│ Global Memory (HBM2e) 40-80 GB ~500 ns ↓ 最慢 │
│ │
│ 带宽对比（GB/s）： │
│ L1 Cache: ~10,000 GB/s │
│ Shared Memory: ~10,000 GB/s │
│ L2 Cache: ~3,500 GB/s │
│ Global Memory: 2,000 GB/s │
│ │
│ 容量对比： │
│ Register: 256 KB（per SM）← 每个线程的"寄存器" │
│ L1 Cache: 128 KB（per SM）← 自动缓存 │
│ L2 Cache: 40 MB（per GPU）← 自动缓存，全 SM 共享 │
│ Global Memory: 40 GB（per GPU）← HBM 显存 │
│ │
└─────────────────────────────────────────────────────────────┘为什么 GPU 延迟高但仍然快？
CPU 的策略：低延迟
- 优化单次访问的响应时间
- 用复杂的分支预测和超线程隐藏延迟
- 适合分支密集、不可预测的程序
GPU 的策略：高带宽 + 大规模并行隐藏延迟
- 一次发起 thousands 个内存请求
- Warp A 在等内存时，调度器切到 Warp B 执行
- 通过"切换"而不是"等待"来隐藏延迟
┌─────────────────────────────────────────────────────────────┐
│ 延迟隐藏技术 │
├─────────────────────────────────────────────────────────────┤
│ │
│ GPU 有 108 个 SM，每个 SM 同时跑多个 Warp： │
│ │
│ A100 每个 SM：2048 个线程 / 32 = 64 个 Warp │
│ A100 总共：108 × 64 = 6912 个 Warp 同时可调度 │
│ │
│ 当 Warp 0 等待内存： │
│ → 调度器立刻切到 Warp 1 执行 │
│ → Warp 1 等待内存时，切到 Warp 2 │
│ → ... │
│ → 永远不闲着！ │
│ │
│ 这就是 GPU 能容忍高延迟的原因： │
│ 足够多的线程 → 永远有其他线程可执行 │
│ │
└─────────────────────────────────────────────────────────────┘第5节：Tensor Core——矩阵乘法的专用硬件
为什么需要 Tensor Core？
传统 CUDA 核心（FP32 FMA）做矩阵乘法：
- 每个 SM 有 16 个 FP32 单元
- 每个单元每个时钟执行 1 次乘加 = 2 FLOP
- 每 SM 每时钟：16 × 2 = 32 FLOP
- A100 108 SM：108 × 32 = 3456 FLOP/时钟
- A100 主频 ~1.41 GHz
- 总算力：3456 × 1.41G ≈ 4.87 TFLOPS（FP32）
Tensor Core 做矩阵乘法：
- 每个 SM 有 4 个 Tensor Core
- 每个 Tensor Core 每个时钟执行 256 FLOP（FP16 矩阵乘）
- 每 SM 每时钟：4 × 256 = 1024 FLOP
- A100 108 SM：108 × 1024 = 110,592 FLOP/时钟
- 总算力：110592 × 1.41G ≈ 156 TFLOPS（FP16 Tensor）
Tensor Core 比传统 CUDA 核心快约 32 倍！Tensor Core 工作原理
┌─────────────────────────────────────────────────────────────┐
│ Tensor Core: D = A × B + C │
├─────────────────────────────────────────────────────────────┤
│ │
│ 输入： │
│ A: M×K 矩阵 B: K×N 矩阵 C: M×N 矩阵（累加） │
│ 输出： │
│ D: M×N 矩阵 = A × B + C │
│ │
│ A100 Tensor Core 支持的形状： │
│ - 16×16×16（FMA 模式） │
│ - 8×16×16（FP16） │
│ - 4×8×16（INT8） │
│ - 1×4×16（INT4） │
│ │
│ 以 16×16×16 为例： │
│ - 一次做 16×16 = 256 个输出元素 │
│ - 每个输出 = 16 次乘 + 15 次加 = 16 FMA │
│ - 总计 4096 FLOP │
│ - 由 4 个 Tensor Core + 16 个 FP32 单元协同完成 │
│ │
└─────────────────────────────────────────────────────────────┘Tensor Core 编程方式
cpp
// CUDA 中使用 Tensor Core 的方式
// 方式 1：cuBLAS 库（最简单，推荐）
#include <cublas_v2.h>
cublasHandle_t handle;
cublasCreate(&handle);
// 调用矩阵乘法，自动使用 Tensor Core
// A: M×K, B: K×N, C: M×N
cublasGemmEx(handle,
CUBLAS_OP_N, CUBLAS_OP_N, // 不转置
M, N, K, // 矩阵维度
&alpha, // 缩放因子
A, CUDA_R_16F, M, // A 矩阵
B, CUDA_R_16F, K, // B 矩阵
&beta, // 累加因子
C, CUDA_R_16F, M, // C 矩阵
CUDA_R_16F, // 计算精度
CUBLAS_GEMM_DEFAULT_TENSOR_OP); // 启用 Tensor Core
// 方式 2：WMMA API（直接操作 Tensor Core）
#include <mma.h>
using namespace nvcuda::wmma;
// 声明矩阵片段
fragment<matrix_a, ...> a_frag;
fragment<matrix_b, ...> b_frag;
fragment<accumulator, ...> c_frag;
// 加载矩阵到片段
load_matrix_sync(a_frag, A, K);
load_matrix_sync(b_frag, B, N);
// 执行矩阵乘法
mma_sync(c_frag, a_frag, b_frag, c_frag);
// 存储结果
store_matrix_sync(C, c_frag, N, mem_layout::row_major);第6节：显存带宽的实际影响
矩阵乘法的带宽分析
┌─────────────────────────────────────────────────────────────┐
│ 矩阵乘法 C = A × B 的带宽需求 │
├─────────────────────────────────────────────────────────────┤
│ │
│ A: M×K, B: K×N, C: M×N │
│ │
│ 数据量： │
│ 读取 A: M×K × 2 字节（FP16） │
│ 读取 B: K×N × 2 字节（FP16） │
│ 写入 C: M×N × 2 字节（FP16） │
│ 总计：2×(M×K + K×N) + M×N × 2 ≈ 2×(M×K + K×N) │
│ │
│ 计算量： │
│ M×N × K 次乘加 = 2×M×N×K FLOP（FP16） │
│ │
│ 算力需求（FLOP vs 带宽）： │
│ Roofline Model: 最大 FLOP = min(算力上限, 带宽 × 算术强度)│
│ │
│ 算术强度 = 计算量 / 数据量 │
│ = 2×M×N×K / (2×M×K + 2×K×N) │
│ ≈ MNK / (MK + KN) │
│ = K / (1 + K/N + K/M) │
│ │
│ 当 K >> M 且 K >> N 时（典型深度学习场景）： │
│ 算术强度 ≈ K（非常大！） │
│ → 带宽充足时，Tensor Core 可以跑满 │
│ │
└─────────────────────────────────────────────────────────────┘Roofline 模型
┌─────────────────────────────────────────────────────────────┐
│ Roofline 性能模型 │
├─────────────────────────────────────────────────────────────┤
│ │
│ 性能 (TFLOPS) │
│ ↑ │
│ 312 ┤ ★ A100 │
│ │ / │
│ 156 ┤─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ Tensor Core│
│ │ /│ FP16 │
│ 20 ┤─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ │
│ │ / │
│ 10 ┤─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ │
│ │ /│ FP32 │
│ 5 ┤─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ ─ │
│ │ / │
│ 0 ┼──────────────────────────────────────────────────→ │
│ 0 200 500 1000 2000 算术强度 │
│ (FLOP/Byte) │
│ │
│ 关键点： │
│ - 算术强度 < 312 时，性能受限于带宽（A100 HBM） │
│ - 算术强度 > 312 时，性能受限于算力 │
│ - 矩阵乘法算术强度很大 → 受限于算力 │
│ │
└─────────────────────────────────────────────────────────────┘升华：GPU 架构的核心记忆点
┌─────────────────────────────────────────────────────────────┐
│ GPU 架构速记卡 │
├─────────────────────────────────────────────────────────────┤
│ │
│ A100：6912 CUDA 核心 = 108 SM × 64 线程/SM │
│ │
│ 调度层次： │
│ Grid → Block → Warp → Thread │
│ (整体) (SM) (32线程) (最小单元) │
│ │
│ 执行模型： │
│ - SIMT：单指令多线程 │
│ - Warp：32 线程为一组，同步执行 │
│ - 分支分化 = 效率损失 │
│ │
│ 内存层次（由快到慢）： │
│ Register → L1/Shared → L2 → Global Memory (HBM) │
│ │
│ Tensor Core： │
│ - 比 CUDA 核心快 32 倍 │
│ - 专门做矩阵乘法 D = A×B+C │
│ - 深度学习训练的加速核心 │
│ │
│ 延迟隐藏： │
│ 6912 个 Warp 随时可调度，永远不闲着 │
│ │
└─────────────────────────────────────────────────────────────┘"AI 可查 vs 必须理解"清单
AI 可查：
✅ A100/H100 的具体核心数、频率、带宽——官网规格表
✅ Tensor Core 支持的具体精度（FP16/BF16/INT8/INT4）
✅ cuBLAS/cuDNN 的具体 API 参数
必须理解：
🔴 SM × Warp × Thread 的三层调度结构
🔴 Warp = 32 线程，同步执行，分支分化 = 效率损失
🔴 GPU 内存层次：Register < L1/Shared < L2 < Global Memory
🔴 为什么 HBM 带宽 2 TB/s 是关键指标
🔴 Tensor Core 比 CUDA 核心快 32 倍，是矩阵乘法的专用硬件
🔴 GPU 用大量 Warp 隐藏内存延迟学习状态：🟡 开始学习