DAY73：阅读Programming Guidelines
DAY73：阅读Programming Guidelines
我们正带领大家开始阅读英文的《CUDA C Programming Guide》,今天是第73天，我们正在讲解CUDA 动态并行，希望在接下来的27天里，您可以学习到原汁原味的CUDA，同时能养成英文阅读的习惯。
关注微信公众号，查看历史信息，可以看到之前的阅读
本文共计310字，阅读时间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
D.4. Programming Guidelines
D.4.1. Basics
The device runtime is a functional subset of the host runtime. API level device management, kernel launching, device memcpy, stream management, and event management are exposed from the device runtime.
Programming for the device runtime should be familiar to someone who already has experience with CUDA. Device runtime syntax and semantics are largely the same as that of the host API, with any exceptions detailed earlier in this document.
The following example shows a simple Hello World program incorporating dynamic parallelism:
This program may be built in a single step from the command line as follows:
$ nvcc -arch=sm_35 -rdc=true hello_world.cu -o hello -lcudadevrtD.4.2. Performance
D.4.2.1. Synchronization
Synchronization by one thread may impact the performance of other threads in the same Thread Block, even when those other threads do not call cudaDeviceSynchronize() themselves. This impact will depend upon the underlying implementation.
D.4.2.2. Dynamic-parallelism-enabled Kernel Overhead
System software which is active when controlling dynamic launches may impose an overhead on any kernel which is running at the time, whether or not it invokes kernel launches of its own. This overhead arises from the device runtime's execution tracking and management software and may result in decreased performance for e.g., library calls when made from the device compared to from the host side. This overhead is, in general, incurred for applications that link against the device runtime library.
本文备注/经验分享：
设备端库提供了一些前几章节曾经说明过的设备,流,事件,传输等支持,这些支持的函数是Host的上的对应版本的功能子集, 而且和Host上的不同也已经在之前的章节说明过了, 这里给出一个演示用的例子子。 这个例子也很简单. GPU上的kernel自己又启动了一个子kernel, 输出了"Hello"，然后父Kernel输出了"World"。最终输出了"Hello World"。 Hello,World是很常见的一种基本演示代码的方式. 注意这里的<<<>>>启动子kernel后立刻检查是否启动成功. 这是一个非常重要的好习惯.建议总是保持这样. 多人没有这个习惯. 这是不好的.包括在论坛上, 从Host上启动kernel的人, 也往往不检查这个. 正常的kernel执行需要检查两步: (1)检查是否能启动. (2)如果能启动, 在后续的同步的时候, 检查执行过程是否有问题. 很多CUDA的书只对kernel检查一次(第二点), 这是不对的. 需要注意的是, 在设备端动态并行的时候, 如果子kernel因为访存之类的原因挂掉,因为父子kernel共享一个CUDA Context执行环境, 此时子kernel挂掉的时候, 因为可能已经破坏父kernel的正常运行环境,例如所需要的显存内容, 会连累父kernel直接一块挂掉的.此时应当从Host上检查父kernel和所有的子kernel做为一个整体, 是否执行成功.任何这个整体之中的单一kernel挂了, 会做为整体反映的. 所以设备上, 和Host上的两步检查稍微还是不同的. 以及, 设备端上的同步也可能会失败(因为有个最大同步深度的问题). 这点也是不同的. 需要注意了.此外, 这段演示代码使用的-arch sm_35的参数，这里的sm_35部分, 实际上是因为当初只有计算能力3.5(在本章节写的最初), 能支持动态并行.请替换成你实际使用的卡的计算能力.例如sm_61, sm_50这种，(6.1是家用Pascal)。 后一章节的后半部分, 则说明了, 使用动态并行可能会带来的性能上的问题. (1)则是1个block中的任何1个thread执行了同步, 则可能会影响整体block，这也是我们之前常说的, "启动是线程自己的行为, 但同步是1个block的行为"，因为例如当前的实现上, block中的任何一个线程要求对子kernel们同步, 则可能会导致该block整体被暂停,然后被切换到显存上冻结.不过本段落也说了, 以后的卡的具体实现, 可能会改变这一点.等等. 这都是之前的内容的强调实际上.然后本章节的后半部分的第二段则是说,如果程序使用动态并行, 链接了设备端的库, 则可能因为设备端的库内含的代码, 需要进行一些, 例如管理任务, 维护动态并行所需要的环境,会带来额外的性能损失.包括你链接了库后, 有的kernel并没有调用过子kernel, 也可能会遭受这点(但一般比较轻微)，所以实际上, 使用动态并行是否100%能带来性能上的正面提升, 则和具体的应用有关. 需要用户自己权衡.这就像纹理引用一样.一个kernel所在的.cu里面有纹理引用,哪怕具体到某个kernel, 它没有使用到该引用,该kernel依然会遭受额外的启动延迟损失.这点有点类似.不过我们以前也常说, 使用纹理等时候一定能带来性能提升, 也要看具体应用.
有不明白的地方，请在本文后留言
或者在我们的技术论坛bbs.gpuworld.cn上发帖