📚 课程目录第2章 GPU架构基础

GPU架构基础：硬件架构、内存层次结构与执行模型

在开始使用 Nsight Compute 定位内核热点之前，必须先建立对 GPU 硬件的基本认知。许多性能分析结论——例如“访存延迟过高”“占用率不足”“Bank Conflict 严重”——本质上都源于硬件架构的特定行为。如果缺乏架构基础，分析工具给出的计数器数值就只是一堆无意义的数字。本节将围绕 NVIDIA GPU 的三大核心主题展开：硬件架构、内存层次结构与执行模型。

2.1 GPU 硬件架构：从 SM 到 CUDA Core

NVIDIA GPU 的基本计算单元是 流式多处理器（Streaming Multiprocessor, SM）。每个 SM 内部包含若干 CUDA Core（即浮点/整数运算单元）、特殊函数单元（SFU）、加载/存储单元（LSU）以及共享内存和寄存器堆。以 Ampere 架构的 GA100 为例，每个 SM 包含 128 个 CUDA Core，而整个 GPU 由多个 SM 组成（如 A100 有 108 个 SM）。

理解 SM 的关键在于：线程并不是独立调度执行的，而是以 Warp（通常 32 个线程）为单位。一个 Warp 中的所有线程在同一时刻执行同一条指令（SIMT 模型），如果发生分支分歧，则不同路径会被串行执行。因此，分析热点时，Warp 级行为往往比线程级行为更重要。

架构启示：Nsight Compute 中的 sm__throughput 和 gpu__compute_memory_throughput 等指标，本质上反映的是 SM 内部执行单元与访存管线的利用率。如果某个内核的算术指令吞吐接近 100%，说明计算单元已饱和，优化方向应转向减少指令数或提高 ILP（指令级并行）。

2.2 内存层次结构：带宽与延迟的权衡

GPU 内存层次从寄存器到全局内存，每一级的容量、带宽和延迟差异巨大。下表总结了主要内存类型及其关键特性：

内存类型	作用域	容量	延迟	带宽	用途
寄存器文件	单个线程	每 SM 256KB（A100）	~1 周期	极高	局部变量、中间计算
共享内存	单个线程块	每 SM 最多 164KB	~20-30 周期	高	块内数据共享、协作计算
L1 缓存	单个 SM	与共享内存可配置（A100 共 192KB）	~30 周期	高	缓存全局内存访问
L2 缓存	整个 GPU	40MB（A100）	~200 周期	中	跨 SM 数据共享、原子操作
全局内存（HBM2）	所有线程	16-80GB	~400-800 周期	~2TB/s（A100）	主数据存储

在 Nsight Compute 中，l1tex__throughput 和 lts__throughput 分别对应 L1/L2 缓存利用率。当 L2 命中率低而全局内存吞吐高时，通常意味着数据局部性差，需要考虑使用共享内存或调整访问模式。

2.3 执行模型：线程层次与调度机制

CUDA 的执行模型采用三层线程层次：Grid → Block → Thread。一个 Grid 包含多个 Block，每个 Block 包含多个 Thread。Block 被调度到 SM 上执行，一个 SM 可以同时驻留多个 Block（受资源限制）。

关键的执行机制包括：

Warp 调度：SM 将 Block 中的线程划分为 Warp（每 32 个连续线程一个 Warp），每个 Warp 独立被调度器发射指令。
占用率（Occupancy）：指 SM 上活跃 Warp 数量与最大可支持 Warp 数量的比值。高占用率有助于隐藏访存延迟，但并非越高越好——过高的占用率可能导致寄存器溢出或缓存抖动。
延迟隐藏：GPU 通过快速切换不同 Warp 来隐藏访存和指令延迟。当某个 Warp 因等待数据而阻塞时，调度器立即切换到其他就绪 Warp。

下面的 SVG 展示了从 Grid 到 Warp 的完整执行流水线，以及内存层次与 SM 的对应关系：

2.4 内存访问模式与合并访问

全局内存访问效率高度依赖 合并访问（Coalesced Access）。当同一个 Warp 中的 32 个线程访问连续的内存地址时，硬件可以将这些访问合并为少数几个内存事务（通常 1-2 个 128 字节事务）。反之，如果线程访问的是分散的地址（如按列访问二维数组），则会产生大量冗余事务，浪费带宽。

在 Nsight Compute 中，l1tex__data_bank_conflicts_pipe_lsu_mem_global_op_ld 和 lts__t_sectors_srcunit_tex_op_read 等指标可以反映合并访问的效率。当 gld_efficiency（全局加载效率）低于 100% 时，说明存在非合并访问。

示例：以下两种访问模式在 A100 上的性能差异可达 8-10 倍：

// 高效：合并访问（连续地址）
__global__ void coalesced(float* out, const float* in, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) out[i] = in[i] * 2.0f;
}

// 低效：非合并访问（跨步访问）
__global__ void strided(float* out, const float* in, int N, int stride) {
    int i = (blockIdx.x * blockDim.x + threadIdx.x) * stride;
    if (i < N) out[i] = in[i] * 2.0f;
}

2.5 常见误区与性能分析要点

在实际使用 Nsight Compute 分析内核时，初学者常犯以下错误：

误区一：只关注占用率。高占用率并不保证高性能。如果内核受访存带宽限制，过高的占用率反而可能导致 L2 缓存抖动。应结合 memory_throughput 和 compute_throughput 综合判断。
误区二：忽略 Warp 分歧。如果 Warp 内线程执行不同分支，所有路径都会被串行执行，指令数成倍增加。Nsight Compute 的 branch_efficiency 指标可以量化此问题。
误区三：将所有数据都放入共享内存。共享内存容量有限，过度使用会降低占用率。应优先优化全局内存访问模式，再考虑使用共享内存。
误区四：忽视寄存器溢出。当每个线程使用的寄存器超过硬件限制时，局部变量会被溢出到本地内存（L1/L2），导致性能骤降。检查 launch__registers_per_thread 和 local_memory_overflow 指标。

实战提示：在 Nsight Compute 的 Source 页签中，可以逐行查看每条指令对应的硬件计数器。当发现某行代码的 stall_long_scoreboard 周期数较高时，说明该指令正在等待全局内存数据返回——此时应考虑使用 __ldg() 或调整数据布局以提高缓存命中率。

2.6 从架构到分析：一个完整示例

假设我们有一个简单的向量加法内核，在 A100 上运行。使用 ncu --set full 分析后，关键指标如下：

Section: GPU Speed Of Light Throughput
    Compute (SM) Throughput:    62.3%
    Memory Throughput:          98.7%

Section: Occupancy
    Block Limit:                32
    Active Warps:               48
    Theoretical Occupancy:      75.0%
    Achieved Occupancy:         71.2%

Section: Memory Workload Analysis
    Memory Throughput:          1.94 TB/s
    L1/TEX Hit Rate:            85.3%
    L2 Hit Rate:                92.1%
    Mem Busy:                   97.5%
    Max Bandwidth:              98.2%

从这些数据可以判断：该内核是访存密集型（Memory Throughput 98.7% 而 Compute 仅 62.3%），L1/L2 命中率较高说明局部性良好，但内存带宽已接近饱和。优化方向应聚焦于减少全局内存访问量（如使用向量化加载 float4）或提高计算密度（如融合操作）。