CUDA编程接口:异步并发执行的概念和API

转载自IT168：http://cuda.it168.com/a2011/1215/1289/000001289321.shtml

【IT168技术】本文主要描述了支持系统中不同层次的异步并发执行的概念和API。

　　相关阅读：

　　CUDA编程接口：分页锁定主机存储器

　　CUDA编程接口：共享存储器实现矩阵相乘

　　CUDA编程接口：运行初始化与设备存储器

　　CUDA编程接口：使用nvcc编译器的兼容性

　　CUDA编程接口：如何用nvcc编译CUDA程序

　　1.主机和设备间异步执行

　　为了易于使用主机和设备间的异步执行，一些函数是异步的：在设备完全完成任务前，控制已经返回给主机线程了。它们是： 内核发射; 设备间数据拷贝函数; 主机和设备内拷贝小于64KB的存储器块时; 存储器拷贝函数中带有Async后缀的; 设置设备存储器的函数调用。

　　程序员可通过将CUDA_LAUNCH_BLOCKING环境变量设置为1来全局禁用所有运行在系统上的应用的异步内核发射。提供这个特性只是为了调试，永远不能作为使软件产品运行得可靠的方式。 当应用通过CUDA调试器或CUDA profiler(cuda-gdb, CUDA Visual Profiler, Parallel Nsight)运行时，所有的内核发射都是同步的。

　　2.数据传输和内核执行重叠

　　一些计算能力1.1或更高的设备可在内核执行时，在分页锁定存储器和设备存储器之间拷贝数据。应用可以通过检查asyncEngineCount 设备属性查询这种能力，如果其大于0，说明设备支持数据传输和内核执行重叠。这种能力目前只支持不涉及CUDA数组和使用cudaMallocPitch()分配的二维数组的存储器拷贝( 见前文，可阅读“相关阅读”中的文章)。

　　3. 并发内核执行

　　一些计算能力2.x的设备可并发执行多个内核。应用可以检查concurrentKernels属性以查询这种能力)(后续文章将介绍)，如果等于1，说明支持。 设备最大可并发执行的内核数目是16。 来自不同CUDA上下文的内核不能并发执行。 使用了许多纹理或大量本地存储器的内核和其它内核并发执行的可能性比较小。

　　4. 并发数据传输

　　在计算能力2.x的设备上，从主机分页锁定存储器复制数据到设备存储器和从设备存储器复制数据到主机分页锁定存储器，这两个操作可并发执行。 应用可以通过检查asyncEngineCount 属性查询这种能力，如果等于2，说明支持。

　　5. 流

　　应用通过流管理并发。流是一系列顺序执行的命令(可能是不同的主机线程发射)。另外，流之间相对无序的或并发的执行它们的命令;这种行为是没有保证的，而且不能作为正确性的的保证(如内核间的通信没有定义)。

　　①创建和销毁

　　可以通过创建流对象来定义流，且可指定它作为一系列内核发射和设备主机间存储器拷贝的流参数。下面的代码创建了两个流且在分页锁定存储器中分配了一个名为hostPtr的浮点数组。