GPU 硬件架构基础：从 SM 到内存层次

要理解 GPU 性能分析工具的输出，首先必须建立对 GPU 硬件的基本认知。GPU 并非“一个巨大的并行核”，而是一组由 流式多处理器（Streaming Multiprocessor，SM） 构成的阵列，每个 SM 内部又包含大量计算单元、寄存器堆和共享内存。性能瓶颈往往不是单一资源耗尽，而是多个硬件子系统之间的配合失衡。

2.1 SM 内部结构：线程执行的物理载体

以 NVIDIA 的 Ampere 架构为例，一个 SM 通常包含以下关键资源：

  • CUDA 核心（FP32/INT32 单元）：执行标量浮点与整数运算。
  • Tensor Core：执行矩阵乘加运算，用于深度学习训练与推理。
  • SFU（特殊函数单元）：处理超越函数（如 sin、exp、倒数）。
  • 寄存器堆（Register File）：每个 SM 拥有 64K 个 32 位寄存器，是访问速度最快的存储。
  • 共享内存（Shared Memory）：SM 内所有线程共享的低延迟存储，容量通常为 128KB~228KB。
  • 线程调度器（Warp Scheduler）：每个 SM 包含 4 个调度器，每个调度器每周期可发射一条指令给一个 warp（32 线程）。

线程在 GPU 上的执行粒度是 warp（32 个线程）。调度器以 warp 为单位分配指令，而 warp 内的线程必须执行相同的指令（SIMT 模型）。如果 warp 内出现分支发散，则不同路径会被串行执行，造成吞吐损失。

关键概念：占用率（Occupancy）= 活跃 warp 数 / SM 最大可容纳 warp 数。高占用率有助于隐藏访存延迟，但并非越高越好——过高的占用率可能导致寄存器溢出或共享内存不足。
GPU 硬件架构与内存层次 SM（流式多处理器） Warp Scheduler 0 每周期发射 1 条指令 Warp Scheduler 1 每周期发射 1 条指令 CUDA 核心 FP32 × 32 INT32 × 32 Tensor Core 矩阵运算 FP16/BF16 SFU / LD/ST 特殊函数 加载/存储 寄存器堆（Register File）— 64K × 32-bit 每个线程最多 255 个寄存器，访问延迟 ~1 周期 共享内存（Shared Memory）— 最大 228 KB 低延迟，SM 内线程间通信，需同步 L1 缓存 / 纹理缓存（可配置） 与共享内存共享 128KB 容量 线程块（Thread Block）驻留在 SM 上 Block 0: 256 线程 | Block 1: 128 线程 warp = 32 线程，SM 最多 64 warp 内存层次结构 寄存器（Register）— 每线程私有 容量最小，速度最快（~1 周期），编译器自动分配 共享内存（Shared Memory）— Block 内共享 延迟 ~20-30 周期，需 __syncthreads() 同步 L2 缓存 — GPU 全局共享 延迟 ~200 周期，所有 SM 共享，容量 40MB（A100） 全局内存（Global Memory / VRAM） 延迟 ~400-800 周期，容量 16-80GB，需通过 PCIe/NVLink 访问 常量内存 只读，广播优化 纹理内存 只读，空间局部性优化 典型带宽：寄存器 > 共享内存 > L2 > 全局内存 带宽差距可达 2 个数量级 数据通路 性能分析关注：SM 利用率、内存带宽、warp 发散、共享内存 bank 冲突

2.2 内核执行模型：从 Grid 到 Thread

CUDA 内核（Kernel）启动时，GPU 以三层结构组织线程：

  1. Grid（网格）：一次内核启动的全部线程集合。
  2. Block（线程块）：Grid 的子集，块内线程可协作（共享内存 + 同步）。
  3. Thread（线程）：最小执行单元，每个线程有独立寄存器状态。

硬件调度器将 Block 分配到 SM 上执行。一个 SM 可同时驻留多个 Block，但受限于寄存器数量、共享内存容量和最大线程数。当 Block 中的线程全部执行完毕后，该 Block 退出，SM 再调度新的 Block 进入。这种调度机制使得 GPU 能自动隐藏延迟——当一个 warp 等待内存访问结果时，调度器可以切换到其他就绪的 warp。

性能分析视角：如果内核的 Block 尺寸设置过小（如 64 线程），可能导致 SM 无法达到足够的占用率来隐藏延迟；如果 Block 过大（如 1024 线程），则可能因寄存器不足导致每个线程只能使用很少的寄存器，迫使编译器进行局部内存溢出（spill）。

2.3 性能瓶颈的常见类型

在 GPU 性能分析中，瓶颈通常可归为以下五类。使用 ncu（NVIDIA Nsight Compute）或 nprof 时，首先判断瓶颈类型，再针对性优化。

瓶颈类型 典型特征 分析指标 优化方向
计算瓶颈 SM 中 ALU/Tensor 利用率高，但内存吞吐低 SM 吞吐率（% of peak） 减少冗余计算、使用快速数学函数、提高 ILP
访存瓶颈 内存吞吐接近峰值，SM 利用率不高 DRAM 吞吐率、L2 命中率 合并访问、使用共享内存、调整数据布局
延迟瓶颈 占用率低，warp 大量处于等待状态 Warp Stall 原因、Occupancy 增加 Block 数、减少寄存器使用、增大并行度
同步瓶颈 __syncthreads() 或原子操作导致大量等待 Barrier 等待周期、原子操作冲突率 减少同步点、使用私有化、避免频繁原子操作
指令发散瓶颈 warp 内分支导致串行执行 Branch Efficiency（分支效率） 重构代码避免 warp 内分支、使用谓词化

2.4 案例：矩阵乘法的瓶颈识别

以经典的 SGEMM（单精度矩阵乘法）为例，假设内核使用朴素的三重循环实现：

__global__ void sgemm_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;
}

使用 ncu --set full 分析后，可能得到以下结果：

  • Memory Throughput：85% 峰值 → 访存瓶颈
  • Compute (SM) Throughput：仅 12%
  • L1/TEX Cache Hit Rate：23%（每次访问都未命中）
  • Warp Stall 原因：Long Scoreboard（等待全局内存返回）占 78%

结论：该内核完全受限于全局内存访问。每个线程每次循环都直接访问全局内存，且相邻线程访问的地址不连续（B 矩阵按列访问），导致无法合并。优化方案是使用 共享内存分块（tiling）：将 A 和 B 的子块加载到共享内存，然后从共享内存中读取数据进行计算，将全局内存访问次数从 O(N³) 降低到 O(N³ / tile_size)。

常见误区：初学者往往认为“GPU 瓶颈一定是计算能力不足”。实际上，在大多数真实内核中，访存瓶颈和延迟瓶颈出现的频率远高于纯计算瓶颈。务必先通过分析工具确认瓶颈类型，再决定优化策略，避免盲目优化计算而忽略真正的瓶颈。

2.5 复盘：性能分析中的关键思维

在后续章节中，我们将深入使用 Nsight Compute 等工具。但在此之前，请记住以下原则：

  1. 先测量，再优化：不要凭直觉猜测瓶颈，用工具获得量化数据。
  2. 关注“瓶颈资源”：GPU 性能受限于最接近 100% 利用率的资源。
  3. 区分“吞吐型”与“延迟型”瓶颈：前者需要减少数据量，后者需要提高并行度。
  4. 理解硬件再谈优化：共享内存、寄存器、warp 调度等硬件特性决定了优化手段的上限。
← 上一章 📖 返回目录 下一章 →