5.7_内存复制

5.7 内存复制

CUDA中含有3种类型的内存——主机内存、设备内存和CUDA数组，并且CUDA中实现了它们之间内存复制的全套函数。对主机与设备间内存复制，CUDA还提供了一个额外的内存复制函数集合来支持锁页主机内存和设备内存或CUDA数组间的异步内存复制。除此之外，还有一个点对点内存复制函数集合来支持GPU间的内存复制。

CUDA运行时和驱动程序API使用不同的方式实现内存复制。对1D内存复制，驱动程序API定义了一个使用强类型参数的函数集。主机到设备、设备到主机和设备到设备内存复制函数是相互独立的。

CResult cuMemcpyHtoD(CUdeviceptr dstDevice, const void *srcHost, size_t ByteCount);  
CResult cuMemcpyDtoH(void *dstHost, CUdeviceptr srcDevice, size_t ByteCount);  
CResult cuMemcpyDtoD(CUdeviceptr dstDevice, CUdeviceptr srcDevice, size_t ByteCount);

与此不同，CUDA运行时倾向于定义接受一个额外的“内存复制类型”参数，这个参数取决于主机内存类型和复制目的地指针。

enumuveaMemcpyKind   
{ CUDAMemcpyHostToHost  $= 0$  -   
cudamemcpyHostToDevice  $\equiv$  1,   
cudamemcpyDeviceToHost  $= 2$  -   
cudamemcpyDeviceToDevice  $= 3$  -   
cudamemcpyDefault  $= 4$    
}；

对于更加复杂的内存复制操作，2种API都使用描述符结构体来指定内存复制参数。

5.7.1 同步内存复制与异步内存复制

因为内存复制的大多数方面（维度、内存类型）是独立于内存复制是异步还是同步的，本节会详细检视其中的不同，下一节会简要介绍异步内存复制的相关内容。

默认的，任何涉及主机内存的复制是同步的：函数直到所有的操作执行完毕，否则不会返回控制权给调用者 [1]。即使是在锁页内存（例如，由函数CUDAMallocHost()分配的内存）上操作，同步内存复制程序必须等待这一操作全部完成，因为应用软件会依赖于这一内存行为 [2]。

出于性能原因，在可能的情况下，我们需要避免同步内存复制。即使在未使用流的情况下，使所有的操作异步的进行会使CPU和GPU并发执行，会提升程序的性能。如果没有其他的操作同时执行，在GPU运行的同时，CPU可以启动更多的GPU操作，像内核启动和其他的内存复制。如果CPU/GPU的并发是我们唯一的目标，我们就没有必要去创建任何的CUDA流，使用NULL流来调用一次异步内存复制就足够了。

虽然涉及主机内存的内存复制默认是同步的，但任何不涉及主机内存的复制（设备与设备、设备与数组之间）全部都是异步的。GPU硬件在内部让这些操作顺序进行，所以对函数来说，没有必要等待GPU结束所有的操作再返回。

异步内存复制函数都用Async()作为后缀。例如，异步主机到设备的内存复制在驱动程序API中的函数是cuMemcpyHtoDAsync()，在CUDA运行时中是CUDAMemcpyAsync()。

实现异步内存复制的硬件随时间更迭已经更新了很多代。最初的支持CUDA的GPU（GeForce 8800 GTX）不含有任何的复制引擎，所以异步内存复制只能使CPU/GPU并发。之后的GPU添加了复制引擎，在SM运行时可以执行1D传输。此后，全功能的复制引擎被添加到GPU中，可以加速2D和3D传输，即使内存复制中包含了等步长布局与CUDA数组中使用的块线性布局的转换。除此之外，早期的CUDA硬件只有一个内存复制引擎，然而在今天，有些时候硬件中有两个复制引擎。而超过两个的复制引擎不会真正起作用。因为单个复制引擎可以使PCIe总线单向达到饱和，为了最大化双向总线性能以及GPU计算和总线传输间的并发，只需要2个复制引擎。

复制引擎的数量可以通过调用函数cuDeviceGetAttribute()并使用标志CU_DEVICE_ATTRIBUTE AsyncENGINE_COUNT查询，或者也可以检查CUDADeviceProp::asyncEngineCount。

5.7.2 统一虚拟寻址

统一虚拟寻址使CUDA可以对地址范围上的内存类型进行推断。因为CUDA寻址范围中包含了设备地址和主机地址，所以在函数CUDAMemcpy()中没有必要指定参数CUDAMemcpyKind。驱动程序API中添加了函数cuMemcpy()，此函数同样可以通过地址推断内存类型。

Cuiresult cuMemcpy(CUdeviceptr dst, CUdeviceptr src, size_t ByteCount);

CUDA运行时中相对应的函数调用被命名为CUDAMemcpy:

cudaError_t cudaMemcpy(void *dst, const void *src, size_t bytes);

5.7.3 CUDA运行时

表5-10总结了CUDA运行时中可用的内存复制函数。

表5-10 内存复制函数 (CUDA运行时)

（续）

1D和2D内存复制函数使用基指针、步长和所需的大小。3D内存复制操作接受一个描述结构体CUDAMemcpy3Dparms，定义如下：

structudaMemcpy3DParams
{
    structudaArray *srcArray;
    structudaPos srcPos;
    structudaPitchedPtr srcPtr;
    structudaArray *dstArray;
    structudaPos dstPos;
    structudaPitchedPtr dstPtr;
    structudaExtent extent;
    enumudaMemcpyKind kind;
};

表5-11总结了CUDAMemcpy3DParams结构体中的每一个成员。CUDAPos和CUDAExtent结构体中的定义如下。

表5-11udaMemcpy3DParams结构体成员

structudaExtent {
    size_t width;
    size_t height;
    size_t depth;
};

5.7.4 驱动程序API

表5- 12总结了驱动程序API中的内存复制函数。

表5-12 内存复制函数（驱动程序API）

函数cuMemcpy3D()设计来实现以往所有内存复制函数的超集。任何的1D、2D或3D内存复制可以在主机、设备或CUDA数组，与复制源和复制目标的任何偏移之间使用。输入结构体CUDA_MEMCDY_3D的成员WidthInBytes、Height和Depth，定义了内存复制的维度：Heigh==0意味着1D内存复制，Depth==0意味着一次2D复制。复制源和复制目标内存类型分别在结构体srcMemoryType和dstMemoryType中给出。

如果不需要，函数cuMemcpy3D()中结构体成员可以被忽略。例如，如果请求1D主机到设备内存复制，则srcPitch、srcHeight、dstPitch和dstHeight全部被忽略；如果srcMemoryType的值是CU_MEMORYTYPE_HOST，srcDevice和srcArray均被忽略。API使用C语言中的惯例用法，为结构体赋值为{0}进行初始化，这也使内存复制被描述的十分简明。绝大多数的其他内存复制函数可以用几行代码实现，例如下面几行代码：

CResult   
my_cuMemcpyHtoD(CUdevice dst, const void \*src,size_t N)   
{ CUDA_MEMCPY_3D cp  $=$  {0}; cp.srcMemoryType  $=$  CU_MEMORYTYPE_HOST; cp.srcHost  $=$  srcHost; cp.dstMemoryType  $=$  CU_MEMORYTYPE_DEVICE; cp.dstDevice  $=$  dst; cp.WidthInBytes  $= \mathbb{N}$  . return cuMemcpy3D(&cp);

[1] 这是因为硬件不能直接访问主机内存，除非它被页锁定并且映射给GPU。对可换页内存的一次异步内存复制可以通过生成另一个CPU线程实现，但是至今为止，CUDA团队选择了回避这一额外的复杂问题。
[2] 当锁页内存被指定到一个同步内存复制操作中，驱动程序会让硬件使用DMA，这通常会提升速度。

Previous5.6_共享内存
Next06_第6章_流与事件