GPU架构与并行计算基础

GPU架构与并行计算基础

当2012年AlexNet在ImageNet挑战赛上一鸣惊人时，恐怕没有多少人预见到这将成为GPU计算时代的序曲。那场竞赛中，研究团队首次使用NVIDIA GPU进行深度卷积神经网络的训练，将图像识别错误率从26%一举降至15%，开启了深度学习革命的序幕。从那一刻起，原本专精于图形渲染的GPU开始承担起通用计算的使命，并在随后十余年间重塑了整个人工智能产业的底层架构。

理解英伟达如何一步步将图形处理器改造为通用算力引擎，必须从最基础的硬件架构设计说起。流式多处理器（Streaming Multiprocessor, SM）作为GPU的基本计算单元，其设计哲学与CPU截然不同——它追求的不是单个线程的极致性能，而是在高并发场景下实现惊人的吞吐量。本章将深入解析SM的微架构细节，探讨SIMT（Single Instruction Multiple Thread）执行模型的工作原理，剖析显存层级架构对AI计算性能的影响，并通过Ampere和Hopper两代架构的具体案例，展示英伟达如何在制程工艺的约束下通过架构创新持续提升算力。

从图形渲染到通用计算的演进之路

GPU的发展历程本身就是一部技术演进的教科书。1999年，英伟达推出GeForce 256，首次将几何变换和光照计算从CPU卸载到专用硬件中执行，这被视为现代GPU的起点。早期的GPU采用固定功能管线（Fixed-Function Pipeline）设计，每个渲染阶段都有专用的硬件电路执行特定任务：顶点着色器处理几何变换，像素着色器负责纹理采样和颜色计算，光栅化引擎将几何图元转换为像素片段。这种设计在当时的图形渲染场景下效率极高，但缺乏灵活性，无法适应不断演进的图形API和新算法需求。

2006年是一个重要的转折点。英伟达在G80架构中引入了统一着色器架构（Unified Shader Architecture），将顶点着色器、几何着色器和像素着色器统一为同一类可编程单元——流处理器（Streaming Processor）。这种设计的革命性在于：硬件不再预先分配给特定渲染阶段，而是根据工作负载动态调度。当场景中几何计算密集时，更多流处理器被用于顶点处理；当纹理和光照计算占主导时，同一批硬件资源又可以重新分配给像素着色任务。G80架构还引入了SIMT执行模型的雏形，为后来的通用计算奠定了基础。

真正让GPU进入通用计算时代的是CUDA（Compute Unified Device Architecture）的推出。2006年，英伟达发布了CUDA 1.0，为开发者提供了在GPU上编写并行程序的统一编程模型。CUDA的核心抽象是将GPU视为一个由大量轻量级线程组成的并行计算机，程序员可以用类似C语言的语法表达并行算法，而无需关心底层的硬件调度细节。这一设计将GPU的通用计算能力从极客圈带入主流工程界。

随后几年，英伟达逐步在硬件中增加通用计算相关的特性。2008年的GT200架构将双精度浮点运算单元加入流处理器，支持IEEE 754-2008标准的完整双精度计算。2010年的Fermi架构是一个重要的里程碑，它引入了真正的内存层次结构——包括L1和L2缓存、聚合内存访问、错误纠正码（ECC）等特性。Fermi还首次支持并行执行多个内核函数（Kernel），显著提升了硬件利用率。

Kepler（2012）、Maxwell（2014）、Pascal（2016）三代架构在通用计算方面持续优化。Pascal引入了统一内存（Unified Memory）和NVLink高速互连技术，并首次在数据中心GPU（P100）中采用HBM2高带宽显存。Volta（2017）架构则是一个转折点，它引入了张量核心（Tensor Core）——专为矩阵乘法运算设计的专用硬件单元，这一创新将深度学习训练的效率提升了数量级。

从历史演进来看，GPU的通用计算之路并非一蹴而就，而是通过硬件和软件的双重迭代逐步实现。早期的图形专用设计反而为通用计算提供了宝贵的经验：大规模并行性、高内存带宽、SIMT执行模型等特性，在深度学习时代都找到了新的用武之地。这种"无心插柳"式的架构演进，恰恰体现了硬件设计中的意外收获。

流式多处理器：GPU的核心执行单元

流式多处理器（Streaming Multiprocessor, SM）是NVIDIA GPU的基本计算单元，每个GPU芯片由多个SM组成。以Ampere架构的A100 GPU为例，它包含108个SM，而最新的Hopper架构H100 GPU则包含132个SM。每个SM本身又是一个高度并行的计算子系统，拥有独立的指令缓存、调度单元、寄存器文件和多个执行单元。理解SM的微架构设计，是掌握GPU并行计算本质的关键。

线程调度机制

GPU的线程调度与传统CPU有着本质的不同。在CPU中，每个核心通常运行一个或两个线程（通过超线程技术），调度器在少数活跃线程之间快速切换，上下文切换的开销相对较高。GPU则采用了一种完全不同的策略：同时维持数万个线程的运行状态，通过硬件层面的线程调度器在极小的开销下在这些线程之间切换。

在CUDA编程模型中，线程被组织成三级层次结构：单个线程（Thread）、线程块（Block）和线程网格（Grid）。一个线程块包含最多1024个线程，这些线程被分配到同一个SM上执行，能够通过共享内存高效通信。不同线程块之间无法直接通信，但可以通过全局内存进行间接数据交换。线程网格由多个线程块组成，覆盖整个GPU。

SM内部的线程调度采用两级仲裁机制。第一级是线程块调度器（Block Scheduler），负责将就绪的线程块分配到空闲的SM上执行。由于线程块的执行时间可能差异很大（取决于具体的工作负载），GPU采用了一种贪婪但公平的调度策略：当某个SM上的线程块因内存访问而停滞时，调度器会立刻将其他就绪的线程块分配到该SM上，保持计算单元始终处于忙碌状态。

第二级调度发生在SM内部，被称为线程束调度器（Warp Scheduler）。在NVIDIA的术语中，一个线程束（Warp）包含32个线程——这是硬件并行执行的基本单位。在SM的每个时钟周期内，线程束调度器从准备好的线程束中选择一个发射指令。由于GPU执行单元的 latency 通常较长（特别是内存访问），一个线程束在等待结果期间可以切换到另一个线程束继续执行，这种技术被称为延迟隐藏（Latency Hiding）。

线程束调度的一个关键参数是调度延迟（Scheduling Latency）。以Ampere架构为例，每个SM有4个线程束调度器，每个调度器在每个时钟周期可以发射一条指令。这意味着理论上每个SM每时钟周期可以发射4条指令。但实际吞吐量取决于多种因素：指令类型（算术、内存访问、控制流）、执行单元利用率、以及线程束调度器的效率。

```python

CUDA线程层次结构示意

代码来源：NVIDIA CUDA Programming Guide

定义一个简单的向量加法kernel

__global__ void vectorAdd(float* a, float* b, float* c, int n) { // 计算全局线程索引 int tid = blockIdx.x * blockDim.x + threadIdx.x;

// 确保不越界 if (tid < n) { c[tid] = a[tid] + b[tid]; } }

// 调用时的线程配置 // 假设向量长度为 N = 1024 int threadsPerBlock = 256; int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;

// 启动kernel vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, N); ```

在上述示例中，`threadsPerBlock`定义了每个线程块的线程数（256），`blocksPerGrid`定义了线程网格包含的线程块数（4）。这意味着总共有1024个线程同时执行，每个线程负责处理向量中的一个元素。这种编程模型将数据级并行（Data-Level Parallelism）直接映射到硬件的线程结构上。

SIMT执行模型深度解析

SIMT（Single Instruction, Multiple Thread，单指令多线程）是NVIDIA GPU的核心执行模型。这一名称借鉴了SIMD（Single Instruction, Multiple Data）的概念，但两者有着重要的区别。SIMD在硬件层面将多个数据元素打包到同一个向量寄存器中，由一条指令同时处理；而SIMT则将每个数据元素映射到独立的线程中，由独立的程序计数器（PC）跟踪每个线程的执行状态。

SIMT模型的关键优势在于线程级并行与数据级并行的统一。从程序员的角度看，他可以像编写串行程序一样为每个数据元素编写处理逻辑，这些逻辑在底层会被并行执行。从硬件的角度看，32个线程组成一个线程束，同一线程束内的所有线程同时执行同一条指令——这又是SIMD的特征。当某个线程因分支 divergence（如if-else语句）而需要执行不同的代码路径时，线程束会被动态划分为多个子线程束，每个子束执行各自的路径。

```cpp // SIMT中的分支分歧示例 __global__ void conditionalKernel(int* data, int* result, int threshold) { int tid = blockIdx.x * blockDim.x + threadIdx.x;

if (data[tid] > threshold) { // 分支1 result[tid] = data[tid] * 2; } else { // 分支2 result[tid] = data[tid] + 1; } } ```

在上述代码中，如果`data`数组中的值既有大于threshold的，也有小于等于threshold的，那么同一线程束内的线程就会产生分歧。此时，GPU硬件会串行执行两个分支，屏蔽掉当前不需要执行路径的线程。虽然这会导致性能下降，但相比纯SIMD架构需要程序员手动向量化，SIMT的自动处理方式大大简化了编程。

为了最小化分支分歧的影响，程序员应当尽可能保证同一线程束内的线程执行相同的代码路径。这通常通过结构化分支（Structuring Branches）来实现：将数据预先按照条件分组，使同一线程块内的线程倾向于走相同的分支。另一种技术是分支predication（Branch Predication），将分支语句转换为predicated指令，由硬件自动处理分歧而非显式分支。

SIMT的硬件实现涉及几个关键组件。程序计数器（PC）寄存器组为每个线程维护独立的PC，尽管通常多个线程共享同一条指令。执行掩码（Active Mask）寄存器记录线程束中哪些线程当前处于活跃状态。分支同步堆栈（Branch Synchronization Stack）用于处理嵌套分支，记录每个分支层次的执行状态。

值得注意的是，从Volta架构开始，NVIDIA引入了独立线程调度（Independent Thread Scheduling）特性，允许同一线程束内的线程独立推进，即使它们暂时处于同步状态。这意味着某些线程可以在其他线程等待时执行其他代码，提供了更细粒度的调度灵活性。这一特性对于某些算法（如某些图算法、复杂控制流代码）可以显著提升效率，但程序员也需要更加小心地管理同步点，避免死锁。

寄存器文件与执行单元设计

寄存器文件（Register File）是SM中最关键的资源之一，直接决定了单个SM能够支持的线程并发数量。在Ampere架构中，每个SM包含65536个32位寄存器，被分配给该SM上执行的所有线程使用。每个线程可以使用的寄存器数量是有限的，当一个线程使用的寄存器超过限制时，超出的部分会被溢出到本地内存（Local Memory），导致性能显著下降。

寄存器分配策略对GPU程序的性能有直接影响。CUDA编译器会分析每个线程函数（Kernel Function）的寄存器使用量，尝试在保持性能的同时最大化并发线程数。编译选项`--maxrregcount`可以手动指定最大寄存器使用量。减少每个线程的寄存器使用可以增加活跃线程束的数量，从而提高内存访问延迟隐藏的能力，但同时可能导致寄存器溢出到慢速的本地内存。

```cuda // 寄存器使用示例 __global__ void registerIntensiveKernel(float* data, int n) { // 声明多个局部变量——每个线程占用更多寄存器 float temp1 = data[threadIdx.x]; float temp2 = temp1 * 2.0f; float temp3 = temp2 + 1.0f; float temp4 = temp3 * 3.0f; float temp5 = temp4 - 1.0f;

data[threadIdx.x] = temp5; } ```

在上述代码中，每个线程声明了5个局部变量`temp1`到`temp5`，这些变量都会分配到寄存器中。如果编译器发现寄存器压力过大，可能会将这些变量溢出到本地内存（实际上存储在全局内存中），显著增加内存访问延迟。

SM的执行单元配置在不同架构间有所差异。以Ampere架构为例，每个SM包含：

• **128个CUDA核心**：每个核心包含一个浮点运算单元（FP32）和一个整数运算单元（INT），可以每时钟周期执行一次单精度浮点运算或整数运算
• **4个张量核心（Tensor Core）**：专门用于矩阵乘法运算，在深度学习中扮演核心角色
• **1个RT核心（RT Core）**：用于光线追踪加速
• **4个SFU（Special Function Units）**：执行超越函数（sin、cos、exp、log等）
• **2个LD/ST单元**：负责加载和存储操作
• **1个分支单元**：处理控制流指令

这种异构执行单元的配置使得SM可以同时处理多种类型的计算任务。当某个线程束执行浮点运算时，CUDA核心处于忙碌状态，但SFU或LD/ST单元可能同时为其他线程束服务。这种粗粒度的多任务机制是GPU实现高指令吞吐量的关键。

显存层级架构与AI计算性能

GPU的内存系统是决定深度学习训练和推理性能的核心因素。与CPU注重低延迟的缓存设计不同，GPU的内存系统针对高带宽进行了优化，这是因为许多GPU工作负载（尤其是深度学习）具有数据访问模式规则、计算密度高的特点。本节将深入分析GPU显存层级架构的设计，以及它如何影响AI计算性能。

内存层级结构总览

GPU的内存系统是一个多层级的复杂结构，从快到慢依次包括：寄存器、本地内存、L1缓存、共享内存、L2缓存、纹理内存（Texture Memory）、常量内存（Constant Memory）和全局内存（Global Memory）。理解每层存储的特性对于编写高性能GPU代码至关重要。

寄存器是访问速度最快的存储，位于执行流水线内部。寄存器访问是零周期延迟的——指令可以在同一个时钟周期内读取寄存器并发射到执行单元。然而，寄存器的数量是严格受限的，每个SM的寄存器文件大小固定。

本地内存在物理上位于全局内存（显存）中，当寄存器不足时用于存储局部变量。本地内存的访问延迟与全局内存相当（数百个时钟周期），是性能敏感代码应当极力避免的。

共享内存是一种用户管理的软件缓存，位于SM附近但独立于L1缓存。共享内存的访问延迟约为L1缓存的2-3倍，但远低于全局内存（可低至1/10）。共享内存通过`__shared__`关键字声明，被同一线程块内的所有线程共享。由于需要手动管理同步，共享内存适合用于线程块内需要重复访问的数据。

L1缓存是每个SM私有的缓存行缓存（Line-based Cache），用于缓存全局内存的访问结果。L1缓存的容量较小（通常为128KB或更少），但访问延迟极低。从Kepler架构开始，L1缓存和共享内存共用一部分物理存储空间，程序员可以选择配置比例。

L2缓存是所有SM共享的缓存，容量较大（Ampere A100有40MB L2缓存）。L2缓存的访问延迟高于L1但远低于全局内存。当多个线程访问相似的数据时，L2缓存可以显著减少对全局内存的访问。L2缓存还支持一级缓存行失效（Cache Line Eviction）的原子操作，对并行编程有重要意义。

全局内存是GPU的"主内存"，容量最大（从8GB到80GB不等），但访问延迟也最高。全局内存通过PCIe总线或NVLink与GPU通信。访问全局内存时，最佳实践是对齐合并访问（Coalesced Access）：同一线程束内的线程访问连续的内存地址，这些访问会被合并成一次或少数几次内存事务。

```cuda // 合并访问与非合并访问对比

// ✅ 合并访问：线程束内线程访问连续地址 __global__ void coalescedAccess(float* data, int n) { int tid = blockIdx.x * blockDim.x + threadIdx.x; if (tid < n) { float value = data[tid]; // 连续访问，合并为少量事务 data[tid] = value * 2.0f; } }

// ❌ 非合并访问：线程束内线程访问分散地址 __global__ void uncoalescedAccess(float* data, int stride, int n) { int tid = blockIdx.x * blockDim.x + threadIdx.x; if (tid < n) { float value = data[tid * stride]; // stride导致分散访问 data[tid * stride] = value * 2.0f; } } ```

在深度学习场景中，合并访问的重要性更加突出。典型的卷积层前向传播会遍历输入特征图，每个像素的相邻通道通常在内存中连续存储。确保内存访问模式与这种数据布局匹配，可以最大化内存带宽利用率。

访问延迟