CUDA学习之第四章：全局内存（二）

4.2 内存管理

4.2.1 内存分配和释放

动态分配全局内存的经典函数：

cuda Error_t cudaMalloc(void ** devPtr, size_t count);

分配了count字节的全局内存给devPtr指向的地址。初始化可以使用如下函数：

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

释放则是经典函数：

cudaError_t cudaFree(void *devPtr);

分配和释放的成本很高，应该多次利用内存少开新的。

4.2.2 内存传输

分配好了就可以传输数据，我们依旧使用之前提到的经典函数：

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

此函数从src复制count字节到dst，kind指定了方向。其取值有四种情况：

cudaMemcpyHostToHost
cudaMemcpyHostToDevice
cudaMemcpyDeviceToHost
cudaMemcpyDeviceToDevice

这个函数大部分情况也是同步的。

4.2.3 固定内存

分配的主机内存默认是可分页的，GPU不能安全地访问这些内存。CUDA驱动程序会先将主机源数据复制到固定内存，再给GPU。
那么我们可以直接分配固定内存，使用如下函数：

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

分配了count字节的主机内存，且是页面锁定的，对设备是可访问的。因为固定内存可以直接被设备访问，其带宽高，但是影响主机的性能。
释放必须使用如下函数：

cudaError_t cudaFreeHost(void *ptr);

对于超过10MB的数据传输，在Fermi设备上使用固定内存会更好。
使用方法：将原本的malloc替换为cudaMallocHost即可。

4.2.4 零拷贝内存

零拷贝内存是主机和设备都可以访问的，其优势有三：

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

使用零拷贝内存时要注意主机和设备的同步，如果同时更改会出问题。
零拷贝内存也是固定内存，映射到设备地址空间中。通过如下函数创建一个到固定内存的映射。

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

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

• cudaHostAllocDefault函数使其与cduaMallocHost函数一致
• cudaHostAllocPortable函数返回能被CUDA上下文使用的固定内存，而不仅是执行内存分配的那个。
• 标志cudaHostAllocWriteCombined返回写结合内存，传输快，但是主机不能有效读取，用作缓冲区很好。
• 标志cudaHostAllocMapped是零拷贝内存的标志，使主机写入和设备读取被映射到地址空间中的主机内存。

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

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

返回了一个在pDevice中的设备指针，可以在设备上被引用来得到固定主机内存。设备不支持的话函数失效，flag留有后用，现在为0。
频繁读写的时候使用零拷贝内存作为设备内存补充会降低性能。
使用方法：先用cudaHostAlloc开一个主机内存，flag设为cudaHostAllocMapped。然后用cudaHostGetDevice函数获取其被映射的设备指针，直接使用设备指针给核函数传参即可。（传回结果仍然使用cudaMemcpy的老一套，就是全局内存。）
零拷贝内存少用，效率很低，而且越大效率越低。

4.2.5 统一虚拟寻址

统一虚拟寻址称为UVA，在CUDA4.0中被引入。这个时候就不用管理设备还是主机了，上一小节零拷贝内存可由此简化为直接cudaHostAlloc一个主机内存，然后直接传参给核函数即可。效率与上一节测试内容中零拷贝内存的效率相同。
主要作用：简化代码（而非效率）

4.2.6 统一内存寻址

也是简化代码用的，创建了一个托管内存池，其中之空间可以用相同的地址提供给主机和设备，其依赖于UVA之支持但是与之不同。统一内存寻址由于将内存和执行空间分离，会提升局部性和性能。
托管内存指的是底层系统自动分配的统一内存，静态分配用__managed__注释（也是全局变量），如下：

__device__ __managed__ int y;

动态分配则用如下函数：

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

分配size字节托管内存，用devPtr返回指针，指针在主机和设备都有效，注意此函数必须于主机端使用。