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

GPU架构基础：GPU硬件架构及其对内核性能的影响

核心原理：从“多核”到“众核”的硬件设计哲学

CPU与GPU的硬件设计目标截然不同。CPU追求低延迟，为此将大量芯片面积用于分支预测、乱序执行和超大缓存；而GPU追求高吞吐量，其设计哲学是“用更多的线程掩盖延迟”。GPU将芯片面积主要分配给计算单元（ALU）和线程调度器，而非缓存与控制逻辑。理解这一根本差异，是后续所有调优决策的出发点。

关键认知：GPU并非“更快的CPU”，而是“规模化的并行计算引擎”。一个典型的现代GPU（如NVIDIA A100）拥有108个SM（Streaming Multiprocessor，流式多处理器），每个SM包含64个FP32 CUDA核心，总计6912个核心。而一颗高端CPU通常只有8~16个核心。这种数量级的差异决定了：GPU擅长处理数据并行任务，而非任务并行或强依赖串行任务。

下图展示了GPU硬件架构的核心层次——从设备级到线程级，以及它们之间的映射关系。这是理解内核性能的“第一性原理”图。

方法流程：线程层次与硬件资源的映射关系

CUDA的线程层次设计为三层：Grid（网格）→ Block（线程块）→ Thread（线程）。硬件执行时，Block被调度到SM上运行，而SM内部将Block进一步划分为Warp（线程束）——每32个连续线程组成一个Warp，这是GPU执行和调度的最小单位。

理解这一映射关系对性能调优至关重要，因为：

Block数量决定了GPU的并行度是否足够。若Block总数少于SM总数，则部分SM空闲，造成资源浪费。
Block大小（即每Block的线程数）影响SM内资源的分配效率。Block过大可能导致寄存器溢出，过小则可能无法充分隐藏延迟。
Warp的收敛性直接影响执行效率。若Warp内32个线程发生分支分歧（if-else），则所有分支路径都会被串行执行，导致吞吐量下降。

一个典型的调优流程如下：

确定目标SM占用率：计算每个SM可容纳的最大线程数（通常为2048），用blockSize除以该值，得到每个SM可容纳的Block数。
计算实际占用率：使用cudaOccupancyMaxPotentialBlockSize API或手动计算，考虑寄存器、共享内存等资源限制。
调整Block大小：通常选择128、256或512的倍数（Warp大小的整数倍），避免出现“尾部Warp”浪费。
验证内存访问模式：确保同一Warp内的线程访问连续内存地址，实现合并访问（Coalesced Access）。

案例/场景：内存层次对矩阵乘法的性能影响

以经典的SGEMM（单精度矩阵乘法）为例，展示内存层次如何决定性能上限。假设计算C = A × B，矩阵大小为4096×4096。

朴素实现中，每个线程直接读取全局内存中的A和B元素。由于全局内存延迟高达400~800个周期，而计算仅需几个周期，GPU必须依赖大量线程切换来隐藏延迟。此时，性能受限于内存带宽而非计算能力。

优化后的实现使用共享内存（Shared Memory）和寄存器分块（Register Tiling）：

// 每个Block负责计算 16×16 的输出子矩阵
// 使用共享内存缓存A和B的对应子块
__global__ void sgemm_shared(float* A, float* B, float* C, int N) {
    __shared__ float As[16][16];
    __shared__ float Bs[16][16];
    
    int row = blockIdx.y * 16 + threadIdx.y;
    int col = blockIdx.x * 16 + threadIdx.x;
    float sum = 0.0f;
    
    for (int k = 0; k < N; k += 16) {
        // 协作加载到共享内存（合并访问）
        As[threadIdx.y][threadIdx.x] = A[row * N + k + threadIdx.x];
        Bs[threadIdx.y][threadIdx.x] = B[(k + threadIdx.y) * N + col];
        __syncthreads();  // 同步，确保加载完成
        
        for (int i = 0; i < 16; i++) {
            sum += As[threadIdx.y][i] * Bs[i][threadIdx.x];
        }
        __syncthreads();  // 防止覆盖
    }
    C[row * N + col] = sum;
}

该优化将全局内存访问次数从2×N次降低到2×N/16次，性能提升可达10~20倍。下表对比了不同优化级别的性能差异：

实现版本	全局内存访问量	共享内存使用	典型性能 (TFLOPS)	瓶颈
朴素实现	2×N³ 次	无	~1.5	内存带宽
共享内存分块	2×N³/16 次	2×16×16×4B	~8.0	计算/同步开销
寄存器分块+向量化	2×N³/64 次	2×32×32×4B	~15.0	寄存器压力

误区/复盘：常见性能陷阱与诊断方法

误区一：盲目追求“满占用率”

许多开发者认为SM占用率越高越好，于是将Block大小设为32或64以增加Block数量。但过小的Block会导致每个Block的计算量不足，无法有效利用共享内存和寄存器分块，反而降低算术强度。正确做法是：在占用率与单线程计算效率之间寻找平衡点。

误区二：忽略Warp分歧的代价

当Warp内线程执行不同分支时，GPU会串行执行所有分支路径。例如，一个if (data[i] > 0)判断若导致Warp内一半线程走true路径、一半走false路径，则执行时间翻倍。调优时应尽量让分支条件基于threadIdx / 32（即Warp ID），保证整个Warp统一决策。

复盘建议：使用Profiler定位瓶颈

使用ncu（NVIDIA Nsight Compute）分析内核时，重点关注以下指标：

Memory Throughput：若接近100%，说明内存带宽饱和，需减少全局访问或提高缓存命中率。
Compute Throughput：若接近100%，说明计算饱和，需优化算法减少冗余计算。
Warp Stall Reasons：查看Warp等待的原因（如Long Scoreboard表示等待全局内存返回），判断是否需要增加并行度或使用预取。
Occupancy：实际占用率与理论值的差距，往往源于寄存器或共享内存超限。

最后，一个常被忽视的硬件细节是L2缓存分区。现代GPU的L2缓存被划分为多个切片（Slice），每个切片服务于特定的内存地址范围。若多个SM同时访问同一L2切片对应的地址，会造成L2切片冲突，降低有效带宽。在编写内核时，可通过调整数据布局或使用cudaMemAdvise等API优化L2访问模式，但这属于进阶调优范畴，需结合具体硬件架构（如Hopper的分布式共享内存）进一步分析。