课程导入:学习目标、课程结构、预备知识、整体框架

欢迎进入《英伟达流处理器阵列底层架构解析》课程。在深入 CUDA 内核调度、Tensor Core 流水线或内存一致性模型之前,我们必须先建立一张清晰的“架构地图”。本章作为全课程的起点,将帮助你明确学习终点、规划路径、补齐必要的前置知识,并从宏观到微观建立英伟达 GPU 流处理器阵列的整体认知框架。

1.1 学习目标:从“会用”到“看懂”

本课程不满足于告诉你“如何写 CUDA 代码”,而是聚焦于代码背后的硬件行为。完成本课程后,你将能够:

  • 拆解流处理器阵列(SM)的内部组成:准确说出 SM 中 CUDA Core、SFU、Tensor Core、LD/ST 单元、寄存器文件、共享内存与调度器之间的数据通路关系。
  • 解释线程调度与执行模型:理解 Warp 的诞生、调度、发射与停滞机制,能够分析分支发散对吞吐率的影响。
  • 追踪内存请求的完整路径:从全局内存访问到 L2 缓存、L1/共享内存、寄存器,描述缓存行失效与合并访问的硬件逻辑。
  • 对比不同架构代际的演进逻辑:从 Volta 的独立线程调度到 Ampere 的第三代 Tensor Core,再到 Hopper 的 DPX 指令,理解每代架构解决的核心瓶颈。
  • 建立性能分析的硬件视角:能够将 Nsight Compute 中的 Stall 原因映射到具体的硬件队列、仲裁器或带宽瓶颈。
学习建议:本课程涉及大量硬件细节,建议每学完一章,回到本节的目标清单,尝试用自己的话复述对应能力是否达成。如果无法清晰解释某个术语(如“Warp 调度器每周期发射 1 条指令”),说明该章节需要重读。

1.2 课程结构:六大模块递进式展开

全课程共 30 章,划分为六个逻辑模块,每个模块解决一个核心问题。下表展示了模块划分与对应章节范围:

模块 章节范围 核心问题
基础框架 第 1-4 章 GPU 整体架构与流处理器阵列的定位
SM 内部微架构 第 5-12 章 CUDA Core、调度器、寄存器与执行流水线
存储层次与数据通路 第 13-18 章 寄存器、共享内存、L1/L2 缓存与全局内存
并行执行与同步 第 19-23 章 Warp 调度、同步原语、协作组与原子操作
现代特性与加速器 第 24-27 章 Tensor Core、RT Core、稀疏性与 DPX 指令
性能分析与实战 第 28-30 章 Nsight 工具链、瓶颈定位与架构敏感优化

每个模块内部章节之间具有强依赖关系。例如,不理解第 6 章的 SIMT 模型,就无法掌握第 20 章的 Warp 同步语义。因此建议按顺序学习,不要跳跃。

1.3 预备知识:你需要具备的基础

本课程面向具备一定 GPU 编程经验、希望深入硬件层的开发者。建议在开课前确认自己已掌握以下内容:

  1. CUDA 编程基础:能编写简单的 kernel,理解 grid、block、thread 三级层次结构,知道 __syncthreads() 的基本用法。
  2. 计算机体系结构常识:了解流水线、缓存、寄存器、指令集的基本概念,知道什么是时钟周期和 IPC。
  3. C/C++ 指针与内存模型:能够区分栈、堆、全局内存,理解指针解引用对应的硬件访存行为。
  4. 基本性能分析经验:使用过 nvprof 或 Nsight Compute 查看过 kernel 的占用率或带宽数据。

如果对上述内容感到陌生,建议先补充相应基础。本课程不会花时间解释“什么是线程块”或“缓存是什么”,而是直接讨论这些概念在英伟达硬件中的具体实现。

常见误区:很多学习者认为“只要熟悉 CUDA 编程就能理解硬件”。实际上,CUDA 编程模型是对硬件的抽象,隐藏了大量细节。例如,编程模型中的“线程”在硬件上以 Warp 为单位执行;编程模型中的“共享内存”在物理上可能与 L1 缓存共用同一块 SRAM。如果只停留在编程层面,你将无法解释为什么某些代码在特定架构上性能骤降。

1.4 整体框架:从 GPU 到 SM 再到执行单元

在开始逐章拆解之前,我们先建立一张全局架构图。这张图将贯穿整个课程,后续所有细节都会挂载到这张图的对应节点上。请仔细观察下图中的层级关系与数据流方向:

英伟达 GPU 流处理器阵列层级架构总览 GPU 芯片 (例如 GA100 / AD102) 包含:GPC (图形处理簇) → TPC (纹理处理簇) → SM (流处理器) 全局内存控制器 (HBM/GDDR)、L2 缓存、PCIe/NVLink 接口 GPC (图形处理簇) 每个 GPC 包含多个 TPC 光栅化单元、ROP 控制(部分型号) TPC (纹理处理簇) 包含 1-2 个 SM(依架构而定) 纹理单元 (Texture Unit) SM (流处理器) — 本课程核心研究对象 Warp 调度器 4 个调度器 每周期发射 1 条 CUDA Core 阵列 FP32/INT32 单元 FP64 单元(部分) SFU / Tensor Core 特殊函数单元 矩阵乘累加加速器 LD/ST 单元 访存指令执行 地址生成与缓存 寄存器文件 256KB 每 SM 32-bit 寄存器 共享内存 / L1 缓存 (SM 内部) 统一架构,可配置容量,低延迟高带宽 L2 缓存 (GPU 全局共享) 所有 SM 共享,缓存一致性点,原子操作驻留 橙色虚线 = 数据通路方向 | 紫色 = 存储层次

图中展示了本课程反复提及的层级关系:GPU → GPC → TPC → SM → 执行单元。请注意,SM 是流处理器阵列的基本单位,也是本课程每一章都会回到的核心节点。在后续章节中,我们将逐层深入:第 2 章先剖析 SM 外部的 GPC/TPC 组织方式;第 5 章开始进入 SM 内部,观察调度器如何从 Warp 池中选取指令;第 13 章则聚焦存储通路,解释共享内存与 L1 缓存的物理统一设计。

这张图也揭示了英伟达架构设计的一个关键哲学:通过层次化复制(GPC/TPC/SM)实现大规模并行,同时通过共享存储(L2)与同步原语维持一致性。理解这个哲学,你就掌握了理解所有具体机制的总钥匙。

1.5 本章小结与学习路径建议

本章我们完成了三件事:明确了学习目标(从编程到硬件)、规划了课程结构(六大模块递进)、确认了预备知识(CUDA 基础与体系结构常识)。同时,通过整体框架图建立了从 GPU 到 SM 的层级认知。

在进入第 2 章之前,请确保你能够不看资料回答以下问题:

  • 一个典型的英伟达 GPU 中,GPC、TPC、SM 三者的包含关系是什么?
  • SM 内部至少包含哪几类执行单元?各自负责什么类型的指令?
  • 共享内存/L1 缓存与 L2 缓存在物理位置和共享范围上有何本质区别?