2. GPU基础架构

要理解GPU资源调度与性能调优，首先必须建立对GPU硬件底层的清晰认知。GPU并非“更快的CPU”，而是一台为大规模数据并行而设计的专用处理器。它的性能来源于数千个简单核心的协同工作，而非少数复杂核心的超高主频。本节将拆解GPU的硬件组成、内存层次结构与并行计算模型，为后续的调度与调优打下根基。

2.1 GPU硬件组成：SM与CUDA Core的层级结构

现代NVIDIA GPU（以Ampere架构为例）的硬件组织遵循清晰的层级关系。最核心的执行单元是CUDA Core（流处理器），它执行整数与浮点运算。多个CUDA Core连同特殊功能单元（SFU）、加载/存储单元（LD/ST）、寄存器堆和共享内存，组成一个流式多处理器（SM，Streaming Multiprocessor）。整个GPU芯片则由多个SM簇、L2缓存、内存控制器和全局内存（显存）构成。

关键硬件单元速查表

硬件单元数量级（A100）主要职责
CUDA Core6912个执行FP32/INT32算术运算
SM（流式多处理器）108个包含CUDA Core、共享内存、调度器
Tensor Core432个执行混合精度矩阵乘加运算
L2缓存40 MB跨SM共享的二级缓存
HBM2e显存80 GB全局内存，带宽可达2 TB/s

每个SM内部包含4个处理块（Processing Block），每个处理块拥有一个Warp调度器和对应的执行单元。Warp是GPU调度的基本单位，由32个线程组成。当一个Warp中的线程执行同一指令时，若发生分支发散，调度器会串行执行各分支路径，导致性能下降。

GPU硬件层级与内存结构示意（A100风格） GPU 芯片（Die） SM 0 处理块0 Warp调度器 32 CUDA Core 处理块1 Warp调度器 32 CUDA Core 处理块2 Warp调度器 32 CUDA Core 处理块3 Warp调度器 32 CUDA Core 共享内存 (Shared Memory) 164KB SM 1 处理块0 Warp调度器 32 CUDA Core 处理块1 Warp调度器 32 CUDA Core 处理块2 Warp调度器 32 CUDA Core 处理块3 Warp调度器 32 CUDA Core 共享内存 164KB … … SM 107 处理块0 Warp调度器 32 CUDA Core 处理块1 Warp调度器 32 CUDA Core 处理块2 Warp调度器 32 CUDA Core 处理块3 Warp调度器 32 CUDA Core 共享内存 164KB L2 缓存（40 MB，所有SM共享） 跨SM数据共享与原子操作的中转站 内存控制器（Memory Controller） HBM2e 接口，512-bit 位宽 内存控制器（Memory Controller） HBM2e 接口，512-bit 位宽 全局内存（HBM2e 显存）— 80 GB，带宽 ~2 TB/s 所有线程均可访问，但延迟高达 ~400-800 周期 图例 处理块 共享内存 L2缓存 全局显存

2.2 内存层次结构：从寄存器到全局内存

GPU的内存层次结构直接决定了数据访问的延迟与带宽，是性能调优的核心杠杆。从硬件物理位置和访问速度来看，可划分为以下层级：

  • 寄存器（Register）：每个线程私有，位于SM内部，访问延迟为0周期。但容量极为有限，每个SM的寄存器总量通常为256KB（A100），若每个线程占用过多寄存器，将直接限制SM上可驻留的线程总数。
  • 共享内存（Shared Memory）：同一线程块（Block）内所有线程共享，位于SM内部，延迟约20-30周期，带宽极高。常用于线程间数据复用与协作，但容量有限（A100为164KB/SM），需手动管理。
  • L1缓存：与共享内存在物理上共用同一块SRAM（可配置比例），用于缓存全局内存访问和局部变量。
  • L2缓存：所有SM共享，容量40MB（A100），延迟约200周期。跨SM的数据共享、原子操作、全局同步都依赖L2。
  • 全局内存（显存）：容量最大（80GB），但延迟高达400-800周期，带宽约2TB/s。所有线程均可访问，但必须通过内存控制器。
  • 常量内存与纹理内存：经过专门优化的只读缓存路径，适用于广播访问或空间局部性强的访问模式。

调优启示：共享内存与寄存器是“快而小”的存储，全局内存是“慢而大”的存储。性能调优的核心思路之一，就是最大化数据在寄存器与共享内存中的复用率，减少对全局内存的重复访问。例如，矩阵分块乘法的经典优化，就是通过将子矩阵加载到共享内存，将全局内存访问次数从O(N³)降低到O(N³/B)。

2.3 并行计算模型：线程层次与Warp执行

CUDA的并行模型采用三层线程层次：网格（Grid）→ 线程块（Block）→ 线程（Thread）。这种层次设计是为了匹配GPU的物理调度单元：线程块被调度到SM上执行，块内的线程通过共享内存协作，块之间则相互独立。

// CUDA线程层次示例
dim3 gridDim(4);      // 网格包含4个线程块
dim3 blockDim(256);   // 每个线程块包含256个线程
kernel<<<gridDim, blockDim>>>(args);

硬件执行时，线程块中的线程被进一步划分为Warp（32线程）。Warp是GPU硬件调度的最小单位，一个Warp中的所有线程在同一时刻执行同一条指令（SIMT模型）。若线程间发生分支发散，则各分支路径会被串行执行，造成部分线程闲置。

Warp执行的关键规则：

  1. 一个Warp中的32个线程共享程序计数器（PC），执行同一指令。
  2. Warp调度器在每个时钟周期选择一个就绪的Warp发射指令，通过线程级并行（TLP）隐藏内存延迟。
  3. 当Warp因访问全局内存而等待时，调度器会切换到其他Warp执行，这就是GPU高延迟容忍能力的来源。
  4. SM上可同时驻留的Warp数量受限于寄存器、共享内存等资源。例如，若每个线程使用64个寄存器，则每个SM最多驻留4096个线程（256KB/64B），即128个Warp。

从编程模型到硬件执行的映射关系如下：

编程抽象硬件映射协作范围同步机制
线程（Thread）CUDA Core上的执行流无（独立执行）无
线程块（Block）SM（一个块只在一个SM上执行）块内线程可协作__syncthreads()
网格（Grid）整个GPU（多个SM）块间独立仅通过全局内存或原子操作
Warp（32线程）SM内的调度单元Warp内隐式同步无显式同步

2.4 案例：矩阵乘法的内存访问优化

以经典的矩阵乘法C = A × B为例，直观展示内存层次对性能的影响。朴素实现中，每个线程直接读取全局内存中的A和B元素：

// 朴素版本：每个线程计算C[i][j]
__global__ void matmul_naive(float *A, float *B, float *C, int N) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    float sum = 0.0f;
    for (int k = 0; k < N; ++k) {
        sum += A[row * N + k] * B[k * N + col];  // 每次循环访问全局内存
    }
    C[row * N + col] = sum;
}

该实现中，每个线程需要访问2N次全局内存，总访问次数为2N³。当N=1024时，全局内存访问量高达2.1亿次，而实际计算量仅为10亿次浮点运算，内存带宽成为瓶颈。

优化策略是使用共享内存分块（Tiling）：将矩阵划分为大小为TILE×TILE的子块，每个线程块负责计算一个输出子块。先将A和B的子块加载到共享内存，再进行计算，从而将全局内存访问次数降低到2N³/TILE。

// 共享内存分块版本（TILE=16）
__global__ void matmul_tiled(float *A, float *B, float *C, int N) {
    __shared__ float As[TILE][TILE];
    __shared__ float Bs[TILE][TILE];
    int row = blockIdx.y * TILE + threadIdx.y;
    int col = blockIdx.x * TILE + threadIdx.x;
    float sum = 0.0f;
    for (int tile = 0; tile < N/TILE; ++tile) {
        // 协作加载到共享内存
        As[threadIdx.y][threadIdx.x] = A[row * N + tile * TILE + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(tile * TILE + threadIdx.y) * N + col];
        __syncthreads();  // 确保加载完成
        for (int k = 0; k < TILE; ++k) {
            sum += As[threadIdx.y][k] * Bs[k][threadIdx.x];
        }
        __syncthreads();  // 防止覆盖共享内存
    }
    C[row * N + col] = sum;
}

实测中，分块版本在A100上的性能可从约1.5 TFLOPS提升至约19 TFLOPS（FP32），提升超过10倍。这充分说明内存访问模式对GPU性能的决定性影响。

2.5 常见误区与调优复盘

误区一：认为线程数越多越好。当线程数超过SM可驻留的上限时，多余的线程会排队等待，并不会提升性能。应通过Occupancy计算器评估每个SM的驻留Warp数，找到资源使用与并行度的平衡点。

误区二：忽略内存合并访问（Memory Coalescing）。当Warp中32个线程访问连续的全局内存地址时，硬件可将其合并为少数几次内存事务。若访问模式为跨步或随机，则会产生大量冗余事务，带宽利用率骤降。例如，按列访问二维数组时，应改为按行访问或使用转置。

误区三：过度使用共享内存而不考虑Bank Conflict。共享内存被划分为32个Bank，若同一Warp中多个线程同时访问同一Bank的不同地址，会产生冲突并串行化。应通过填充（Padding）或调整数据布局来避免。

复盘总结：GPU性能调优的本质，是在理解硬件资源约束的前提下，设计出高并行度、高数据复用率、低访问延迟的执行方案。核心检查清单如下：

  • 每个SM的Occupancy是否达到50%以上？
  • 全局内存访问是否满足合并访问条件？
  • 共享内存是否有效复用了加载的数据？是否存在Bank Conflict？
  • 是否存在Warp发散？能否通过数据重排或分支优化消除？
  • 寄存器使用量是否限制了线程驻留数？

掌握GPU硬件组成与内存层次，是后续学习资源调度器原理、并发流优化、以及CUDA图（CUDA Graph）等高级主题的基石。只有深刻理解SM如何执行Warp、数据如何在各级存储间流动，才能在设计并行算法时做出正确的权衡决策。

← 上一章 📖 返回目录 下一章 →