计算机组成原理：GPU架构、并行计算、内存层次结构等

发布于：2025-09-06 ⋅ 阅读:(13) ⋅ 点赞:(0)

好的，这是一个非常精彩且深入的话题。我们将从计算机组成原理的视角，深入剖析CUDA的底层机制，涵盖GPU架构、并行计算模型和内存层次结构。

一、为什么需要GPU？—— 设计与目标的根本差异

要理解CUDA和GPU，首先要明白它与CPU在设计哲学上的根本区别：

  • CPU (Central Processing Unit)：** latency-oriented** (延迟导向)。旨在以尽可能快的速度完成单个复杂任务。它拥有强大的ALU（算术逻辑单元）、复杂的控制逻辑和大型缓存（Cache）来减少指令和数据的等待时间，擅长处理分支预测、乱序执行等复杂逻辑。CPU是“全能冠军”，但核数较少。
  • GPU (Graphics Processing Unit)：** throughput-oriented** (吞吐量导向)。旨在同时处理海量的、相互独立的简单任务（如处理屏幕上数百万个像素）。它将更多的晶体管用于计算单元而非缓存和控制逻辑。GPU是“人海战术”，拥有成千上万个计算核心。

这种差异导致了截然不同的架构：

特性 CPU GPU
设计目标 低延迟，强通用性 高吞吐，数据并行
核心数 几个到几十个（强大多核） 成千上万个（精简小核）
缓存 巨大，用于减少延迟 较小，主要用于服务数据流
控制逻辑 非常复杂（分支预测、乱序执行） 相对简单
适用场景 操作系统、应用程序逻辑、数据库 图形渲染、科学计算、深度学习

二、GPU硬件架构（以NVIDIA现代GPU为例）

GPU是一个层次化的并行处理器。我们以NVIDIA的Fermi/Ampere架构为例来解析其核心组成部分。

1. 宏观架构：GPU -> Streaming Multiprocessors (SMs)

一张GPU芯片由多个流多处理器（SM） 组成。SM是GPU的核心计算单元，相当于CPU的一个“大核”，但数量要多得多（例如，一个GA102芯片有84个SM）。你的CUDA程序最终就是在这些SM上被执行的。

2. 微观架构：深入一个SM内部

每个SM内部又包含以下关键组件：

  • CUDA Cores： 最基本的计算单元。每个CUDA Core包含一个FP32单元（单精度浮点）和一个INT32单元（整数运算）。注意：CUDA Core是NVIDIA的营销术语，本质上就是一个能够执行浮点和整数运算的标量处理器（Scalar Processor）。
  • Tensor Cores (现代GPU)： 专为执行矩阵乘加操作（D = A * B + C）设计的特殊计算单元，速度极快，是深度学习训练和推理的核心。
  • Warp Schedulers ( warp调度器 )： 这是SM的大脑，也是GPU高效并发的关键。它负责从指令流中取出指令，并分发给下面的执行单元。一个SM有多个Warp Scheduler。
  • Dispatch Units ( 分发单元 )： 每个Warp Scheduler连接多个Dispatch Units，可以在一个时钟周期内分发多条指令。
  • Register File ( 寄存器文件 )： SM上速度最快的存储器。每个线程都有自己的私有寄存器。寄存器资源是有限的，线程越多，每个线程能分到的寄存器就越少。
  • Shared Memory / L1 Cache： SM内部的一块可由程序员控制的高速、片上缓存。它由SM上的所有线程块共享，是实现线程间通信和减少全局内存访问的关键。它与L1 Cache共享物理硬件，可以通过配置来划分大小。
  • Load/Store Units： 负责处理从内存到寄存器的数据加载和存储操作。
  • 特殊功能单元（SFU）： 执行一些特殊的数学运算，如正弦、余弦、平方根等。

三、CUDA并行计算模型：软件视角

CUDA提供了一种层次化的并行编程模型，它与GPU的硬件层次紧密对应。

1. 核心概念：Kernel（内核函数）

Kernel是在GPU上执行的并行函数。在主机（CPU）代码中调用，在设备（GPU）上执行。

2. 线程层次结构：Grid -> Block -> Thread

这是CUDA编程模型的核心抽象，它定义了如何组织成千上万个线程来解决问题。

  • Thread（线程）： 最基本的执行单元。每个线程都独立地执行kernel函数，拥有自己的程序计数器和寄存器。
  • Block（线程块）： 一组线程的集合。这些线程会被分配到一个SM上执行。块内的线程可以：
    • 通过 __syncthreads() 函数进行同步。
    • 通过 Shared Memory 进行高效通信。
  • Grid（网格）： 所有线程块的集合，共同完成一个Kernel函数的执行。

当你启动一个Kernel时，需要指定Grid和Block的维度（Dimensions）：

// 定义一个Kernel
__global__ void myKernel(float *data) {
    // 每个线程的工作
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    data[idx] *= 2;
}

int main() {
    ...
    // 启动Kernel！
    // <<<网格维度，线程块维度>>>
    // 这里定义了一个包含N个线程块的Grid，每个Block有256个Threads。
    myKernel<<<N, 256>>>(data);
    ...
}

3. 硬件映射：软件模型如何映射到硬件上？

  • 一个Grid -> 在一个GPU设备上执行。
  • 一个Block -> 被分配到一个SM上执行。一个SM可以同时处理多个Block，具体数量取决于SM的资源（寄存器、Shared Memory）。
  • Warp（线程束）： 这是GPU执行的基本单位，也是最关键的概念。
    • 一个Block在SM上执行时，会被进一步划分为多个Warp。
    • 每个Warp包含32个连续的线程（这是NVIDIA GPU的硬件规定）。
    • Warp是SM调度和执行的基本单元。Warp Scheduler每次调度一个Warp（32个线程）去执行同一条指令。

4. SIMT（单指令多线程）模型

CPU是SISD（单指令单数据）或SIMD（单指令多数据）。 GPU采用的是SIMT（Single-Instruction, Multiple-Thread） 模型。

  • SIMT vs SIMD：
    • SIMD： 一条指令操作一个向量（例如4个float），所有操作是锁步的，作为一个整体。
    • SIMT： 一条指令被一个Warp（32个线程）同时执行。每个线程有自己的寄存器状态和自己的数据地址。
  • SIMT的优势与挑战：
    • 优势： 编程模型更灵活，你像写标量程序一样编写每个线程的行为，硬件负责将它们组织成并行执行。
    • 挑战：分支分化（Branch Divergence）。
      • 如果Warp内的线程在执行if-else语句时走了不同的分支，那么SM必须串行化地执行所有分支路径，禁用那些不满足条件的线程。
      • 例如： if (threadIdx.x % 2 == 0) { A } else { B }。这个Warp会先执行所有偶数线程的A代码（同时奇数线程被禁用），然后再执行所有奇数线程的B代码（偶数线程被禁用）。这会导致性能急剧下降。

四、GPU内存层次结构：性能的关键

GPU拥有复杂的内存层次，理解它是优化CUDA程序性能的重中之重。其设计原则是：让最频繁的访问发生在最快的内存上。

从上到下，速度变慢，容量变大：

  1. Register（寄存器）：
    • 位置： 在SM片上。
    • 速度： 最快。
    • 作用： 每个线程私有。用于存储局部变量和中间结果。
    • 注意： 资源有限。使用过多寄存器会限制SM上可同时活跃的线程数（Warp数）。
  2. Shared Memory（共享内存）：
    • 位置： 在SM片上。
    • 速度： 非常快，接近寄存器。
    • 作用： 一个Block内的所有线程共享。用于实现块内线程协作、数据复用，是减少全局内存访问的最关键武器。
    • 注意： 需要程序员显式管理（__shared__变量）。通常用作可编程的缓存（Software-Managed Cache）。
  3. L1 Cache / Read-Only Cache：
    • 位置： SM片上。
    • 速度： 很快。
    • 作用： 自动缓存局部数据，减少访问全局内存的延迟。
  4. L2 Cache（二级缓存）：
    • 位置： 所有SM共享，在芯片上。
    • 速度： 比L1慢。
    • 作用： 缓存所有SM访问全局内存的数据，是全局内存的缓冲区。
  5. Global Memory（全局内存）：
    • 位置： GPU板载的DRAM（如GDDR6, HBM2e）。
    • 速度： 很慢（相对于片上内存），延迟高。
    • 作用： GPU的主内存，CPU和GPU都可以访问（通过PCIe）。所有数据最初都在这里。
    • 关键概念：合并访问（Coalesced Access）
      • GPU为了高效访问全局内存，希望一个Warp（32个线程）的内存访问能合并成尽可能少的（理想情况下是1次）事务。
      • 合并访问： 如果Warp中的线程访问连续的、对齐的内存块（例如，thread0访问addr0, thread1访问addr1, ...），则访问会被合并，效率最高。
      • 非合并访问： 如果线程随机地、分散地访问内存，会导致多次内存事务，性能极差。
  6. Constant Memory & Texture Memory：
    • 它们是Global Memory的一部分，但有特殊的缓存机制。
    • Constant Memory： 用于存储只读数据。有专用的Constant Cache，适合所有线程读取相同常数值的场景。
    • Texture Memory： 专为图形纹理访问模式（具有空间局部性的访存）优化，也适用于一些科学计算。

总结：编写高效CUDA程序的核心要点

  1. 最大化并行度： 启动足够多的Thread和Block来隐藏内存延迟，让SM始终保持忙碌。
  2. 优化内存访问：
    • 尽力实现全局内存的合并访问。
    • 积极使用Shared Memory来缓存数据，减少对全局内存的访问。
  3. 避免分支分化： 确保同一个Warp内的线程尽可能走相同的执行路径。
    • Occupancy（占用率）**： 衡量SM上活跃的Warp数与最大可能支持的Warp数之比。较高的占用率有助于隐藏延迟，但并非越高越好，需要与寄存器/SMEM的使用做权衡。

理解这些底层机制，你就能从“程序员”升级为“优化专家”，真正释放GPU的强大计算潜力。

网站公告

今日签到

点亮在社区的每一天
去签到

热门文章

最新发布