关于CUDA零拷贝内存的实验

零拷贝介绍和主要函数

通常来说，之不能直接访问设备变量，同时设备也不能直接访问主机变量。但是有一个例外：另拷贝内存。主机和设备都可以访问零拷贝内存。

GPU线程可以直接访问零拷贝内存。在CUDA核函数中使用另拷贝内存有以下优势：

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

当使用零拷贝内存来共享主机和设备间的数据时，必须同步主机和设备间的内存访问，同时更改主机和设备的零拷贝内存中的数据将导致不可预料的后果

零拷贝内存是固定（不可分页）内存，该内存映射到设备地址空间中。使用下函数创建一个到固定内存的映射（和锁页内存相似）：

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

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

• cudaHostAllocDefault：使cudaHostAlloc函数的行为与cudaMallocHost函数一致
• cudaHostAllocPortable：可以返回能被CUDA上下文所使用的固定内存，而不仅是执行内存分配的那一个
• cudaHostAllocWriteCombined：返回写结合内存，该内存可以再某些系统配置上通过PCIe总线上更快地传输，但是他在大多数主机上不能被有效地读取。因此，写结合内存对缓冲区来说是一个很好的选择，该内存通过设备使用映射的固定内存或主机到设备的传输。
• cudaHostAllocMapped：另拷贝内存，该标志返回可以实现主机写入和设备读取被映射到设备地址空间的主机内存。

分配好零拷贝内存后，就可以使用下列函数获取映射到固定内存的设备指针了：

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

该函数返回一个在pDevice中的设备指针，该指针可以在设备上被引用以访问映射得到的固定主机内存。如果设备不支持映射得到的固定内存，则该函数失效。flag保留，始终为0。

在进行频繁的读写操作时，使用另拷贝内存作为设备内存的补充将显著降低性能。因为每一次映射到内存的传输必须经过PCIe总线。与全局内存相比，延迟也显著增加（集成架构在后面讨论）。

在使用零拷贝内存时，需要检查设备是否支持固定内存映射

cudaDevieProp的canMapHostMemory成员是一个bool类型值，true表示支持固定内存映射

以下是主要的代码，包含全局内存计算和零拷贝内存计算两个大块。

#include "arrayTools.h"#include "cudaCode.h"#include "MyTimer.h"template<class T>__global__ void warmup(T *A, T *B, T *C, const int n){int tid = blockDim.x*blockIdx.x + threadIdx.x;if (tid >= n)return;C[tid] = A[tid] + B[tid];}template<class T>__global__ void sumVecOnDeviceZeroCopy(T *A, T *B, T *C, const int n){int tid = blockDim.x*blockIdx.x + threadIdx.x;if (tid >= n)return;C[tid] = A[tid] + B[tid];}int main(int argc, char *argv[]){int dev = 0;CHECK(cudaSetDevice(dev));cudaDeviceProp deviceProp;CHECK(cudaGetDeviceProperties(&deviceProp, dev));if (!deviceProp.canMapHostMemory){printf("Device %d(%s) does not support mapping CPU host memory!\n", dev, deviceProp.name);CHECK(cudaDeviceReset());exit(EXIT_SUCCESS);}printf("Useing Device %d: %s ", dev, deviceProp.name);int ipower = 16;if (argc > 1) ipower = atoi(argv[1]);int nElem = 1 << ipower;int nBytes = nElem * sizeof(float);if (ipower >= 18)printf("nElem=%f MB\n", nElem / (1024.0*1024.0));elseprintf("nElem=%f KB\n", nElem / 1024.0);float *h_A = (float*)malloc(nBytes);float *h_B = (float*)malloc(nBytes);float *cpuRes = (float*)malloc(nBytes);float *gpuRes = (float*)malloc(nBytes);initialInt(h_A, nElem);initialInt(h_B, nElem);sumVecOnHost<float>(h_A, h_B, cpuRes, nElem);dim3 block(512, 1);dim3 grid((nElem + block.x - 1) / block.x);float *d_A, *d_B, *d_C;CHECK(cudaMalloc((void**)&d_A, nBytes));CHECK(cudaMalloc((void**)&d_B, nBytes));CHECK(cudaMalloc((void**)&d_C, nBytes));warmup<float> << <grid, block >> > (d_A, d_B, d_C, nElem);CHECK(cudaDeviceSynchronize());Timer t1;t1.start();CHECK(cudaMemcpy(d_A, h_A, nBytes, cudaMemcpyHostToDevice));CHECK(cudaMemcpy(d_B, h_B, nBytes, cudaMemcpyHostToDevice));Timer t;t.start();sumVecOnDevice<float> << <grid, block >> > (d_A, d_B, d_C, nElem);CHECK(cudaDeviceSynchronize());t.stop();CHECK(cudaMemcpy(gpuRes, d_C, nBytes, cudaMemcpyDeviceToHost));t1.stop();printf("sumVecOnDevice spend %lf ms(kernel+copy), %lf(kernel), <<<%d, %d>>>\n", t1.getElapsedTimeInMilliSec(), t.getElapsedTimeInMilliSec(), grid.x, block.x);compareVec(gpuRes, cpuRes, 0, nElem);CHECK(cudaFree(d_A));CHECK(cudaFree(d_B));//cudaFree(d_C);free(h_A);free(h_B);free(gpuRes);//申请固定内存unsigned int flags = cudaHostAllocMapped;CHECK(cudaHostAlloc((void**)&h_A, nBytes, flags));CHECK(cudaHostAlloc((void**)&h_B, nBytes, flags));CHECK(cudaHostAlloc((void**)&gpuRes, nBytes, flags));//初始化输入数据initialInt(h_A, nElem);initialInt(h_B, nElem);//获取零拷贝内存的设备指针CHECK(cudaHostGetDevicePointer((void**)&d_A, h_A, 0));CHECK(cudaHostGetDevicePointer((void**)&d_B, h_B, 0));CHECK(cudaHostGetDevicePointer((void**)&d_C, gpuRes, 0));//主机端计算sumVecOnHost<float>(h_A, h_B, cpuRes, nElem);t1.start();//零拷贝内存计算（与设备端计算看起来没有不同，只需要注意传入指针即可）sumVecOnDeviceZeroCopy<float> << <grid, block >> > (d_A, d_B, d_C, nElem);CHECK(cudaDeviceSynchronize());t1.stop();printf("sumVecOnDeviceZeroCopy spend %lf ms, <<<%d, %d>>>\n", t1.getElapsedTimeInMilliSec(), grid.x, block.x);compareVec<float>(cpuRes, gpuRes, 0, nElem);//释放固定内存CHECK(cudaFreeHost(h_A));CHECK(cudaFreeHost(h_B));CHECK(cudaFreeHost(gpuRes));free(cpuRes);CHECK(cudaDeviceReset());system("pause");return 0;}

实验测试

在TX1开发板（集成架构，即cpu和gpu集成在一个芯片上，并且在物理地址上共享内存）上的测试结果如下所示：

1、使用简单的矢量加法运算

template<class T>__global__ void sumVecOnDeviceZeroCopy(T *A, T *B, T *C, const int n){int tid = blockDim.x*blockIdx.x + threadIdx.x;if (tid >= n)return;C[tid] = A[tid] + B[tid];}

2、使用跨块加法运算

template<class T>__global__ void sumVecOnDeviceZeroCopy(T *A, T *B, T *C, const int n){int tid = blockDim.x*blockIdx.x + threadIdx.x;if (tid >= n-blockDim.x*3)return;C[tid] = A[tid] + A[tid+blockDim.x*1] + A[tid+blockDim.x*2] + B[tid] + B[tid+blockDim.x*1] + B[tid+blockDim.x*2];// if (tid >= n)return;// C[tid] = A[tid] + B[tid];} 

3、使用滤波式加法运算

template<class T>__global__ void sumVecOnDeviceZeroCopy(T *A, T *B, T *C, const int n){int tid = blockDim.x*blockIdx.x + threadIdx.x;if (tid >= n-1 && tid<=0)return;C[tid] = A[tid] + A[tid+1] + A[tid-1] + B[tid] + B[tid+1] + B[tid-1];// if (tid >= n)return;// C[tid] = A[tid] + B[tid];}

可以看出在第1和3中种情况下零拷贝内存和全局内存的纯计算时间在内存量不大的情况下几乎没有区别，而包含了数据拷贝在内的时间后，零拷贝内存具有极大地优势。在内存使用量特别大的时候，零拷贝内存的计算时间表现出了劣势，但在整体上仍具有优势。

而在跨块加法运算时，零拷贝内存的纯计算时间比全局内存大了许多，但考虑整体时，零拷贝内存仍有不可替代的优势。
离散架构下的测试没有做，因为只在内存极小的时候有意义，且跟PCIe传输速度和主机设备的性能关系都比较大，而且提升不大，只是会编程方便，可以在几十字节几百字节时考虑使用。

集成架构和离散架构

这两种架构是常见的异构计算系统架构。

集成架构：cpu和gpu集成在一个芯片上，并且在物理内存上共享内存。在这种架构中，由于无需在PCIe总线上备份，所以零拷贝内存在性能和可编程性方面可能更好一些。

离散架构：gpu和gpu是分离的，物理内存上也是分离的。数据需要通过PCIe总线进行交互，数据拷贝的耗时和延迟通常都是代价比较大的操作。因此这种架构下，另拷贝内存只在特殊情况下有优势（例如内存量较小的时候，数据复制函数的开销可能快赶上数据处理的开销了）

由于通过映射的固定内存在主机和设备之间是共享的，所以必须同步内存访问来避免任何潜在的数据冲突，这种数据冲突一般是有多线程异步访问相同的内存而引起的。

注意：

零拷贝内存还有另一种实现方式，就是直接对malloc申请的内存使用cudaHostRegister进行标志位设置（设置成固定内存），然后使用cudaHostGetDevicePointer获取其设备指针

代码如下：

#include "arrayTools.h"​#include "cudaCode.h"#include "MyTimer.h"​template<class T>__global__ void warmup(T *A, T *B, T *C, const int n){    int tid = blockDim.x*blockIdx.x + threadIdx.x;    if (tid >= n)return;    C[tid] = A[tid] + B[tid];}​template<class T>__global__ void sumVecOnDeviceZeroCopy(T *A, T *B, T *C, const int n){    int tid = blockDim.x*blockIdx.x + threadIdx.x;    if (tid >= n)return;    C[tid] = A[tid] + B[tid];}​template<class T>__global__ void sumVecOnDeviceZeroCopyRegister(T *A, T *B, T *C, const int n){    int tid = blockDim.x*blockIdx.x + threadIdx.x;    if (tid >= n)return;    C[tid] = A[tid] + B[tid];}int main(){    int nElem=1<<20;    int nBytes=nElem*sizeof(float);    float * d_A, * d_B, * d_C;    float * h_A, * h_B, * gpuRes;    h_A=(float*)malloc(nBytes);    h_B=(float*)malloc(nBytes);    initialInt(h_A, nElem);    initialInt(h_B, nElem);    gpuRes=(float*)malloc(nBytes);        dim3 block(512, 1);    dim3 grid((nElem + block.x - 1) / block.x);​    //将h_A、h_B、h_C三块内存锁定并获取对应的设备地址，进行计算，这样避免了从一开始就得cudaHostAlloc申请零拷贝内存的不便    CHECK(cudaHostRegister(h_A, nBytes, cudaHostRegisterMapped));    CHECK(cudaHostGetDevicePointer((void **)&d_A, h_A, 0));    CHECK(cudaHostRegister(h_B, nBytes, cudaHostRegisterMapped));    CHECK(cudaHostGetDevicePointer((void **)&d_B, h_B, 0));    CHECK(cudaHostRegister(gpuRes, nBytes, cudaHostRegisterMapped));    CHECK(cudaHostGetDevicePointer((void **)&d_C, gpuRes, 0));    memset(gpuRes, 0, nBytes);    sumVecOnDeviceZeroCopyRegister<float> << <grid, block >> > (d_A, d_B, d_C, nElem);    CHECK(cudaDeviceSynchronize());        compareVec(gpuRes, cpuRes, 0, nElem);    return 0;}

实践发现：

1、在Windows7+CUDA8.0+GTX1050（Pascal）上测试结果与直接用cudaHostAlloc申请的固定内存运行速度基本无二（奇怪的是好像用vs插件Performance Analysis无法分析这块操作），NVVP（NVIDIA Visual Profiler）能够分析

2、在Ubuntu16.04+CUDA8.0+TX1（Maxwell）上测试结果却是无法实现错误为

Error: test.cu: 47, code: 71, reason: operation not supportedsh: 1: pause: not foundError: test.cu: 48, code: 11, reason: invalid argumentsh: 1: pause: not foundError: test.cu: 49, code: 71, reason: operation not supportedsh: 1: pause: not foundError: test.cu: 50, code: 11, reason: invalid argumentsh: 1: pause: not found

其中code后面是错误码，reason后面是错误解释。看了两眼没看懂具体意思，去查了下cudaHostRegister和cudaHostGetDevicePointer的返回值分别为

 * \return * ::cudaSuccess, * ::cudaErrorInvalidValue, * ::cudaErrorMemoryAllocation, * ::cudaErrorHostMemoryAlreadyRegistered * \notefnerr

和

 * \return * ::cudaSuccess, * ::cudaErrorInvalidValue, * ::cudaErrorMemoryAllocation * \notefnerr

再查错误码的头文件为driver_types.h，查到如下两个错误码为71和11的解释

/** * This error indicates the attempted operation is not supported * on the current system or device. */cudaErrorNotSupported                 =     71,​/** * This indicates that one or more of the parameters passed to the API call * is not within an acceptable range of values. */cudaErrorInvalidValue                 =     11,

看下注释就明白了，cudaHostRegister调用造成了cudaErrorNotSupported错误，原因是这个系统或设备上不支这个操作。造成了获取到的内存未设置成固定内存，而在后面cudaHostGetDevicePointer又用到了，输入不对，造成了cudaErrorInvalidValue错误，内存不在合理的可接受范围。

疑问1：cudaHostRegister的return中没有写有cudaErrorNotSupported返回值，却返回了这个，为什么？

疑问2：Ubuntu16.04+CUDA8.0+TX1（Maxwell）不支持cudaHostRegister，究竟是那个system或者device不支持呢？

猜想1：估计notefnerr表示可能还有其他错误吧。。。

猜想2：Linux系统一般应该是支持页锁定的，将已分配的内存设为页锁定这类操作应该是不存在难度的。而CUDA8.0显然不是原因，那么有理由相信cudaHostRegister这个操作是Pascal支持的操作而Maxwell不支持。手头没有别的卡了，就不验证了。当然也不排除是通过别的代码设置可以解决这个问题的可能。

对了，查阅资料过程中，发现有个人遇到类似的问题，是关于统一内存的问题（看来很多操作都是和架构有关的或者操作系统有关的，以后需要多注意）

cudaMallocManaged() returns “operation not supported”

但并不是一个问题

https://stackoverflow.com/questions/23600403/cudamallocmanaged-returns-operation-not-supported

帖子中说道，这类情况应该是统一内存的问题，应符合3个条件才能使用。

J.1.4. System RequirementsUnified Memory has three basic requirements:•a GPU with SM architecture 3.0 or higher (Kepler class or newer)•a 64-bit host application and operating system, except on Android•Linux or Windows

官方文档如是说：

K.1.1. System Requirements

Unified Memory has two basic requirements:

• a GPU with SM architecture 3.0 or higher (Kepler class or newer)
• a 64-bit host application and non-embedded operating system (Linux, Windows, macOS)

Read more at: http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#ixzz4yetZxUH3 

Follow us: @GPUComputing on Twitter | NVIDIA on Facebook

这里有些关于固定内存和统一内存的资料，引用下

http://blog.csdn.net/LIYUAN123ZHOUHUI/article/details/53672656

附一张存储器带宽图，非官方，可能不准，且每种不同设备不一样，仅做大致参考