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

GPU架构基础：GPU与CPU的差异，英伟达GPU的硬件架构

核心原理：为什么GPU与CPU走的是两条路

CPU与GPU的根本差异，源于它们被设计来解决不同性质的问题。CPU面向延迟优化，需要快速响应单个复杂任务；而GPU面向吞吐量优化，追求在单位时间内处理海量并行的小任务。这种目标差异直接体现在芯片的晶体管分配上：CPU将大量面积用于控制单元（分支预测、乱序执行）和大容量缓存，而GPU则将绝大部分晶体管用于算术逻辑单元（ALU），仅保留极简的控制逻辑和小型缓存。

关键对比：CPU的核心数量少但单核能力强，适合执行分支密集、依赖关系复杂的串行逻辑；GPU拥有数千个轻量级核心，适合执行数据并行、计算密集且分支简单的任务。例如，一个4核CPU与一个拥有3072个CUDA核心的GPU相比，后者在浮点运算峰值上往往高出数十倍。

英伟达GPU硬件架构：从SM到CUDA核心

英伟达GPU的基本计算单元是流式多处理器（Streaming Multiprocessor, SM）。一个GPU由多个SM组成，例如Ampere架构的A100拥有108个SM，而消费级的RTX 3060拥有28个SM。每个SM内部包含若干CUDA核心（即浮点/整数运算单元）、共享内存、寄存器堆、线程调度器和特殊功能单元（SFU）。

CUDA核心是执行浮点运算的最小硬件单元，但它并不等同于CPU核心——它没有独立的指令获取与解码能力，必须依赖SM中的调度器统一分发指令。一个SM中的CUDA核心以32个为一组构成一个Warp，这是英伟达GPU线程调度的基本单位。当程序发出一个线程块（Block）时，SM会将其中的线程划分为多个Warp，每个Warp内的32个线程在同一时刻执行同一条指令（SIMT模式）。

硬件层级	典型配置（以Ampere架构为例）	作用
GPU 芯片	108 个 SM（A100）	整体计算资源池
SM（流式多处理器）	128 个 CUDA 核心 / SM	独立执行线程块，拥有调度器与共享内存
Warp	32 个线程 / Warp	硬件调度的最小单位，SIMT 执行
CUDA 核心	1 个 FP32 单元 / 核心	执行单精度浮点/整数运算

内存层次：数据流动的阶梯

英伟达GPU的内存层次结构决定了数据访问的效率，理解这一层次是优化性能的关键。从最靠近计算单元到最远，依次是：寄存器 → 共享内存 → L1/L2缓存 → 全局内存（显存）。访问速度从几个时钟周期到数百个时钟周期不等，差异巨大。

寄存器（Register）：每个线程私有的存储空间，速度最快，但容量极小（每个线程通常最多255个寄存器）。寄存器溢出会迫使数据写入本地内存，导致性能骤降。
共享内存（Shared Memory）：同一线程块内所有线程可访问的片上存储，延迟约为全局内存的1/20。常用于线程间数据交换和复用，是GPU优化中最灵活的手动管理缓存。
L1/L2缓存：硬件自动管理的缓存，L1与共享内存在物理上共用同一块SRAM（可配置比例），L2缓存则被整个GPU共享。
全局内存（Global Memory）：即显存（VRAM），容量最大（8~80GB），但延迟高达数百周期。所有线程均可访问，必须通过合并访问（Coalesced Access）模式提升带宽利用率。

性能直觉：如果全局内存访问延迟为400个时钟周期，而算术运算只需4个周期，那么一次未命中的内存访问浪费的时间足够执行100次浮点运算。因此，减少全局内存访问次数比减少计算量往往更能提升性能。

案例：矩阵乘法的内存层次利用

以经典的矩阵乘法 C = A × B 为例，朴素实现中每个线程直接读取全局内存中的A和B元素，会产生大量重复访问。利用共享内存进行分块（Tiling）优化后，每个线程块先将所需的数据子块加载到共享内存，再从共享内存中读取参与计算，大幅减少全局内存访问。

// 分块矩阵乘法核心逻辑（示意）
__global__ void matmul_tiled(float *A, float *B, float *C, int N) {
    __shared__ float As[TILE][TILE];
    __shared__ float Bs[TILE][TILE];
    int bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;
    float sum = 0.0f;

    for (int t = 0; t < N / TILE; ++t) {
        // 协作加载数据到共享内存
        As[ty][tx] = A[by * TILE + ty][t * TILE + tx];
        Bs[ty][tx] = B[t * TILE + ty][bx * TILE + tx];
        __syncthreads();  // 确保所有线程加载完成

        for (int k = 0; k < TILE; ++k)
            sum += As[ty][k] * Bs[k][tx];
        __syncthreads();  // 防止覆盖共享内存
    }
    C[by * TILE + ty][bx * TILE + tx] = sum;
}

该案例展示了如何通过共享内存将全局内存访问次数从 O(N³) 降低到 O(N³/TILE)。实际测试中，TILE=32时，A100上的性能可从约1.2 TFLOPS提升至19.5 TFLOPS，提升超过16倍。

误区与复盘：常见理解偏差

误区1：CUDA核心越多，性能一定越高。实际上，性能还受限于内存带宽、调度效率、分支发散和占用率。如果程序是内存密集型（Memory-bound），增加核心数不会带来线性提升。

误区2：共享内存越大越好。共享内存与L1缓存共用片上SRAM，若将共享内存配置得过大，L1缓存空间会被压缩，可能导致全局内存访问的缓存命中率下降。应根据程序访存模式动态配置。

误区3：忽略Warp发散。当Warp内32个线程执行不同分支时，所有分支会串行执行，导致性能成倍下降。例如，若一个Warp中一半线程走if路径、一半走else路径，实际执行时间为两条路径之和。在设计算法时应尽量让同一Warp内的线程保持执行路径一致。

复盘要点：在编写CUDA程序时，应首先判断程序是计算密集型还是访存密集型。对于前者，关注算术强度与指令混合；对于后者，重点优化数据布局以实现合并访问，并利用共享内存减少全局内存往返。同时，通过调整网格（Grid）和线程块（Block）的维度，确保SM的占用率足够高，以隐藏内存访问延迟。