英伟达的GPU和CUDA

一. 前言

早期的 GPU 具有固定功能管线，只能执行特定的图形渲染任务，如几何处理和像素填充。这些操作都是预定义的，不能编程。显存（VRAM）主要存储图形数据，如纹理、顶点数据、帧缓冲区、深度缓冲区等。在这一阶段，GPU 并没有执行可编程代码的能力，也没有类似 CPU 的线程模型。显存仅用于存储图形数据，没有执行通用计算任务的线程空间概念。

后来GPU 进入通用计算（GPGPU）时代

1. 可编程着色器的引入：

   • 随着可编程着色器的引入（例如，顶点着色器和像素着色器），GPU 开始能够运行自定义代码。这些着色器程序允许开发者编写小型程序（着色器代码），并由 GPU 并行执行。
   • 虽然这些程序具有一定的通用性，但它们的设计仍然高度依赖图形渲染任务。

2. 通用计算的出现（GPGPU）：

   • GPGPU（General-Purpose computing on GPUs）技术的出现使得 GPU 能够被用于非图形计算任务。CUDA、OpenCL 等框架的引入为 GPU 提供了执行通用计算的能力。
   • 在此之后，显存不仅仅存储图形数据，还用于存储通用计算的输入数据、输出数据以及程序代码（如 CUDA 核函数）。

3. 线程模型的引入：

   • 随着 CUDA 和 OpenCL 等编程模型的引入，GPU 的线程模型变得可编程。每个 CUDA 核函数的调用都会在 GPU 上创建大量的轻量级线程，这些线程在显存中分配空间来存储中间数据、结果和其他必要信息。
   • 这些线程在 GPU 的执行单元上并行执行。虽然这些线程没有 CPU 线程那样的独立性和灵活性，但它们能够高效地处理大规模并行任务。

二. GPU和8051以及其他CPU的对比

以前比较简单的MCU比如8051，是哈佛结构，程序和临时中间数据分别存在ROM和RAM里，因为它的ROM具有直接寻址的能力，所以可以直接运行程序，而ROM中的数据不能改变，因此不能存放临时数据，就得存到RAM里。

对于冯诺伊曼结构的CPU来说，内存中通常创建一个进程，进程里有一部分是放置不可变的程序代码、常量数据和初始化数据，这段空间相当于ROM，而其他的一些临时数据，比如动态数据、临时数据、堆栈放在线程里，这部分空间就相当于RAM。

对GPU来说，也有点类似，以前整个显存都存放的临时数据，比如需要处理的图像数据之类的。但GPU具备通用计算的能力之后，显存有一部分用来存放CUDA的核函数或者OpenCL 内核，这地方就类似于ROM，CUDA核函数还会创建很多空间来存放临时数据，比如很多网格（grid），每个网格又包含很多线程块（block），每个线程块可以被多个线程共用。

三. GPU中的线程和CPU中的线程的对比

说到这里就奇怪了，明明就是一个核函数创建的grid和block，为什么会有多线程。CPU中的线程和GPU中的线程不一样，为了发挥GPU的并行计算能力，一个核函数被拆成了好多线程。而CPU中一个函数还是一个线程。但是GPU不能像CPU很全能，有各种计算单元，可以同时创建不同进程，每个进程管理不同线程。GPU 的执行模型在某种程度上可以看作是一个进程管理所有线程。

• CPU 线程和 GPU 线程 是不同的概念。CPU 线程是重量级的，独立性强，而 GPU 线程是轻量级的，设计用于大规模并行执行。
• GPU 可以将一个函数分成多个线程执行，这是为了并行处理大量数据。每个线程执行相同的代码，但处理不同的数据片段。
• SM（流式多处理器） 是 GPU 内的固定计算单元组合。每个 SM 包含多个 CUDA 核心，可以并行执行多个线程块，利用硬件资源进行高效计算。

四. CPU在管理GPU计算中的作用

GPU虽然具有了一定执行程序的能力，但是仍然需要CPU来主导。CPU的角色是负责启动和管理 GPU 的计算任务。它调用 CUDA API，将核函数（代码）发送到 GPU，将相应的配置（如 Grid 和 Block 的大小、数据指针等）传递给 GPU，并启动核函数的执行。一旦 GPU 接收到核函数的执行命令，CUDA 调度器会根据 GPU 的架构将核函数分成多个线程块（block），并在 GPU 的多个流处理器（Streaming Multiprocessors, SMs）上并行执行。每次 CUDA 调用只会启动一个核函数，所有线程都会执行同一个核函数的代码。核函数的执行是并行的，但它们执行相同的代码，只是处理不同的数据。

在 CUDA 编程中，Grid 和 Block 的大小并不是固定的，而是在运行时由你定义的，在核函数执行时动态创建的。它们的配置取决于任务的并行性需求、数据大小以及 GPU 硬件资源

例如这段代码：kernel<<<numBlocks, numThreadsPerBlock>>>(...);

这里的 numBlocks 是 Grid 中的 Block 数量，而 numThreadsPerBlock 是每个 Block 中的线程数。

尽管 CUDA 的执行模型通常聚焦于单个核函数的大量并行线程，但在某些情况下，你可以通过不同的 CUDA 流（Streams）启动多个核函数，它们可能在同一时间并行执行。例如，你可以在一个流中启动一个核函数，在另一个流中启动另一个核函数。如果 GPU 有足够的资源（如足够的 SM），这些核函数可以并行运行。

五. GPU中的SM（Streaming Multiprocessor, 流式多处理器）是什么？

• SM 的定义：

  • 流式多处理器（SM） 是 GPU 内部的一个关键组件。一个 GPU 通常包含多个 SM，每个 SM 内部包含多个 CUDA 核心、寄存器、共享内存、调度器等。
  • 每个 SM 可以并行执行多个线程块（Block），这些线程块内的线程共享 SM 内的资源（如共享内存、寄存器等）。
  • SM 内的这些单元不随着任务的变化而变化，它们是 GPU 架构的基础，负责执行核函数中的并行计算任务。

• SM 的结构和工作方式：每个 SM 都有一组固定的计算单元（如 CUDA 核心、特殊功能单元等），这些计算单元共同执行分配给 SM 的线程块。在每个 SM 内，线程按照 SIMD（单指令多数据）或 SIMT（单指令多线程）模型执行，这意味着多个线程同时执行相同的指令，但处理不同的数据。

  • CUDA 核心（CUDA Cores）：
    • 这些是执行整数和浮点运算的基本计算单元。每个 SM 包含多个 CUDA 核心，能够并行处理多个线程的计算任务。
    • CUDA 核心执行基本的算术运算（如加法、乘法、逻辑运算等）。在一次时钟周期内，多个 CUDA 核心可以并行执行来自不同线程的相同指令。

  • 特殊功能单元（SFUs, Special Function Units）：

    • SFUs 处理一些特殊的数学运算，如三角函数、平方根、指数等。与 CUDA 核心相比，SFUs 更专注于复杂的数学操作。
    • 每个 SM 通常有几个 SFU，用于加速特殊数学函数的计算。

  • 载入/存储单元（LD/ST Units）：

    • 这些单元负责处理内存的加载和存储操作。它们从显存中读取数据或将计算结果写回显存。
    • 载入/存储单元使得线程可以高效地访问全局内存、共享内存和常量内存等不同的存储空间。

  • 寄存器文件：

    • 每个 SM 包含大量的寄存器，供其管理的线程使用。每个线程都有自己的寄存器，用于存储局部变量和中间计算结果。
    • 这些寄存器非常高速，线程在执行过程中主要依赖寄存器进行计算。

  • 共享内存（Shared Memory）：

    • 共享内存是一个高速缓存，供同一线程块（Block）中的线程共享。共享内存允许线程之间快速交换数据，减少对全局内存的访问需求。
    • 共享内存是可编程的，并且可以配置为不同大小，以适应不同的计算任务需求。

  • 硬件调度器：

    • SM 内的硬件调度器负责管理和调度线程块的执行。它决定哪些线程在当前时钟周期内执行，并分配 CUDA 核心和其他资源。
    • 调度器使用轮询或其他策略来最大化硬件利用率，避免执行单元的闲置。

  • 纹理和常量缓存（Texture and Constant Cache）：

    • SM 还包含一些特定用途的缓存，如纹理缓存和常量缓存。纹理缓存用于加速图形处理中的纹理采样，而常量缓存用于快速访问常量内存中的数据。

六. GPU中的存储管理

前面说到CUDA核函数还会创建很多空间来存放临时数据，比如很多网格（grid），每个网格又包含很多线程块（block），每个线程块可以被多个线程共用。线程除了共用线程块之外，还需要共享内存，但是这个共享内存在流处理器（SM）中，可以理解为共享内存是在 Block 内，但它是由 SM 提供的硬件资源，物理上位于 GPU 的流处理器（SM）中的。而且不同 Block 之间的共享内存是隔离的，无法相互访问。除了共享内存（Shared Memory）外，GPU 的每个线程块（Block）和线程还可以访问其他几种类型的内存。以下是 CUDA 内存层次结构的详细信息：

1. 寄存器（Registers）

• 寄存器 是最接近 CUDA 核心的存储单元，每个线程都有一组专用的寄存器，用于存储局部变量和中间计算结果。也是每个SM都有的，不能共用。
• 寄存器的访问速度最快，但数量有限。寄存器空间的使用直接影响线程的并行度（即每个 SM 上能够同时调度的线程数）。

2. 本地内存（Local Memory）

• 本地内存 是针对每个线程的私有内存空间，用于存储超过寄存器容量的局部变量或数组。
• 尽管名为“本地内存”，它实际上位于全局内存（显存）中，因此访问速度较慢。本地内存主要用于处理寄存器溢出的数据。

3. 共享内存（Shared Memory）

• 共享内存的大小是有限的，通常可以在核函数启动时配置。
• 每个 Block 内的线程可以共享一个专用的共享内存空间（Shared Memory）。共享内存是一个高效的、低延迟的存储区域，位于每个 SM 内，供同一 Block 的所有线程访问。
• 共享内存的设计初衷是为了让 Block 内的线程可以高效地合作。例如，多个线程可以共享中间计算结果、缓存常用数据，或者进行同步操作。
• Block 内的多个线程需要协同工作。通过共享内存，线程可以快速地交换数据，而不必通过全局内存（显存），从而减少数据传输的开销，提高计算性能。例如，在矩阵乘法中，多个线程可以通过共享内存协同处理同一块数据，以减少重复的全局内存访问。

4. 全局内存（Global Memory）

• 全局内存 是所有线程（无论是同一 Block 内还是不同 Block 内）都可以访问的内存空间，通常是 GPU 的显存（VRAM）。
• 全局内存的容量最大，但访问速度较慢，具有高延迟。全局内存适用于存储大型数据集，如输入数据、输出数据和模型参数。
• 因为访问速度慢，通常使用共享内存或寄存器来缓存从全局内存读取的数据。
• 核函数，以及那种需要处理的图形数据都放在全局内存里

5. 常量内存（Constant Memory）

• 常量内存 是一个只读内存空间，适用于存储在整个核函数执行过程中不变的数据。
• 常量内存是共享的，并且所有线程都可以读取，但只能由 CPU 写入。常量内存的读取速度较快，适合用于存储小型、常用的数据（如模型参数、配置信息等）。

6. 纹理内存（Texture Memory）

• 纹理内存 是一种专门用于优化图像处理的只读内存空间。
• 纹理内存包括特殊的缓存机制，用于加速非对齐或随机访问的读取操作，特别适合图像处理和采样操作。

这些内存类型构成了 CUDA 的内存层次结构，各有不同的访问速度和使用场景。合理利用这些内存可以显著优化 CUDA 程序的性能。