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的浮点数组。

cudaStream_t stream[2];

for (int i = 0; i < 2; ++i)

cudaStreamCreate(&stream[i]);

float* hostPtr;

cudaMallocHost((void**)&hostPtr, 2 * size);

 　　下面的代码定义的每个流是一个由一次主机到设备的传输，一次内核发射，一次设备到主机的传输组成的系列。

for (int i = 0; i < 2; ++i){ cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size, size, cudaMemcpyHostToDevice, stream[i]); MyKernel<<<100, 512, 0, stream[i]>>> (outputDevPtr + i * size, inputDevPtr + i * size, size); cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size, size, cudaMemcpyDeviceToHost, stream[i]); }

　　每个流将它的hostPtr输入数组的部分拷贝到设备存储器数组inputdevPtr，调用MyKernel()内核处理inputDevPtr，然后将结果outputDevPtr传输回hostPtr同样的部分。后文描述了例子中的流如何依赖设备的计算能力重叠。必须注意为了使用重叠hostPtr必须指向分页锁定主机存储器。

　　调用cudaStreamDestroy()来释放流。

for (int i = 0; i < 2; ++i)

 cudaStreamDestroy(stream[i]);

 　　cudaStreamDestroy()等待指定流中所有之前的任务完成，然后释放流并将控制权返回给主机线程。

 

　　② 默认流

　　内核启动和没有使用流参数的主机设备间数据拷贝，或者等价地将流参数设为0，此时发射到默认流。因此顺序执行。

　　③显式同步

　　有很多方法显式的在流之间同步。

　　cudaDeviceSynchronize()直到前面所有流中的命令都执行完。

　　cudaStreamSynchronize()以某个流为参数，强制运行时等待该流中的任务都完成。可用于同步主机和特定流，同时允许其它流继续执行。

　　cudaStreamWaitEvent()以一个流和一个事件为参数(后文将介绍)，使得在调用cudaStreamWaitEvent()后加入到指定流的所有命令暂缓执行直到事件完成。流可以是0，此时在调用cudaStreamWaitEvent()后加入到所有流的所有命令等待事件完成。

　　cudaStreamQuery()用于查询流中的所有之前的命令是否已经完成。

　　为了避免不必要的性能损失，这些函数最好用于计时或隔离失败的发射或存储器拷贝。

　　④隐式同步

　　如果是下面中的任何一种情况，来自不同流的两个命令也不能并发：分页锁定主机存储器分配，设备存储器分配，设备存储器设置，设备之间存储器拷贝，默认流中调用的任何CUDA命令， F.4.1节描述的一级缓存/共享存储器之间配置切换。

　　对于支持并发内核执行的设备，任何需要依赖检测以确定内核发射是否完成的操作：

　　1）只有来自CUDA上下文中的任何流中的所有的前面的内核启动的线程块开始执行，才能够开始执行;

　　2）会阻塞CUDA上下文中后面任何流中所有的内核发射直至被检测的内核发射完成。

　　需要依赖检测的操作包括同一个流中的一些其它类似被检查的发射的命令和流中的任何cudaStreamQuery()调用。因此，应用应当遵守这些指导以提升潜在的内核并发执行：

　　1）所有独立操作应当在依赖操作之前发出，

　　2）任何类型同步尽量延后。

　　⑤重叠行为

　　两个流的重叠执行数量依赖于发射到每个流的命令的顺序和设备是否支持数据传输和内核执行重叠、并发内核执行、并发数据传输。

　　例如，在不支持并发数据传输的设备上，前面例程的两个流并没有重叠，因为发射到流1的从主机到设备的存储器拷贝在发射到流0的从设备到主机的存储器拷贝之后，因此只有发射到流0的设备到主机的存储器拷贝完成它才开始。如果代码重写成如下方式(同时假设设备支持数据传输和内核执行重叠)。

for (int i = 0; i < 2; ++i)

cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size, size, cudaMemcpyHostToDevice, stream[i]);

for (int i = 0; i < 2; ++i)

MyKernel<<<100, 512, 0, stream[i]>>> (outputDevPtr + i * size, inputDevPtr + i * size, size);

for (int i = 0; i < 2; ++i)

cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size, size, cudaMemcpyDeviceToHost, stream[i]);

　　此时发射到流1的从主机到设备的存储器拷贝和发射到流0的内核执行重叠。

　　在支持并发数据传输的设备上，前文例程的两个流重叠：发射到流1的从主机到设备的存储器拷贝和发射到流0的设备到主机的存储器拷贝，甚至和发射到流0的内核执行(假设设备支持数据传输和内核执行重叠)。但是内核执行不可能重叠，因为发射到流1的第二个内核执行在发射到流0的设备到主机的存储器拷贝之后，因此会被阻塞直到发射到流0的内核执行完成。如果代码被重写成上面的样子，内核执行就重叠了(假设设备支持并发内核执行)，因为发射到流1的第二个内核执行在发射到流0的设备到主机的存储器拷贝之前。然而在这种情况下，发射到流0的设备到主机的存储器只和发射到流1的内核执行的最后一个线程块重叠，这只占总内核执行时间的一小部分。

 

　　6.事件

　　通过在应用的任意点上异步地记载事件和查询事件是否完成，运行时提供了精密地监测设备运行进度和精确计时。当事件记载点前面，事件指定的流中的所有任务或者指定流中的命令全部完成时，事件被记载。只有记载点之前所有的流中的任务/命令都已完成，0号流的事件才会记载。

　　①创建和销毁

　　下面的代码创建了两个事件：

cudaEvent_t start, stop; cudaEventCreate(&start); cudaEventCreate(&stop);

 　　以下面的方式销毁它们：

cudaEventDestroy(start); cudaEventDestroy(stop);

 　　②过去的时间

　　节建立的事件可以用下面的方式给3.2.5.5.1节的代码计时：

cudaEventRecord(start, 0); for (int i = 0; i < 2; ++i)｛ cudaMemcpyAsync(inputDev + i * size, inputHost + i * size, size, cudaMemcpyHostToDevice, stream[i]); MyKernel<<<100, 512, 0, stream[i]>>> (outputDev + i * size, inputDev + i * size, size); cudaMemcpyAsync(outputHost + i * size, outputDev + i * size, size, cudaMemcpyDeviceToHost, stream[i]); cudaEventRecord(stop, 0); cudaEventSynchronize(stop); float elapsedTime; cudaEventElapsedTime(&elapsedTime, start, stop);

 　　7.同步调用

　　直到设备真正完成任务，同步函数调用的控制权才会返回给主机线程。在主机线程执行任何其它CUDA调用前，通过调用cudaSetDeviceFlags()并传入指定标签(参见参考手册)可以指定主机线程的让步，阻塞，或自旋状态。

        更多内容请点击：

        CUDA专区：http://cuda.it168.com/

        CUDA论坛：http://cudabbs.it168.com/