CUDA Memory Copy

CUDA memcpy

CUDA Memory Copy

CUDA内存拷贝

在CUDA程序的内存数据拷贝中包含以下几种情况：Host2Device、Device2Host和Device2Device（Host2Host即正常程序中的copy）。在刚开始编写CUDA程序时对cudaArray和其他DeviceArray的copy存在较大的疑惑，在此对CUDA memcpy进行一个尽量详细的记录。

cudaArray以及绑定纹理对象、表面对象的拷贝

以2D图像数组为例进行说明。

Host2Device

cudaMemcpy2DToArray(
	cudaArray_t dst,
	size_t wOffset,
	size_t hOffset,
	const void* src,
	size_t spitch,
	size_t width,
	size_t height,
	cudaMemcpyKind kind
);
cudaMemcpy2DToArray(
	_cuda_array, 0, 0,
	_host.data,
	_host.cols * _host.elemSize(),
	_host_cols * _host.elemSize(),
	_host.rows,
	cudaMemcpyHostToDevice
);

使用函数cudaMallocPitch和cudaMemcpy2D来使用二维数组。C/C++中二维数组内存分配是转化为一维数组，连贯紧凑，每次访问数组中的元素都必须从数组首元素开始遍历；而CUDA中分配的二维数组内存保证数组每一行首元素的地址值都按照 256 或 512 的倍数对齐，提高访问效率，但使得每行末尾元素与下一行首元素地址可能不连贯，使用指针寻址时要注意考虑尾部。

cudaMallocPitch传入存储器指针**devPtr，偏移值的指针*pitch，数组行字节数widthByte，数组行数height。函数返回后指针指向分配的内存（每行地址对齐到 AlignByte 字节，为 256B 或 512B），偏移值指针指向的值为该行实际字节数=sizeof(datatype) * width + alignByte - 1) / alignByte）。

cudaMemcpy2D传入目标存储器的指针*dst，目标存储器行字节数dpitch，源存储器指针*src，源存储器行字节数spitch，数组行字节数widthByte，数组行数height，拷贝方向 kind。这里要求存储器行字节数不小于数组行字节数，多出来的部分就是每行尾部空白部分。

我们在使用cudaMemcpy2DToArray时，第一个存放的为cudaArray_t的参数，个人理解为指向GPU中某一个内存地址的指针，但采用cudaArray_t的格式对其进行封装。在利用cv::Mat对二维图像进行host到device的拷贝时，一般pitch的大小和行字节数相等。

对pitch的理解，个人认为和C++中的内存对齐类似，提高了读写速率。

Device2Host

cudaMemcpy2DFromArray(
	void* dst,
	size_t dpitch,
	cudaArray_const_t src,
	size_t wOffset,
	size_t hOffset,
	size_t width,
	size_t height,
	cudaMemcpyKind kind
);
cudaMemcpy2DFromArray(
	_host.data,
	_host.cols * _host.elemSize(),
	_cuda_array,
	0, 0,
	_host.cols * _host.elemSize(),
	_host.rows,
	cudaMemcpyDeviceToHost
);

Device到Host拷贝和Host拷贝到Device基本一致，仅仅是参数传入的位置进行了改变。

Device2Device

一般也利用cudaMemcpy2DToArray的方式，将host的data地址改成GPU中指向某个分配的数据块指针即可。

cudaTexture_t

当我们获取到cudaTexture_t的句柄时，但并不能直接获得cudaArray_t的句柄，可以使用desc的方式对cudaArray_t数据进行获取。

cudaResourceDesc desc;
cudaGetTextureObjectResourceDesc(&desc, _cuda_texobj);
cudaMemcpy2DFromArray(
	_host.data,
	_host.cols * _host.elemSize(),
	desc.res.array.array,
	0, 0,
	_host.cols * _host.elemSize(),
	_host.rows,
	cudaMemcpyDeviceToHost
);

使用这种方式就能够不去获取cudaArray_t的句柄，通过cudaTextureObject_t就可以直接获取到CUDA上存储的数组并下载到Host上。

cudaArray和cuda对象的拷贝

cudaArray_t和CUDA malloc对象之间的拷贝，也可以使用cudaMemcpy2DFromArray、cudaMemcpy2DToArray的方式来实现，同样在编写好对应指针地址、pitch大小等参数后，利用cudaMemcpyDeviceToDevice进行拷贝。

通常，拷贝采用Async的异步拷贝方式进行，并采用cudaStreamSynchronize的方式完成同步。这是因为CUDA指令基本都是异步进行，程序不会等待CUDA调用结束再进行下一步操作。但拷贝的指令若不实用Async异步方式则会让CPU进行等待，造成最终程序时长过大。

Reference

Memory Alignment For CUDA - Fang’s Notebook