关于CUDA 6.0统一寻址的讨论

[_梦魇花葬]

大家好！我是CSDN CUDA版块新任版主！希望能够尽自己的力量帮助大家解决有关CUDA编程的问题，并且会定期发布一些问题同大家一起讨论，共同学习，共同进步！大家有问题可以发帖留言也可以发邮件到我的邮箱[e-mail]。
在大家的共同努力下，祝愿CSDN的CUDA版块能够越办越好！
下面进入此次讨论的主题！
———————————————分割线————————————————————
大家知道，Nvidia公司发布的CUDA6开发包拥有一个新特性，就是“统一内存寻址”，那究竟统一内存寻址有什么特殊的地方呢？我们编写CUDA代码跟以前有什么区别呢？现在拥有的GPU架构能够很好的支持吗？我们带着这些问题开始我们的话题。
从名字上看，统一内存寻址就是将CPU端的内存同GPU显存统一起来，使得程序猿在编写代码的时候不用明显的使用诸如cudaMalloc或者cudaMemcpy等操作显存的指令，而能够在Kernel函数中直接使用定义的变量。
例子如下：
原始代码（CUDA6.0文档）

__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;
}

统一寻址的代码：

__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产品所前进的一个方向，势必会改变我们的编程方式，因此，这一话题还是有重大的意义的，希望大家热烈讨论，在讨论中真正掌握这一技术！
谢谢大家~~！

[PerfectHacker]

未来的热点，先点再收藏

[大_猫]

我在大学的时候做过CUDA编程。记得最深的就头像的3D建模、渲染。用CPU跑能把电脑卡死。 采用CUDA 跑得飞起。感觉爽歪歪。当时学CUDA，学校的课都不听，拿本书自己在那里啃。 经常在实验室编程到深夜。一直觉得这东西高大上。 出来工作了就一直没碰过CUDA了。唉，都忘得差不多了。

[_梦魇花葬]

https://developer.nvidia.com/cuda-gpus这个网址可以查看支持CUDA的GPU GTX 520是可以支持CUDA的 ~

[木子超同学]

这个对显卡有什么要求没？是特定的显卡才能使用，还是都可以，比如gtx 520.

[nadleeh]

引用 6 楼 sparrow986831 的回复:

感谢版主们支持~共同讨论 共同进步~~嘿嘿

楼主把你的核弹借我玩几天行么

[fu3210]

cuda是什么呀 求科普

[_梦魇花葬]

引用 15 楼 u012648302 的回复:

cuda是什么呀 求科普

语出维基百科： CUDA（Compute Unified Device Architecture，统一计算架构[1]）是由NVIDIA所推出的一种集成技术，是该公司对于GPGPU的正式名称。通过这个技术，用户可利用NVIDIA的GeForce 8以后的GPU和较新的Quadro GPU进行计算。亦是首次可以利用GPU作为C-编译器的开发环境。NVIDIA营销的时候[2]，往往将编译器与架构混合推广，造成混乱。实际上，CUDA可以兼容OpenCL或者自家的C-编译器。

[psyangqi]

UVA的性能怎么样？ 楼主讲从软件 来实现统一寻址 的意义是：物理上 还是在host端device端都申请了存储 并且 copy 如果 每次只访问一个变量 而不是一个数组 ，访问效率怎么样？

[木子超同学]

引用 14 楼 sparrow986831 的回复:

https://developer.nvidia.com/cuda-gpus这个网址可以查看支持CUDA的GPU GTX 520是可以支持CUDA的 ~

我是想问cuda 6.0的统一寻址对显卡有什么要求，是否要求最新的显卡才可以？

[bob76012]

我还在用5.5的路过 不知有对资料转移加速吗? 因为我在处理数值分析会常用CPU和GPU协同处理 在记忆体拷贝过去在拷贝回来 数据一大很耗时

[meatball1982]

这个不错，mark一下，需要好好学习。

[_梦魇花葬]

引用 20 楼 bob76012 的回复:

我还在用5.5的路过 不知有对资料转移加速吗? 因为我在处理数值分析会常用CPU和GPU协同处理 在记忆体拷贝过去在拷贝回来 数据一大很耗时

现在的GPU即使是使用了内存同意寻址，也只是从软件层面加速，省去了程序员的工作，而实际的工作方式还是要将数据从CPU传递到GPU，期待Nvidia下一代GPU吧，也许会从硬件层面有所突破~~是数据传输的性能提高~~

[Sedna1]

引用 22 楼 sparrow986831 的回复:

[quote=引用 20 楼 bob76012 的回复:] 我还在用5.5的路过 不知有对资料转移加速吗? 因为我在处理数值分析会常用CPU和GPU协同处理 在记忆体拷贝过去在拷贝回来 数据一大很耗时

现在的GPU即使是使用了内存同意寻址，也只是从软件层面加速，省去了程序员的工作，而实际的工作方式还是要将数据从CPU传递到GPU，期待Nvidia下一代GPU吧，也许会从硬件层面有所突破~~是数据传输的性能提高~~[/quote]谢谢了，对我们还在使kepler架构的有一定指导借鉴。

[_梦魇花葬]

引用 19 楼 llcchh012 的回复:

[quote=引用 14 楼 sparrow986831 的回复:] https://developer.nvidia.com/cuda-gpus这个网址可以查看支持CUDA的GPU GTX 520是可以支持CUDA的 ~

我是想问cuda 6.0的统一寻址对显卡有什么要求，是否要求最新的显卡才可以？[/quote] 要求计算能力在3.0以上的GPU才可以支持~

[_梦魇花葬]

引用 26 楼 bob76012 的回复:

[quote=引用 25 楼 sparrow986831 的回复:] [quote=引用 19 楼 llcchh012 的回复:] [quote=引用 14 楼 sparrow986831 的回复:] https://developer.nvidia.com/cuda-gpus这个网址可以查看支持CUDA的GPU GTX 520是可以支持CUDA的 ~

我是想问cuda 6.0的统一寻址对显卡有什么要求，是否要求最新的显卡才可以？[/quote] 要求计算能力在3.0以上的GPU才可以支持~[/quote] 如何得知当下GPU计算能力?[/quote] 可以运行CUDA自带的例子，在sample/1_Utilities/deviceQuery中，如果成功安装了CUDA，则会在终端中查看GPU计算能力~

[bob76012]

引用 25 楼 sparrow986831 的回复:

[quote=引用 19 楼 llcchh012 的回复:] [quote=引用 14 楼 sparrow986831 的回复:] https://developer.nvidia.com/cuda-gpus这个网址可以查看支持CUDA的GPU GTX 520是可以支持CUDA的 ~

我是想问cuda 6.0的统一寻址对显卡有什么要求，是否要求最新的显卡才可以？[/quote] 要求计算能力在3.0以上的GPU才可以支持~[/quote] 如何得知当下GPU计算能力?

[weitherd]

引用 18 楼 psyangqi 的回复:

UVA的性能怎么样？ 楼主讲从软件 来实现统一寻址 的意义是：物理上 还是在host端device端都申请了存储 并且 copy 如果 每次只访问一个变量 而不是一个数组 ，访问效率怎么样？

Unified Memory or Unified Virtual Addressing? CUDA has supported Unified Virtual Addressing (UVA) since CUDA 4, and while Unified Memory depends on UVA, they are not the same thing. UVA provides a single virtual memory address space for all memory in the system, and enables pointers to be accessed from GPU code no matter where in the system they reside, whether its device memory (on the same or a different GPU), host memory, or on-chip shared memory. It also allows cudaMemcpy to be used without specifying where exactly the input and output parameters reside. UVA enables “Zero-Copy” memory, which is pinned host memory accessible by device code directly, over PCI-Express, without a memcpy. Zero-Copy provides some of the convenience of Unified Memory, but none of the performance, because it is always accessed with PCI-Express’s low bandwidth and high latency. UVA does not automatically migrate data from one physical location to another, like Unified Memory does. Because Unified Memory is able to automatically migrate data at the level of individual pages between host and device memory, it required significant engineering to build, since it requires new functionality in the CUDA runtime, the device driver, and even in the OS kernel. The following examples aim to give you a taste of what this enables. UVA不拷贝，统一寻址拷贝

[resolutexjh]

对于下一段程序中的 __device__ __managed__ int ret[1000]; 该怎么理解？是在设备端（或全局）定义了一个使用统一内存的变量吗？

[企鹅的外层世界]

平台均为cuda 6.5，使用开普勒架构和麦克斯韦的显卡，同统一寻址的效率有多大的差别？

[search111]

统一虚拟地址UVA，最底层的时候，gpu 计算时候数据是不是还要从cpu 复制到gpu的globle memory