CUDA_C_NOTES [4]

CH04 全局内存

4.1 CUDA内存模型概述

在现有的硬件存储子系统下， 必须依靠内存模型获得最佳的延迟和带宽。 CUDA内存模结合了主机和设备的内存系统， 展现了完整的内存层次结构， 使你能显式地控制数据布以优化性能.s

4.1.1 内存层次结构的优点

两种不同类型的局部性:

• 时间局部性：时间局部性认为如果一个数据位置被引用， 那么该数据在较短的间周期内很可能会再次被引用， 随着时间流逝， 该数据被引用的可能性逐渐降低
• 空间局部性：空间局部性认为如果一个内存位置被引用， 则附近的位置也可能会被引用

现代计算机使用不断改进的低延迟低容量的内存层次结构来优化性能。 这种内存层次结构仅在支持局部性原则的情况下有效。 一个内存层次结构由具有不同延迟、 带宽容量的多级内存组成。 通常， 随着从处理器到内存延迟的增加， 内存的容量也在增加。

CPU和GPU的主存都采用的是DRAM（动态随机存取存储器），而低延迟内存（如CPU一级缓存）使用的则是SRAM（静态随机存取存储器）。内存层次结构中最大且最慢的级别通常使用磁盘或闪存驱动来实现。在这种内存层次结构中，当数据被处理器频繁使用时，该数据保存在低延迟、低容量的存储器中；而当该数据被存储起来以备后用时，数据就存储在高延迟、大容量的存储器中。这种内存层次结构符合大内存低延迟的设想。

GPU和CPU内存模型的主要区别是， CUDA编程模型能将内存层次结构更好地呈现给用户， 能让我们显式地控制它的行为.

4.1.2 CUDA内存模型

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

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

在CPU内存层次结构中， 一级缓存和二级缓存都是不可编程的存储器。

CUDA内存模型提出了多种可编程内存的类型:

• 寄存器 (register)
• 共享内存 (shared memory)
• 本地内存 (local memory)
• 常量内存（constant memory）
• 纹理内存 ()
• 全局内存(global memory)

一个核函数中的线程都有自己私有的本地内存。 一个线程块有自己的共享内存， 对同一线程块中所有线程都可见， 其内容持续线程块的整个生命周期。 所有线程都可以访问全局内存。 所有线程都能访问的只读内存空间有： 常量内存空间和纹理内存空间。 > 全局内存、 常量内存和纹理内存空间有不同的用途。 纹理内存为各种数据布局提供了不同的寻址模式和滤波模式。 对于一个应用程序来说， 全局内存、 常量内存和纹理内存中的内容具有相同的生命周期.

4.1.2.1 寄存器

寄存器是GPU上运行速度最快的内存空间。

核函数中声明的一个没有其他修饰符的自变量， 通常存储在寄存器中。 在核函数声明的数组中， 如果用于引用该数组的索引是常量且能在编译时确定， 那么该数组也存储在寄存器中。

寄存器变量对于每个线程来说都是私有的， 一个核函数通常使用寄存器来保存需要频 繁访问的线程私有变量。 寄存器变量与核函数的生命周期相同。 一旦核函数执行完毕， 就不能对寄存器变量进行访问了。

寄存器是一个在SM中由活跃线程束划分出的较少资源:

在Fermi架构中，每个线程最多有63个寄存器； 在Kepler架构中，每个线程最多有255个寄存器；

在核函数中使用较少的寄存器将使在SM上有更多的常驻线程块。 每个SM上并发线程块越多，使用率和性能就越高

如果一个核函数使用了超过硬件限制数量的寄存器， 则会用本地内存替代多占用的寄 存器。

4.1.2.2 本地内存（local memory）

编译器可能存放到本地内存中的变量有：

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

“本地内存”这一名词是有歧义的： 溢出到本地内存中的变量本质上与全局内存在同一 块存储区域， 因此本地内存访问的特点是高延迟和低带宽， 并且如在本章后面的4.3节中所描述的那样， 本地内存访问符合高效内存访问要求.

4.1.2.3 共享内存

在核函数中使用如下修饰符修饰的变量存放在共享内存中： __shared__ 因为共享内存是片上内存， 所以与本地内存或全局内存相比， 它具有更高的带宽和更 低的延迟。 它的使用类似于CPU一级缓存， 但它是可编程的。

每一个SM都有一定数量的由线程块分配的共享内存。 因此， 必须非常小心不要过度使用共享内存， 否则将在不经意间限制活跃线程束的数量。

共享内存在核函数的范围内声明， 其生命周期伴随着整个线程块。 当一个线程块执行结束后， 其分配的共享内存将被释放并重新分配给其他线程块。

共享内存是线程之间相互通信的基本方式。 一个块内的线程通过使用共享内存中的数 据可以相互合作。 访问共享内存必须同步使用如下调用， 该命令是在之前章节中介绍过的CUDA运行时调用： void __syncthreads();

该函数设立了一个执行障碍点， 即同一个线程块中的所有线程必须在其他线程被允许 执行前达到该处。 为线程块里所有线程设立障碍点， 这样可以避免潜在的数据冲突。

SM中的一级缓存和共享内存都使用64KB的片上内存， 它通过静态划分， 但在运行时 可以通过如下指令进行动态配置： cudaError_t cudaFuncSetCacheConfig(const void* func, enum cadaFuncCache cacheConfig)

4.1.2.4 常量内存

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

常量变量必须在全局空间内和所有核函数之外进行声明。 对于所有计算能力的设备， 都只可以声明64KB的常量内存。 常量内存是静态声明的， 并对同一编译单元中的所有核函数可见。

核函数只能从常量内存中读取数据。（不能往常量内存中写数据） 因此， 常量内存必须在主机端使用下面的函数来 初始化： cudaError_t cudaMemoryToSymbol(const void* symbol, const void* src, size_t count)

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

线程束中的所有线程从相同的内存地址中读取数据时， 常量内存表现最好。

举个例子， 数学公式中的系数就是一个很好的使用常量内存的例子， 因为一个线程束中所有的线程使用相同的系数来对不同数据进行相同的计算。 如果线程束里每个线程都从不同的地址空间读取数据， 并且只读一次， 那么常量内存中就不是最佳选择， 因为每从一个常量内存中读取一次数据， 都会广播给线程束里的所有线程。

4.1.2.5 纹理内存

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

4.1.2.6 全局内存

全局内存是GPU中最大、 延迟最高并且最常使用的内存。 global指的是其作用域和生命周期。 它的声明可以在任何SM设备上被访问到， 并且贯穿应用程序的整个生命周期。

一个全局内存变量可以被静态声明或动态声明。 你可以使用如下修饰符在设备代码中 静态地声明一个变量：

__device__

在第2章的2.1节中， 你已经学习了如何动态分配全局内存。 在主机端使用cuda-Malloc 函数分配全局内存， 使用cudaFree函数释放全局内存。 然后指向全局内存的指针就会作为 参数传递给核函数。 全局内存分配空间存在于应用程序的整个生命周期中， 并且可以访问 所有核函数中的所有线程。 从多个线程访问全局内存时必须注意。 因为线程的执行不能跨 线程块同步， 不同线程块内的多个线程并发地修改全局内存的同一位置可能会出现问题， 这将导致一个未定义的程序行为。

优化内存事务对于获得最优性能来说是至关重要的。 当一个线程束执行内存加载/ 存储时， 需要满足的传输数量通常取决于以下两个因素：

• 跨线程的内存地址分布
• 每个事务内存地址的对齐方式

对于一个给定的线程束内存请求， 事务数量和数据吞吐率是由设备的计算能力来确定 的。 对于计算能力为1.0和1.1的设备， 全局内存访问的要求是非常严格的。 对于计算能力高于1.1的设备， 由于内存事务被缓存， 所以要求较为宽松。 缓存的内存事务利用数据局部性来提高数据吞吐率。

4.1.2.7 GPU缓存

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

• 一级缓存
• 二级缓存
• 只读常量缓存
• 只读纹理缓存

每个SM都有一个一级缓存， 所有的SM共享一个二级缓存。 一级和二级缓存都被用来在存储本地内存和全局内存中的数据， 也包括寄存器溢出的部分。对Fermi GPU和Kepler K40或其后发布的GPU来说， CUDA允许我们配置读操作的数据是使用一级和二级缓存，还是只使用二级缓存。

在GPU上只有内存加载操作可以被缓存，内存存储操作不能被缓存。 每个SM也有一个只读常量缓存和只读纹理缓存， 它们用于在设备内存中提高来自于各自内存空间内的读取性能。

4.1.2.8 CUDA变量声明总结

4.1.2.9 静态全局内存

4.2 内存管理

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

• 分配和释放设备内存
• 在主机和设备之间传输数据

4.2.1 内存分配和释放

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

你可以在主机上使用下列函数分配全局内存：

cudaError_t cudaMalloc(void **devPrt, size_t count);

这个函数在设备上分配了count字节的全局内存， 并用devptr指针返回该内存的地址。

你需要用从主机上传输的数据来填充所分配的全局内存， 或用下列函数将其初始 化:

cudaError_t cudaMemset(void *devPtr, int value, size_t count);

这个函数用存储在变量value中的值来填充从设备内存地址devPtr处开始的count字节。

一旦一个应用程序不再使用已分配的全局内存， 那么可以以下代码释放该内存空间： cudaError_t cudaFree(void *devPtr);

这个函数释放了devPtr指向的全局内存， 该内存必须在此前使用了一个设备分配函数 （如cudaMalloc） 来进行分配。 否则， 它将返回一个错误cudaErrorInvalidDevicePointer。 如果地址空间已经被释放， 那么cudaFree也返回一个错误。

4.2.2 内存传输

一旦分配好了全局内存， 你就可以使用下列函数从主机向设备传输数据：

cudaError_t cudaMemory(void *dst, const void *src, size_t count, enum cudaMemcpyKind kind)

这个函数从内存位置src复制了count字节到内存位置dst。 变量kind指定了复制的方向， 可以有下列取值：

• cudaMemcpyHostToHost
• cudaMemcpyHostToDevice
• cudaMemcpyDeviceToHost
• cudaMemcpyDeviceToDevice

CUDA编程的一个基本原则应是尽可能地减少主机与设备之间的传输.

4.2.3 固定内存

分配的主机内存默认是pageable（可分页） ， 它的意思也就是因页面错误导致的操 作， 该操作按照操作系统的要求将主机虚拟内存上的数据移动到不同的物理位置。 虚拟内存给人一种比实际可用内存大得多的假象， 就如同一级缓存好像比实际可用的片上内存大得多一样。

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

CUDA运行时允许你使用如下指令直接分配固定主机内存： cudaError_t cudaMallocHost(void **devPtr, size_t count);

这个函数分配了count字节的主机内存， 这些内存是页面锁定的并且对设备来说是可 访问的。 由于固定内存能被设备直接访问， 所以它能用比可分页内存高得多的带宽进行读写。 然而， 分配过多的固定内存可能会降低主机系统的性能， 因为它减少了用于存储虚拟内存数据的可分页内存的数量， 其中分页内存对主机系统是可用的。

主机与设备间的内存传输

与可分页内存相比， 固定内存的分配和释放成本更高， 但是它为大规模数据传输提供 了更高的传输吞吐量

4.2.4 零拷贝内存

通常来说， 主机不能直接访问设备变量， 同时设备也不能直接访问主机变量。 但有一个例外： 零拷贝内存。 主机和设备都可以访问零拷贝内存。

GPU线程可以直接访问零拷贝内存。 在CUDA核函数中使用零拷贝内存有以下几个优 势。

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

当使用零拷贝内存来共享主机和设备间的数据时， 你必须同步主机和设备间的内存访 问， 同时更改主机和设备的零拷贝内存中的数据将导致不可预知的后果。

零拷贝内存是固定（不可分页） 内存， 该内存映射到设备地址空间中。 你可以通过下列函数创建一个到固定内存的映射：

cudaError_t cudaHostAlloc(void **pHost, size_t count, unsigned int flags);

这个函数分配了count字节的主机内存， 该内存是页面锁定的且设备可访问的。 用这 个函数分配的内存必须用cudaFreeHost函数释放。 flags参数可以对已分配内存的特殊属性 进一步进行配置：

- cudaHostAllocDefault
- cudaHostAllocPortable
- cudaHostAllocWriteCombined
- cudaHostAllocMapped

cudaHostAllocDefault函数使cudaHostAlloc函数的行为与cudaMallocHost函数一致。

设置cudaHostAllocPortable函数可以返回能被所有CUDA上下文使用的固定内存， 而不仅是执 行内存分配的那一个。

标志cudaHostAllocWriteCombined返回写结合内存， 该内存可以在某些系统配置上通过PCIe总线上更快地传输， 但是它在大多数主机上不能被有效地读取。因此， 写结合内存对缓冲区来说是一个很好的选择， 该内存通过设备使用映射的固定内存或主机到设备的传输。

零拷贝内存的最明显的标志是cudaHostAllocMapped， 该标志返回， 可以实现主机写入和设备读取被映射到设备地址空间中的主机内存。

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

cudaError_t cudaHostGetDevicePointer(void **pDevice, void *pHost, unsigned int flags);

该函数返回了一个在pDevice中的设备指针， 该指针可以在设备上被引用以访问映射得到的固定主机内存。 如果设备不支持映射得到的固定内存， 该函数将失效。 flag将留作以后使用。 现在， 它必须被置为0。

在进行频繁的读写操作时， 使用零拷贝内存作为设备内存的补充将显著降低性能。 因为每一次映射到内存的传输必须经过PCIe总线。 与全局内存相比， 延迟也显著增加。

零拷贝内存

有两种常见的异构计算系统架构： 集成架构和离散架构。

在集成架构中， CPU和GPU集成在一个芯片上， 并且在物理地址上共享主存。 在这种架构中， 由于无须在PCIe总线上备份， 所以零拷贝内存在性能和可编程性方面可能更佳。

对于通过PCIe总线将设备连接到主机的离散系统而言， 零拷贝内存只在特殊情况下有优势。

因为映射的固定内存在主机和设备之间是共享的， 你必须同步内存访问来避免任何潜在的数据冲突， 这种数据冲突一般是由多线程异步访问相同的内存而引起的。

注意不要过度使用零拷贝内存。 由于其延迟较高， 从零拷贝内存中读取设备核函数可能很慢。

4.2.5 统一虚拟寻址