CUDA全局内存

1. 简介

CPU 和 GPU 的主存都采用的是 DRAM（动态随机存取存储器），而低延迟内存（如 CPU 一级缓存）使用的则是 SRAM（静态随机存取存储器）。CPU 和 GPU 在内存层次结构设计中都使用相似的准则和模型。GPU 和 CPU 内存模型的主要区别是，CUDA 编程模型能将内存层次结构更好地呈现给用户，能让我们显式地控制它的行为。

2. CUDA 内存模型

对于程序员来说，一般有两种类型的存储器：

可编程的：需要显式地控制哪些数据存放在可编程内存中；
不可编程的：无法人为决定数据的存放位置，程序将自动生成存放位置以获得良好的性能。

在 CPU 内存层次中，一级缓存和二级缓存都是不可编程的存储器。另一方面，CUDA 内存模型提出了多种可编程内存的类型：

寄存器：GPU 上运行速度最快的内存空间。核函数中声明的一个没有其他修饰符的自变量，通常存储在寄存器中。寄存器变量对于每个线程来说都是私有的，寄存器变量的生命周期与核函数相同。
共享内存：一个线程块有自己的共享同存，对同一线程块中所有线程都可见，其内容持续线程块的整个生命周期。
本地内存：一个核函数中的线程都有自己私有的本地内存。
常量内存：所有线程都可以访问的只读内存空间。
纹理内存：所有线程都可以访问的只读内存空间。纹理内存为各种数据布局提供了不同的寻址模式和滤波模式。
全局内存：所有线程都可以访问全局内存。

寄存器

寄存器是一个在 SM 中由活跃线程束划分出的较少资源。在 Fermi GPU 中，每个线程限制最多拥有 $63$ 个寄存器。Kepler GPU 将该限制扩展至每个线程可拥有 $255$ 个寄存器。在核函数中使用较少的寄存器将使在 SM 上有更多的常驻线程块。每个 SM 上并发线程块越多，使用率和性能就越高。如果一个核函数使用了超过硬件限制数量的寄存器，则会使用本地内存替代多占用的寄存器。

nvcc 编译器使用启发式策略来最小化寄存器的使用，以避免寄存器溢出。我们也可以在代码中为每个核函数显式地加上额外的信息来帮助编译器进行优化：

__global__ void
__launch_bounds__ (maxThreadsPerBlock, minBlocksPerMultiprocessor)
kernel(...) {
	// your kernel body
}

其中，maxThreadsPerBlock 指出了每个线程块可以包含的最大线程数，这个线程块由核函数来启动。minBlockPerMultiprocessor 是可选参数，指明了在每个 SM 中预期的最小的常驻线程块数量。对于给定的核函数，最优的启动边界会因主要架构的版本不同而有所不同。可以使用 -maxrregcount 编译器选项来指定一个编译单元里所有核函数使用的寄存器的最大数量。

本地内存

核函数中符合存储在寄存器中但不能进入被该核函数分配的寄存器空间中的变量将溢出到本地内存中。编译器可能存放到本地内存中的变量有：

在编译时使用未知索引引用的本地数组
可能会占用大量寄存器空间的较大本地结构体或数组
任何不满足核函数寄存器限定条件的变量

本地内存访问符合高效内存访问要求，对于计算能力 $2.0$ 及以上的 GPU 来说，本地内存数据也是存储在每个 SM 的一级缓存和每个设备的二级缓存中。

共享内存

在核函数中使用如下修饰符修饰的变量存放在共享内存中：

1	__shared__

因为共享内存是片上内存，所以与本地内存或全局内存相比，它具有更高的带宽和更低的延迟。它的使用类似于 CPU 一级缓存，但它是可编程的。每一个 SM 都有一定数量的由线程块分配的共享内存。共享内存在核函函数范围内声明，其生命周期伴随着整个线程块。**共享内存是线程之间相互通信的基本方式。**一个块内的线程通过使用共享内存中的数据可以相互合作。访问共享内存必须同步使用如下调用：

1	void __syncthreads();

该函数设立了一个执行障碍点，即同一个线程块中的所有线程必须在其它线程被允许执行前达到该处。

SM 中的一级缓存和共享内存者使用 $64$ KB 的片上内存，它通过静态划分，但在运行时可以通过如下指令进行动态配置：

1	cudaError_t cudaFuncSetCacheConfig(const void* func, enum cudaFuncCache cacheConfig);

此函数在每个核函数的基础上配置了片上内存划分，为 func 指定的核函数设置了配置。支持的缓存配置如下：

cudaFuncCachePreferNone:	没有参考值（默认）
cudaFuncCachePreferShared:	建议48KB的共享内存和16KB的一级缓存
cudaFuncCachePreferL1:		建议48KB的一级缓存和16KB的共享内存
cudaFuncCachePreferEqual:	建议相同尺寸的一级缓存和共享内存，都是32KB

Fermi 设备支持前三种配置，Kepler 设备支持以上所有配置。

常量内存

常量内存驻留在设备内存中，并在每个 SM 专用的常量缓存中缓存。常量变量用如下修饰符来修饰：

1	__constant__

常量变量必须在全局空间内和所有核函数之外进行声明。对于所有计算能力的设备，都只可以声明 $64$ KB 的常量内存。**常量内存是静态声明的，并对同一编译单元中的所有核函数可见。**核函数只能从常量内存中读取数据，因此常量内存必须在主机端使用下面的函数来初始化：

1	cudaError_t cudaMemcpyToSymbol(const void* symbol, const void* src, size_t count);

这个函数将 count 个字节从 src 指向的内存复制到 symbol 指向的内存中，这个变量存放在设备的全局内存或常量内存中。在大多数情况下，这个函数是同步的。

**每从一个常量内存中读取一次数据，都会广播给线程束中所有线程。**因此线程束中的所有线程从相同的内存地址中读取数据时，常量内存表现最好。

纹理内存

纹理内存驻留在设备内存中，并在每个 SM 的只读缓存中缓存。纹理内存是一种通过指定的只读缓存访问的全局内存。

只读缓存包括硬件滤波的支持，它可以将浮点插入作为读过程的一部分来执行。
纹理内存是对二维空间局部性的优化，所以线程束里使用纹理内存访问二维数据的线程可以达到最优性能。

硬件滤波是啥？二维空间局部性优化又是啥？

全局内存

全局内存是 GPU 中最大、延迟最高并且最常用的内存。global 指的是其作用域和生命周期。它的声明可以在任何 SM 设备上被访问到，并且贯穿应用程序的整个生命周期。一个全局内存变量可以被静态声明或动态声明，可以使用如下修饰符在设备代码中静态地声明一个变量：

1	__device__

在主机端使用 cudaMalloc 函数分配全局内存，使用 cudaFree 函数释放全局内存。全局内存分配空间存在于应用程序的整个生命周期中，并且可以访问所有核函数中的所有线程。从多个线程访问全局内存时必须注意，操作的原子性。

全局内存常驻于设备内存中，可以通过 $32$ 字节、 $64$ 字节或 $128$ 字节的内存事务进行访问。这些内存事务必须自然对齐。当一个线程束执行内存加载/存储时，需要满足的传输数量通常取决于以下两个因素：

跨线程的内存地址分布

每个事务内存地址的对齐方式

GPU 缓存

跟 CPU 缓存一样，GPU 缓存是不可编程的内存。在 GPU 上有 $4$ 种缓存：

一级缓存：粒度为 $128$ 字节
二级缓存：粒度为 $32$ 字节
只读常量缓存：粒度为 $32$ 字节
只读纹理缓存：粒度为 $32$ 字节

每个 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	一个线程	线程
本地	片外	-	R/W	一个线程	线程
共享	片上	n/a	R/W	块内所有线程	块
全局	片外	-	R/W	所有线程 + 主机	主机配置
常量	片外	Yes	R	所有线程 + 主机	主机配置
纹理	片外	Yes	R	所有线程 + 主机	主机配置

其中 + 表示只在计算能力 $\mathrm{2.x}$ 的设备上进行缓存。

设备变量

在主机端声明的设备变量只是作为一个标识符，并不是设备全局内存的变量地址。因此，不能在主机端的设备变量中使用运算符 &，因为它只是一个在 GPU 上表示物理位置符号。

有一种例外，可以直接从主机引用 GPU 内存：CUDA 固定内存。主机代码和设备代码都可以通过简单的指针引用直接访问固定内存。

文件作用域

在 CUDA 编程中，需要控制主机和设备这两个地方的操作。一般情况下，设备核函数不能访问主机变量，并且主机函数也不能访问设备变量，即使这些变量在同一文件作用域中被声明。CUDA 运行时 API 能够访问主机和设备变量。

3. 内存管理

CUDA 编程的内存管理与 C 语言的类似，需要程序员显式地管理主机和设备之间的数据移动。随着 CUDA 版本的升级，NVIDIA 正系统地实现主机和设备内存空间的统一。

统一内存是 CUDA 编程模型的一个组件，在 CUDA 6.0 中首次引入，它定义了一个托管内存空间，在该空间中所有处理器都可以看到具有公共地址空间的单个连贯内存映像。

内存分配和释放

CUDA 编程模型假设了一个包含一个主机和设备的异构系统，每一个异构系统都有自己独立的内存空间。核函数在设备内存空间中运行，CUDA 运行时提供以下函数以分配和释放设备内存：

函数	说明
`cudaError_t cudaMalloc(void **devPtr, size_t count);`	在设备上分配 `count` 字节的全局内存。执行失败则返回 `cudaErrorMemory Allocation`。
`cudaError_t cudaMemset(void *devPtr, int value, size_t count);`	用存储在变量 `value` 中的值来填充从设备内存地址 `devPtr` 处开始的 `count` 字节。
`cudaError_t cudaFree(void *devPtr)`	释放全局内存变量占用的空间。执行失败则返回错误 `cudaErrorInvalidDevicePointer`。

设备内存的分配和释放操作成本较高，所以应用程序应重利用设备内存，以减少对整体性能的影响。

内存传输

一旦分配好了全局内存，就可以使用以下函数在主机和设备之间传输数据：

函数	说明
`cudaError_t cudaMemcpy(void dst, const void src, size_t count, enum cudaMemcpyKind kind)`	指定从内存位置 `src` 复制 `count` 字节到内存位置 `dst`，变量 `kind` 指定了复制的方向。

kind 指定的复制方向有以下四种：

cudaMemcpyHostToHost

cudaMemcpyHostToDevice

cudaMemcpyDeviceToHost

cudaMemcpyDeviceToDevice

如果指针 dst 和 src 与 kind 指定的复制方向不一致，那么 cudaMemcpy 的行为就是未定义行为。这个函数在大多数情况下都是同步的。下图为 Fermi C2050 GPU 与 CPU 之间的传输带宽图：

上图说明，主机和设备间的数据传输会降低应用程序的整体性能。因此，CUDA 编程的一个基本原则应是尽可能减少主机和设备之间的传输。

固定内存

分配的主机内存默认是可分页的。GPU 不能在可分页主机内存上安全地访问数据，因为当主机操作系统在物理位置上移动该数据时，它无法控制。当从可分页内存传输数据到设备内存时，CUDA 驱动程序首先分配临时页面锁定的或固定的主机内存，将主机源数据复制到固定内存中，然后从固定内存传输数据给设备内存。

分页是一种内存管理技术，用于将进程所需的内存空间分割成固定大小的页面（或称为页框），并将这些页面映射到物理内存中的页面框。每个页面的大小通常是固定的，例如4KB或8KB。

CUDA 运行时允许你使用如下指令直接分配/释放固定主机内存：

函数	说明
`cudaError_t cudaMallocHost(void **devPtr, size_t count);`	分配 `count` 字节的固定内存，设备可以直接访问。
`cudaError_t cudaFreeHost(void *ptr);`	固定内存必须通过该函数进行释放。

零拷贝内存

通常来说，主机不能直接访问设备变量，同时设备也不能直接访问主机变量。但有一个例外：零拷贝内存，主机和设备都可以访问零拷贝内存。GPU 线程可以直接访问零拷贝内存，在 CUDA 核函数中使用零拷贝内存有以下几个优势：

当设备内存不足时可利用主机内存
避免主机和设备间的显式数据传输
提高 PCIe 传输率

当使用零拷贝内存来共享主机和设备间的数据时，必须同步主机和设备间的内存访问。零拷贝内存是固定内存，该内存映射到设备地址空间中。可以通过以下函数创建一个到固定内存的映射：

函数	说明
`cudaError_t cudaHostAlloc(void **pHost, size_t count, unsigned int flags);`	分配 `count` 字节的主机内存，该内存是页面锁定的且设备可访问的。

释放该函数分配的内存使用 cudaFreeHost 函数。flags 参数可以对已分配内存的特殊属性进一步进行配置：

cudaHostAllocDefault：使 cudaHostAlloc 函数的行为与 cudaMallocHost 函数一致。
cudaHostAllocPortable：可以返回能被所有 CUDA 上下文使用的固定内存，而不仅是执行内存分配的那一个。
cudaHostAllocWriteCombined：返回写结合内存，该内存可以在某些系统配置上通过 PCIe 总线上更快地传输，但是它在大多数主机上不能被有效地读取。
cudaHostAllocMapped：该标志返回可以实现主机写入和设备读取被映射到设备地址空间中的主机内存。

可以使用下列函数获取映射到固定内存的设备指针：

函数	说明
`cudaError_t cudaHostGetDevicePointer(void *pDevice, void pHost, unsigned int flags);`	返回一个在 `pDevice` 中的设备指针，该指针可以在设备上被引用以访问映射得到的固定主机内存。

如果想共享主机和设备端的少量数据，零拷贝内存可能会是一个不错的选择，因为它简化了编程并且有较好的性能。对于由 PCIe 总线连接的离散 GPU 上的更大数据集来说，零拷贝内存不是一个好的选择，它会导致性能的显著下降。

有两种常见的异构计算系统架构：

集成架构：在集成架构中，CPU 和 GPU 集成在一个芯片上，并且在物理地址上共享内存。在这种架构中，由于无须在 PCIe 总线上备份，所以零拷贝内存在性能和可编程性方面可能更佳。
离散架构：通过 PCIe 总线将设备连接到主机，零拷贝内存只在特殊情况下有优势。

统一虚拟寻址

计算能力 $2.0$ 及以上版本的设备支持一种特殊的寻址方式，称为统一虚拟寻址（UVA）。UVA 在 CUDA $4.0$ 中被引入，支持 $64$ 位 Linux 系统。有了 UVA，主机内存和设备内存可以共享同一个虚拟地址空间，如下图所示：

有 UVA 之前，需要管理哪些指针指向主机和哪些指针指向设备内存。有了 UVA 后，由指针指向的内存空间对应用程序代码来说是透明的。通过 UVA，由 cudaHostAlloc 分配的固定主机具有相同的主机和设备指针，可以将返回的指针直接传递给核函数。

统一内存寻址

在 CUDA $6.0$ 中，引入了「统一内存寻址」这一新特性，它用于简化 CUDA 编程模型中的内存管理，统一内存中创建了一个托管内存池，内存池中已分配的空间可以用相同的内存地址（即指针）在 CPU 和 GPU 上进行访问。底层系统在统一内存空间中自动在主机和设备之间进行数据传输。这种数据传输对应用程序是透明的。

统一内存寻址依赖于 UVA 的支持。二者是完全不同的技术：

UVA 为系统中所有处理器提供了一个单一的虚拟内存地址空间。
统一内存寻址会自动将数据从一个物理位置转移到另一个位置。

和零拷贝内存不同，后者是在主机内存中进行分配的，受到 PCIe 总线上访问零拷贝内存的影响，核函数的性能将具有较高的延迟。

托管内存指的是由底层系统自动分配的统一内存，与特定于设备的分配内存可以互操作，它们都是使用 cudaMalloc 程序创建的。可以在核函数中使用两种类型的内存：由系统控制的托管内存，以及由应用程序明确分配和调用的未托管内存。所有在设备内存上有效的 CUDA 操作也同样适用于托管内存，其主要区别是主机也能够引用和访问托管内存。

托管内存可以被静态分配也可以被动态分配，可以通过添加 __managed__ 注释，静态声明一个设备变量作为托管变量。但这个操作只能在文件范围和全局范围内进行，该变量可以从主机或设备代码中直接被引用。

1	__device__ __managed__ int y;

还可以使用下述 CUDA 运行时函数动态分配托管内存：

函数	说明
`cudaError_t cudaMallocManaged(void **devPtr, size_t size, unsigned int flags=0);`	分配 `size` 字节的托管内存，并用 `devPtr` 返回一个指针，该指针在所有设备和主机上都是有效的。使用托管内存的程序可以利用自动数据传输和重复指针消除功能。

在 CUDA $6.0$ 中，设备代码不能调用 cudaMallocManaged 函数，所有的托管内存必须在主机端动态声明或者在全局范围内静态声明。

4. 内存访问模式

大多数设备端数据访问都是从全局内存开始的，并且多数GPU应用程序容易受内存带宽的限制。因此，最大限度地利用全局内存带宽是调控核函数性能的基本。

对齐与合并访问

全局内存通过缓存来实现加载/存储，如下图所示：

核函数的内存请求通常是在 DRAM 设备和片上内存间以 $128$ 字节或 $32$ 字节内存事务来实现的。所有对全局内存的访问都会通过二级缓存，也有许多访问会通过一级缓存，这取决于访问类型和 GPU 架构。

若两级缓存都被用到，那么内存访问是由一个 $128$ 字节的内存事务实现的。
若只使用了二级缓存，那么这个内存访问是由一个 $32$ 字节的内存事务实现的。

对于全局缓存架构，如果允许使用一级缓存，那么可在编译时选择启用或禁用一级缓存。

在优化应用程序时，需要注意设备内存访问的两个特性：

对齐内存访问：当设备内存事务的第一个地址是用于事务服务的缓存粒度的偶数位时，就会出现对齐内存访问。运行非对齐的加载会造成带宽浪费。
合并内存访问：当一个线程束中全部 $32$ 个线程访问一个连续的内存块时，就会出现合并内存访问。

全局内存读取

在 SM 中，数据通过以下 $3$ 种缓存/缓冲路径进行传输，具体使用何种方式取决于引用于哪种类型的设备内存：

一级和二级缓存
常量缓存
只读缓存

一二级缓存是默认路径，想要通过其他两种路径传递数据需要应用程序显式地说明，但要想提升性能还要取决于访问模式。全局内存加载操作是否会通过一级缓存取决于两个因素：设备的计算能力、编译器选项。

默认情况下，在 Fermi 设备上对于全局内存载可以用一级缓存，在 K40 及以上 GPU 中禁用。以下标志通知编译器禁用/启用一级缓存：-Xptxas -dlcm=cg、-Xptxas -dlcm=ca。

CPU 一级缓存和 GPU 一级缓存之间的差异：CPU 一级缓存优化了时间和空间局部性，GPU 一级缓存是专为空间局部性而不是为时间局部性设计的。频繁访问一个一级缓存的内存位置不会增加数据留在缓存中的概率。

只读缓存

只读缓存有两种方式可以指导内存通过只读缓存进行读取：

使用函数 __ldg
在间接引用的指针上使用修饰符 __restrict__，帮助 nvcc 编译器识别无别名指针（即专门用来访问特定数组的指针）

全局内存写入

内存的存储操作相对简单。一级缓存不能用在 Fermi 或 Kepler GPU 上进行存储操作，在发送到设备内存之前存储操作只通过二级缓存。存储操作在 $32$ 个字节段的粒度上被执行，内存事务可以同时被分为一段、两段或四段。

结构体数组与数组结构体

数组结构体（AoS）和结构体数组（SoA）是两种常见的数据组织方式。用 SoA 模式存储数据充分利用了 GPU 的内存带宽，由于没有相同字段元素的交叉存取，GPU 上的 SoA 布局提供了合并内存访问，并且可以对全局内存实现更高效的利用。

许多并行编程范式，尤其是 SIMD 型范式，更倾向于使用 SoA。

性能调整

优化设备内存带宽利用率有两个目标：

对齐及合并内存访问，以减少带宽的浪费
足够的并发内存操作，以隐藏内存延迟
- 增加每个线程中执行独立内存操作的数量
- 对核函数启动的执行配置进行实验，以充分体现每个 SM 的并行性

最大化带宽利用率 - 影响设备内存操作性能的因素主要有两个：

有效利用设备 DRAM 和 SM 片上内存之间的字节移动：为避免设备内存带宽的浪费，内存访问模式应是对齐和合并的
当前的并发内存操作数：可通过以下两点实现最大化当前存储器操作数
- 展开，每个线程产生更多的独立内存访问
- 修改核函数启动的执行配置来使每个 SM 有更多的并行性

5. 内存带宽

一般有如下两种类型的带宽：

理论带宽：是当前硬件可以实现的绝对最大带宽。对禁用 ECC 的 Fermi M2090 来说，理论上设备内存带宽的峰值为 $177.6$ GB/s。
有效带宽：核函数实际达到的带宽，它是测量带宽，可以以下公式计算：

$有效带宽（GB/s）= \frac{(读字节数 + 写字节数) \times 10^{-9}}{运行时间}$