2. GPU硬件基础

要真正理解GPU加速计算,必须从硬件底层出发。CPU与GPU的协同设计并非简单的“把任务丢给GPU”,而是需要理解GPU的架构哲学——它如何通过大规模并行隐藏延迟、如何组织线程、如何管理内存带宽。本章将系统拆解GPU的架构演进脉络、流式多处理器(SM)的内部结构、内存层次设计、线程调度机制以及SIMT执行模型,帮助你在后续编写CUDA/OpenCL内核时,能够从硬件行为反推代码优化方向。

GPU硬件架构核心链路图 CPU 主机端 x86/ARM 多核处理器 PCIe Gen4/5 根复合体 页表管理 / 驱动栈 PCIe 总线 双向 32 GB/s (Gen4 x16) GPU 设备端 GigaThread 全局调度器 线程块分发 / 负载均衡 L2 缓存 (40-96 MB) 所有SM共享 / 原子操作 SM 0 (流式多处理器) CUDA核心 ×128 FP32 SFU 特殊函数 LD/ST 访存单元 共享内存 / L1 (256 KB) 低延迟 32 银行并行访问 4× Warp 调度器 / 时钟周期 SM 1 ... SM N CUDA核心 ×128 SFU LD/ST 共享内存 / L1 全局内存 (GDDR6X / HBM2e / HBM3) 带宽 1-3 TB/s · 延迟 400-800 周期 · 需合并访问 常量缓存 / 纹理路径 广播优化 / 空间局部性 统一内存 (UM) 页迁移 / 零拷贝 数据路径 控制路径 存储层次 注:Ampere/Ada 架构 SM 含 128 FP32 + 4 Tensor Core

2.1 GPU架构演进:从固定管线到通用并行计算

GPU的演进史本质上是从“专用图形加速器”向“通用并行处理器”的转变过程。早期GPU(2000年前后)采用固定功能管线,顶点着色和像素着色由专用硬件单元完成,程序员无法改变其行为。2001年NVIDIA GeForce3引入可编程顶点着色器,2006年NVIDIA GeForce 8800(G80架构)首次将统一着色器(Unified Shader)引入GPU,所有着色器单元既能处理顶点也能处理像素,这为通用计算(GPGPU)铺平了道路。

从G80到Hopper/Blackwell,GPU架构经历了几个关键里程碑:

架构代号 发布年份 关键特性 对计算的影响
G80 (Tesla) 2006 统一着色器、SIMT执行模型 首次支持CUDA,GPGPU起点
Fermi 2010 真L1/L2缓存、ECC、并发内核 计算能力2.x,支持C++特性
Kepler 2012 SMX单元、动态并行、Hyper-Q 能效比大幅提升,Grid Management Unit
Maxwell 2014 共享内存只读路径、独立L2分区 功耗优化,适合移动端
Pascal 2016 统一内存、NVLink、16nm工艺 支持页迁移,多GPU互联
Volta/Turing 2017/2018 Tensor Core、独立线程调度 AI训练加速,INT8/FP16混合精度
Ampere 2020 第三代Tensor Core、MIG 多实例GPU,弹性计算
Hopper/Ada 2022 Transformer引擎、DPX指令 大模型训练优化,光追与计算融合

架构演进的核心驱动力是并行度与通用性的平衡。早期GPU的并行度来自大量简单核心,但缺乏通用编程模型;现代GPU则在保持高并行度的同时,增加了缓存层次、分支预测、Tensor Core等专用单元,使得GPU不仅能做图形渲染,还能高效运行深度学习、科学计算、数据库等通用负载。

架构演进的关键判断方法:评估一个GPU架构是否适合你的计算任务,不能只看峰值FLOPS。需要关注三个维度:计算密度(每平方毫米的FLOPS)、内存带宽与容量(决定数据搬运效率)、编程模型支持(如CUDA版本、统一内存、动态并行)。例如,Volta引入的独立线程调度解决了Warp内分支发散的性能损失,但代价是芯片面积增加——如果你的负载是高度规整的矩阵运算,这一特性收益有限。

2.2 SM/CU核心结构:GPU的基本计算单元

流式多处理器(Streaming Multiprocessor,SM)是GPU的核心计算单元,相当于CPU中的“核中核”。一个现代GPU(如A100)包含108个SM,每个SM内部又包含多个处理块。理解SM的内部结构是编写高效内核的前提,因为线程调度、资源分配、访存优化都发生在SM层面。

以NVIDIA Ampere架构的GA100 SM为例,其内部结构如下:

  • 4个处理块(Processing Block):每个处理块包含32个FP32 CUDA核心、16个INT32核心、4个Tensor Core(第三代)、4个SFU(特殊函数单元)、4个LD/ST(加载/存储)单元。
  • Warp调度器:每个SM有4个Warp调度器,每个调度器每时钟周期可以发射1条指令到其对应的处理块。这意味着SM每周期最多执行4条不同的Warp指令。
  • 寄存器文件:SM拥有64K个32位寄存器(共256KB),这些寄存器以Warp为单位分配给线程。每个线程最多可使用255个寄存器,但实际分配取决于内核需求和占用率。
  • 共享内存/L1缓存:Ampere架构中共享内存和L1缓存统一为192KB(可配置为共享内存或缓存),A100上为164KB可配置。共享内存以32个银行(Bank)组织,每个银行4字节宽,支持无冲突并行访问。
  • 统一L1/纹理缓存:用于缓存全局内存访问和纹理采样,减少对L2的访问压力。

在AMD的CDNA架构中,对应的单元称为CU(Compute Unit),每个CU包含2个SIMD单元(每个SIMD有32个ALU),共享一个LDS(Local Data Share,相当于共享内存)。虽然命名不同,但设计哲学一致:用大量简单核心配合高带宽内存,通过线程级并行隐藏延迟。

实战经验:在CUDA编程中,__shared__ 变量映射到SM的共享内存,__constant__ 映射到常量缓存。共享内存的访问速度比全局内存快约100倍(延迟约20-30周期 vs 400-800周期),但容量有限。当你的内核需要频繁复用数据时,应优先考虑将数据加载到共享内存中。例如,矩阵乘法中,将A和B的子块加载到共享内存,可以避免对全局内存的重复访问,性能提升可达10倍以上。

2.3 内存层次:从寄存器到全局内存的完整路径

GPU的内存层次设计遵循“容量越大、速度越慢”的原则,但与传统CPU缓存层次不同,GPU的内存层次更强调带宽而非延迟。GPU通过大量并发线程隐藏内存延迟,因此即使全局内存延迟高达数百周期,只要线程数足够多,计算单元就不会空闲。

完整的内存层次从快到慢排列如下:

内存类型 位置 容量 延迟 带宽 作用域
寄存器文件 SM内部 256KB/SM ~1周期 极高 线程私有
共享内存/L1 SM内部 164-228KB/SM ~20-30周期 ~128B/周期 线程块内共享
常量缓存 SM内部 64KB/SM ~20周期(广播时更快) 广播优化 全局只读
L2缓存 GPU芯片 40-96MB ~200-300周期 ~2-4TB/s 所有SM共享
全局内存 显存(HBM/GDDR) 16-80GB ~400-800周期 1-3TB/s 所有线程
主机内存 CPU侧 64-512GB ~1000+周期(PCIe) 32-64GB/s CPU+GPU

内存访问的合并(Coalescing)是GPU性能的关键。当Warp中的32个线程访问全局内存时,硬件会将这32个访问合并为尽可能少的缓存行事务。理想情况下,如果32个线程访问连续地址,只需1个128字节的事务;如果访问分散,则可能需要32个独立事务,带宽利用率下降32倍。

常见误区:很多初学者认为“共享内存越大越好”,但实际上共享内存容量与SM上可驻留的线程块数量成反比。每个线程块占用的共享内存越多,SM能同时容纳的线程块就越少,可能导致占用率下降,无法隐藏内存延迟。例如,一个SM有164KB共享内存,如果每个线程块使用32KB,则最多5个线程块驻留;如果每个线程块使用64KB,则最多2个线程块。在优化时,需要权衡共享内存使用量与线程块数量,通常建议共享内存使用量不超过SM容量的50%,以保证至少2个线程块驻留。

2.4 线程调度与SIMT执行模型

GPU的线程调度基于Warp(线程束)概念——32个线程组成一个Warp,这是硬件调度的最小单位。Warp中的所有线程在同一时刻执行同一条指令(SIMT,Single Instruction Multiple Thread),但每个线程有自己的寄存器状态和程序计数器,因此可以独立分支。

SIMT与传统的SIMD(如CPU的AVX)有本质区别:SIMD要求数据在向量寄存器中对齐,而SIMT允许每个线程访问任意地址。SIMT的灵活性来自硬件调度器的复杂逻辑——每个Warp调度器维护多个Warp的就绪状态,当某个Warp因内存访问而阻塞时,调度器立即切换到另一个就绪的Warp,从而实现延迟隐藏。

线程调度的完整流程如下:

  1. 内核启动:CPU通过驱动将内核代码和参数传递给GPU,GigaThread引擎创建线程网格(Grid)。
  2. 线程块分发:GigaThread将线程块(Block)分发到各个SM。分发策略基于SM的资源可用性(寄存器、共享内存、线程槽位)。
  3. Warp划分:每个线程块内的线程被划分为多个Warp(线程块大小除以32,向上取整)。例如,256线程的块分为8个Warp。
  4. Warp调度:SM中的Warp调度器每周期选择1个就绪Warp发射指令。就绪条件包括:指令无数据依赖、操作数寄存器已就绪、无内存屏障阻塞。
  5. 执行与切换:当Warp执行访存指令时,由于延迟较长,调度器切换到其他Warp执行计算指令。这种切换开销几乎为零,因为寄存器状态已保存在物理寄存器文件中。

分支发散(Branch Divergence)是SIMT模型的主要性能杀手。当Warp内不同线程执行不同分支时,硬件会串行执行每个分支路径,并屏蔽不满足条件的线程。例如,一个Warp中16个线程走if路径,16个线程走else路径,则两条路径串行执行,性能减半。

分支发散优化策略:在编写内核时,应尽量让Warp内的线程执行相同路径。常用方法包括:
1. 数据重排:将相同条件的数据分组,使每个Warp处理同质数据。
2. 使用位运算替代分支:例如 result = (a > 0) ? a : 0 可改写为 result = a * (a > 0)。
3. 循环展开:减少循环控制指令的分支开销。
4. Warp聚合:在Volta及以后架构中,独立线程调度允许Warp内不同线程执行不同指令,但代价是调度器复杂度增加,实际收益有限。

2.5 实战案例:矩阵乘法的硬件适配优化

矩阵乘法(C = A × B)是GPU计算的经典基准,也是理解硬件特性的最佳案例。我们以一个具体的优化过程来展示如何将硬件知识转化为性能提升。

背景与约束:假设矩阵大小为4096×4096(单精度浮点),运行在A100 GPU上(108个SM,每个SM 164KB共享内存,峰值FP32约19.5 TFLOPS)。初始实现为朴素的三重循环,每个线程计算C的一个元素。

初始性能:朴素实现中,每个线程需要读取A的一行和B的一列,共2×4096次全局内存访问。总访问量为4096×4096×2×4096×4字节 ≈ 550GB,而A100的全局内存带宽约1.5TB/s,理论最短时间约0.37秒。但实际运行时间为2.1秒,效率仅17%。瓶颈在于:全局内存访问未合并(每个线程访问不同行/列),且没有利用共享内存复用数据。

优化步骤:

// 步骤1:分块(Tiling)——利用共享内存
#define TILE_SIZE 32
__global__ void matmul_tiled(float* A, float* B, float* C, int N) {
    __shared__ float As[TILE_SIZE][TILE_SIZE];
    __shared__ float Bs[TILE_SIZE][TILE_SIZE];
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    float sum = 0.0f;
    
    for (int k = 0; k < N; k += TILE_SIZE) {
        As[threadIdx.y][threadIdx.x] = A[row * N + k + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(k + threadIdx.y) * N + col];
        __syncthreads();
        
        for (int i = 0; i < TILE_SIZE; i++) {
            sum += As[threadIdx.y][i] * Bs[i][threadIdx.x];
        }
        __syncthreads();
    }
    C[row * N + col] = sum;
}

结果对比:

实现版本 全局内存访问量 运行时间 效率 瓶颈分析
朴素三重循环 550 GB 2.1 s 17% 未合并访问,无数据复用
32×32分块 550 GB / 32 = 17.2 GB 0.42 s 85% 共享内存复用,合并访问
分块+向量化(float4) 17.2 GB / 4 = 4.3 GB 0.31 s 95% 减少指令数,提高内存吞吐

复盘经验:分块大小32是硬件友好的选择——32对应Warp大小,确保每个Warp的线程访问连续的共享内存地址,避免银行冲突。进一步使用float4向量化后,每个线程一次加载16字节,减少LD/ST指令数量,提升内存子系统利用率。最终性能接近理论峰值,说明硬件资源已被充分利用。

调优检查清单:当你的GPU内核性能不达标时,按以下顺序排查:
1. 全局内存访问是否合并?(检查访问地址是否连续)
2. 共享内存是否有银行冲突?(检查同一Warp访问的地址是否映射到不同银行)
3. 占用率是否足够?(使用CUDA Occupancy Calculator计算)
4. 是否有分支发散?(分析Warp内线程的分支路径)
5. 是否使用了向量化加载(float4/int4)?
6. 是否过度使用同步(__syncthreads)导致线程阻塞?

2.6 常见误区与总结复盘

在GPU硬件基础的学习中,以下误区最为常见,需要特别注意:

  • 误区一:认为GPU核心数越多性能越高。实际上,GPU的FLOPS需要足够的数据并行度才能发挥。如果内核的线程数少于SM可容纳的Warp总数,大量核心将空闲。通常需要至少数万个线程才能充分利用现代GPU。
  • 误区二:忽略内存带宽瓶颈。很多计算任务(如向量加法、数据过滤)是内存密集型而非计算密集型。对于这类任务,优化计算指令数毫无意义,应聚焦于减少内存访问量或提高访问效率。
  • 误区三:过度使用共享内存。共享内存虽然快,但容量有限。如果数据复用率不高(如每个数据只被使用一次),使用共享内存反而增加同步开销和资源占用。
  • 误区四:忽视Warp调度的影响。当Warp数量不足时,调度器无法有效隐藏延迟。例如,一个SM最多驻留64个Warp(2048线程),如果内核只启动少量线程块,每个SM可能只有几个Warp,内存延迟将直接暴露。
  • 误区五:认为所有GPU架构相同。不同代际的GPU在共享内存大小、寄存器数量、Tensor Core支持、NVLink带宽等方面差异巨大。代码在A100上优化良好,迁移到RTX 4090(Ada架构)可能需要重新调整分块大小和向量化宽度。

总结:GPU硬件基础的核心在于理解“并行度换延迟”的设计哲学。SM是计算的基本单元,内存层次决定了数据搬运的成本,SIMT模型规定了线程的执行方式。在实际开发中,应始终从硬件行为出发思考优化方向:先确保内存访问合并,再考虑共享内存复用,最后调整计算指令和调度策略。掌握这些硬件原理,你才能在CPU-GPU协同设计中做出正确的架构决策——例如,哪些任务适合放在GPU执行(数据并行、计算密集),哪些任务应保留在CPU(逻辑控制、小规模数据、I/O操作),以及如何设计数据流以减少PCIe传输开销。