DAY63：阅读Execution Environment

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

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

本文共计845字，阅读时间20分钟

D.2. Execution Environment and Memory Model

D.2.1. Execution Environment

The CUDA execution model is based on primitives of threads, thread blocks, and grids, with kernel functions defining the program executed by individual threads within a thread block and grid. When a kernel function is invoked the grid's properties are described by an execution configuration, which has a special syntax in CUDA. Support for dynamic parallelism in CUDA extends the ability to configure, launch, and synchronize upon new grids to threads that are running on the device.

D.2.1.1. Parent and Child Grids

A device thread that configures and launches a new grid belongs to the parent grid, and the grid created by the invocation is a child grid.

The invocation and completion of child grids is properly nested, meaning that the parent grid is not considered complete until all child grids created by its threads have completed. Even if the invoking threads do not explicitly synchronize on the child grids launched, the runtime guarantees an implicit synchronization between the parent and child.

Figure 12. Parent-Child Launch Nesting

D.2.1.2. Scope of CUDA Primitives

On both host and device, the CUDA runtime offers an API for launching kernels, for waiting for launched work to complete, and for tracking dependencies between launches via streams and events. On the host system, the state of launches and the CUDA primitives referencing streams and events are shared by all threads within a process; however processes execute independently and may not share CUDA objects.

A similar hierarchy exists on the device: launched kernels and CUDA objects are visible to all threads in a thread block, but are independent between thread blocks. This means for example that a stream may be created by one thread and used by any other thread in the same thread block, but may not be shared with threads in any other thread block.

D.2.1.3. Synchronization

CUDA runtime operations from any thread, including kernel launches, are visible across a thread block. This means that an invoking thread in the parent grid may perform synchronization on the grids launched by that thread, by other threads in the thread block, or on streams created within the same thread block. Execution of a thread block is not considered complete until all launches by all threads in the block have completed. If all threads in a block exit before all child launches have completed, a synchronization operation will automatically be triggered.

D.2.1.4. Streams and Events

CUDA Streams and Events allow control over dependencies between grid launches: grids launched into the same stream execute in-order, and events may be used to create dependencies between streams. Streams and events created on the device serve this exact same purpose.

Streams and events created within a grid exist within thread block scope but have undefined behavior when used outside of the thread block where they were created. As described above, all work launched by a thread block is implicitly synchronized when the block exits; work launched into streams is included in this, with all dependencies resolved appropriately. The behavior of operations on a stream that has been modified outside of thread block scope is undefined.

Streams and events created on the host have undefined behavior when used within any kernel, just as streams and events created by a parent grid have undefined behavior if used within a child grid.

D.2.1.5. Ordering and Concurrency

The ordering of kernel launches from the device runtime follows CUDA Stream ordering semantics. Within a thread block, all kernel launches into the same stream are executed in-order. With multiple threads in the same thread block launching into the same stream, the ordering within the stream is dependent on the thread scheduling within the block, which may be controlled with synchronization primitives such as __syncthreads().

Note that because streams are shared by all threads within a thread block, the implicit NULL stream is also shared. If multiple threads in a thread block launch into the implicit stream, then these launches will be executed in-order. If concurrency is desired, explicit named streams should be used.

Dynamic Parallelism enables concurrency to be expressed more easily within a program; however, the device runtime introduces no new concurrency guarantees within the CUDA execution model. There is no guarantee of concurrent execution between any number of different thread blocks on a device.

The lack of concurrency guarantee extends to parent thread blocks and their child grids. When a parent thread block launches a child grid, the child is not guaranteed to begin execution until the parent thread block reaches an explicit synchronization point (e.g. cudaDeviceSynchronize()).

While concurrency will often easily be achieved, it may vary as a function of deviceconfiguration, application workload, and runtime scheduling. It is therefore unsafe to depend upon any concurrency between different thread blocks.

D.2.1.6. Device Management

There is no multi-GPU support from the device runtime; the device runtime is only capable of operating on the device upon which it is currently executing. It is permitted, however, to query properties for any CUDA capable device in the system.

本文备注/经验分享：

今天这部分主要是说的在设备端使用动态并行(DP)的时候, 所涉及到的多流, 同步, kernel启动, events等等之类的,和在CPU端使用它们的异同。实际上, CUDA尽量被设计成你的老经验能够被重复使用, 之前我们所说过的从CPU上使用的这些,基本上动态并行也差不多, 手册的这个章节, 主要是说了不同的方面.在阅读本章节之前, 用户应当重新复习一下之前的Host上的相关内容. 以避免疑惑. 本章节的这些内容, 可以看成是略有差别的Host上的Runtime API的子集.还是很容易使用的.首先动态并行引入了父kernel和子kernel的概念, (精确的说, 这里是父grid和子grid, grid是指kernel的一次启动, 为了方便说明, 下面我们可能随时混用它们)，这个概念是新的. 因为从Host上启动kernel, 是用的CPU代码, 此时并不存在嵌套的两层kernel启动的概念.只有CPU--->一次kernel启动的grid, 这样的对应关系.而这里则需要考虑类似: 父kernel->所启动的下一层kernel们->再下一层的kernels(如果有)->再再下一层 如此反复. 这里的父子关系是相对的. 例如A->B->C->D假设连续的嵌套启动了4个kernel，那么A是B的父kernel, B也是C的父kernel，同时B是A的子kernel, C也是B的子kernel，因此你看到并不单独存在绝对的父或者子kernel, 而是一个相对关系.这种关系主要是会带来同步和一致性之类的方面的问题.目前的CUDA实现中, 父kernel的子kernel(们)结束之前, 它不能结束.因为实际中可能动态并行的时候, 父和子kernel在同时执行(特别是都启动规模较小的时候很容易这样), 一旦子kernel提前执行完毕(很有可能), 父kernel将自动等待它(们).而子kernel又可能有子kernel, 这样实际上,最初层次的kernel(通过CPU端启动的那个), 实际上会等待自己的所有的子kernel和后代的kernel们.这点需要注意. OpenCL也是这样的. 这可能是因为你无法从CPU上跟踪后续被动态启动的kernels, 所以设计成这样.这是第一段的部分. 请注意的是, 你不需要手工使用一次最后的cudaDeviceSynchronize()去等待(该函数是能在设备端使用的Runtime API的子集的其中一个函数),这种等待是自动的.以及, 这里的父kernel和子kernel的同时执行关系不能被保证, 和实际运行时刻的情况有关. 第二段则是说, 一些基本的CUDA中的event或者stream之类的对像, 在动态并行时候的存活周期，这点和CPU端的CUDA Runtime API略微不同. 主要体现在: (1)Host上的event/stream, 在你创建它们之后, 手工销毁之前(以及,出错之前), 是一直有效的. 而设备端的存活周期很短, 只在一个父kernel的创建了它们的特定block中有效. 一旦该block结束, 这些会被自动释放(是的, 能在设备端创建和使用它们, 也是该Runtime API的子集中的函数). (2)一些这种创建它们的函数, 可能会和CPU端的常规Runtime API有区别. 例如stream只能被创建为non-blocking的.你可以理解成设备端的CUDA Runtime API的子集, 功能有限. (3)这些对像使用的地方也有限制. 只能在创建它们的父kernel的特定block中的线程中使用.不能给该父kernel的其他block使用.也不能传给host上使用.也不能传给该block创建的任何子kernel使用.同时event这种, 也没有计时之类的功能.使用限制颇多. 但它们的基本功能还是存在的. 例如streams可以让动态并行, 能同时启动多个彼此执行无关的子kernel. 在设备允许的时候, 就可以同时让这些子kernel执行.例如A->B0, A->B1, A->B2, A->B3，如果它们使用了设备端的streams, 同时硬件此时的情况运行,则B0,B1,B2,B3有可能同时执行(请参考之前的host上的, 用stream进行concurrent kernels执行的章节，和这个情况类似，只不过现在时从GPU上自己给自己启动的. 而不是CPU控制启动的). 关于同步这点, 还需要注意的是,启动子kernel的时候, 任何一个block里的任何一个线程都可以启动.也就是说, 启动是线程级别的，而同步则是该block级别的. 1个block里面的任何1个线程, 都可以对本block中其他伙伴线程启动的子kernel们进行同步.实际上, 只要本block中有任何1个线程启动的子block还在进行, 本block里的所有线程都不能结束.其次, 设备端, 可用cudaDeviceSynchronize()进行中途的同步(而不是结尾的时刻, 自动的隐式同步). 这个时候也是以block为整体单位进行的, 主要涉及到:当前硬件已经没有资源执行子kernel了(例如所有的SM都被父kernel的blocks占满)，此时如果1个block中的1个线程要求进行cudaDeviceSynchronize(), 等待之前启动的所有子kernels们，为了避免死锁, 那么硬件只能强迫性的临时释放出来一些资源,来执行子kernel的blocks.例如此时可能会将该要求进行同步等待的父kernel的该block, 强制性的切换出SM，将它的寄存器中的值, shared memory中的值, 其他的所有状态, 都临时保存到显存.空出资源来, 执行子kernel，等该block中的某线程所要求同步等待的子kernel执行完成后,硬件再自动的从显存中, 读取被临时挂起的该block, 继续从等待完成的点开始执行.所以说, 启动子kernel是线程的行为, 但等待子kernel, 将是block级别的行为.这点需要注意了.

关于event, 这里也需要说明一下. event在设备端只能被用在cudaStreamWaitEvent上,该函数的作用和CPU上的版本一样倒是.允许协调多个子kernel, 在不同的一个block所创建的steams中, 之间的依赖关系.注意CUDA无论是设备端还是CPU端的stream都是顺序执行的.所以想要能有子kernel之间无关执行的并行性, 和这种情况下部分的有关性(依赖关系), 只能多个streams + cudaStreamWaitEvent().这点和OpenCL不同, 来自OpenCL的用户知道, 设备端的默认队列(等于CUDA的默认流),是默认就是乱序执行的, 想要任何顺序或者依赖关系, 必须依靠这样.而CUDA默认的就是顺序的, 但可以通过手工多个流不保持顺序 + Event.这两种设计都可以.CUDA的版本更好一点. 等于默认就有强制性的barrier存在, 保证顺序关系, 比较易用,不容易出错.然后你可以随时通过多流和event消除这种强制性. 关于ordering这里, 还需要说一点.刚才说了, 如果中途要求进行同步等待, 会*可能*导致父kernel的block被切换到显存冻结状态, 暂停执行(例如资源不足的时候).如果中途不要求进行等待, 则父kernel的block可以一直继续往下执行, 而此时,没有足够资源被执行的子kernel, 可能一直得不到执行.所以手册这里强调了:如果父kernel的block不进行同步等待, 则子kernel可能一直没有机会得到执行(没有资源啊)，只有父kernel进行了cudaDeviceSynchronize()之类的操作, 子kernel才可能会开始执行.这是为何这么奇怪的情况存在的原因.其实你会看到, 这种中途的等待, 导致父kernel的block被切换到显存冻结状态, 需要一定量的显存空间，所以同步等待的层数越多, 例如连续3层kernel都因为等待被切换blocks出去到显存了,这种消耗越大.所以CUDA默认有一定的限制(但是可以通过手工扩大它)，这个后续有讲.用户可以认为, 一般的显存紧张的时候, 不太适合使用太多层数的动态并行.或者换句话说, 大显存的卡,在使用动态并行的时候具有天生的优势. 这点需要注意了，因为我们论坛很多来访的用户, 还是在使用只有1GB显存的卡，目前根据情况看, 是GTX750Ti之类的卡居多.这种卡在使用动态并行的时候稍微需要注意一点了，显存的确太少.一旦自己的应用上去就使用了0.8-9个GB, 能干其他用途的就不多了.此外, 根据目前的资源, 这种block被切换出去的情况, 只会发生在你手工要求同步等待的时候,到是不会自动的.所以在kernel里面使用互斥锁之类的实现(还记得之前的章节提供过一个范例的锁的实现吗),还算是比较安全的.不用担心出现在父kernel的多个blocks的线程们之间,某个block当前正持有锁的所有权的线程, 因为突然莫名的被切换出去, 冻结执行,导致其他锁有需要等待锁的线程和block整体卡死.但在父kernel和子kernel之间使用同一个锁的需要需要注意了,任何时候不建议在持有锁的同时, 进行同步等待操作. 这很危险，你很可能会立刻遭遇TDR超时(Windows上)..或者kernel永远不能结束.类似这种问题. 最后的段落则说, 目前DP不支持跨卡操作.这代表了：

（1）你不能在其他卡上启动子kernel，

(2)event不具有跨卡的同步功能.但是这两点均可以使用Host上的CUDA Runtime API解决.Host上的版本的, 还是全功能的. 只是动态并行的时候不可以.

动态并行直接套用了原本的CPU上的CUDA Runtime API用法.还是节省了大量的学习成本的.这样你只需要学习不同的地方即可.成本很低(然后就能在设备端动态的kernel启动kernel了. 而不需要CPU控制)，因为很多应用, 实话说,很难将结果临时传给CPU, 然后CPU在处理这些临时结果, 决定是否启动下个kernel, 什么样的启动配置来启动kernel的时候,能给GPU同时见缝插针的塞入其他活. 很难有活的.此时GPU往往会只能闲置.如果一台8卡的系统, 这种情况一多,浪费的GPU时间将很可观.动态并行消除了需要GPU暂停, 等待CPU处理和发任务的时间,至少在很多场合能让GPU持续忙碌.还是很有意义的.

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

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