第二十七课：GPU 架构学习总结与统一分析框架

导语：最后一课将把前面所有内容压缩成一套统一分析框架，并给出面向 GPU 架构与 RTL 设计的后续实践路线。

学习 GPU 架构的最终目标，不是记住多少型号、核心数量和术语，而是面对一个新问题时，能够回答：

这个任务的并行性在哪里？
线程怎样映射到硬件？
数据存放在哪里？
真正限制性能的资源是什么？
应当优化算法、数据流，还是硬件结构？

前面二十六课可以归纳为五个层次。

算法层
  ↓
并行编程模型
  ↓
SM 微架构
  ↓
整芯片存储与互连
  ↓
系统与多 GPU

一、GPU 的统一抽象

可以把 GPU 理解成：

通过保存大量线程状态、快速切换可执行 Warp，以高吞吐率执行规则数据并行任务的处理器。

其核心结构是：

大量线程
   ↓
组成 Warp
   ↓
多个 Warp 驻留在 SM
   ↓
Warp Scheduler 选择 Ready Warp
   ↓
CUDA Core / Tensor Core / LSU 等执行
   ↓
通过 Cache、L2 和显存获得数据

GPU 与 CPU 的根本差异不只是核心数量。

CPU 的主要目标是：

让少量线程尽快完成

GPU 的主要目标是：

让大量线程总体完成得更快

因此：

CPU：优化单线程延迟
GPU：优化整体吞吐率

二、GPU 性能的基本公式

任何 GPU Kernel 的时间都可以抽象为几个因素共同决定：

T\approx \max \left( T_{\text{compute}}, T_{\text{memory}}, T_{\text{communication}} \right) + T_{\text{latency exposed}} + T_{\text{sync}} + T_{\text{launch}}

其中：

$T_{\text{compute}}$ ：算术执行时间；
$T_{\text{memory}}$ ：数据搬运时间；
$T_{\text{communication}}$ ：多 GPU 通信时间；
$T_{\text{latency exposed}}$ ：无法隐藏的访存或依赖延迟；
$T_{\text{sync}}$ ：Barrier、Atomic 等同步成本；
$T_{\text{launch}}$ ：Kernel 启动与调度开销。

判断性能瓶颈，就是找出其中最大的部分。

三、看到一个算法时，先问哪五个问题

1. 有多少独立任务

例如向量加法：

每个元素相互独立

并行度很高。

串行递推：

x_i=f(x_{i-1})

存在强依赖，并行度低。

2. 线程之间是否需要通信

例如：

Elementwise：
基本不需要通信

Reduction：
需要逐步合并

Matrix Multiply：
需要共享和复用 Tile

Softmax：
需要行内 Max 和 Sum Reduction

通信需求决定是否需要：

Warp Shuffle；
Shared Memory；
Barrier；
Atomic；
多级 Kernel。

3. 数据能否重复使用

例如矩阵乘法中，一个 Tile 可以反复参与计算，所以算术强度很高。

向量加法中，一个元素通常只读一次、计算一次、写一次，数据复用低。

AI= \frac{\text{Operations}} {\text{Memory Bytes}}

算术强度是判断 GPU 算法性能的重要指标。

4. 访问地址是否规则

GPU 喜欢：

相邻线程
访问相邻地址

GPU 不喜欢：

相邻线程
通过随机索引访问完全不同的位置

规则访存有利于：

合并事务；
Cache；
预取；
内存分区均衡；
高显存带宽。

5. 不同线程工作量是否一致

如果一个 Warp 中：

Lane 0：处理 2 个元素
Lane 1：处理 4 个元素
Lane 2：处理 2000 个元素

其他 Lane 会长时间等待。

GPU 的高性能依赖：

同一 Warp 中的线程
执行相近的控制流和工作量

四、看到一张 GPU 架构图应该怎么看

不要先数 CUDA Core。

建议按照以下顺序。

第一层：找到计算基本单元

看：

SM 有多少个；
SM 如何组成 GPC、TPC 或其他分区；
是否有不同类型的计算 Tile。

SM 决定通用计算资源的总体规模。

第二层：观察 SM 内部

重点看：

Warp Scheduler 数量
Register File
FP32 / INT 单元
Tensor Core
Load/Store Unit
SFU
Shared Memory / L1

需要回答：

每周期能发射多少指令；
有多少独立执行流水线；
一个 Warp 如何映射到这些流水线；
寄存器和 Shared Memory 容量多大。

第三层：观察存储层次

按距离 SM 的远近分析：

寄存器
Shared Memory / L1
L2 Cache
HBM / GDDR
Host Memory

每下降一级，通常：

容量增大；
延迟增加；
每字节能耗增加；
软件控制能力下降。

第四层：观察内存分区

找出：

L2 Cache Slice
Memory Partition
Memory Controller
HBM/GDDR Channel

需要考虑：

地址如何分散到不同分区；
总带宽如何形成；
是否存在分区热点；
ROP、Atomic 等结构位于哪里。

第五层：观察片上互连

计算能力再强，也需要片上网络输送数据。

需要判断：

SM 到 L2 如何连接；
L2 与内存控制器如何组织；
多 Die 之间如何通信；
是否可能出现片上网络拥塞。

第六层：观察专用硬件

例如：

Tensor Core；
RT Core；
Texture Unit；
Rasterizer；
ROP；
Video Engine；
Copy Engine；
TMA。

专用单元通常针对高频固定操作，以降低面积和能耗。

五、看到一个 CUDA Kernel 应该怎么分析

可以使用以下统一流程。

第一步：确定并行映射

回答：

一个 Thread 负责什么？
一个 Warp 负责什么？
一个 Block 负责什么？
一个 Grid 覆盖什么？

例如矩阵乘法：

一个 Thread：计算若干输出元素
一个 Warp：计算一个输出子 Tile
一个 Block：计算一个较大 C Tile
一个 Grid：覆盖整个输出矩阵

如果无法清楚回答这些问题，说明线程映射尚未设计成熟。

第二步：画出数据流

明确数据从哪里来，到哪里去：

Global Memory
    ↓
Shared Memory
    ↓
Register
    ↓
ALU / Tensor Core
    ↓
Register
    ↓
Global Memory

记录每个数据：

读取几次；
复用几次；
是否写入中间数组；
是否能保留在寄存器或 Shared Memory。

第三步：分析访存模式

检查一个 Warp 内地址：

连续？
对齐？
跨多少 Cache Line？
是否随机 Gather？
是否存在 Scatter？
是否存在 Partition Camping？

然后检查 Shared Memory：

是否产生 Bank Conflict？
是否可以广播？
是否需要 Padding？

第四步：分析控制流

检查：

分支条件是否随 Lane 变化；
循环次数是否一致；
是否存在提前退出；
是否有大规模 Predicate 屏蔽；
是否存在不规则任务长度。

核心问题是：

同一个 Warp 的有效 Lane 有多少？

第五步：分析依赖关系

区分：

线程间数据依赖
单线程指令依赖
Warp 间同步依赖
Block 间依赖

例如：

sum = sum * a + b;

存在单线程长依赖链。

Reduction 则存在跨线程合并依赖。

第六步：分析资源占用

检查：

每线程寄存器；
每 Block Shared Memory；
Block 线程数；
Warp 数；
Barrier 资源；
Cluster 资源。

资源决定一个 SM 能同时驻留多少 Block 和 Warp。

第七步：判断理论瓶颈

通过算术强度和数据流初步判断：

计算受限
带宽受限
访存延迟受限
依赖延迟受限
同步受限
任务规模受限

之后再用 Profiler 验证。

第八步：实际测量

使用：

Nsight Systems：
确定时间花在哪里

Nsight Compute：
确定 Kernel 内部为什么慢

最终优化标准永远是：

\text{kernel duration decreases}

而不是某个单独指标是否更漂亮。

六、六类常见性能瓶颈

1. 计算吞吐率受限

表现：

Compute Throughput 高
Tensor/FP32/FP64/SFU 某条管线接近满载
Memory Throughput 相对较低

优化方向：

Tensor Core；
更低精度；
减少冗余运算；
使用 FMA/MMA；
平衡不同执行流水线；
改善指令级并行。

典型算子：

大型 GEMM；
高复用卷积；
Tensor Core Attention。

2. 显存带宽受限

表现：

DRAM Throughput 接近上限
计算吞吐率不高
访问相对连续

优化方向：

减少字节数；
Kernel Fusion；
低精度；
寄存器或 Shared Memory 复用；
避免中间张量；
压缩权重。

典型算子：

Vector Add；
LayerNorm；
RMSNorm；
大型 Reduction；
解码阶段 GEMV。

3. 访存延迟受限

表现：

DRAM 带宽没有跑满
Long Scoreboard 高
Eligible Warp 少
访存随机或细碎

优化方向：

重排数据；
改善局部性；
合并事务；
增加 Memory-Level Parallelism；
减少指针链；
使用更合适的稀疏格式。

典型算子：

SpMV；
图计算；
Embedding Lookup；
随机 Gather。

4. 指令依赖受限

表现：

Compute 和 Memory 都不高
Wait 较高
存在长串行依赖链

优化方向：

多累加器；
循环展开；
软件流水；
提前加载；
增加独立指令；
改变算法结构。

5. 同步与冲突受限

表现：

Barrier 高
Atomic 竞争明显
Shared Memory Conflict
不同 Warp 工作时间差异大

优化方向：

Warp 级原语；
缩小同步范围；
Warp Aggregated Atomic；
分层归约；
Padding；
重构任务分配；
降低热点地址竞争。

6. 并行度与任务规模受限

表现：

Grid 很小
Waves Per SM 很低
计算和内存利用率都低
Kernel 很短

优化方向：

Batch；
Kernel Fusion；
合并小任务；
Persistent Kernel；
CUDA Graph；
CPU 执行小任务。

七、不要混淆带宽受限与延迟受限

两者都可能表现为线程在等待内存，但含义不同。

带宽受限

像高速公路上所有车道都满了：

请求很多
地址规则
显存持续满载

优化目标是减少总数据量。

延迟受限

像每辆车都需要走完全不同的小路：

请求少而分散
等待时间长
总带宽却未充分利用

优化目标是改善访问规则和并行请求数量。

八、Occupancy 的正确地位

Occupancy 是：

\frac{\text{Active Warps}} {\text{Maximum Warps}}

它的作用是提供更多可切换 Warp，以隐藏延迟。

但 Occupancy 不是性能本身。

Occupancy 高：
可能拥有更多候选 Warp

Occupancy 低：
可能仍然有足够 Warp 隐藏延迟

高性能矩阵乘法可能因为大量寄存器分块而只有中等 Occupancy，却能使 Tensor Core 接近满载。

因此：

Occupancy 是手段，不是目标。

九、GPU 优化中的三个核心交换

几乎所有 GPU 优化都可以归结为以下交换。

1. 用空间换时间

例如：

更多寄存器
更多 Shared Memory
    ↓
减少 Global Memory 访问

代价是 Occupancy 下降。

2. 用计算换数据搬运

例如 FlashAttention：

重新计算局部分数
    ↓
避免保存和读取 N×N 矩阵

现代 GPU 算力增长快，数据搬运相对昂贵，因此这种交换经常有效。

3. 用通用性换效率

例如：

CUDA Core → 通用
Tensor Core → 矩阵专用
RT Core → 求交专用
Texture Unit → 采样专用

ASIC 则把这种交换推向更极端。

十、几个典型算子的统一比较

算子	主要并行对象	主要复用	典型瓶颈
Vector Add	元素	很少	显存带宽
GEMM	输出 Tile	A/B Tile	计算吞吐
Convolution	输出像素和通道	输入、权重	计算或带宽
Reduction	输入元素	较少	带宽、同步
Softmax	行	输入行	带宽、归约、SFU
LayerNorm	Token 行	输入行	带宽、归约
FlashAttention	Q/K/V Tile	Q、K、V Tile	Tensor Core、片上资源
SpMV	非零元素	很少	随机访存延迟
GEMV	输出元素	输入向量	权重带宽
MoE	Token—Expert	Expert 权重	通信、负载均衡

十一、从算法映射到 GPU 的统一模板

面对一个新算法，可以按下面顺序设计。

1. 选择并行维度

找出：

Batch
Token
Channel
Spatial
Matrix Row/Column
Nonzero Element

中的哪些维度可以并行。

2. 选择线程粒度

决定：

一个线程计算一个元素
一个线程计算多个元素
一个 Warp 处理一行
一个 Block 处理一个 Tile

粒度过小：

启动和调度开销高；
数据复用差。

粒度过大：

寄存器过多；
并行 Block 减少；
负载不均衡。

3. 设计片上数据复用

决定哪些数据放入：

Register
Shared Memory
L1
L2

优先复用最昂贵的数据。

4. 设计协作和同步

选择：

Warp Shuffle；
Shared Memory；
__syncthreads()；
Atomic；
多 Kernel；
Cooperative Groups。

原则是使用能满足正确性的最小同步范围。

5. 处理边界和不规则情况

例如：

矩阵尺寸不是 Tile 整数倍；
最后一个 Warp 不完整；
序列长度不同；
稀疏行长度不均；
Mask；
Padding。

边界处理不应破坏主路径效率。

6. 评估资源和性能

估算：

\text{Registers per Thread}

\text{Shared Memory per Block}

\text{Blocks per SM}

AI=\frac{\text{FLOPs}}{\text{Bytes}}

然后再进入实现和测量。

十二、从算法映射到 ASIC 的统一框架

GPU 和 ASIC 的分析方法有大量共通之处。

首先画出算法数据流：

输入
  ↓
局部缓冲
  ↓
计算阵列
  ↓
归约或激活
  ↓
输出缓冲

然后回答：

主要运算是什么？
数据能复用多少次？
片上 SRAM 需要多大？
外部 DRAM 每次搬运多少？
数据路径需要多宽？
是否需要多级流水？
不规则控制如何处理？
精度和位宽如何选择？
计算阵列利用率是多少？
带宽能否持续供数？

十三、GPU 与 NPU 的本质差异

GPU 的核心机制是：

保存大量线程上下文
动态选择 Ready Warp
通过线程切换隐藏延迟

NPU 的核心机制通常是：

提前规划 Tensor 数据流
DMA 搬运大块数据
矩阵阵列持续计算
尽量避免细粒度动态调度

因此：

特性	GPU	NPU
灵活性	高	中等
控制方式	线程和指令	张量和数据流
不规则任务	相对较好	通常较弱
矩阵能效	高	通常更高
软件生态	成熟	依赖厂商
硬件控制开销	较高	较低
开发适应性	强	面向特定算子

十四、CPU、GPU、FPGA、ASIC 如何选择

平台	最适合的场景	主要优势	主要限制
CPU	复杂控制、小任务、串行逻辑	灵活、低延迟	并行吞吐率有限
GPU	大规模规则并行、AI训练	通用性与吞吐率平衡	功耗高、控制成本高
FPGA	定制流水、低批量、快速迭代	可重构、接口灵活	频率和开发效率有限
ASIC	固定算法、大规模部署	最高能效和密度	成本高、修改困难
NPU	规则神经网络推理	AI能效高	算子和控制灵活性有限

选择平台时不能只比较 TOPS。

还要考虑：

算法是否稳定
部署数量
功耗限制
延迟要求
开发周期
精度要求
软件生态
接口需求

十五、什么时候适合使用 GPU

优先考虑 GPU 的情况：

算法仍在快速迭代；
需要支持多种模型；
任务规模大；
并行性强；
数据中心已有 GPU；
需要成熟 AI 软件生态；
开发时间有限；
功耗不是最严格限制。

十六、什么时候值得做 FPGA 或 ASIC

更适合定制硬件的情况：

算法和数据流已经稳定；
部署量足够大；
功耗要求严格；
实时延迟要求高；
数据类型和精度可固定；
通用 GPU 的控制和存储成本过高；
算法具有大量特殊结构；
需要与传感器或片上系统紧密集成。

你的 SNN、在线学习和专用 CNN 加速器项目，就属于适合从 GPU 数据流逐渐抽象到 ASIC 的任务。

十七、GPU RTL 项目的推荐实践路线

最适合当前阶段的项目不是立即实现完整 CUDA GPU，而是完成一个可验证、可综合的小型 SIMT Core。

建议目标：

4 Lane
4 个驻留 Warp
16 个寄存器/线程
单个 SM
整数 ISA
Global Memory
Shared Memory
Barrier
Predicate
简单分支发散

Phase 0：规格冻结

明确：

Warp Size；
最大 Warp 数；
寄存器数量；
指令编码；
存储器接口；
支持的 Kernel；
分支和 Barrier 语义。

输出：

architecture_spec.md
isa_spec.md
memory_model.md

Phase 1：软件模拟器

用 Python 或 C++ 实现：

Instruction Decoder
Warp State
Register File
Active Mask
Memory
Barrier
Divergence Stack

先让软件模型正确运行：

Vector Add；
ReLU；
Shared Memory Reverse；
Reduction。

Phase 2：单 Warp RTL

实现：

Instruction Fetch；
Decoder；
Vector Register File；
4-Lane ALU；
Writeback；
EXIT。

先不支持长延迟和调度。

Phase 3：多 Warp 与调度

加入：

Warp Context Table；
Round-Robin Scheduler；
多 Warp PC；
Active Mask。

验证 Warp 可以交错执行。

Phase 4：Global Memory 与 Scoreboard

加入：

LDG；
STG；
固定延迟 Memory Model；
请求队列；
Scoreboard。

运行 Vector Add。

Phase 5：Shared Memory 与 Barrier

加入：

Banked Shared Memory；
LDS/STS；
Bank Conflict；
Block Barrier。

运行 Block Reduction。

Phase 6：分支发散

加入：

Predicate Register；
True/False Mask；
Reconvergence Stack；
部分 Lane EXIT。

运行分支测试。

Phase 7：综合与性能统计

加入硬件计数器：

总周期
发射指令数
ALU Busy Cycle
Memory Stall Cycle
Barrier Stall Cycle
Eligible Warp 数
Bank Conflict 数

这些计数器相当于简化版 Nsight Compute。

十八、GPU RTL 中最值得研究的三个问题

如果需要形成论文或研究项目，以下方向比单纯“实现一颗 GPU”更有研究价值。

1. Warp 调度策略

比较：

Round-Robin；
Greedy-Then-Oldest；
Memory-Aware；
Criticality-Aware；
Cache-Conscious。

观察对不同 Kernel 的影响。

2. 寄存器与 Occupancy

研究：

静态寄存器分配；
动态寄存器分配；
Register File 分区；
Operand Collector；
Bank Conflict；
Register Cache。

3. 存储与稀疏数据流

研究：

Memory Coalescer；
Shared Memory Bank；
Cache；
稀疏 Gather；
事件驱动调度；
原子聚合。

这些方向与你的 AI 芯片和 SNN 项目关联较强。

十九、后续阅读顺序

不需要一次性阅读所有文档，可以按照以下顺序。

第一组：CUDA 编程模型

重点阅读：

CUDA C++ Programming Guide
CUDA Best Practices Guide
CUDA Occupancy 相关章节

目标：

建立 Thread、Warp、Block、Memory 的准确语义。

第二组：指令和编译器

阅读：

PTX ISA
CUDA Binary Utilities
nvcc Compiler Driver 文档

配合：

nvcc -ptx
cuobjdump -sass
nvdisasm

目标：

从 CUDA C++ 追踪到真实指令。

第三组：性能分析

阅读：

Nsight Compute Profiling Guide
Nsight Systems User Guide
Roofline 相关资料

实际分析：

Vector Add；
Reduction；
Matrix Multiply；
Softmax。

第四组：架构白皮书

建议按演进顺序选择性阅读：

Fermi
Kepler
Volta
Ampere
Hopper
Blackwell

不需要逐字阅读。

重点寻找：

SM 结构；
Tensor Core；
Cache；
数据搬运；
多 GPU 互连；
新增硬件解决了什么瓶颈。

第五组：开源 GPU RTL

阅读开源项目时重点观察：

Warp Context
Scheduler
Scoreboard
Register File
LSU
Memory Coalescer
Barrier
SIMT Stack

不要一开始陷入完整工程的外设、驱动和工具链细节。

先沿一条指令的执行路径阅读：

Fetch
→ Decode
→ Schedule
→ Register Read
→ Execute
→ Writeback

二十、常见错误认识汇总

错误一：CUDA Core 越多就一定越快

实际还受：

频率；
指令吞吐；
显存；
Cache；
调度；
软件；
功耗；

影响。

错误二：Occupancy 越高越好

Occupancy 只表示驻留 Warp 数量，不表示执行单元一定忙碌。

错误三：GPU 会自动并行任何程序

程序必须有足够独立任务，并正确映射到 Thread、Warp 和 Block。

错误四：线程越多性能越高

如果显存已经饱和，增加线程没有意义。

线程过多还可能增加资源竞争。

错误五：Shared Memory 一定比 Global Memory 快

Shared Memory 本身低延迟，但：

需要额外搬运；
需要同步；
可能 Bank Conflict；
占用有限资源。

只有存在足够复用时才值得使用。

错误六：分支一定很慢

如果 Warp 内线程条件一致，分支不会产生发散。

真正的问题是同一 Warp 走不同路径。

错误七：稀疏一定比稠密快

稀疏可能引入索引、随机访存和负载不均衡。

错误八：Tensor Core 会自动加速所有浮点代码

只有适合矩阵块运算、数据类型和布局满足要求时，才能使用 Tensor Core。

错误九：Profiler 中某个 Stall 很高，就一定是首要问题

只有当 Scheduler 缺少 Eligible Warp 时，Stall 才可能暴露为实际性能损失。

错误十：峰值算力等于真实算力

真实性能由：

\text{algorithm mapping} + \text{dataflow} + \text{hardware utilization}

决定。

二十一、最终统一检查表

面对一个 GPU 问题，可以依次检查。

算法

并行度是否足够？
数据依赖在哪里？
计算复杂度是多少？

线程映射

Thread、Warp、Block 各自处理什么？
Warp 内工作量是否一致？

访存

是否合并？
是否对齐？
是否存在随机 Gather/Scatter？
是否有复用？

片上存储

哪些数据放寄存器？
哪些放 Shared Memory？
是否 Bank Conflict？

控制流

是否分支发散？
是否有提前退出？
是否需要 Predicate？

同步

是否需要 Barrier？
能否缩小到 Warp 范围？
Atomic 是否存在热点？

资源

寄存器多少？
Shared Memory 多少？
理论 Occupancy 多少？
是否 Spill？

性能

计算受限？
带宽受限？
延迟受限？
同步受限？
并行度不足？

测量

Duration 是否改善？
数据结果是否正确？
是否只优化了指标而没有优化时间？

二十二、整个课程最核心的十句话

GPU 通过大量线程和 Warp 切换隐藏延迟。
Warp 是调度单位，Thread 是编程和数据语义单位。
分支发散的本质是同一 Warp 的 Lane 需要执行不同路径。
Global Memory 性能主要取决于 Warp 的地址模式，而不只是单次访问延迟。
Shared Memory 的价值来自数据复用，不只是因为它速度快。
Occupancy 的价值是提供足够的 Eligible Warp，不是追求100%。
Tensor Core 的价值是用专用矩阵数据通路提高计算密度和能效。
现代 GPU 优化越来越重视数据搬运，而不只是增加算术单元。
高性能算法往往通过 Tiling、Fusion 和异步流水减少 HBM 访问。
GPU 架构分析最终都可以归结为：

\boxed{ \text{parallelism} + \text{data reuse} + \text{scheduling} + \text{memory hierarchy} + \text{hardware utilization} }

课程结束后的合理目标

完成这套课程后，应当具备以下能力：

能解释 Thread、Warp、Block 和 SM 的关系；
能理解 GPU 如何隐藏访存延迟；
能分析分支发散、数据冒险和同步；
能理解寄存器、Shared Memory、Cache 和显存层次；
能使用 Roofline 判断计算和带宽瓶颈；
能理解 GEMM、卷积、Reduction、Softmax 和 Norm 的 GPU 数据流；
能解释 Tensor Core、RT Core 和低精度计算；
能阅读 PTX、SASS 和基础性能报告；
能从顶层理解完整 GPU 芯片结构；
能规划一颗简化 SIMT GPU 的 RTL 实现。

这套 GPU 架构基础到进阶课程至此结束。下一阶段最有效的学习方式，不是继续增加概念章节，而是实际完成一个4-Lane SIMT GPU，或者选择一个 CUDA Kernel，逐步完成源码、PTX、SASS和 Nsight Compute 的完整分析闭环。