DAY74：阅读Runtime
我们正带领大家开始阅读英文的《CUDA C Programming Guide》,今天是第73天，我们正在讲解CUDA 动态并行，希望在接下来的27天里，您可以学习到原汁原味的CUDA，同时能养成英文阅读的习惯。
关注微信公众号，查看历史信息，可以看到之前的阅读
本文共计590字，阅读时间15分钟
前情回顾：
DAY61：关于Cooperative Groups
DAY65:阅读Device-Side Kernel Launch
DAY67：阅读阅读Events、Synchronization和Device Management
DAY69：阅读API Errors and Launch Failures
DAY71：阅读Device-side Launch from PTX
DAY72：阅读Toolkit Support for Dynamic Parallelism
DAY73：阅读Programming Guidelines
D.4.3. Implementation Restrictions and Limitations
Dynamic Parallelism guarantees all semantics described in this document, however, certain hardware and software resources are implementation-dependent and limit the scale, performance and other properties of a program which uses the device runtime.
D.4.3.1. Runtime
D.4.3.1.1. Memory Footprint
The device runtime system software reserves memory for various management purposes, in particular one reservation which is used for saving parent-grid state during synchronization, and a second reservation for tracking pending grid launches. Configuration controls are available to reduce the size of these reservations in exchange for certain launch limitations. See Configuration Options, below, for details.
The majority of reserved memory is allocated as backing-store for parent kernel state, for use when synchronizing on a child launch. Conservatively, this memory must support storing of state for the maximum number of live threads possible on the device. This means that each parent generation at which cudaDeviceSynchronize() is callable may require up to 150MB of device memory, depending on the device configuration, which will be unavailable for program use even if it is not all consumed.
D.4.3.1.2. Nesting and Synchronization Depth
Using the device runtime, one kernel may launch another kernel, and that kernel may launch another, and so on. Each subordinate launch is considered a new nesting level, and the total number of levels is the nesting depth of the program. The synchronization depth is defined as the deepest level at which the program will explicitly synchronize on a child launch. Typically this is one less than the nesting depth of the program, but if the program does not need to call cudaDeviceSynchronize() at all levels then the synchronization depth might be substantially different to the nesting depth.
The overall maximum nesting depth is limited to 24, but practically speaking the real limit will be the amount of memory required by the system for each new level (see Memory Footprint above). Any launch which would result in a kernel at a deeper level than the maximum will fail. Note that this may also apply to cudaMemcpyAsync(), which might itself generate a kernel launch. See Configuration Options for details.
By default, sufficient storage is reserved for two levels of synchronization. This maximum synchronization depth (and hence reserved storage) may be controlled by callingcudaDeviceSetLimit() and specifying cudaLimitDevRuntimeSyncDepth. The number of levels to be supported must be configured before the top-level kernel is launched from the host, in order to guarantee successful execution of a nested program. Calling cudaDeviceSynchronize() at a depth greater than the specified maximum synchronization depth will return an error.
An optimization is permitted where the system detects that it need not reserve space for the parent's state in cases where the parent kernel never calls cudaDeviceSynchronize(). In this case, because explicit parent/child synchronization never occurs, the memory footprint required for a program will be much less than the conservative maximum. Such a program could specify a shallower maximum synchronization depth to avoid over-allocation of backing store.
D.4.3.1.3. Pending Kernel Launches
When a kernel is launched, all associated configuration and parameter data is tracked until the kernel completes. This data is stored within a system-managed launch pool.
The launch pool is divided into a fixed-size pool and a virtualized pool with lower performance. The device runtime system software will try to track launch data in the fixed-size pool first. The virtualized pool will be used to track new launches when the fixed-size pool is full.
The size of the fixed-size launch pool is configurable by calling cudaDeviceSetLimit() from the host and specifying cudaLimitDevRuntimePendingLaunchCount.
本文备注/经验分享：
这章节主要是说, 使用动态并行所需要付出的设备上的代价, 以及, 对这些代价的注意事项。
我们之前的章节, 曾经说过, 父kernel可以选择对它的子kernel们进行同步,这种同步目前是在block级别上完成的: 任何block内部的线程要求进行同步, 都可能导致这个线程所在的block被整体暂时冻结,将状态保存到显存上, 空出来它之前占用的SM上的地方, 让给子kernel(们)运行.那么很显然的, 这种临时的保存, 需要付出占用显存的代价.所以这里为何说, 这里有"memory footprint".(手册还有一定的剩余章节的. 计算能力后面那里的区别还是需要说一下的)。 常见的显存占用代价有两种, 一种是同步深度导致的代价, 另外一种是尚未启动的子kernel列表所占用的空间的代价.这两个都很好理解. 前者你知道如果父kernel和子kernel是相对的概念, 父kernel可以等待自己的子kernel而被交换到显存冻结,子kernel又可能又自己的子kernel, 做为自己的子kernel的父kernel, 它在等待的时候也可能被冻结到显存,这样一层一层的冻结, 每层都需要一定的显存容量空间.所以存在一个同步层次(或者说深度)的概念, 需要在进行的同步越深, 则需要提前保留出来的显存容量越大(很好理解吧?). 而本章节为了说明同步深度的概念, 首先引入了嵌套深度的概念, 这种嵌套的深度(nesting depth), 不妨用图片表示一下: Host -> Kernel A -> Kernel B -> Kernel C -> Kernel D ... 这么一个调用层次, 请注意总是存在一个从Host启动的最外层的kernel的,然后这个最外层的kernel, 然后才能利用动态并行, 再次启动kernel.根据你看到的NV自家的文档不同,有的时候, 这个最外层的kernel被叫做深度0，它直接启动的子kernel叫深度1，子kernel的子kernel叫深度2，也就是一个0, 1, 2, 3...N这样的层次概念.但是很令人遗憾的是, NV目前还有另外一种叫法, 也是官方文档的,就是最外层那个kernel叫层次1, 下一个叫层次2.如此类推.所以这导致了实际上存在从0开始, 和从1开始的两种深度的叫法.这两种叫法都很常见. 当用户无法确定自己需要的深度是那种做为起点的时候(例如你只看到了本手册做为文档, 而没有看当年的GTC),我建议总是保留1个深度做为余量.不过这里可以直接说的是, 手册这章节里面的同步深度是从最外层就算层次1开始的.这样你通过本章节的, 设定过的最大同步深度数值(例如8), 实际上只能用到第7层子kernel上.这点需要注意了. (如果你总是+1的设定, 则可以无视这点)，然后同步深度和嵌套深度本章节还说明了, 其实是两回事, 这里也需要注意.因为有些父kernel实际上可以不等待自己的子kernel. 例如这种调用流程: Host -> Kernel A -> 启动Kernel B和C Host -> 读取kernel B和C的结果, 然后进行kernel D的启动这种做法, 因为Host上才需要kernel B和C的结果,而通过kernel A可能只需要负责动态并行启动B和C的时候的参数准备, 并不需要自己处理结果,因为kernel A做为B和C的父kernel, 它可以不显式的用cudaDeviceSynchronize()等待自己的子kernel结束,因为它不需要中途等待她们的结果做下一步的判断---这工作在Host上会后续完成, 所以它没有中途需要被暂停出来的说法,而没有中途暂停冻结起来到显存, 因此也不需要占用额外的显存容量,所以这里实际上并不需要任何同步深度.实际这样在动态并行时候, 所需要付出的显存占用代价较小.
类似的还有A->B->C->D->E->F, 这样6层，嵌套深度虽然较深, 但实际的同步深度可能在0到5层中间的任何值，A,B,C,D,E做为相对的父kernel们, 如果他们都一个等一个下来, 那么同步深度是5, 你需要指定5的同步深度，而如果一个算法最后, A,B,C,D,E这5个父kernel没有任何人需要等待, 你可以直接设定成0层的同步深度。这样会减少预先保留的kernel交换空间，这点也需要注意.注意还有一点, 不同档次的卡的SM个数不同,能同时在SM上驻留的blocks最大数量也不同,所以实际上, 同样的需要预留出来给相同深度层次的同步用的空间,实际也不同, 高档的卡往往显存大, 但它的SM往往也多,所以每层同步付出的保留空间代价也大.而低档的卡往往显存很小, 例如只有1GB, 但它的SM也相应的较少,所以每层同步付出的空留空间也小.所以有的人可能感觉, 在某些卡上, 每增加一层父kernel的同步都需要额外占用, 例如100MB显存,她如果换成小显存的卡, 那么光动态并行, 例如10层kernel就占用光了显存, 啥都不用干了?实际上不是的, 小显存的卡因为自己规格弱小, 所以每层的代价也较低, 例如20MB,所以还是可以正常使用动态并行的. 这里需要注意了.本章节可能在后面, 有个万年不变的每层会占用150MB显存的说法, 这个说法是当年针对K20和K40说的,万年没有更新了, 请以你的现在手头的卡的实际为准.不用担心小显存的卡如此巨大的动态并行同步开销. 然后本章节的后一部分, 还有一个pending launch count, 也就是等待启动中的最大子kernel数量限制和资源占用问题.动态并行本身会维持一个需要被动态启动的子kernel列表的,这个列表本身也占用空间, 需要在使用动态并行的时候预留出来.这个默认的大小大约是2000个左右等待启动的子kernel. 一般足够用了. 如果不够用了, 可以手工通过本章节的最后一句的说法, 来扩大它. 这个占用的空间比同步所需要的预留空间要小. 所以可以适当的扩大较大.注意根据NV的文档不同, 存在两种说法,一种是本手册的说法, 可以预先分配固定的等待启动的kernel的所需数量的固定空间, 这个空间超过后, 就再也不能启动子kernel了.另外一种则是GTC上的NV的说法, 如果当前等待启动的子kernel数量超过了这个固定分配的大小, 则会使用另外一个慢速的方式, 存在额外超出的需要等待启动的子kernel, 依然可以继续启动的. 只是性能会降低.后者还当年配合了想当生动的图例说明这两种(子kernel数量超了, 和子kernel数量没超)时候的动态并行启动性能变化.我建议保持手册的说法, 这样万一将来实现(CUDA实现, 或者AMD的HIP兼容实现)发生了改变,使用未定义的行为存在对未来的兼容性的风险. 注意需要说明的一点是,之前曾经说过, 子kernel完全无法保证任何执行进度的,只有当父kernel调用cudaDeviceSynchronize()等它的时候, 它才可能一定会开始执行, 所以当用户跟踪自己的还有多少子kernel每有启动的时候, 例如用户的最外层父kernel的10个线程, 每个启动了10个子kernel,则应当总是假设这100个子kernel一定都没有执行.只有当用户明确的调用cudaDeviceSynchronize()的时候, 才能总是假定kernel此时方能执行, 以及, 必须此时假定父kernel一定可能会被冻结1个或者多个或者全部blocks.但实际上的情况可能要比这宽泛, 例如父kernel的blocks和线程规模很小, 子kernel可能立刻就执行了.而父kernel再等待子kernel的时候, 也可能并没有被切换出去.使用后者这种假设增加了死锁的风险, 特别是父kernel正在持有一把和子kernel共享的锁的当前所有权的时候,所以用户如果要做这种宽泛的假设, 则一定要结合自己的当前的卡的实际情况,并不能假设任何未来的兼容性.(虽然根据目前的NV的暗示, 将来会引入动态的kernel任务切换,一个kernel的blocks可能再GPU被反复的调度进来SM, 调度出去的,但用户现在不应当假设这点)，论坛上的用户代码的BUG都是用户写的太随意造成的.
有不明白的地方，请在本文后留言
或者在我们的技术论坛bbs.gpuworld.cn上发帖