GPU基础：GPU架构、CUDA核心、显存层次、与CPU的差异

在AI算力集群中，GPU是承担并行计算任务的核心硬件。理解GPU的底层架构、执行单元组织方式以及显存层次，是后续理解CUDA编程模型、性能调优和集群调度策略的前提。本节从硬件视角出发，梳理GPU的关键组成，并对比CPU的设计哲学差异。

2.1 GPU架构总览：从SM到CUDA Core

现代数据中心GPU（如NVIDIA A100、H100、AMD MI250等）均采用流式多处理器（Streaming Multiprocessor, SM）作为基本计算单元。一个GPU芯片由多个SM组成，每个SM内部又包含若干计算核心、调度器、寄存器文件和共享内存。

以NVIDIA Ampere架构的A100为例，其内部结构可概括为：

• GPC（Graphics Processing Cluster）：GPU顶层分组，每个GPC包含多个SM和光栅化单元（AI计算中光栅化单元不参与计算）。
• SM（Streaming Multiprocessor）：执行计算的基本硬件单元，A100每个SM包含64个FP32 CUDA Core、32个FP64 Core、4个Tensor Core（第三代）。
• CUDA Core：执行整数和浮点运算的标量核心，每个时钟周期可执行一条指令。注意，CUDA Core并非独立的处理器，而是SM内部高度并行的执行通道。
• Tensor Core：专为矩阵乘加运算设计的专用硬件，支持FP16、BF16、INT8等混合精度计算，是AI训练和推理吞吐量的关键。
• Warp：NVIDIA GPU的调度单位，32个线程组成一个Warp，SM以Warp为单位进行指令发射和执行。

一个典型的SM内部数据流如下：

指令调度器（4个） → 分发到不同执行单元
   ├── FP32 CUDA Core（64个） → 标量浮点/整数运算
   ├── INT32 Core（64个） → 整数运算（与FP32共享寄存器）
   ├── Tensor Core（4个） → 矩阵乘加（混合精度）
   ├── SFU（特殊函数单元） → 倒数、平方根、三角函数等
   └── 共享内存（L1 Cache合并，A100为192KB/SM）

这种“SM内多执行单元 + 多SM并行”的组织方式，决定了GPU适合处理数据并行任务：同一指令作用于大量数据元素，而控制流简单。

2.2 CUDA核心与线程层次模型

CUDA Core是程序员直接面对的硬件抽象，但实际编程时并不直接操作单个Core，而是通过线程层次模型来组织计算。CUDA的线程层次与硬件映射关系如下：

一个典型的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];
    }
}

// 启动配置：256个线程/Block，共N/256个Block
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

当Block被调度到SM上时，SM会将其中的线程划分为多个Warp（每个Warp 32线程），然后由Warp调度器依次发射指令。如果Warp内线程发生分支分歧（即不同线程走不同代码路径），硬件会串行执行各分支路径，导致性能下降——这是GPU编程中需要避免的典型问题。

2.3 显存层次：寄存器、共享内存、全局内存

GPU的存储层次直接影响数据访问效率和编程模型。理解各层级的容量、延迟和可见性，是优化AI算子性能的基础。

2.3.1 存储层次结构

2.3.2 关键设计要点

• 共享内存与L1合并：在NVIDIA架构中，共享内存和L1 Cache共用同一块物理SRAM，可通过API配置比例（如A100支持0-164KB共享内存）。共享内存由程序员显式控制，用于Block内数据复用和线程间通信。
• 全局内存访问模式：全局内存以32字节、64字节或128字节的段为单位访问。当Warp内32个线程访问连续地址时，硬件可合并为少数几次内存事务，实现高带宽利用。若访问地址分散（非合并访问），则会产生多次事务，带宽利用率大幅下降。
• 常量内存与纹理内存：常量内存适合所有线程读取同一地址的场景（如标量参数），纹理内存适合二维/三维局部性访问模式，但在AI计算中较少使用。

一个典型的共享内存优化示例——矩阵分块乘法：

// 使用共享内存分块：每个Block计算一个TILE×TILE的子矩阵
__global__ void matMulShared(float* A, float* B, float* C, int N) {
    __shared__ float As[TILE][TILE];
    __shared__ float Bs[TILE][TILE];
    
    int bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;
    int row = by * TILE + ty;
    int col = bx * TILE + tx;
    
    float sum = 0.0f;
    for (int k = 0; k < N; k += TILE) {
        As[ty][tx] = A[row * N + k + tx];
        Bs[ty][tx] = B[(k + ty) * N + col];
        __syncthreads();  // 确保数据加载完成
        for (int i = 0; i < TILE; i++) {
            sum += As[ty][i] * Bs[i][tx];
        }
        __syncthreads();  // 防止覆盖共享内存
    }
    C[row * N + col] = sum;
}

该模式将全局内存访问次数从O(N³)降低到O(N³/TILE)，显著提升访存效率。

2.4 与CPU的架构差异：设计哲学的对立

CPU与GPU的差异根植于设计目标的不同：CPU追求低延迟和复杂控制流处理，GPU追求高吞吐和大量并行计算。

一个直观的比喻：CPU像一位精通多领域知识的专家，能处理复杂多变的任务；GPU像一支由数千名普通工人组成的队伍，每人只做简单重复的操作，但合在一起能完成海量工作。

2.5 典型场景：AI训练中的GPU工作负载

以一次典型的深度学习训练迭代为例，GPU的工作流程如下：

1. 数据加载：CPU从磁盘/网络读取训练样本，进行预处理（裁剪、归一化），并传输到GPU全局内存。
2. 前向传播：GPU执行卷积、矩阵乘法、激活函数等算子。卷积运算被分解为矩阵乘法（im2col或隐式GEMM），由Tensor Core高效执行。
3. 损失计算：计算预测值与真实值的误差，通常涉及逐元素运算和归约操作。
4. 反向传播：计算梯度，同样以矩阵乘法和卷积为主，计算量约为前向的2-3倍。
5. 参数更新：优化器（如SGD、Adam）更新权重，涉及逐元素乘加和状态变量维护。

在整个过程中，GPU的利用率取决于：

• 算子是否足够并行（矩阵维度是否足够大）
• 数据是否在显存中连续布局（NHWC vs NCHW）
• CPU数据预处理和GPU计算是否重叠（使用CUDA Stream实现异步）

2.6 常见误区与性能陷阱

2.7 小结

GPU通过大量简单核心和SIMT执行模式实现高吞吐并行计算，其架构核心是SM、CUDA Core和Tensor Core的层次化组织。显存层次中，寄存器、共享内存和全局内存各有不同的容量、延迟和编程模型，合理利用共享内存和合并访问是性能优化的关键。与CPU相比，GPU牺牲了单核复杂度和分支处理能力，换取了数十倍的并行度和内存带宽。在AI集群中，理解这些硬件特性有助于正确选择算子实现方式、配置训练参数（如batch size、混合精度策略），并为后续学习CUDA编程、多卡通信和集群调度打下基础。