1. 为什么要优化Device ↔ Host的内存传递

设备内存和 GPU 之间的峰值理论带宽（例如，NVIDIA Tesla V100 上的峰值理论带宽为 898 GB/s）远高于主机内存和设备内存之间的峰值理论带宽（PCIe x16 Gen3 上的 16 GB/s）。因此，为了获得最佳的整体应用程序性能，最大限度地减少主机和设备之间的数据传输非常重要，即使这意味着在 GPU 上运行内核与在主机 CPU 上运行内核相比，这些内核不会显示任何加速。
注意: 尽量减少主机和设备之间的数据传输，即使这意味着在设备上运行某些内核，与在主机 CPU 上运行这些内核相比，这些内核的性能提升不会提高。

2. 优化内存传递的几种方式

2.1. 使用Host上的Pinned Memor存储并使用DeviceCopy

Page-locked 或 pinned 内存有更好的访问效率。

注意：

不应过度使用固定内存。过度使用可能会降低整体系统性能，因为固定内存是稀缺资源，但很难事先知道多少是太多。
分配系统内存的耗时是较慢的。

2.2. 计算与传输的重叠

Host & Device 与内存传输并行执行 e.g. cudacopy + cuda algorithm 与 cpu algorithm 并行
多stream执行重叠的传输和计算 注意:
1. 需要GPU支持计算、传输并发
2. 出现Host 与 Device 双向传输的并发，需要注意GPU是否为每个方向传输准备了专门的硬件
  e.g.

2.3. 零拷贝内存形式

注意： 为Device映射一片Host内存。这导致GPU与CPU访问内存时出现频繁的映射。不建议在当前GPU上使用。

float *a_h, *a_map;
...
cudaGetDeviceProperties(&prop, 0);       // 检查设备驱动是否支持零拷贝形式的内存
if (!prop.canMapHostMemory)
    exit(0);
cudaSetDeviceFlags(cudaDeviceMapHost);         // 将设备设置成零拷贝模式，导致申请Host内存时可以同时申请一块Device内存做映射
cudaHostAlloc(&a_h, nBytes, cudaHostAllocMapped);       // 开辟一块Host内存，同时引用cudaHostAllocMapped启用零拷贝
cudaHostGetDevicePointer(&a_map, a_h, 0);          // 通过Host内存指针获取device内存指针
kernel<<<gridSize, blockSize>>>(a_map);
cudaFreeHost(a_h);           // 释放零拷贝的内存

2.4. 统一虚拟寻址UVA (Unified Virtual Addressing)

使用 UVA 时，主机内存和所有已安装受支持设备的设备内存共享一个虚拟地址空间。

注意： 使用pinned memory方式申请的内存仍然保持Host与Device不同的地址指针。

2.5. 统一内存UM（Unified Managed Memory）

另见本栏专题文章

4.0.0 【NVIDIA-GPU-CUDA】数据传输优化 —— Device ↔ Host