2、GPU基础架构：从图形处理器到通用计算，SIMT/SIMD模型、流处理器、线程束（Warp）与波前（Wavefront）概念

GPU（Graphics Processing Unit）最初是为图形渲染而设计的专用处理器，其核心任务是对大量顶点、像素进行并行变换与着色。然而，正是这种“为大规模并行而生”的硬件基因，使得GPU在21世纪初被逐步改造为通用计算引擎（GPGPU，General-Purpose computing on GPU）。理解GPU的基础架构，必须从三个关键维度入手：执行模型（SIMT/SIMD）、硬件执行单元（流处理器）以及线程调度组织（Warp/Wavefront）。这三者共同构成了现代GPU并行计算的核心骨架，也是后续深入学习显存层次、缓存一致性、Tensor Core等高级主题的基石。

2.1 从图形渲染到通用计算：GPU 的“变形记”

GPU 的早期架构完全围绕图形流水线设计：顶点着色器（Vertex Shader）处理几何变换，像素着色器（Pixel Shader）处理片段着色，两者都是典型的“数据并行”任务——每个顶点或像素的操作相互独立，且执行相同的计算逻辑。这种“同构大规模并行”的特性，恰好与 SIMD（Single Instruction, Multiple Data）模型高度契合。早期的 GPU 使用固定功能单元，后来逐步演变为可编程着色器，但本质上仍然是“宽 SIMD”引擎：一条指令驱动多个 ALU 对不同的数据执行相同操作。

2006 年，NVIDIA 推出 CUDA（Compute Unified Device Architecture），首次将 GPU 的并行计算能力以通用编程接口的形式开放给开发者。这一变革的关键在于：GPU 不再仅仅是“图形加速卡”，而是一个“大规模并行协处理器”。硬件上，统一着色器架构（Unified Shader Architecture）将顶点和像素着色器合并为通用的流处理器（Streaming Processor, SP），使得同一批硬件单元既可以处理图形任务，也可以运行通用计算内核（Kernel）。AMD 随后也推出了类似的 Stream Processor 架构和 OpenCL 支持。

从硬件角度看，GPU 与 CPU 的设计哲学截然不同：CPU 追求低延迟（Latency-Oriented），通过大容量缓存、分支预测、乱序执行来加速单线程；GPU 则追求高吞吐（Throughput-Oriented），通过数千个 ALU 和极简的控制逻辑，以“人多力量大”的方式换取整体计算速度。这种差异直接体现在芯片面积分配上：CPU 将大量晶体管用于缓存和控制，而 GPU 将超过 80% 的晶体管用于 ALU 阵列和寄存器堆。

2.2 SIMT 与 SIMD：看似相同，实则迥异

SIMD（单指令多数据）是经典的并行执行模型，在 CPU 的 MMX/SSE/AVX 指令集和早期的 GPU 向量处理器中广泛使用。其核心特征是：一条指令同时作用于多个数据元素，数据以向量形式打包，硬件执行单元以固定宽度（如 128 位、256 位、512 位）并行处理。程序员需要显式地将数据组织为向量类型，并调用对应的向量指令。例如，使用 AVX-512 指令一次可以处理 16 个单精度浮点数。

SIMT（单指令多线程）是 NVIDIA 提出的执行模型，它是对 SIMD 的“线程化”抽象。在 SIMT 模型中，程序员编写的是标量线程代码（每个线程处理一个数据元素），硬件自动将多个线程打包为一个 Warp（通常 32 个线程），并让整个 Warp 在同一时刻执行同一条指令。关键区别在于：

• 编程视角：SIMD 需要程序员手动管理向量宽度和数据布局；SIMT 中程序员只需编写标量代码，硬件自动完成线程打包。
• 分支处理：SIMD 遇到分支时通常需要“掩码”操作（predication），所有通道执行所有路径但用掩码屏蔽结果；SIMT 则允许 Warp 内的线程独立维护程序计数器（PC），但实际执行时仍以锁步（lock-step）方式进行——如果分支发散，则串行执行各分支路径。
• 灵活性：SIMT 允许线程有独立的寄存器状态和内存地址，因此可以处理不规则的数据访问模式（如稀疏矩阵）；SIMD 则要求数据在内存中连续对齐。

需要特别指出的是，SIMT 并非 SIMD 的替代品，而是 SIMD 的“超集”——SIMT 硬件底层仍然使用 SIMD 风格的宽执行单元（如 32 个 FP32 ALU 并行工作），但通过线程抽象提供了更高的编程灵活性和硬件利用率。AMD 的 Wavefront 概念与 NVIDIA 的 Warp 本质相同，只是 Wavefront 通常包含 64 个线程（RDNA 架构之前），而 Warp 固定为 32 个线程。

2.3 流处理器（SP）与流多处理器（SM）：硬件执行单元

流处理器（Streaming Processor, SP）是 GPU 中最基本的计算单元，本质上是一个标量 ALU（算术逻辑单元），能够执行整数、浮点运算（FP32/FP64）以及逻辑操作。在 NVIDIA 的架构中，SP 也被称为 CUDA Core；在 AMD 的架构中，则称为 Stream Processor 或 Shader Processor。单个 SP 的性能并不突出，但一个 GPU 中集成了数千个 SP，通过大规模并行实现高吞吐。

流多处理器（Streaming Multiprocessor, SM）是 GPU 硬件架构中的“计算集群”层级。一个 SM 包含多个 SP（例如，NVIDIA Ampere 架构的 GA102 芯片中，每个 SM 包含 128 个 FP32 CUDA Core）、共享内存（Shared Memory）、寄存器堆（Register File）、Warp 调度器（Warp Scheduler）以及指令分发单元（Dispatch Unit）。SM 是 GPU 调度的基本单位——一个线程块（Thread Block）被调度到某个 SM 上执行，SM 负责管理该线程块内所有 Warp 的生命周期。

以 NVIDIA A100（Ampere 架构）为例，其硬件规格如下：

• 108 个 SM，每个 SM 包含 64 个 FP32 CUDA Core + 64 个 INT32 Core（可混合执行）
• 每个 SM 有 4 个 Warp 调度器，每个调度器每周期可发射 1 条指令
• 每个 SM 拥有 256KB 寄存器堆（64K 个 32 位寄存器）
• 每个 SM 配置 164KB 可配置共享内存/L1 缓存

这种“SM + SP”的两级结构，使得 GPU 能够灵活地分配计算资源：一个线程块可以包含数百个线程，这些线程被划分为多个 Warp，由 SM 内的调度器分时发射。SM 的寄存器堆是线程上下文切换的关键——每个线程的寄存器状态保存在 SM 的寄存器堆中，切换 Warp 时只需更换寄存器索引，无需保存/恢复现场，因此上下文切换开销几乎为零。

2.4 Warp 与 Wavefront：线程调度的核心粒度

Warp（NVIDIA）和 Wavefront（AMD）是 GPU 线程调度的最小硬件单位。一个 Warp 包含 32 个线程（Wavefront 为 64 个线程），这些线程在硬件上以锁步方式执行——即同一时刻执行同一条指令。Warp 的引入解决了两个关键问题：

1. 指令发射效率：如果每个线程单独调度，指令发射带宽将成为瓶颈。将 32 个线程打包为一个 Warp，调度器只需每周期发射 1 条指令，即可驱动 32 个 ALU 并行工作。
2. 延迟隐藏：当某个 Warp 因访存或同步而阻塞时，调度器可以立即切换到另一个就绪的 Warp，无需保存/恢复上下文（因为寄存器已预先分配）。这种“零开销切换”使得 GPU 能够用“线程并行”掩盖“流水线延迟”。

Warp 的执行过程遵循“单指令多线程”原则：所有线程共享同一个程序计数器（PC），但每个线程有独立的寄存器状态和内存地址。当遇到分支指令（如 if (tid % 2 == 0)）时，Warp 会执行“分支发散”处理：

• 硬件计算分支掩码（Active Mask），标记哪些线程走 true 路径，哪些走 false 路径
• 先执行 true 路径的指令，此时 false 路径的线程被“禁用”（不写回结果）
• 再执行 false 路径的指令，true 路径的线程被禁用
• 最后合并两条路径，继续执行后续指令

这种串行执行分支路径的方式，会导致性能惩罚——如果分支两边的指令数量相同，则 Warp 的执行时间翻倍。因此，在 GPU 编程中，应尽量避免 Warp 内的分支发散，或者将分支条件设计为“Warp 级别一致”（即整个 Warp 的线程都走同一路径）。

2.5 案例拆解：矩阵乘法的 Warp 级优化

矩阵乘法（C = A × B）是 GPU 性能测试的经典基准，也是理解 Warp 调度和内存合并访问的最佳案例。假设我们有一个 1024×1024 的单精度矩阵乘法，在 NVIDIA A100 上运行，初始实现（朴素版本）的性能可能只有理论峰值的 10% 左右，而经过 Warp 级优化后可以达到 60% 以上。

初始实现（Naive Kernel）：

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

该内核的问题在于：每个线程直接访问全局内存中的 A 和 B，且访问模式不满足“合并访问”（Coalesced Access）条件——相邻线程访问的 B 矩阵地址不连续（B[k*N + col] 中 col 是线程索引，但 B 的行跨度是 N，导致同一 Warp 的 32 个线程访问的地址间隔为 N×4 字节）。这会导致全局内存带宽利用率极低。

优化步骤：

1. 使用共享内存（Shared Memory）分块：将矩阵划分为 32×32 的块（Block），每个线程块加载一个子矩阵到共享内存，然后从共享内存中读取数据。共享内存的访问延迟远低于全局内存（约 20-30 周期 vs 400-600 周期）。
2. 确保合并访问：在加载 A 和 B 到共享内存时，让相邻线程访问相邻地址。例如，A_tile[threadIdx.y][threadIdx.x] = A[row * N + col]，此时同一 Warp 的线程（threadIdx.x 连续）访问的地址是连续的。
3. 避免 Bank Conflict：共享内存被划分为 32 个 Bank（每 Bank 4 字节），如果同一 Warp 的多个线程访问同一 Bank 的不同地址，会产生冲突，导致串行化。通过填充（Padding）或调整索引方式可以避免。
4. 循环展开（Loop Unrolling）：每个线程计算多个输出元素（如 4×4 子块），减少循环开销，提高指令级并行。

优化后的关键参数对比：

复盘经验：矩阵乘法的优化过程充分体现了 GPU 架构的三个核心约束：全局内存带宽（通过共享内存分块缓解）、共享内存 Bank 冲突（通过 Padding 解决）、以及 Warp 调度效率（通过循环展开提高指令级并行）。最终性能从 1.2 TFLOPS 提升到 9.8 TFLOPS（A100 理论峰值 19.5 TFLOPS），证明了“理解硬件才能写出高性能代码”这一原则。

2.6 常见误区与总结

在学习和使用 GPU 基础架构时，以下误区值得特别注意：

• 误区一：将 GPU 线程等同于 CPU 线程。GPU 线程是“轻量级”的，创建和切换开销极低，但单个线程的性能远低于 CPU 线程。GPU 的优势在于“数量”而非“质量”。
• 误区二：认为 Warp 内线程可以独立分支。虽然 SIMT 模型允许线程有独立 PC，但硬件执行时仍以 Warp 为单位锁步。分支发散会导致串行执行，应尽量避免。
• 误区三：忽略占用率（Occupancy）。占用率 = 活跃 Warp 数 / 最大 Warp 数。如果寄存器或共享内存使用过多，占用率下降，延迟隐藏能力减弱，即使 ALU 数量充足也无法达到峰值性能。
• 误区四：混淆 SIMD 和 SIMT。SIMD 是数据级并行（DLP），SIMT 是线程级并行（TLP）+ 数据级并行的混合。理解这一区别有助于选择合适的并行化策略。

总结而言，GPU 的基础架构可以概括为“一个模型、两级硬件、一种调度”：SIMT 执行模型提供了编程抽象，SM/SP 两级硬件提供了物理执行能力，Warp/Wavefront 调度机制则连接了软件与硬件。掌握这三者的关系，是深入理解 GPU 性能优化、内存层次、以及未来架构演进（如 Tensor Core、Ray Tracing Core）的必备基础。在实际工作中，无论是使用 CUDA、OpenCL 还是 ROCm，开发者都应时刻关注“Warp 如何映射到硬件”、“内存访问是否合并”、“分支是否发散”这三个核心问题。