cuda的统一内存寻址 cudaMalloc VS cudaMemcpy

 大家知道，Nvidia公司发布的CUDA6开发包拥有一个新特性，就是“统一内存寻址”，那究竟统一内存寻址有什么特殊的地方呢？我们编写CUDA代码跟以前有什么区别呢？现在拥有的GPU架构能够很好的支持吗？我们带着这些问题开始我们的话题。
从名字上看，统一内存寻址就是将CPU端的内存同GPU显存统一起来，使得程序猿在编写代码的时候不用明显的使用诸如cudaMalloc或者cudaMemcpy等操作显存的指令，而能够在Kernel函数中直接使用定义的变量。
    例子如下：
原始代码（CUDA6.0文档）

C/C++ code

?

1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22

__global__  void  AplusB( int  *ret,  int  a,  int  b) {    ret[threadIdx.x] = a + b + threadIdx.x; }   int  main() {    int  *ret;    //**************************************    cudaMalloc(&ret, 1000 *  sizeof ( int ));    AplusB<<<1, 1000>>>(ret, 10, 100);    //**************************************    int  *host_ret = ( int  *) malloc (1000 *  sizeof ( int ));    cudaMemcpy(host_ret, ret, 1000 *  sizeof ( int ), cudaMemcpyDefault);      for ( int  i = 0; i < 1000; i++)      printf ( "%d: A + B = %d\n" , i, host_ret[i]);      free (host_ret);    cudaFree(ret);    return  0; }

统一寻址的代码：

C/C++ code

?

1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19

__global__  void  AplusB( int  *ret,  int  a,  int  b) {    ret[threadIdx.x] = a + b + threadIdx.x; }   int  main() {    int  *ret;    //***********************************************    cudaMallocManaged(&ret, 1000 *  sizeof ( int ));    AplusB<<<1, 1000>>>(ret, 10, 100);    //***********************************************    cudaDeviceSynchronize();    for ( int  i = 0; i < 1000; i++)      printf ( "%d: A + B = %d\n" , i, ret[i]);      cudaFree(ret);    return  0; }

    从上面不同的代码可以看出，统一寻址后的代码更简洁，使用了函数cudaMallocManaged()开辟一块存储空间，无论是在Kernel函数中还是main函数中，都可以使用这块内存，达到了统一寻址的目的。
    这里有一个需要注意的地方就是main函数在调用Kernel函数之后，使用了一个同步函数。仔细思考后就会有所领悟——既然这块存储空间既可以被Kernel函数访问，也可以被main函数访问，为了解决访问冲突的问题，因此使用了同步函数，使得在Kernel改变变量的值后，main函数才能使用该变量。
    上面只是简单的举了个例子，但已经把统一寻址的问题描述的很清楚：
    1. 开辟空间所使用的新函数；
    2. 变量访问同步。
    在CUDA6.0的文档中更加仔细的描述了相关的问题，包括流和多GPU条件下使用统一寻址的问题等，其主要针对的还是变量访问同步，大家有兴趣可以参看文档。
    同AMD的APU不同，即使是Kepler架构的GPU，也不是真正意义上从硬件手段实现“统一内存寻址”，只是从软件上解决的一种方式，Nvidia公司的新一代GPU产品——Maxwell架构，才真正是从硬件层面实现这一特性。而速度，还有待测试与研究~~
  既然“统一内存寻址”是GPU产品所前进的一个方向，势必会改变我们的编程方式，因此，这一话题还是有重大的意义的，希望大家热烈讨论，在讨论中真正掌握这一技术！