首页
学习
活动
专区
工具
TVP
发布
社区首页 >专栏 >DAY66:阅读Streams

DAY66:阅读Streams

作者头像
GPUS Lady
发布2018-08-17 14:42:17
5910
发布2018-08-17 14:42:17
举报
文章被收录于专栏:GPUS开发者GPUS开发者

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

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

本文共计282字,阅读时间10分钟

D.3.1.2. Streams

Both named and unnamed (NULL) streams are available from the device runtime. Named streams may be used by any thread within a thread-block, but stream handles may not be passed to other blocks or child/parent kernels. In other words, a stream should be treated as private to the block in which it is created. Stream handles are not guaranteed to be unique between blocks, so using a stream handle within a block that did not allocate it will result in undefined behavior.

Similar to host-side launch, work launched into separate streams may run concurrently, but actual concurrency is not guaranteed. Programs that depend upon concurrency between child kernels are not supported by the CUDA programming model and will have undefined behavior.

The host-side NULL stream's cross-stream barrier semantic is not supported on the device (see below for details). In order to retain semantic compatibility with the host runtime, all device streams must be created using the cudaStreamCreateWithFlags() API, passing the cudaStreamNonBlocking flag. The cudaStreamCreate() call is a host-runtime- only API and will fail to compile for the device.

As cudaStreamSynchronize() and cudaStreamQuery() are unsupported by the device runtime, cudaDeviceSynchronize() should be used instead when the application needs to know that stream-launched child kernels have completed.

D.3.1.2.1. The Implicit (NULL) Stream

Within a host program, the unnamed (NULL) stream has additional barrier synchronization semantics with other streams (see Default Stream for details). The device runtime offers a single implicit, unnamed stream shared between all threads in a block, but as all named streams must be created with the cudaStreamNonBlocking flag, work launched into the NULL stream will not insert an implicit dependency on pending work in any other streams.

本文备注/经验分享:

动态并行时候的, 设备端使用的stream.我们常用的stream是指Host上的, 但动态并行的时候, 因为CUDA和OpenCL相反, 设备端的流总是顺序的,也就是同一个流里面的kernel总是串行执行的,所以如果需要在设备端, 通过动态并行启动多个能同时执行的kernel, 则必须在设备端使用多流.这也是论坛上面, 为何很多"我只想使用同一个流, 却需要让里面的多个kernel能同时并行执行"之类的问题,无法回答一样.这是因为对于同一个流, 无论从Host上, 或者从设备上, 都只能在里面串行的启动流的.无论是常规启动, 还是动态并行设备端启动, 都做不到.很多有这种疑问的用户, 因为OpenCL和这里是反的.默认的OpenCL里的设备端流(队列), 是自然就并行的(乱序)的. 而CUDA里面,默认的设备端的流是顺序的。所以需要乱序(嗯嗯), 你需要手工像本章那样, 单独的创建多个设备端的流, 给动态并行使用。而OpenCL则需要单独的通过事件(实际上是barrier)进行约束,用户可以分别理解成, CUDA的动态并行, 默认提供的功能较少, 但用起来简单。而OpenCL的实现, 看上去功能很多, 但对用户造成的负担较大(用户需要理解顺序和乱序两种队列(流), 同时用barrier控制住自己不需要的特性),不是那么易用, 但会给人带来更加精细的控制力的感觉.这两种哪种都可以.CUDA的这种好处是, 用户可以完全在没有相关知识(本章节的多流, 例如只在动态并行里面使用默认流, 或者干脆不指定任何流---这种等于默认流)的情况下,就能写出功能正常, 逻辑完备的应用了动态并行的kernel来.这其实是CUDA的一贯特色, 入门很容易,然后后期用户感觉不满足, 可以逐渐的再增强自己的CUDA知识. 而OpenCL用户则有一个陡峭的入门曲线,例如CUDA用户培训3天就能上岗,OpenCL不培训个一个月整, 直接上岗, 会处处碰到坑和问题.但好处是OpenCL是可以后期基本不用再怎么学习了.这种鲜明的特色对比, 实际上贯彻了这整本手册的全部章节.读者可以不时的体验到. 再说一下动态并行的时候, 设备端使用的stream和host端的主要区别.这个区别主要体现在创建, 使用, 和同步上. 我分别说一下这三点: Host上对流的创建, 可以创建一个普通流, 也可以创建一个非阻塞流.后者这种流, 不对默认流进行隐式的自动同步.而Host上的默认流, 目前也有两种, 一种是经典默认流, 另外一种是新型的per-thread的默认流,(后者需要编译的时候, 使用--default-stream per-thread编译, 用户如果不知道这点, 则可以认为只有经典默认流,这也是CUDA的易用性的体现, 你不知道, 完全不妨碍; 知道了, 有可能利用这点提升性能),但是动态并行的时候, 设备端的流则不同,首先的确也分成默认存在的流, 和需要手工创建的流.但是手工创建的流必须使用非阻塞标志创建, 理由请参考本章节.这是为了维持和Host上的代码语义一致.因为动态并行的时候, 并没有经典的带有自动同步的默认流. 这是第一点, 创建上来说. 关于第二点, 在动态并行的时候, 设备端的流的使用, 注意你能使用的函数, 比较少, 因为之前曾经说过, 设备端的CUDA Runtime API,只是全功能的Host上的CUDA Runtime API的一个子集.像是流查询, 流同步之类的函数(cudaStreamQuery/Synchronize)都不能使用,因为动态并行的CUDA Runtime, 是一个简化版本,就没有这些东西。这点实际上用户很容易发现. 因为试图使用这些设备端的动态并行时候, 超出使用范围的API, 编译的时候都会报错.不过虽然只是简化版本的, 但对于动态并行的时候, 一般也够用了(不够用你可以考虑传统的从Host上启动kernel嘛! 当然, 方便性和可能的性能优势就没有了). 第三点则是同步上.请注意, 动态并行的同步只有1种, 就是cudaDeviceSynchronize(),这将等待当前父kernel的当前block(请注意当前父kernel可能也是其他kernel的子kernel, 父-子kernel是一个相对的概念, 之前说过)中, 所启动的所有子kernel结束. 这里需要注意两点:

(1)这种同步等待是阻塞的, 也就是说, 一旦父kernel的某个block中有线程执行了这种同步(设备端的cudaDeviceSynchornize()函数), 则会导致阻塞住执行, 也就是暂停了. (2)并不存在Host上的polling操作(轮询/查询),一旦你决定同步, 就需要作出会被暂停的风险(也可能不会, 看当时的硬件资源和kernel执行情况),而不能像Host上那样, 只是查询一下启动的kernel们是否还是执行中, 或者已经执行完了, 这点做不到的. 这两点的特别行为主要是收限于: 动态并行的时候, 父-子kernel, 均是通过GPU执行的; 而Host上启动, host线程却是由CPU来执行的, CPU不会再执行kernel. 这导致了前者(动态并行的时候), 父kernel轮询子kernel的执行状况可能无意义 --- 父kernel有时间浪费GPU来polling子kernel的状态, 不如直接暂停它(的blocks), 全力执行子kernel更好. 而传统的Host上启动kernel, CPU很可能命令GPU干活后, 就无所事事了, 所以轮询提供一下也不错. 这点需要注意. 所以虽然你看CUDA的动态并行设计比较简单(只是一个cuda runtime api的子集), 但还是考虑过性能很多的, 没有必要的特性, 不提供, 也简化了用户. 和这个相反的是, OpenCL的"设备端"可以同时指GPU和CPU, 也就是可能存在OpenCL上, CPU同时要负责执行kernel, 还要负责调度的情况。 此时就很尴尬了。特别是任何对自己的轮询, 特别是非常频繁无节制的, 均将影响CPU自身上的kernel执行性能. 特别的, 还影响UI交互的实时性.所以OpenCL用户经常不知道为何, 性能就剧烈下降了("我似乎也没有执行神马操作啊, 怎么会这样")。

而为了解决这种问题, OpenCL又引入了CPU设备的拆分(fission)功能,允许将CPU拆分出来(例如7/8)做为设备端执行kernel, 其他1/8用来当成Host端执行调度.但这极大的增加了代码逻辑的复杂性,开发团队苦不堪言,所以之前为何说, CUDA能3天入门, OpenCL需要30天. 这里也可以做为另外一个辅佐的例子,大致这样. 最后需要说明的是, 动态并行时候的默认流(也叫0流, 或者未命名流), 行为和Host端上的任何一种默认流都不同.这样实际上存在3种默认流。 (1)经典默认流. CUDA从最开始就提供这种. 带有全局自动隐式同步功能(per device); (2)新默认流. 需要用户要求启用. 这种是每个Host线程单独一个的. 共享NULL或者0流的称呼; (3)动态并行时候的默认流. 这有点类似(2), 但实际上是每个block都有一个的. 同样所有的这些默认流, 共享一个名字称呼。 用户在使用的时候, 需要这点不同。

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

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

本文参与 腾讯云自媒体分享计划,分享自微信公众号。
原始发表:2018-08-10,如有侵权请联系 cloudcommunity@tencent.com 删除

本文分享自 GPUS开发者 微信公众号,前往查看

如有侵权,请联系 cloudcommunity@tencent.com 删除。

本文参与 腾讯云自媒体分享计划  ,欢迎热爱写作的你一起参与!

评论
登录后参与评论
0 条评论
热度
最新
推荐阅读
目录
  • D.3.1.2. Streams
  • D.3.1.2.1. The Implicit (NULL) Stream
领券
问题归档专栏文章快讯文章归档关键词归档开发者手册归档开发者手册 Section 归档