CUDA C编程（八）CUDA内存管理

  CUDA编程的内存管理与C语言的类似，需要程序员显式管理主机和设备之间的数据移动。随着CUDA版本的升级，NVIDIA正系统地实现主机和设备内存空间的统一，但对于大多数应用程序来说，仍需要手动移动数据。对于CUDA内存管理来说，工作重点在于如何使用CUDA函数来显式地管理内存和数据移动，主要是两个方面：分配和释放设备内存；在主机和设备之间传输数据。为了达到最优性能，CUDA提供了在主机端准备设备内存地函数，并且显式地向设备传输数据和从设备中获取数据。
内 存 分 配 和 释 放
  CUDA编程模型假设了一个包含一个主机和一个设备地异构系统，每一个异构系统都有自己独立地内存空间。核函数在设备内存空间中运行，CUDA运行时提供函数以分配和释放设备内存。可以使用cudaError_t cudaMalloc(void **devPtr, size_t count)函数分配全局内存，这个函数在设备上分配了count字节的全局内存，并用devptr指针返回该内存的地址。所分配的内存支持任何变量类型，包含整型、浮点类型变量、布尔类型等。如果cudaMalloc函数执行失败则返回cudaErrorMemoryAllocation。在已分配的全局内存中的值不会被清除。你需要用从主机上传输的数据来填充所分配的全局内存，或用下列函数将其初始化：cudaError_t cudaMemcpy(void *devPtr, int value, size_t count),这个函数用存储在变脸value中的值来填充设备内存地址devPtr处开始的count个字节。一旦一个应用程序不再使用已分配的全局内存，那么可以使用以下代码释放该内存空间：cudaError_t cudaFree(void *devPtr);这个函数释放了devPtr指向的全局内存，该内存必须在此前使用了一个设备分配函数（cudaMalloc）来进行分配。否则，它将返回一个错误cudaErrorInvalidDevicePointer。如果地址空间已经被释放，那么cudaFree也返回一个错误。设备内存的分配和释放操作成本较高，所以应用程序应重利用设备内存，以减少对整体性能的影响。

内 存 传 输
  一旦分配好了全局内存，就可以使用cudaError_t cudaMemcpy(void *dst,const void *src,size_t count,enum cudaMemcpyKind kind)函数从主机向设备传输数据，这个函数从内存weizhisrc复制了count字节到内存位置dst。变量kind指定了复制的方向，可以有下列取值；cudaMemcpyHostToHost、cudaMemcpyHostToDevice、cudaMemcpyDeviceToHost、cudaMemcpyDeviceToDevice。
固 定 内 存
  分配的主机内存默认是pageable（可分页的），它的意思也就是因页面错误导致的操作，该操作按照操作系统的要求将主机虚拟内存上的数据移动到不同的物理位置。虚拟内存给人一种给比实际可用内存大得多的假象，就如同一级缓存好像比实际可用的片上内存大得多一样。
  GPU不能在可分页内存上安全地访问数据，因为当主机操作系统在物理位置上移动该数据时，它无法控制。当从可分页内存传输数据到设备内存时，CUDA驱动程序首先分配临时页面锁定的或固定的主机内存，将主机源数据复制到固定内存中，然后从固定内存传输数据给设备内存，如下图所示：

  CUDA运行时允许使用如下指令直接分配固定主机内存：cudaError_t cudaMallocHost(void **devPtr,size_t count);这个函数分配了count字节的主机内存，这些内存是页面锁定的并且对设备来说是可访问的。由于固定内存能被设备直接访问，所以它能用比分页内存高得多的带宽进行读写。然而，分配过多的固定内存可能会降低主机系统的性能，因为它减少了用于存储虚拟内存数据的可分页内存的数量，其中分页内存对主机系统是可用的。固定主机内存被絮通过下列指令来释放：cudaError_t cudaFreeHost(void *ptr);
  与可分页内存相比，固定内存的分配和释放成本更高，但是它为大规模数据传输提供了更高的传输吞吐量。相对于可分页内存，使用固定内存获得的加速取决于设备计算能力。将许多小的传输批处理为一个更大的传输能提高性能，因为他减少了单位传输消耗。主机和设备之间的数据传输有时可以与内核执行重叠，我们应该尽可能地减少或重叠主机和设备间的数据传输。
零拷贝内存
  通常来说，主机不能直接访问设备变量，同时设备也不能直接访问主机内存。但有个例外：零拷贝内存。主机和设备都可以访问零拷贝内存。GPU线程可以直接访问零拷贝内存。在CUDA核函数中使用零拷贝内存有以下几个优势：1.当设备内存不足时可利用主机内存；2.避免主机和设备间的显式数据传输；3.提高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总线将设备连接到主机的离散系统而言，零拷贝内存只在特殊情况下有优势。因为映射的固定内存再主机和设备之间是共享的，你必须同步内存访问来避免任何潜在的数据冲突，这种数据冲突一般是由多线程异步访问相同的内存而引起的。注意不要过渡使用零拷贝内存，由于其延迟较高，从零拷贝内存中读取设备核函数可能很慢。
统 一 虚 拟 寻 址
  计算能力为2.0及以上版本的设备支持一种特殊的寻址方式，称为统一虚拟寻址(UAV)。UAV 4.0中被引入，支持64位Linux系统。有了UAV，主机内存和设备内存可以共享一个虚拟地址空间，如下图所示：

  有UVA之前，我们需要管理哪些指针指向主机内存和哪些指针指向设备内存。有了UAV，由指针指向的内存空间对应用程序来说是透明的。通过UVA，由cudaHostAlloc分配的固定主机内存具有相同的主机和设备指针。因此，可以将返回的指针直接传递给核函数。有了UVA，无须获取设备指针或管理物理上数据完全相同的两个指针
统 一 内 存 寻 址
  在CUDA6.0中，引入了“统一内存寻址”这一特性，它用于简化CUDA编程模型中的内存管理。统一内存中创建了一个托管内存池，内存池中已分配的空间可以用相同的内存地址（即指针）在CPU和GPU上进行访问。底层系统在统一内存空间中自动在主机和设备之间进行数据传输。这种数据传输对应用程序是透明的，这大大简化了程序代码。统一内存寻址依赖于UVA的支持，但它们是完全不同的技术。UVA为系统中的所有处理器提供了一个单一的虚拟内存地址空间。但是，UVA不会自动地将数据从一个物理位置转移到另一个位置，这是统一内存寻址的一个特有功能。统一内存寻址提供了一个“单指针到数据”模型，在概念上它类似于零拷贝内存，但是零拷贝内存在主机内存中进行分配，因此，由于受到在PCIe总线上访问零拷贝内存的影响，核函数的性能将具有较高的延迟。另一方面，统一内存寻址将内存和执行空间分离，因此可以根据需要将数据透明地传输到主机或设备上，以提高局部性和性能。
  托管内存指的是由底层系统自动分配的统一内存，与特定于设备的分配内存可以互操作，如它们的创建都使用cudaMalloc程序。因此，我们可以在核函数中使用两种类型的内存：由系统控制的托管内存，以及由应用程序明确分配和调用的未托管内存。所以在设备内存上有效的CUDA操作也同样用于托管内存。其主要区别是主机也能够引用和访问托管内存。托管内存可以被静态分配也可以被动态分配，可以通过添加__managed__注释，静态声明一个设备变量作为托管变量。但这个操作只能在文件范围和全局范围内进行。该变量可以从主机或设备代码中直接被引用：__device__ __managed__ int y;还可以使用下述的CUDA运行时函数动态分配托管内存：cudaError_t cudaMallocManaged(void **dePtr,size_t size,unsigned int flags=0);这个函数分配size字节的托管内存，并用devPtr返回一个指针。该指针在所有设备和主机上都是有效的。使用托管内存的程序行为与使用未托管内存的程序副本行为在功能上是一致的。但是，使用托管内存的程序可以利用自动数据传输和重复指针消除功能。