CUDA编程笔记之内存管理-拓冰建站

CUDA 中的内存管理主要包括

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

内存分配与释放

Device端

cudaError_t cudaMalloc(void **devPtr, 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）来进行分配。
设备内存的分配和释放操作成本较高，所以应用程序应重利用设备内存，以减少对整体性能的影响。

内存传输

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

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

这个函数从内存位置src复制了count字节到内存位置dst。变量kind指定了复制的方向，可以有下列取值：
在这里插入图片描述
图所示为CPU内存和GPU内存间的连接性能。从图中可以看到GPU芯片和板载GDDR5 GPU内存之间的理论峰值带宽非常高，对于Fermi C2050 GPU来说为144GB/s。CPU和GPU之间通过PCIe Gen2总线相连，这种连接的理论带宽要低得多，为8GB/s（PCIe Gen3总线最大理论限制值是16GB/s）。这种差距意味着如果管理不当的话，主机和设备间的数据传输会降低应用程序的整体性能。因此，CUDA编程的一个基本原则应是尽可能
地减少主机与设备之间的传输。
在这里插入图片描述

固定内存

GPU不能在可分页主机内存上安全地访问数据，因为当主机操作系统在物理位置上移动该数据时，它无法控制。当从可分页主机内存传输数据到设备内存时，CUDA驱动程序首先分配临时页面锁定的或固定的主机内存，将主机源数据复制到固定内存中，然后从固定内存传输数据给设备内存，如图左边部分所示：
在这里插入图片描述
CUDA运行时允许你使用如下指令直接分配固定主机内存

cudaError_t cudaMallocHost(void **devPtr, size_t count);

这个函数分配了count字节的主机内存，这些内存是页面锁定的并且对设备来说是可访问的。由于固定内存能被设备直接访问，所以它能用比可分页内存高得多的带宽进行读写。然而，分配过多的固定内存可能会降低主机系统的性能，因为它减少了用于存储虚拟内存数据的可分页内存的数量，其中分页内存对主机系统是可用的。
__ (可以理解为，固定内存是稀缺资源， GPU占用过多会影响CPU的性能) __
在这里插入图片描述

固定主机内存必须通过下述指令来释放：

cudaError_t cudaFreeHost(void *ptr);

主机与设备之间的内存传输

与可分页内存相比，固定内存有着更高的分配与释放成本，但是它为大规模数据传输提供了更高的传输吞吐量。
将许多小的传输批处理为一个更大的传输能提高性能，因为它减少了单位传输消耗。（可参考归并中的循环展开）。主机和设备之间的数据传输有时可以与内核执行重叠。

零拷贝内存

后续介绍
零拷贝内存。主机和设备都可以访问零拷贝内存。

统一内存寻址

在CUDA 6.0中，引入了“统一内存寻址”这一新特性，它用于简化CUDA编程模型中的内存管理。统一内存中创建了一个托管内存池，内存池中已分配的空间可以用**相同的内存地址（即指针）**在CPU和GPU上进行访问。底层系统在统一内存空间中自动在主机和设备之间进行数据传输。这种数据传输对应用程序是透明的。
托管内存指的是由底层系统自动分配的统一内存，与特定于设备的分配内存可以互操作，如它们的创建都使用cudaMalloc程序。因此，你可以在核函数中使用两种类型的内存：由系统控制的托管内存，以及由应用程序明确分配和调用的未托管内存。所有在设备内存上有效的CUDA操作也同样适用于托管内存。其主要区别是主机也能够引用和访问托管内存。
托管内存可以被静态分配也可以被动态分配。可以通过添加__managed__注释，静态声明一个设备变量作为托管变量。但这个操作只能在文件范围和全局范围内进行。该变量可以从主机或设备代码中直接被引用：

__device__ __managend__ int y;

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

cudaError_t cudaMallocManagend(void **devPtr, size_t size, unsigned int flags = 0);

这个函数分配size字节的托管内存，并用devPtr返回一个指针。该指针在所有设备和主机上都是有效的。使用托管内存的程序行为与使用未托管内存的程序副本行为在功能上是一致的。但是，使用托管内存的程序可以利用自动数据传输和重复指针消除功能。
注: 在CUDA 6.0中，设备代码不能调用cudaMallocManaged函数。所有的托管内存必须在主机端动态声明或者在全局范围内静态声明。

内存访问模式

大多数设备端数据访问都是从全局内存开始的，并且多数GPU应用程序容易受内存带宽的限制。为了在读写数据时达到最佳的性能，内存访问操作必须满足一定的条件。CUDA执行模型的显著特征之一就是指令必须以线程束为单位进行发布和执行。存储操作也是同样。在执行内存指令时，线程束中的每个线程都提供了一个正在加载或存储的内存地址。在线程束的32个线程中，每个线程都提出了一个包含请求地址的单一内存访问请求，它并由一个或多个设备内存传输提供服务。根据线程束中内存地址的分布，内存访问可以被分成不同的模式。

对齐与合并访问

设备内存访问的两个特性：

对齐内存访问
合并访问内存

在内存访问中， 所有对全局内存的访问都会通过二级缓存，也有许多访问会通过一级缓存，这取决于访问类型和GPU架构。如果这两级缓存都被用到，那么内存访问是由一个128字节的内存事务实现的。如果只使用了二级缓存，那么这个内存访问是由一个32字节的内存事务实现的。对全局内存缓存其架构，如果允许使用一级缓存，那么可以在编译时选择启用或禁用一级缓存。
一级、二级缓存都用到： 128字节的事务
只用到二级缓存：32字节的事务
在这里插入图片描述
内存对齐：
当设备内存事务的第一个地址是用于事务服务的缓存粒度的偶数倍时（32字节的二级缓存或128字节的一级缓存），就会出现对齐内存访问。运行非对齐的加载会造成带宽浪费。（ps：比如内存事务的第一个地址为256，正好是偶数倍。如果为100就不是偶数倍）。
内存访问合并：
当一个线程束中全部的32个线程访问一个连续的内存块时，就会出现合并内存访问。（一个线程束访问连续的内存地址）。

对齐合并内存访问的理想状态是线程束从对齐内存地址开始访问一个连续的内存块。为了最大化全局内存吞吐量，为了组织内存操作进行对齐合并是很重要的。图4-7描述了对齐与合并内存的加载操作。在这种情况下，只需要一个128字节的内存事务从设备内存中读取数据。图4-8展示了非对齐和未合并的内存访问。在这种情况下，可能需要3个128 字节的内存事务来从设备内存中读取数据：一个在偏移量为0的地方开始，读取连续地址之后的数据；一个在偏移量为256的地方开始，读取连续地址之前的数据；另一个在偏移量为128的地方开始读取大量的数据。注意在内存事务之前和之后获取的大部分字节将不能被使用，这样会造成带宽浪费。
在这里插入图片描述

一般来说，需要优化内存事务效率：用最少的事务次数满足最多的内存请求。事务数量和吞吐量的需求随设备的计算能力变化。

全局内存读取

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

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

内存加载访问模式

内存加载可以分为两类：

缓存加载（启用一级缓存）
没有缓存的加载（禁用一级缓存）

内存加载的访问模式特点如下：
有缓存与没有缓存：如果启用一级缓存，则内存加载被缓存
对齐与非对齐：如果内存访问的第一个地址是32字节的倍数，则对齐加载
合并与未合并：如果线程束访问一个连续的数据块则合并加载

缓存加载
缓存加载操作经过一级缓存，在粒度为128字节的一级缓存行上由设备内存事务进行传输。缓存加载可以分为对齐/非对齐及合并/非合并。（排列组合）

情形1：对齐与合并内存访问
线程束中所有线程请求的地址都在128 字节的缓存行范围内。完成内存加载操作只需要一个128字节的事务。总线的使用率为100%，在这个事务中没有未使用的数据。
在这里插入图片描述

情形2：访问是对齐的，引用的地址不是连续的线程ID，而是128 字节范围内的随机值。
由于线程束中线程请求的地址仍然在一个缓存行范围内，所以只需要一个128字节的事务来完成这一内存加载操作。总线利用率仍然是100%，并且只有当每个线程请求在128字节范围内有4个不同的字节时，这个事务中才没有未使用的数据。（一个线程束32个线程，只有当每个线程的请求为4
字节时才能100%占用）
在这里插入图片描述

情形3：线程束请求32个连续4个字节的非对齐数据元素。（合并连续但不对齐）
在全局内存中线程束的线程请求的地址落在两个128字节段范围内。因为当启用一级缓存时，由SM执行的物理加载操作必须在128个字节的界线上对齐，所以要求有两个128字节的事务来执行这段内存加载操作。总线利用率为50%，并且在这两个事务中加载的字节有一半是未使用的。
(一级缓存的内存事务最小单位为128字节)
在这里插入图片描述

情形4：线程束中所有线程请求相同的地址。
因为被引用的字节落在一个缓存行范围内，所以只需请求一个内存事务，但总线利用率非常低。如果加载的值是4字节的，则总线利用率是4字节请求/128字节加载＝3.125%。
在这里插入图片描述

情形5：线程束中线程请求分散于全局内存中的32个4字节地址。（一个线程束32个线程，每个线程请求的地址分散）
线程束请求的字节总数仅为128个字节，但地址要占用N个缓存行（0＜N≤32，每个线程一个）。完成一次内存加载操作需要申请N次内存事务。
在这里插入图片描述

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

没有缓存的加载
没有缓存的加载不经过一级缓存，它在内存段的粒度上（32个字节）而非缓存池的粒度（128个字节）执行。这是更细粒度的加载，可以为非对齐或非合并的内存访问带来更好的总线利用率。
和上述一级缓存类似，不经过一级缓存也有一下几种情形（只是事件单位从128字节变为32字节）

情形1：对齐与合并内存访问
128个字节请求的地址占用了4个内存段，总线利用率为100%。
在这里插入图片描述

情形2：内存访问是对齐的且线程访问是不连续的，在128个字节的范围内随机进行。
只要每个线程请求唯一的地址，那么地址将占用4个内存段，并且不会有加载浪费。这样的随机访问不会抑制内核性能。
在这里插入图片描述

情形3：线程束请求32个连续的4字节元素但加载没有对齐到128个字节的边界。
请求的地址最多落在5个内存段内，总线利用率至少为80%。与这些类型的请求缓存加载相比，使用非缓存加载会提升性能，这是因为加载了更少的未请求字节。
在这里插入图片描述

情形4：线程束中所有线程请求相同的数据。
地址落在一个内存段内，总线的利用率是请求的4字节/加载的32字节＝12.5%，在这种情况下，非缓存加载性能也是优于缓存加载的性能。
在这里插入图片描述

情形5：线程束请求32个分散在全局内存中的4字节字。
由于请求的128个字节最多落在N个32字节的内存分段内而不是N个128个字节的缓存行内，所
以相比于缓存加载，即便是最坏的情况也有所改善。
在这里插入图片描述

只读缓存

只读缓存的加载粒度是32个字节。通常，对分散读取来说，这些更细粒度的加载要优于一级缓存。
有两种方式可以指导内存通过只读缓存进行读取：

使用函数__ldg
在间接引用的指针上使用修饰符
示例：

你可以使用内部函数__ldg来通过只读缓存直接对数组进行读取访问：

你也可以将常量__restrict__修饰符应用到指针上。这些修饰符帮助nvcc编译器识别无别名指针（即专门用来访问特定数组的指针）。nvcc将自动通过只读缓存指导无别名指针
的加载。