DAY67：阅读阅读Events、Synchronization和Device Management

我们正带领大家开始阅读英文的《CUDA C Programming Guide》,今天是第67天，我们正在讲解CUDA C语法，希望在接下来的33天里，您可以学习到原汁原味的CUDA，同时能养成英文阅读的习惯。

关注微信公众号，查看历史信息，可以看到之前的阅读

本文共计413字，阅读时间12分钟

D.3.1.3. Events

Only the inter-stream synchronization capabilities of CUDA events are supported. This means that cudaStreamWaitEvent() is supported, but cudaEventSynchronize(), cudaEventElapsedTime(), and cudaEventQuery() are not. As cudaEventElapsedTime() is not supported, cudaEvents must be created via cudaEventCreateWithFlags(), passing thecudaEventDisableTiming flag.

As for all device runtime objects, event objects may be shared between all threads withinthe thread-block which created them but are local to that block and may not be passed to other kernels, or between blocks within the same kernel. Event handles are not guaranteed to be unique between blocks, so using an event handle within a block that did not create it will result in undefined behavior.

D.3.1.4. Synchronization

The cudaDeviceSynchronize() function will synchronize on all work launched by any thread in the thread-block up to the point where cudaDeviceSynchronize() was called. Note that cudaDeviceSynchronize() may be called from within divergent code (see Block Wide Synchronization).

It is up to the program to perform sufficient additional inter-thread synchronization, for example via a call to __syncthreads(), if the calling thread is intended to synchronize with child grids invoked from other threads.

D.3.1.4.1. Block Wide Synchronization

The cudaDeviceSynchronize() function does not imply intra-block synchronization. In particular, without explicit synchronization via a __syncthreads() directive the calling thread can make no assumptions about what work has been launched by any thread other than itself. For example if multiple threads within a block are each launching work and synchronization is desired for all this work at once (perhaps because of event-based dependencies), it is up to the program to guarantee that this work is submitted by all threads before calling cudaDeviceSynchronize().

Because the implementation is permitted to synchronize on launches from any thread in the block, it is quite possible that simultaneous calls to cudaDeviceSynchronize() by multiple threads will drain all work in the first call and then have no effect for the later calls.

D.3.1.5. Device Management

Only the device on which a kernel is running will be controllable from that kernel. This means that device APIs such as cudaSetDevice() are not supported by the device runtime. The active device as seen from the GPU (returned from cudaGetDevice()) will have the same device number as seen from the host system. The cudaDeviceGetAttribute() call may request information about another device as this API allows specification of a device ID as a parameter of the call. Note that the catch-all cudaGetDeviceProperties() API is not offered by the device runtime - properties must be queried individually.

本文备注/经验分享：

今天的章节首先引入了动态并行时候的, 设备端的Event，虽然说设备端的时间不具有轮询(polling)和阻塞(blocking)同步功能,也不具有计时功能,更没有Host上的Event的特色的跨设备同步(Stream不能做这个),也就是说它基本上是一个非常简化的CPU上的CUDA Runtime的event的功能的子集,它在动态并行的时候, 在设备端上只具有在同一个block中的streams之间的互相barrier功能(cudaStreamWaitEvent),但是对于很多应用动态并行的代码来说, 它也是必须的, 因为它将原本设备上的动态并行时候, 启动的kernel的两种顺序: 纯串行(1个流多个子kernel), 和纯并行(多个流), 给变成了图(Graph)状的, 例如某父kernel需要启动3个子kernel来完成一个任务, 这3个子kernel分别叫A,B,C, 其中C需要使用A,B的输出结果,而A,B之间却无所谓顺序.此时, 用户可以选择只使用1个流, 例如: A -> B -> C的启动,也可以选择B -> A -> C的启动顺序(因为A,B之间无所谓顺序)，这种就是用昨天的章节, 使用1个流进行动态并行就能实现的.但是无论哪种, 均实际上浪费了显卡的可能能力, 例如A,B的启动规模都很小, 显卡此时又有足够的资源,如果能同时启动A,B, 等它们都完成后, 再启动C，将在可能的情况下, 能将C的启动延迟从原本的A执行时间 + B执行时间,变成A执行时间, 和B执行时间中的, 较大的一个即可. 而两者较大的一个时间, 要小于两者的总和(相差了一个较小的时间)，也就是说: 原本: 流: A->B->C 现在: 流1: A->Event同步后执行C 流2: B --------------| 也就是说, 现在可以使用2个流, 同时用一个event约束C, 让他等待B完成后再启动,因为C本来和A都在流1里面, 本来就会等待A完成, 现在使用了event, 增加了等待B即可(cudaEventRecord + cudaStreamWaitEvent), 而在C等待A,B都完成的时候, A,B又可以同时进行，(因为A,B在2个流里) 这样就实现了原本的正确的逻辑执行顺序(先A和B, 再C), 同时又能表达A,B的并行性.需要注意的是, CUDA总是易用的, 你也完全可以不使用event这里. 例如原本的只有一个流的A->B->C, 但这样会付出更多的执行时间. 也就是说, 不懂本章节的内容, 不妨碍你写出正确的代码,但是只是可能会执行的慢一些, 例如别人只需要一张1080即可, 而你却需要一张Titan-V才能完成同样的任务(夸张了. 只是举个例子). 类似的, 用户也完全可以另外一种方式, 不利用event: 第一次动态并行调用: 流1: A-------\ 执行cudaDeviceSynchronize() 流2: B-------/ 第二次动态并行调用: 流1或者流2: C 这样也是可以的, 即先只启动A,B, 然后父kernel暂停, 进行同步等待, 等待完成后, 恢复父kernel运行, 再动态并行启动kernel C，这样也可以只用上一章节(stream)的知识内容即可.但同样的, 也需要付出额外的代价, 性能可能没有直接用event连续启动3个子kernel好.所以阅读了本章节的内容, 还是很有用的, CUDA这东西设计的有意思, 你完全可以软件上懂得很少+买一张好卡,也可以软件上懂得多+稍微差一点的卡.但是如果读者试图将CUDA做为一种业务能力, 还是应当阅读理解会所有章节地.本章节的event, 实际上, 能将动态并行的时候, 子任务(子kernel们的启动)的集合,当成一个有向无环图(DAG)来处理, 通过event来处理每个点和边之间的关系.因为很多复杂的互相逻辑上依赖的任务, 实际上都是有向无环图.这是一个很重要的数据结构.包括你购买多卡系统, 例如8卡的1080, 看上去每张卡都是一样的,可以直接的均衡的将所有的任务平均每个卡1/8即可简单处理, 似乎并不需要图这种结构.但实际上, 每张卡的位置不同, 散热条件不同(这将影响性能), PCI-E拓扑上的位置也不同(这也将影响性能--特别是传输和P2P),这将导致实际上你的8卡1080系统实际上等效的变成了8张不同性能的卡的系统.这种情况下, 每个卡上的同样的kernel执行, 执行的时间和结束的时间都可能不同,直接简单粗暴的1/8每卡, 将实际上等于你受限于你散热条件和插槽位置最不好的那张1080.而在更实际的应用中, CPU上的调度, 多组kernel之间的大小, 执行复杂度不同, 使得你几乎连完美的1/8切分都做不到.此时应当考虑DAG.而考虑将任务组描述成DAG, 则无论在动态并行设备端的启动, 还是在Host上启动, 你将都需要event.而Host上, event更多增加了多卡之间的互相协调能力, 非常适合使用.但event在动态并行的时候, 被局限在同一个卡上, 暂时不能这样做.(但可以协调卡内的多个子kernel任务)。 实际上, 本章节最后的末尾段落, 提到动态并行和多卡的时候, 说到, 你不能跨卡通过动态并行启动kernel的.主要是因为一种考虑: 多卡上的同名kernel可能具有不同的binary(cubin)表示. 这主要是因为不同的卡可能具有不同的计算能力, 不一定二进制兼容. 而如果动态并行同时还要解决这个问题(例如对PTX的JIT使用), 就设计的太复杂了. 所以干脆没有，其实类似本章节的结尾段落说的, 你连cudaSetDevice()这种切换当前卡的操作也不能进行的. 所有的动态并行都必须针对当前卡.而当前卡, 则是父kernel正在执行的这个卡.此外, 动态并行的时候, 对kernel的计时也不能被支持, 这也是仅限从Host上能使用的特性.这样做也很容易理解: 动态并行的时候, 往往有父kernel之类(或者其他伙伴子kernel)在同时运行,而我们对时间的测量往往代表kernel的性能,但是在这种情况下, 每个kernel的执行时间往往会被延长. 甚至变得毫无规律，此时继续测试, 不能反映一个kernel的真实性能, 因此干脆不提供为好.(你依然可以从Host上, 对一次父kernel启动和它的所有子kernel们的整体时间计时的, 此时反映了这个任务集合的整体性能,还是有意义的)

最后在本章节的中间部分, 还谈到的动态并行时候, 父kernel对子kernel(们)的唯一同步方式的说明，因为另外两大同步方式(对Stream和Event的轮询或者阻塞同步)都不能使用.这唯一的cudaDeviceSynchronize()就是我们之前说的, "同步是block的行为".具体本章节表现如下: (1)任何一个block中的任何一个线程, 调用了cudaDeviceSynchronize()进行阻塞同步的时候, 将等待该调用前的, 本block建立的所有流中的所有子kernels们完成. 不能有选择的等待1个, 或者几个(如果需要有选择的等待1个或者几个, 你需要单独用动态并行启动这几个+阻塞同步一次, 然后重复这个过程即可，但这样可能会造成性能上的损失)，这是为何说同步是block的行为的原因. (2)block中的任何一个线程调用了cudaDeviceSynchronize(), 都有可能导致本block整体(所有的线程执行状态, 当前使用的寄存器内容, 当前的shared memory内容), 被交换到global memory上进行状态冻结, 暂停运行. 也就是说, 一旦使用了这唯一的同步方式, 任何1个线程将会影响当前block整体(这个之前实际上稍微提到过). 所以这就是我们为何经常说的, 动态并行的时候, "子kernel的启动是单独的线程的行为, 而同步则是block整体的行为"的原因.

有不明白的地方，请在本文后留言

或者在我们的技术论坛bbs.gpuworld.cn上发帖