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字节的内存事务实现的。对全局内存缓存其架构，如果允许使用一级缓存，那么可以在编译时选择启用或禁用一级缓存。
内存对齐：
当设备内存事务的第一个地址是用于事务服务的缓存粒度的偶数倍时（32字节的二级缓存或128字节的一级缓存），就会出现对齐内存访问。运行非对齐的加载会造成带宽浪费。 （ps： 比如内存事务的第一个地址为256， 正好是偶数倍。 如果为100就不是偶数倍）。
内存访问合并：
当一个线程束中全部的32个线程访问一个连续的内存块时，就会出现合并内存访问。（一个线程束访问连续的内存地址）。
对齐合并内存访问的理想状态是线程束从对齐内存地址开始访问一个连续的内存块。为了最大化全局内存吞吐量，为了组织内存操作进行对齐合并是很重要的。图4-7描述了对齐与合并内存的加载操作。在这种情况下，只需要一个128字节的内存事务从设备内存中读取数据。图4-8展示了非对齐和未合并的内存访问。在这种情况下，可能需要3个128 字节的内存事务来从设备内存中读取数据：一个在偏移量为0的地方开始，读取连续地址之后的数据；一个在偏移量为256的地方开始，读取连续地址之前的数据；另一个在偏移量为128的地方开始读取大量的数据。注意在内存事务之前和之后获取的大部分字节将不能被使用，这样会造成带宽浪费。
一般来说，需要优化内存事务效率：用最少的事务次数满足最多的内存请求。事务数量和吞吐量的需求随设备的计算能力变化。