📚 课程目录第2章 GPU硬件架构

GPU硬件架构：SM、CUDA核心、内存层次与线程调度硬件基础

要理解英伟达GPU运行时内核（如CUDA、TensorRT、驱动层调度器）的行为，首先必须建立对底层硬件的准确认知。GPU并非“许多CPU的简单堆叠”，而是一种大规模并行吞吐机，其设计目标是用大量线程的并行执行来隐藏访存延迟，而非像CPU那样依赖大缓存和分支预测来降低延迟。本节将拆解GPU硬件架构的四个核心支柱：流式多处理器（SM）、CUDA核心、内存层次结构，以及线程调度硬件基础。

2.1 流式多处理器（SM）：GPU的基本计算单元

GPU由多个图形处理簇（GPC）组成，每个GPC内部包含若干流式多处理器（Streaming Multiprocessor, SM）。SM是GPU中真正执行计算指令的硬件单元，相当于GPU的“心脏”。以英伟达Ampere架构（A100）为例，一颗GA100芯片包含108个SM（A100启用108个），每个SM内部又细分为四个处理块（Processing Block），每个处理块拥有一个独立的调度器（Scheduler）和对应的执行资源。

SM内部的关键组件包括：

CUDA核心（FP32单元）：执行单精度浮点与整数运算。
Tensor Core：执行矩阵乘加运算（如FP16、BF16、INT8），是AI推理与训练的核心加速器。
特殊功能单元（SFU）：执行超越函数（如sin、cos、倒数、平方根）。
加载/存储单元（LD/ST）：负责访问全局内存或共享内存。
寄存器堆（Register File）：容量极大（A100每个SM有64K个32位寄存器），用于存放线程的局部变量。
共享内存（Shared Memory）：SM内所有线程可访问的高速缓存，容量可配置（A100为164KB/SM）。
线程束调度器（Warp Scheduler）：负责将线程束（Warp）分派到执行单元。

  关键概念：线程束（Warp）是GPU硬件调度的最小单位，通常由32个线程组成。一个SM中的线程束调度器每次选择一个就绪的Warp，将其指令发送到对应的执行单元。如果某个Warp因访存而阻塞，调度器会立即切换到另一个就绪的Warp，这种机制称为零开销上下文切换，是GPU隐藏延迟的核心手段。

2.2 CUDA核心与Tensor Core：不同精度的执行引擎

CUDA核心本质上是单精度浮点单元（FP32 ALU），每个时钟周期可执行一次FP32运算（FMA）。在Ampere架构中，SM内共有128个FP32单元（4个处理块×32个）。此外，每个处理块还包含16个INT32单元，用于整数运算。值得注意的是，在Ampere及后续架构中，FP32和INT32单元是分离的，这意味着整数运算不会占用浮点流水线。

Tensor Core则是专门为矩阵乘加运算设计的硬件加速器。以A100的第三代Tensor Core为例，每个Tensor Core每个时钟周期可执行4×4×4的矩阵乘加操作。对于FP16输入，A100的Tensor Core峰值算力可达312 TFLOPS（稀疏化后为624 TFLOPS），而FP32 CUDA核心的峰值算力仅为19.5 TFLOPS。这种数量级的差距解释了为什么深度学习推理和训练必须依赖Tensor Core。

执行单元	数量/SM	精度	每周期操作	典型用途
FP32 CUDA核心	128	FP32	1 FMA	通用计算、非矩阵运算
INT32 核心	64	INT32	1 整数运算	地址计算、逻辑判断
Tensor Core	16	FP16/BF16/INT8	4×4×4 MMA	矩阵乘法、卷积、Transformer
SFU	16	FP32	1 超越函数	倒数、平方根、三角函数

编程启示：在编写CUDA内核时，应尽量将计算组织为矩阵乘加形式（如使用cuBLAS、cuDNN或CUTLASS），以充分利用Tensor Core。纯标量循环（如逐元素操作）只会使用CUDA核心，无法发挥GPU的最大算力。

2.3 内存层次结构：从寄存器到全局内存的延迟阶梯

GPU的内存层次设计遵循“容量越大，延迟越高”的原则。理解各级内存的延迟和带宽，是优化内核性能的前提。下表展示了A100的关键内存参数：

内存层级	容量	延迟（周期）	带宽	作用域
寄存器	64K×32bit/SM	~1	极高	单个线程
共享内存	164KB/SM	~20	~19TB/s（聚合）	线程块（CTA）
L1缓存	与共享内存共用	~30	同共享内存	SM内
L2缓存	40MB（A100）	~200	~5.5TB/s	整个GPU
HBM2e显存	40GB/80GB	~400-600	~2TB/s	所有SM共享

在编程模型中，全局内存（Global Memory）对应显存（HBM），是所有线程均可访问的“主存”。访问全局内存的延迟高达数百周期，因此必须通过合并访问（Coalesced Access）来最大化带宽利用率——即让同一Warp的32个线程访问连续的内存地址，硬件会将这32个访问合并为少数几个缓存行事务。

共享内存（Shared Memory）是片上SRAM，延迟远低于全局内存。它被划分为32个存储体（Bank），每个存储体宽度为4字节。当多个线程同时访问不同存储体时，可以并行完成；但如果多个线程访问同一存储体的不同地址，就会发生存储体冲突（Bank Conflict），导致访问串行化，性能下降数倍。

2.4 线程调度硬件：Warp与SIMT执行模型

GPU采用单指令多线程（SIMT）执行模型。硬件以Warp（32线程）为单位进行调度。每个SM的4个Warp调度器，每个时钟周期可以发射一条指令到一个执行单元组。调度器维护着多个Warp的就绪状态，当某个Warp因等待内存返回而阻塞时，调度器会无开销地切换到另一个就绪Warp。

线程块（CTA/Block）是软件层面的调度单位。GPU的GigaThread引擎负责将线程块分配到SM上。一个SM可以同时驻留多个线程块（取决于资源限制），但同一时刻SM上所有线程块中的线程总数不能超过最大线程数（A100为2048线程/SM，即64个Warp）。

  硬件资源与Occupancy的关系：每个SM的寄存器堆和共享内存是有限的。如果每个线程使用过多寄存器（例如超过128个），则SM上能同时驻留的线程数就会减少，导致Occupancy（占用率）下降。低Occupancy意味着可用于隐藏延迟的Warp数量不足，可能无法有效掩盖访存延迟。通常建议Occupancy至少达到50%（即32个Warp/SM）以保证较好的延迟隐藏能力。

2.5 案例：从硬件视角分析一个简单的向量加法内核

假设我们执行以下CUDA内核：

__global__ void vecAdd(float* A, float* B, float* C, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) {
        C[i] = A[i] + B[i];
    }
}

当启动配置为gridDim=ceil(N/256)，blockDim=256时，硬件执行流程如下：

GigaThread引擎将线程块（每个256线程）分发到不同的SM。假设SM有足够资源，每个SM可驻留8个线程块（2048线程/SM）。
每个线程块被拆分为8个Warp（256/32）。SM的4个调度器各分配2个Warp。
每个Warp执行LDG（加载全局）指令时，32个线程访问A[i]。若i连续，则合并为一次或两次内存事务。
在等待数据返回期间（约400周期），调度器切换到其他Warp执行加法或加载B[i]。
数据返回后，Warp恢复执行，完成加法并执行STG（存储全局）指令。

如果N很大（如1000万），则整个内核的执行时间主要由内存带宽决定，而非计算延迟。此时，合并访问和高Occupancy是性能的关键。

2.6 常见误区与复盘

误区一：认为“线程越多越好”。实际上，如果每个线程使用过多寄存器或共享内存，SM能驻留的线程数反而减少。例如，若每个线程使用128个寄存器，则SM最多驻留512线程（16个Warp），Occupancy仅25%，延迟隐藏能力严重不足。应通过__launch_bounds__或maxrregcount限制寄存器使用量。

误区二：忽略存储体冲突。在共享内存中，如果线程访问的地址满足bank = (threadIdx.x * 2) % 32，则线程0和线程16会访问同一存储体，产生2路冲突。应尽量设计为bank = threadIdx.x的连续访问模式，或使用填充（padding）技术。

误区三：将全局内存视为“大缓存”。全局内存没有硬件预取机制，且延迟极高。应主动使用共享内存进行数据复用（如矩阵分块），而非依赖L2缓存。L2缓存主要用于跨SM的数据共享和原子操作，不应作为主要的数据复用手段。

复盘总结：GPU性能优化的本质是最大化有效并行度与最小化内存延迟暴露。硬件层面，SM提供了大量执行单元和Warp调度机制来隐藏延迟；软件层面，开发者需要合理配置线程块大小、寄存器使用量和共享内存分配，确保Occupancy充足，并设计合并的内存访问模式。理解SM内部结构、内存层次和Warp调度机制，是编写高性能CUDA内核的基础，也是后续理解TensorRT、CUDA Graph等运行时优化技术的硬件前提。