GPU硬件基础：GPU架构演进、SM/CU核心结构、内存层次、线程调度与SIMT执行模型

要真正理解GPU加速计算，必须从硬件底层出发。CPU与GPU的协同设计并非简单的“把任务丢给GPU”，而是需要理解GPU的架构哲学——它如何通过大规模并行隐藏延迟、如何组织线程、如何管理内存带宽。本章将系统拆解GPU的架构演进脉络、流式多处理器（SM）的内部结构、内存层次设计、线程调度机制以及SIMT执行模型，帮助你在后续编写CUDA/OpenCL内核时，能够从硬件行为反推代码优化方向。

2.1 GPU架构演进：从固定管线到通用并行计算

GPU的演进史本质上是从“专用图形加速器”向“通用并行处理器”的转变过程。早期GPU（2000年前后）采用固定功能管线，顶点着色和像素着色由专用硬件单元完成，程序员无法改变其行为。2001年NVIDIA GeForce3引入可编程顶点着色器，2006年NVIDIA GeForce 8800（G80架构）首次将统一着色器（Unified Shader）引入GPU，所有着色器单元既能处理顶点也能处理像素，这为通用计算（GPGPU）铺平了道路。

从G80到Hopper/Blackwell，GPU架构经历了几个关键里程碑：

架构演进的核心驱动力是并行度与通用性的平衡。早期GPU的并行度来自大量简单核心，但缺乏通用编程模型；现代GPU则在保持高并行度的同时，增加了缓存层次、分支预测、Tensor Core等专用单元，使得GPU不仅能做图形渲染，还能高效运行深度学习、科学计算、数据库等通用负载。

2.2 SM/CU核心结构：GPU的基本计算单元

流式多处理器（Streaming Multiprocessor，SM）是GPU的核心计算单元，相当于CPU中的“核中核”。一个现代GPU（如A100）包含108个SM，每个SM内部又包含多个处理块。理解SM的内部结构是编写高效内核的前提，因为线程调度、资源分配、访存优化都发生在SM层面。

以NVIDIA Ampere架构的GA100 SM为例，其内部结构如下：

• 4个处理块（Processing Block）：每个处理块包含32个FP32 CUDA核心、16个INT32核心、4个Tensor Core（第三代）、4个SFU（特殊函数单元）、4个LD/ST（加载/存储）单元。
• Warp调度器：每个SM有4个Warp调度器，每个调度器每时钟周期可以发射1条指令到其对应的处理块。这意味着SM每周期最多执行4条不同的Warp指令。
• 寄存器文件：SM拥有64K个32位寄存器（共256KB），这些寄存器以Warp为单位分配给线程。每个线程最多可使用255个寄存器，但实际分配取决于内核需求和占用率。
• 共享内存/L1缓存：Ampere架构中共享内存和L1缓存统一为192KB（可配置为共享内存或缓存），A100上为164KB可配置。共享内存以32个银行（Bank）组织，每个银行4字节宽，支持无冲突并行访问。
• 统一L1/纹理缓存：用于缓存全局内存访问和纹理采样，减少对L2的访问压力。

在AMD的CDNA架构中，对应的单元称为CU（Compute Unit），每个CU包含2个SIMD单元（每个SIMD有32个ALU），共享一个LDS（Local Data Share，相当于共享内存）。虽然命名不同，但设计哲学一致：用大量简单核心配合高带宽内存，通过线程级并行隐藏延迟。

2.3 内存层次：从寄存器到全局内存的完整路径

GPU的内存层次设计遵循“容量越大、速度越慢”的原则，但与传统CPU缓存层次不同，GPU的内存层次更强调带宽而非延迟。GPU通过大量并发线程隐藏内存延迟，因此即使全局内存延迟高达数百周期，只要线程数足够多，计算单元就不会空闲。

完整的内存层次从快到慢排列如下：

内存访问的合并（Coalescing）是GPU性能的关键。当Warp中的32个线程访问全局内存时，硬件会将这32个访问合并为尽可能少的缓存行事务。理想情况下，如果32个线程访问连续地址，只需1个128字节的事务；如果访问分散，则可能需要32个独立事务，带宽利用率下降32倍。

2.4 线程调度与SIMT执行模型

GPU的线程调度基于Warp（线程束）概念——32个线程组成一个Warp，这是硬件调度的最小单位。Warp中的所有线程在同一时刻执行同一条指令（SIMT，Single Instruction Multiple Thread），但每个线程有自己的寄存器状态和程序计数器，因此可以独立分支。

SIMT与传统的SIMD（如CPU的AVX）有本质区别：SIMD要求数据在向量寄存器中对齐，而SIMT允许每个线程访问任意地址。SIMT的灵活性来自硬件调度器的复杂逻辑——每个Warp调度器维护多个Warp的就绪状态，当某个Warp因内存访问而阻塞时，调度器立即切换到另一个就绪的Warp，从而实现延迟隐藏。

线程调度的完整流程如下：

1. 内核启动：CPU通过驱动将内核代码和参数传递给GPU，GigaThread引擎创建线程网格（Grid）。
2. 线程块分发：GigaThread将线程块（Block）分发到各个SM。分发策略基于SM的资源可用性（寄存器、共享内存、线程槽位）。
3. Warp划分：每个线程块内的线程被划分为多个Warp（线程块大小除以32，向上取整）。例如，256线程的块分为8个Warp。
4. Warp调度：SM中的Warp调度器每周期选择1个就绪Warp发射指令。就绪条件包括：指令无数据依赖、操作数寄存器已就绪、无内存屏障阻塞。
5. 执行与切换：当Warp执行访存指令时，由于延迟较长，调度器切换到其他Warp执行计算指令。这种切换开销几乎为零，因为寄存器状态已保存在物理寄存器文件中。

分支发散（Branch Divergence）是SIMT模型的主要性能杀手。当Warp内不同线程执行不同分支时，硬件会串行执行每个分支路径，并屏蔽不满足条件的线程。例如，一个Warp中16个线程走if路径，16个线程走else路径，则两条路径串行执行，性能减半。

2.5 实战案例：矩阵乘法的硬件适配优化

矩阵乘法（C = A × B）是GPU计算的经典基准，也是理解硬件特性的最佳案例。我们以一个具体的优化过程来展示如何将硬件知识转化为性能提升。

背景与约束：假设矩阵大小为4096×4096（单精度浮点），运行在A100 GPU上（108个SM，每个SM 164KB共享内存，峰值FP32约19.5 TFLOPS）。初始实现为朴素的三重循环，每个线程计算C的一个元素。

初始性能：朴素实现中，每个线程需要读取A的一行和B的一列，共2×4096次全局内存访问。总访问量为4096×4096×2×4096×4字节 ≈ 550GB，而A100的全局内存带宽约1.5TB/s，理论最短时间约0.37秒。但实际运行时间为2.1秒，效率仅17%。瓶颈在于：全局内存访问未合并（每个线程访问不同行/列），且没有利用共享内存复用数据。

优化步骤：

// 步骤1：分块（Tiling）——利用共享内存
#define TILE_SIZE 32
__global__ void matmul_tiled(float* A, float* B, float* C, int N) {
    __shared__ float As[TILE_SIZE][TILE_SIZE];
    __shared__ float Bs[TILE_SIZE][TILE_SIZE];
    int row = blockIdx.y * TILE_SIZE + threadIdx.y;
    int col = blockIdx.x * TILE_SIZE + threadIdx.x;
    float sum = 0.0f;
    
    for (int k = 0; k < N; k += TILE_SIZE) {
        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 < TILE_SIZE; i++) {
            sum += As[threadIdx.y][i] * Bs[i][threadIdx.x];
        }
        __syncthreads();
    }
    C[row * N + col] = sum;
}

结果对比：

复盘经验：分块大小32是硬件友好的选择——32对应Warp大小，确保每个Warp的线程访问连续的共享内存地址，避免银行冲突。进一步使用float4向量化后，每个线程一次加载16字节，减少LD/ST指令数量，提升内存子系统利用率。最终性能接近理论峰值，说明硬件资源已被充分利用。

2.6 常见误区与总结复盘

在GPU硬件基础的学习中，以下误区最为常见，需要特别注意：

• 误区一：认为GPU核心数越多性能越高。实际上，GPU的FLOPS需要足够的数据并行度才能发挥。如果内核的线程数少于SM可容纳的Warp总数，大量核心将空闲。通常需要至少数万个线程才能充分利用现代GPU。
• 误区二：忽略内存带宽瓶颈。很多计算任务（如向量加法、数据过滤）是内存密集型而非计算密集型。对于这类任务，优化计算指令数毫无意义，应聚焦于减少内存访问量或提高访问效率。
• 误区三：过度使用共享内存。共享内存虽然快，但容量有限。如果数据复用率不高（如每个数据只被使用一次），使用共享内存反而增加同步开销和资源占用。
• 误区四：忽视Warp调度的影响。当Warp数量不足时，调度器无法有效隐藏延迟。例如，一个SM最多驻留64个Warp（2048线程），如果内核只启动少量线程块，每个SM可能只有几个Warp，内存延迟将直接暴露。
• 误区五：认为所有GPU架构相同。不同代际的GPU在共享内存大小、寄存器数量、Tensor Core支持、NVLink带宽等方面差异巨大。代码在A100上优化良好，迁移到RTX 4090（Ada架构）可能需要重新调整分块大小和向量化宽度。

总结：GPU硬件基础的核心在于理解“并行度换延迟”的设计哲学。SM是计算的基本单元，内存层次决定了数据搬运的成本，SIMT模型规定了线程的执行方式。在实际开发中，应始终从硬件行为出发思考优化方向：先确保内存访问合并，再考虑共享内存复用，最后调整计算指令和调度策略。掌握这些硬件原理，你才能在CPU-GPU协同设计中做出正确的架构决策——例如，哪些任务适合放在GPU执行（数据并行、计算密集），哪些任务应保留在CPU（逻辑控制、小规模数据、I/O操作），以及如何设计数据流以减少PCIe传输开销。