GPU 体系结构：SIMT、Warp、SM 与吞吐优先设计 / GPU Architecture with SIMT, Warps, and Streaming Multiprocessors

📅 创建时间：2026-07-20 🏷️ 标签：#GPU #SIMT #Warp #SM #CUDA 📚 前置知识：[[04-memory-hierarchy]] 📚 相关知识：[[/02-systems-and-performance/02-computer-architecture-and-hardware/03-gpu-architecture]]

1. GPU 为什么拥有大量计算单元

CPU 需要高效处理复杂控制流、操作系统、中断和低延迟任务，因此每个核心包含复杂的预测、乱序和缓存结构。

GPU 面对的是大量相似任务：

text

为每个像素着色
为每个向量元素做运算
为矩阵中的大量输出元素做乘加

GPU 减少单线程控制资源，把更多芯片面积用于算术单元，通过同时维护大量线程隐藏等待延迟。

2. SIMT 执行模型

SIMT 是 Single Instruction, Multiple Threads：

编程时看到许多独立线程
硬件把线程分组执行
同组线程通常执行同一条指令

NVIDIA GPU 通常以 32 个线程组成一个 Warp。

text

Warp 0：Thread 0 ... Thread 31
Warp 1：Thread 32 ... Thread 63

它与 SIMD 相似，但编程抽象是线程，硬件负责将线程映射到执行通道。

3. SM：GPU 的基本执行组织

Streaming Multiprocessor（SM）可以理解为 GPU 上负责调度和执行线程块的一组资源，通常包含：

Warp 调度器
FP32/INT 执行单元
Tensor Core
寄存器文件
Shared Memory / L1 Cache
Load/Store 单元

一个 GPU 包含多个 SM。线程块会被分配到某个 SM 上执行，通常不会在执行中迁移到其他 SM。

4. 用大量线程隐藏延迟

当 Warp A 等待显存数据时，SM 可以切换到已经就绪的 Warp B：

text

Warp A：等待内存 ───────────┐
Warp B：执行计算 ████████   │ 交错执行
Warp C：准备执行     ███████│

这种切换由硬件完成，成本远低于操作系统线程上下文切换。

前提是 SM 上有足够多的活跃 Warp。如果单个线程使用过多寄存器或单个 Block 使用过多 Shared Memory，可同时驻留的 Block 数会下降。

5. 分支发散

同一个 Warp 中的线程走不同分支时，硬件可能需要分别执行各条路径：

cpp

if (threadIdx.x % 2 == 0)
    pathA();
else
    pathB();

概念上变成：

text

先让偶数线程执行 A，奇数线程等待
再让奇数线程执行 B，偶数线程等待

分支本身不是禁止项，真正的问题是同一 Warp 内分支不一致且两条路径工作量较大。

6. GPU 内存层次

text

寄存器          每线程私有，最快，容量有限
Shared Memory   每 Block 共享，可编程管理
L1/L2 Cache     硬件缓存
Global Memory   GPU 显存，容量大、延迟高
Host Memory     CPU 内存，通常需通过 PCIe/NVLink 访问

GPU 优化的核心经常是：

合并全局内存访问
把重复使用的数据放入 Shared Memory
控制寄存器使用
减少 CPU 与 GPU 间的数据往返

7. 合并访存

如果一个 Warp 中相邻线程访问相邻地址，硬件可以把访问合并成较少的内存事务：

text

线程 0 → a[0]
线程 1 → a[1]
...
线程 31 → a[31]   ✓ 连续、容易合并

跨步或随机访问会浪费带宽：

text

线程 0 → a[0]
线程 1 → a[1024]
线程 2 → a[2048]  ✗ 分散访问

8. Tensor Core 与专用数据通路

现代 GPU 不只有通用浮点单元，还包含针对矩阵乘加设计的 Tensor Core。它们以块为单位执行：

text

D = A × B + C

使用 FP16、BF16、TF32、FP8 等低精度格式可以显著提高吞吐，但必须考虑数值范围、累加精度和算法稳定性。

9. GPU 不适合什么

数据规模很小，启动开销占主导
分支高度不规则
每个任务工作量差异巨大
大量指针追逐和随机访问
CPU-GPU 传输比计算本身更贵
必须频繁进行全局同步

GPU 的核心优势不是“单个核心更快”，而是让海量相似工作以高吞吐方式执行。

历史 GPU 架构补充

本页保留课程合并前的 GPU 架构说明。

当前版本增加 SOL、功耗、虚拟化、拓扑和工具流程。

SM

SM 包含调度、执行、Register、Shared Memory 与 Cache 资源。

Kernel 的资源使用决定驻留能力。

Warp

Warp 以共同指令执行线程。

就绪 Warp 用于隐藏延迟。

发散

同一 Warp 不同分支路径会分阶段执行。

分区和专用 Kernel 是候选方案。

Register

线程私有且快速。

过多使用降低驻留或导致 Spill。

Shared Memory

Block 内共享，适合 Tile 与复用。

注意 Bank、容量和 Barrier。

Global Memory

容量大，访问应尽量合并。

线程索引与布局共同决定事务。

Cache

L1/TEX 和 L2 减少部分重复访问。

命中率需要结合字节和时间。

Occupancy

Occupancy 是驻留比例，不是性能分数。

足够隐藏延迟后继续提高可能无益。

Warp Stall

常见等待：

Memory；
Barrier；
Dependency；
Execution Unit；
Atomic；
Instruction Fetch。

Stall 指标需要结合源码和指令。

SOL

Memory SOL 高提示接近内存路径上限。

Compute SOL 高提示接近计算路径上限。

两者低可能是延迟和并行度问题。

Tensor Core

针对特定矩阵和精度。

检查布局、Tile、累加与数值。

Copy Engine

传输与计算能否重叠取决于 Pinned Memory、Stream 和硬件。

多 GPU

考虑 PCIe、NVLink、P2P、NUMA 和集合通信。

功耗

频率受功耗与温度限制。

长任务记录稳态而非短时峰值。

共享环境

多进程与容器会引入队列、Context 和显存竞争。

端到端延迟包含等待。

Nsight Systems

查看 CPU/GPU、传输、Kernel、Stream、同步和 Idle。

Nsight Compute

查看单 Kernel 的 Warp、SM、Memory 和指令。

先系统后 Kernel。

实验记录

GPU；
Driver；
Toolkit；
Compute Capability；
输入；
Grid/Block；
Register；
Shared Memory；
Stream；
Profile；
正确性。

历史参数边界

SM 数、Warp、Cache、Tensor Core 和指标随架构变化。

旧型号阈值不能直接外推。

当前课程映射

当前文章补充：

资源细节；
SOL；
Tensor Core；
Copy Engine；
多 GPU；
功耗；
虚拟化；
Profile；
实验和完成标准。

阅读完成标准

应能：

解释 SM/Warp；
分析发散；
理解内存层次；
判断 Occupancy；
解释 SOL；
分析拓扑；
使用 Nsight；
限定历史结论。

归档验证清单

架构与设备记录；
Driver/Toolkit 记录；
Kernel 配置记录；
资源使用记录；
传输计入时间；
Profile 原始文件；
CPU 参考结果；
数值容差；
功耗与温度；
多 GPU 拓扑；
共享环境说明；
当前文档入口；
旧指标版本边界；
构建命令；
输入校验和；
原始样本。所有结论都绑定记录的 GPU 架构。

核心总结

GPU 用大量线程和执行单元换取吞吐。
Warp 是理解分支和访存行为的关键层次。
SM 资源限制决定同时活跃的线程数量。
合并访存、数据复用和分支一致性直接影响效率。
Tensor Core 是矩阵计算的专用高吞吐通路。

下一篇：[[06-cuda-programming-model]]

原体系结构与硬件 / Original Architecture

原并行计算 / Original Parallel Computing

GPU 体系结构：SIMT、Warp、SM 与吞吐优先设计 / GPU Architecture with SIMT, Warps, and Streaming Multiprocessors

1. GPU 为什么拥有大量计算单元

2. SIMT 执行模型

3. SM：GPU 的基本执行组织

4. 用大量线程隐藏延迟

5. 分支发散

6. GPU 内存层次

7. 合并访存

8. Tensor Core 与专用数据通路

9. GPU 不适合什么

历史 GPU 架构补充

SM

Warp

发散

Register

Shared Memory

Global Memory

Cache

Occupancy

Warp Stall

SOL

Tensor Core

Copy Engine

多 GPU

功耗

共享环境

Nsight Systems

Nsight Compute

实验记录

历史参数边界

当前课程映射

阅读完成标准

归档验证清单

核心总结

GPU 体系结构：SIMT、Warp、SM 与吞吐优先设计 / GPU Architecture with SIMT, Warps, and Streaming Multiprocessors ​

1. GPU 为什么拥有大量计算单元 ​

2. SIMT 执行模型 ​

3. SM：GPU 的基本执行组织 ​

4. 用大量线程隐藏延迟 ​

5. 分支发散 ​

6. GPU 内存层次 ​

7. 合并访存 ​

8. Tensor Core 与专用数据通路 ​

9. GPU 不适合什么 ​

历史 GPU 架构补充 ​

SM ​

Warp ​

发散 ​

Register ​

Shared Memory ​

Global Memory ​

Cache ​

Occupancy ​

Warp Stall ​

SOL ​

Tensor Core ​

Copy Engine ​

多 GPU ​

功耗 ​

共享环境 ​

Nsight Systems ​

Nsight Compute ​

实验记录 ​

历史参数边界 ​

当前课程映射 ​

阅读完成标准 ​

归档验证清单 ​

核心总结 ​

GPU 体系结构：SIMT、Warp、SM 与吞吐优先设计 / GPU Architecture with SIMT, Warps, and Streaming Multiprocessors

1. GPU 为什么拥有大量计算单元

2. SIMT 执行模型

3. SM：GPU 的基本执行组织

4. 用大量线程隐藏延迟

5. 分支发散

6. GPU 内存层次

7. 合并访存

8. Tensor Core 与专用数据通路

9. GPU 不适合什么

历史 GPU 架构补充

SM

Warp

发散

Register

Shared Memory

Global Memory

Cache

Occupancy

Warp Stall

SOL

Tensor Core

Copy Engine

多 GPU

功耗

共享环境

Nsight Systems

Nsight Compute

实验记录

历史参数边界

当前课程映射

阅读完成标准

归档验证清单

核心总结