CUDA Unified Memory统一内存使用注意

  CUDA 6.0推出了统一内存寻址（Unified Memory）的方式，可以大大简化CUDA程序的编写。在这之前，CUDA程序的写作方式一般是这样的：

[cpp]  view plain  copy

1. float *h_a; //主机内存指针  
2. float *d_a; //设备内存指针  
3.   
4. cudaMallocHost  为h_a分配内存  
5. cudaMalloc      为d_a分配内存  
6.   
7. 主机对h_a内容进行填充  
8.   
9. cudaMemcpy h_a 到 d_a  
10. kernel <<<M,N,0,stream0>>>(d_a); //设备对数据进行处理  
11. cudaMemcpy d_a 到 h_a  
12.   
13. 主机对h_a结果进行使用  
14.   
15. cudaFreeHost(h_a);  
16. cudaFree(d_a);  

        这主要有两个问题：需要为主机和设备（GPU）分别定义指针和分布malloc内存，然后在使用前后需要在CPU和GPU进行显式的拷贝，代码冗长，维护难度较大。

Unified Memory大大简化了CUDA程序的编写，使用Unified Memory的CUDA编程方式如下：

[cpp]  view plain  copy

1. float *a;  
2. cudaMallocManaged 为a分配内存  
3.   
4. 主机对a内容进行填充  
5. kernel <<<M,N,0,stream0>>>(a); //设备对数据进行处理  
6. cudaDeviceSynchronize(); //等待GPU执行完成， 有多种方式  
7. 主机对a结果进行使用  
8. cudaFree(a);  

        可见，一方面，只需要定义一个指针即可在主机和设备端通用；另一方面，去掉了显式的内存拷贝，大大缩短了代码数量和编写、维护难度。

在CUDA编程使用Unified Memory的时候有几个注意点：

Unified Memory has three basic requirements:

1, a GPU with SM architecture 3.0 or higher (Kepler class or newer)
2, a 64-bit host application and operating system, except on Android
3, Linux or Windows

        这个限制也会对用户的应用场合进行一些限制，一方面是只能是程序只能在Kepler及以后的显卡才能运行，另一方面必须工作在64位的系统上。考虑到Kepler的显卡已经普及了三四年了，大部分gtx6xx系列的显卡已经就是Kepler架构，现在64位系统也已经非常普遍，这些限制应该还可以接受。

附：

1，cuda的release比debug快的多

2，使用流和cudaMemcpyAsync 的性能比使用unified memory的性能要高，我的一个程序中前者比后者高10%。（An important point is that a carefully tuned CUDA program that uses streams and cudaMemcpyAsync to efficiently overlap execution with data transfers may very well perform better than a CUDA program that only uses Unified Memory. Understandably so: the CUDA runtime never has as much information as the programmer does about where data is needed and when! CUDA programmers still have access to explicit device memory allocation and asynchronous memory copies to optimize data management and CPU-GPU concurrency. ）

参考：

https://en.wikipedia.org/wiki/CUDA

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

https://devblogs.nvidia.com/parallelforall/unified-memory-in-cuda-6/

转载来自：http://blog.csdn.net/u013701860/article/details/51140762