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

GPU架构基础：从硬件结构理解Nsight Compute的分析视角

要真正读懂Nsight Compute的剖析报告，首先需要建立对GPU硬件架构的清晰认知。GPU并非“许多CPU核心的简单堆叠”，而是一种以大规模并行吞吐为设计目标的特殊处理器。NVIDIA GPU采用SIMT（单指令多线程）执行模型，其硬件层次结构直接决定了性能分析工具所呈现的每一个指标的含义。

核心硬件层次：从GPU到线程

以现代NVIDIA数据中心GPU（如A100、H100）为例，其硬件组织自上而下可分为四个关键层级：

GPU芯片（Device）：包含全部计算资源，由多个图形处理簇（GPC）组成，GPC之间通过高速互连总线通信。
流式多处理器（SM）：GPU的基本计算单元，每个SM包含独立的计算核心、缓存和调度硬件。一个GPU通常有数十到上百个SM。
处理块（Processing Block）：SM内部进一步划分为4个处理块，每个处理块包含一组CUDA核心、特殊函数单元（SFU）和寄存器文件。
CUDA核心：执行浮点/整数运算的最小执行单元，一个SM内通常有128个CUDA核心。

关键认知：Nsight Compute中报告的“SM占用率”“计算吞吐量”等指标，本质上是针对SM内部资源利用情况的度量。不理解SM内部结构，就无法定位性能瓶颈究竟出在哪个子系统中。

SM内部执行管线与Warp调度

SM内部最核心的调度机制围绕Warp展开。Warp是GPU执行指令的最小单位，包含32个并行线程。当CUDA程序启动内核时，线程块（Block）被分配到SM上，SM再将线程块中的线程划分为多个Warp。每个SM包含4个指令调度器，每个调度器负责管理一组Warp，每个周期可以发射一条指令。

这种设计使得GPU能够通过线程级并行来隐藏内存访问延迟——当一个Warp因等待数据而暂停时，调度器立即切换到另一个就绪的Warp执行计算。Nsight Compute中报告的“Warp Stall”原因（如“等待固定内存”“等待共享内存”等）正是对这一调度行为的量化呈现。

内存层次与数据局部性

GPU内存层次的设计目标是最大化数据复用率，减少对显存的访问。理解各级存储的容量与延迟差异，是优化内核的关键：

存储层级	容量（A100）	访问延迟	作用域	Nsight Compute 关注指标
寄存器文件	256KB / SM	~1 周期	单个线程	寄存器溢出（Spill）
共享内存 / L1	192KB / SM（可配置）	~30 周期	线程块内	共享内存 Bank 冲突
L2 缓存	40MB	~200 周期	GPU 全局	L2 命中率、扇出
HBM 显存	40GB	~400 周期	GPU 全局	DRAM 吞吐量、带宽利用率

实操建议：在Nsight Compute的“Memory Workload Analysis”页面中，如果发现L1/TEX Hit Rate低于50%，通常意味着数据局部性较差，应考虑使用共享内存进行数据分块（Tiling）或调整访问模式以提高缓存命中率。

案例：矩阵乘法内核的硬件映射

以经典的sgemm（单精度矩阵乘法）为例，观察硬件如何执行CUDA代码：

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

当启动配置为blockDim(16, 16)时，每个线程块包含256个线程，被划分为8个Warp。SM的4个调度器将这8个Warp分配到不同的处理块中执行。每个线程负责计算输出矩阵中的一个元素，内层循环每次迭代需要从全局内存读取两个浮点数。

从Nsight Compute的视角分析此内核，会观察到以下典型瓶颈：

全局内存负载过高：每个线程每次迭代执行2次全局内存读取，而计算仅1次乘加运算，导致Memory Throughput远高于Compute Throughput。
L2缓存命中率低：由于矩阵B的列访问模式导致缓存行利用率不足，L2命中率通常低于30%。
Warp Stall集中在“等待全局内存”：在Warp State统计中，“Long Scoreboard”占比可能超过60%。

优化方向是使用共享内存分块：将A和B的16×16子块加载到共享内存，使每个子块被16×16线程复用16次，从而将全局内存访问量减少16倍。优化后Nsight Compute中Memory Throughput显著下降，而Compute Throughput提升至接近峰值。

常见误区与复盘

误区一：认为“GPU核心数越多，程序一定越快”。实际上，如果内核受内存带宽限制，增加SM数量不会带来线性加速。Nsight Compute的Speed of Light图可以清晰显示当前内核是计算密集型还是内存密集型。

误区二：忽略Warp发散的影响。当Warp内32个线程执行不同分支时，GPU会串行执行每个分支路径，导致指令吞吐量下降。例如，若代码中if (threadIdx.x % 2 == 0)，则Warp内一半线程执行分支A，另一半执行分支B，总执行时间翻倍。Nsight Compute的Branch Efficiency指标低于100%即表明存在发散。

复盘要点：在分析任何内核时，建议遵循以下流程：首先查看Speed of Light确认瓶颈类型（计算/访存/延迟）；然后检查Occupancy确认SM资源利用是否充分；接着深入Warp State分析线程等待的具体原因；最后结合Memory Workload Analysis定位数据访问模式问题。这一分析路径能够帮助开发者从硬件底层理解性能瓶颈，而非盲目尝试各种优化技巧。