CUDA内存模型详解：锁页内存、统一寻址、CPU/GPU交互

> 最近开始研究异构架构下的通信实现，主要是`Nvidia GPU`环境下。因此涉及到CUDA内存的读写等。故写本文总结相关内容。 [TOC] ## Nvidia GPU物理结构 > 本章以最新的 `Ampere`（安培）架构为例，目前安培架构在`Tesla`系列只有`A100`一个产品（2020.12），我们以它为例讲解GPU物理架构。 ### GPU 整体结构  上图显示了一个完整的 GA100 GPU结构，共有128个SMs。 `NVIDIA GA100 GPU` 由多个 [G PC ]^(GPU 处理簇)、[TPC]^(纹理处理簇)、[SMs]^(流式多处理器) 、显存`HBM2` 、L2 Cache 和 `HBM2 内存控制器`组成。 GA100 全功能 GPU 包括以下单元： - 8 个`GPC` 。其中每个GPC含 8 个 TPC ；每个TPC含 2 个 SMs 。因此每个GPU含有`8 * 8 * 2 = 128`个SMs。 - 每个SM含有64 个 FP32 CUDA 核 。故GPU共8192 个 FP32 CUDA 个核。 - 每个SM含有4 个第三代张量核心 。故GPU含 512 个第三代张量核心。 - 6 个 HBM2 内存，12个 512 位内存控制器。 用于张量计算的 __A100__ 包括以下单元： - 7 个 `GPC` 。每个GPC含 7 或 8 个 TPC ；每个 TPC含有2个SMs 。因此最多 16 个 SMs / GPC ， GPU最多 108 个 SMs。 - 每个SM含有64 个 FP32 CUDA 核 。故GPU共6912 个 FP32 CUDA 个核。 - 每个SM含有4 个第三代张量核心 。故GPU含 432 个第三代张量核心。 - 5 个 HBM2 内存， 10 个 512 位内存控制器 ### GPU SM结构 {:max-height: 800px;} A100中的SM包含4个`Tensor Core`，此外还有192KB的L1 Cache/Shared Memory。 每个`Tensor Core`中含有16K个32位寄存器。 新的异步复制指令将数据直接从全局内存加载到共享内存中，可以选择绕过一级缓存，并且不需要使用中间寄存器文件（ RF ）。 ### GPU runtime 资源分配 GPU `kernel<<>>` 运行时指定Block数、Thread数。 Block对应SM单元，每个SM包含4个Tensor Core，Tensor Core包含一个warp调度器，用来每次调度一组thread。 A100下，每个SM最多可以有64个warp/32个Thread Blocks，各warp任务不断切换。每个warp有32个thread，故每个SM最多2048个thread。 ## CUDA内存模型 对于程序员来说，一般有两种类型的存储器： - 可编程的：你需要显式地控制哪些数据存放在可编程内存中 - 不可编程的：你不能决定数据的存放位置，程序将自动生成存放位置以获得良好的性能 在CPU内存层次结构中，一级缓存和二级缓存都是不可编程的存储器。另一方面，CUDA内存模型提出了多种可编程内存的类型： - 寄存器 - 共享内存 - 本地内存 - 常量内存 - 纹理内存 - 全局内存 下图为这些内存空间的层次结构，每种都有不同的作用域、生命周期和缓存行为。一个[Kernel]^(核函数)中的[Thread]^(线程)都有自己私有的`本地内存`。一个[Block]^(线程块)有自己的`共享内存`，对同一线程块中所有[Thread]^(线程)都可见，其内容持续Block的整个生命周期。所有Thread都可以访问`全局内存`。 所有Thread都能访问的只读内存空间有：`常量内存空间`和`纹理内存空间`。`全局内存`、`常量内存`和`纹理内存空间`有不同的用途。`纹理内存`为各种数据布局提供了不同的寻址模式和滤波模式。对于一个应用程序来说，`全局内存`、`常量内存`和`纹理内存`中的内容具有相同的生命周期。 {:max-height:400px} ### 寄存器(Register) 每个SM包含成千上万个32位寄存器，当[kernel]^(核函数)启动时，这些寄存器会被分配到指定的线程中。在SM中，寄存器是速度最快，也是数量最多的存储资源。 例如：Kepler架构中（SM 3.0）的SMX包含65536个寄存器，容量总共256KB，而纹理寄存器只有48KB。 寄存器变量对于每个`Thread`来说都是私有的，一个`Kernel`通常使用寄存器来保存需要频繁访问的`Thread`私有变量。寄存器变量与`Kernel`的生命周期相同。一旦`Kernel`执行完毕，就不能对寄存器变量进行访问了。 Kepler GPU将该限制扩展至每个线程可拥有255个寄存器。在核函数中使用较少的寄存器将使在SM上有更多的常驻线程块。每个SM上并发线程块越多，使用率和性能就越高。 ### 本地内存(Local Memory) [Kernel]^(核函数)中符合存储在寄存器的变量不能被分配寄存器空间时，将溢出到`本地内存`中。编译器可能存放到本地内存中的变量有： - 在编译时使用未知索引引用的本地数组 - 可能会占用大量寄存器空间的较大本地结构体或数组 - 任何不满足核函数寄存器限定条件的变量 `本地内存`这一名词是有歧义的：溢出到本地内存中的变量本质上与`全局内存`在同一块存储区域，因此本地内存访问的特点是高延迟和低带宽。本地内存数据也是存储在每个SM的一级缓存和每个设备的二级缓存中，因此也可以满足高效内存访问要求。 ### 全局内存(Global Memory) 全局内存是GPU中最大、延迟最高并且最常使用的内存。global指的是其作用域和生命周期。它的声明可以在任何SM设备上被访问到，并且贯穿应用程序的整个生命周期。 一个全局内存变量可以被静态声明或动态声明。你可以使用如下修饰符在设备代码中静态地声明一个变量： ```cpp __device__ ``` ### 共享内存(shared memory) 在[Kernel]^(核函数)中使用如下修饰符修饰的变量存放在共享内存中： ```cpp __shared__ ``` 因为共享内存是片上内存，所以与本地内存或全局内存相比，它具有更高的带宽和更低的延迟。它的使用类似于CPU一级缓存，但它是可编程的。 每一个SM都有一定数量的由线程块分配的共享内存。因此，必须非常小心不要过度使用共享内存，否则将在不经意间限制活跃线程束的数量。 共享内存在核函数的范围内声明，其生命周期伴随着整个线程块。当一个线程块执行结束后，其分配的共享内存将被释放并重新分配给其他线程块。 共享内存是线程之间相互通信的基本方式。一个块内的线程通过使用共享内存中的数据可以相互合作。访问共享内存必须同步使用如下调用，该命令是在之前章节中介绍过的CUDA运行时调用： ```cpp void syncthreads(); ``` SM中的一级缓存和共享内存都使用64KB的片上内存（Fermi/Kepler），它通过静态划分，但在运行时可以通过如下指令进行动态配置： ```cpp cudaError_t cudaFunSetCacheConfig(const void* func, enum cudaFuncCache cacheConfig); ``` 这个函数在每个核函数的基础上配置了片上内存划分，为func指定的核函数设置了配置。支持的缓存配置如下： ```cpp cudaFuncCachePreferNone; // 没有参考值（默认） cudaFuncCachePreferShared; // 48KB 共享内存；16KB L1 Cache cudaFuncCachePreferL1; // 48KB L1 cache；16KB 共享内存 cudaFuncCachePreferEqual1; // 相同尺寸，都是32KB ``` Fermi设备支持前三种配置，Kepler设备支持以上所有配置。 ### 常量内存(constant memory) 常量内存驻留在设备内存中，并在每个SM专用的常量缓存中缓存。常量变量用如下修饰符来修饰： ```cpp __constant__ ``` 常量变量必须在全局空间内和所有核函数之外进行声明。对于所有计算能力的设备，都只可以声明64KB的常量内存。常量内存是静态声明的，并对同一编译单元中的所有核函数可见。 核函数只能从常量内存中读取数据。因此，常量内存必须在主机端使用下面的函数来初始化： ```cpp cudaError_t cudaMemcpyToSymbol(const void* symbol, const void* src, size_t count); ``` 这个函数将count个字节从src指向的内存复制到symbol指向的内存中，这个变量存放在设备的`全局内存`或`常量内存`中。在大多数情况下这个函数是同步的。 [warp]^(线程束)中的所有线程从相同的内存地址中读取数据时，常量内存表现最好。举个例子，数学公式中的系数就是一个很好的使用常量内存的例子，因为一个线程束中所有的线程使用相同的系数来对不同数据进行相同的计算。如果线程束里每个线程都从不同的地址空间读取数据，并且只读一次，那么常量内存中就不是最佳选择，因为每从一个常量内存中读取一次数据，都会广播给线程束里的所有线程。 ### 纹理内存(texture memory) 纹理内存驻留在设备内存中，并在每个SM的只读缓存中缓存。纹理内存是一种通过指定的只读缓存访问的全局内存。只读缓存包括硬件滤波的支持，它可以将浮点插入作为读过程的一部分来执行。纹理内存是对二维空间局部性的优化，所以线程束里使用纹理内存访问二维数据的线程可以达到最优性能。对于一些应用程序来说，这是理想的内存，并由于缓存和滤波硬件的支持所以有较好的性能优势。然而对于另一些应用程序来说，与全局内存相比，使用纹理内存更慢。 ### GPU 缓存 跟CPU缓存一样，GPU缓存是不可编程的内存。在GPU上有4种缓存： - 一级缓存 - 二级缓存 - 只读常量缓存 - 只读纹理缓存 每个SM都有一个一级缓存，所有的SM共享一个二级缓存。一级和二级缓存都被用来在存储本地内存和全局内存中的数据，也包括寄存器溢出的部分。对Fermi GPU和Kepler K40或其后发布的GPU来说，CUDA允许我们配置读操作的数据是使用一级和二级缓存，还是只使用二级缓存。 在CPU上，内存的加载和存储都可以被缓存。但是，在GPU上只有内存加载操作可以被缓存，内存存储操作不能被缓存。 每个SM也有一个只读常量缓存和只读纹理缓存，它们用于在设备内存中提高来自于各自内存空间内的读取性能 ### CUDA变量声明总结 表：CUDA变量和类型修饰符 | 修饰符 | 变量名称 | 存储器 | 作用域 | 生命周期 | | --- | | | float var | 寄存器 | 线程 | 线程 | | | float var[100] | 本地 | 线程 | 线程 | | \_\_shared\_\_ | float var(*) | 共享 | 块 | 块 | | \_\_device\_\_ | float var(*) | 全局 | 全局 | 应用程序 | | \_\_constant\_\_ | float var(*) | 常量 | 全局 | 应用程序 | 表：设备存储器的重要特征 | 存储器 | 片上/片外 | 缓存 | 存取 | 范围 | 生命周期 | | --- | | 寄存器 | 片上 | n/a | R/W | 一个线程 | 线程 | | 本地 | 片外 | Yes | R/W | 一个线程 | 线程 | | 共享 | 片上 | n/a | R/W | 块内所有线程 | 块 | | 全局 | 片外 | Yes | R/W | 所有线程+主机 | 主机配置 | | 常量 | 片外 | Yes | R | 所有线程+主机 | 主机配置 | | 纹理 | 片外 | Yes | R | 所有线程+主机 | 主机配置 | ## Reference [1] Ronny Krashinsky, Olivier Giroux, Stephen Jones, et al. NVIDIA Ampere Architecture In-Depth. [EB/OL]. 2020[2020-05-14]. [https://developer.nvidia.com/blog/nvidia-ampere-architecture-in-depth/]() [2] John Cheng. CUDA C编程权威指南 (高性能计算技术丛书) [M/OL] [http://reader.epubee.com/books/mobile/ff/ffb9b417f6f199c88f4076c66f33d1eb/text00084.html](). 北京: 机械工业出版社, 2017 [3] Nicholas Wilt. CUDA 专家手册，GPU编程权威指南. [M] 北京: 机械工业出版社, 2014