GPU架构基础：从SM到Tensor Core的硬件视角

深度学习训练的性能优化，本质上是对GPU硬件资源的调度与匹配。如果只停留在框架API层面调参，很难解释为什么某些算子在不同型号GPU上表现迥异。因此，理解英伟达GPU的底层架构——从流式多处理器（SM）到CUDA核心、Tensor Core，再到显存层次——是进阶优化的第一块基石。

核心原理：GPU的并行层次与执行模型

英伟达GPU采用SIMT（单指令多线程）执行模型。一个kernel启动后，线程以Grid → Block → Warp的层次组织。硬件调度与执行的基本单位是Warp（通常32个线程），而承载Warp执行的核心物理单元是SM（Streaming Multiprocessor）。

以Ampere架构（A100）和Hopper架构（H100）为例，两者的关键区别在于：

• SM数量与规模：A100拥有108个SM，H100提升至132个SM，且每个SM的调度带宽和寄存器文件更大。
• Tensor Core代际：Ampere为第三代Tensor Core，支持TF32和BF16；Hopper为第四代，新增FP8支持，并引入异步Warp专用执行（TMA，Tensor Memory Accelerator）。
• 显存带宽：H100的HBM3带宽达3.35TB/s，远超A100的2TB/s，但显存层次结构逻辑一致。

下图展示了从GPU全局视角到SM内部细节的架构关系，以及数据在显存层次中的流动路径：

方法流程：如何将算子映射到硬件资源

在深度学习训练中，一个典型的矩阵乘法（GEMM）算子，其执行流程可以拆解为以下步骤：

1. 数据分块（Tiling）：将大矩阵切分为适合SM内共享内存大小的子块（如128×128）。
2. 加载到共享内存：通过__syncthreads()同步，将子块从全局显存加载到SM的共享内存，避免重复访问HBM。
3. 寄存器级计算：每个线程从共享内存读取数据到寄存器，执行乘加操作。若使用Tensor Core，则数据需按特定布局（如ldmatrix指令）加载到Tensor Core的寄存器片段。
4. 写回：计算结果从寄存器写回共享内存，再批量写回全局显存。

以下是一个简化的CUDA代码片段，展示如何显式利用共享内存和Tensor Core（伪代码示意）：

// 伪代码：使用Tensor Core的矩阵乘片段
__global__ void matmul_tensor_core(float* A, float* B, float* C) {
    __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
    
    // 1. 协作加载数据到共享内存
    load_tile_to_shared(A, As, ...);
    load_tile_to_shared(B, Bs, ...);
    __syncthreads();
    
    // 2. 使用warp级矩阵指令（如mma.sync）
    // 每个warp计算一个16x16x16的矩阵块
    mma_sync(C_frag, As_frag, Bs_frag, C_frag);
    
    // 3. 写回全局内存
    store_tile_to_global(C, C_frag, ...);
}

理解这一流程后，优化方向就清晰了：最大化共享内存复用、减少线程同步开销、确保数据布局满足Tensor Core对齐要求。

案例/场景：A100与H100上的GEMM性能差异

假设我们要训练一个Transformer模型，其中自注意力机制的QKV投影涉及大规模GEMM。在A100上，使用TF32精度（Tensor Core支持）比FP32快约8倍；而在H100上，使用FP8精度（第四代Tensor Core原生支持）比TF32再快约2倍。

实际测试中，在H100上使用FP8训练GPT-3级别模型，相比A100的FP16训练，吞吐量提升约1.8倍，但需要额外处理缩放因子（scaling factor）以避免精度溢出。这要求框架层（如PyTorch的torch.fp8）和自定义CUDA内核的紧密配合。

误区/复盘：常见硬件理解偏差