DAY65:阅读Device-Side Kernel Launch

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

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

本文共计507字,阅读时间15分钟

D.3. Programming Interface

D.3.1. CUDA C/C++ Reference

This section describes changes and additions to the CUDA C/C++ language extensions for supporting Dynamic Parallelism.

The language interface and API available to CUDA kernels using CUDA C/C++ for Dynamic Parallelism, referred to as the Device Runtime, is substantially like that of the CUDA Runtime API available on the host. Where possible the syntax and semantics of the CUDA Runtime API have been retained in order to facilitate ease of code reuse for routines that may run in either the host or device environments.

As with all code in CUDA C/C++, the APIs and code outlined here is per-thread code. This enables each thread to make unique, dynamic decisions regarding what kernel or operation to execute next. There are no synchronization requirements between threads within a block to execute any of the provided device runtime APIs, which enables the device runtime API functions to be called in arbitrarily divergent kernel code without deadlock.

D.3.1.1. Device-Side Kernel Launch

Kernels may be launched from the device using the standard CUDA <<< >>> syntax:

kernel_name<<< Dg, Db, Ns, S >>>([kernel arguments]);
  • Dg is of type dim3 and specifies the dimensions and size of the grid
  • Db is of type dim3 and specifies the dimensions and size of each thread block
  • Ns is of type size_t and specifies the number of bytes of shared memory that is dynamically allocated per thread block for this call and addition to statically allocated memory. Ns is an optional argument that defaults to 0.
  • S is of type cudaStream_t and specifies the stream associated with this call. The stream must have been allocated in the same thread block where the call is being made. S is an optional argument that defaults to 0.

D.3.1.1.1. Launches are Asynchronous

Identical to host-side launches, all device-side kernel launches are asynchronous with respect to the launching thread. That is to say, the <<<>>> launch command will return immediately and the launching thread will continue to execute until it hits an explicit launch-synchronization point such as cudaDeviceSynchronize(). The grid launch is posted to the device and will execute independently of the parent thread. The child grid may begin execution at any time after launch, but is not guaranteed to begin execution until the launching thread reaches an explicit launch-synchronization point.

D.3.1.1.2. Launch Environment Configuration

All global device configuration settings (e.g., shared memory and L1 cache size as returned from cudaDeviceGetCacheConfig(), and device limits returned from cudaDeviceGetLimit()) will be inherited from the parent. That is to say if, when the parent is launched, execution is configured globally for 16k of shared memory and 48k of L1 cache, then the child's execution state will be configured identically. Likewise, device limits such as stack size will remain as-configured.

For host-launched kernels, per-kernel configurations set from the host will take precedence over the global setting. These configurations will be used when the kernel is launched from the device as well. It is not possible to reconfigure a kernel's environment from the device.

本文备注/经验分享:

今天这章节比较简单, 类似之前的Host上的启动配置章节.首先上去的段落说明了, 在CUDA C里面的调用, 和之前的Host上的调用是非常相似的; 而能在device端使用的CUDA Runtime API函数(Device Runtime), 也非常相似Host上的CUDA Runtime API, 只不过是一个功能的子集.注意手册后面还会介绍如何从PTX中调用. 但这里没说.这适合那些使用较大的篇幅的PTX代码写东西的人. 规避了动态并行调用的时候, 临时将代码切换成CUDA C的麻烦.因为常规的使用PTX, 要么是完全PTX, 要么是CUDA C/C++里面嵌套(inline)PTX,以前的常规运算还好办, 需要动态并行的时候, PTX也必须要作出改进. 否则不能完全PTX, 而只能选择后者--也就是基本的一个空壳的CUDA C的kernel, 里面除了动态并行启动kernel的地方都是PTX. 这样很不方便.而具体的CUDA C里面的动态并行的<<<>>>语法, 和Host上的基本一致.

需要注意这里的两个地方:

第三个参数, 动态分配的shared memory大小(还记得动态和静态分配的shared memory吗? 这里的动态和动态并行无关), 依赖于硬件上的shared memory大小配置.很多计算能力上是可调的, 另外一些计算能力不可调.2.x和Kepler上是可调的(64KB一共大小, 可选16KB Shared + 48KB L1, 或者48KB shared + 16KB L1),然后到了Maxwell开始, 变得不可调, 固定64KB shared memory.到了Pascal,维持了这点, 依然是64KB固定大小.等到了7.X又改了回去, 变成多种大小可调(最大到96KB)。

本章节强调了, 因为动态并行的时候, 你只能进行当前配置好的(从Host上)大小下的shared memory动态分配,而不能进行大小配置修改(这点很好理解, 假设父kernel要求大shared memory, 然后子kernel要求将shared memory调小, 而动态并行很又可能是父子kernel同时在执行中的, 那么这种调节后, 父kernel还怎么运行),所以不能进行大小配置修改的.因为目前9.2支持的计算能力中, 只有Kepler和7.X是可调的,因此这个特性只会对3.x和7.x的用户造成困扰,而对5.X和6.x的用户来说, 和Host上毫无区别的.另外一点则是说, 动态并行启动的kernel是相对父kernel异步执行的.异步这里就是说, 无固定的顺序,可能包括父kernel执行完毕后, 然后子kernel; 也可能子kernel先执行完毕后, 然后再父kernel; 或者同时执行.所以如果需要使用动态并行, 然后还需要同时严格的前后执行顺序的话.你只能选择从Host启动一个实际上不干活的kernel(一个wrapper kernel),里面启动一堆需要使用动态并行的子kernel们, 然后在需要严格的前后顺序的某两个kernel后,使用stream控制一下,也就是说这种情况下, 你不能直接从host上直接启动一个上就真正干活的kernel的. 而是启动一个控制者.论坛前几天的有人问动态并行, 虽然她的那个问题, 最终无法用动态并行提速,但是她的做法(启动了一个wrapper kernel), 还是很值得参考的.

这是今天的主要内容. 非常简单.其实和Host上几乎完全一样.特定的计算能力(Maxwell/Pascal)更是几乎毫无差别.但需要补充的是:

从Device上动态并行启动kernel的代价, 在启动只有少量几个kernel的时候, 性能不一定比host好.例如很多计算能力上的动态并行启动kernel的启动延迟都要比host上的高.我们还没有测试7.X上的情况. 欢迎用户自行测试反馈.

(2)从Host上适合启动一个较大blocks/线程规模的kernel, 而从device上适合用动态并行启动较多的小规模的kernel. 后者在这种情况下还是很有利的. 往往能提升性能.注意这里特别适合kernel编译出来的代码体积较小(可以用cuobjdump看), 或者干脆只有1个kernel, 在用动态并行递归调用的情况.这种情况对L1 I-Cache有利(大约在Maxwell到Pascal上在8KB-16KB之间, 不是一个精确的测试)。

此外还需要补充的是:

1.启动较多体积较大的kernel, 又可能造成I-Cache的负担, 最终结果(性能提升或者下降), 则需要以实际为准.

2.计算能力7.X上, 需要使用动态分配才能用到96KB的shared memory(静态分配的, 也就是写死在代码里的, 依然是收到48KB/Block的shared memory大小限制), 但在动态并行的时候的动态shared memory分配, 结果会如何.目前尚不清楚. 手册也没说.实际上, 手册这章节自从有了动态并行(大约在CUDA 5到CUDA 6左右),都是基本上按照章节更新的,例如多了unified memory, 则只增加unified memory章节, 而对其他章节基本改动不大.因此新增了计算能力7.X后, 只里的使用稍微需要注意一下.

类似的, 在动态并行里面使用async的cudaMemcpy*(), 原本的手册在没有unified memory的时候,只是说明, 只能进行device memory(显存) -> device memory的操作.但实际上有了unified memory后, 可以支持unified memory的存储, 而不一定是显存. 但是动态并行章节依然没有更新.所以这种情况, 用户需要思考一下在新卡上的使用范围(例如是否扩大了?), 特别是对于有新卡的用户, 以及, 我们即将喜迎下一代卡的问世(大约在1-2个月内),所以在阅读本手册的时候, 不妨考虑一下很近或者已经到来的未来.

3.启动是异步的里面的说法, 说唯一明确的执行关系, 就是在父kernel在执行cudaDeviceSynchronize()等待子kernel的时候,只有这个时候的执行关系才是明确的(父kernel的block肯定暂停执行, 子kernel肯定开始执行).这个实际上上次说过了. 因此这里就不再重复了. 但用户需要注意一下.

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

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

原文发布于微信公众号 - 吉浦迅科技(gpusolution)

原文发表时间:2018-08-08

本文参与腾讯云自媒体分享计划,欢迎正在阅读的你也加入,一起分享。

发表于

我来说两句

0 条评论
登录 后参与评论

相关文章

来自专栏Golang语言社区

高可用性、负载均衡的mysql集群解决方案

一、mysql的市场占有率 二、mysql为什么受到如此的欢迎 三、mysql数据库系统的优缺点 四、网络服务器的需求 五、什么是mysql的集群 六、什么是负...

4945
来自专栏编程一生

Redis和消息队列使用实战

1123
来自专栏Coding01

跟着《架构探险》学轻量级微服务架构 (二)

上一篇主要简单搭建了 Spring Boot 框架,写了一个简单的路由/hello,Spring Boot 的其它功能根据后续的学习,再不断完善,接下来我们开始...

442
来自专栏一名叫大蕉的程序员

合格的配置中心应有的素养No.76

最近在看配置中心的一些设计,好像基本都是五花八门,主要看的是还是携程的 Apollo 这个开源的配置中心项目。一直以来都觉得配置中心很重要,因为这对于灰度发布,...

1698
来自专栏熊二哥

Hadoop快速入门

传说中的Hadoop,我终于来对着你唱"征服"了,好可爱的小象,!J 总的来说,hadoop的思路比较简单(map-reduce),就是将任务分开进行,最后汇总...

4666
来自专栏大魏分享(微信公众号:david-share)

上酒桌,今儿咱保证把容器SDN的这事一把唠清楚!

容器的SDN 很多人都说2017将是容器年,大卫也这么认为。但在很长一点时间里,容器与虚拟化都是相互依存,相互补充的问题。 之前笔者发表过一篇文章,放开眼界,看...

2665
来自专栏数据和云

Oracle 12.2新特性掌上手册 - 第五卷 RAC and Grid

编辑手记:RAC是Oracle最重要的高可用架构之一,具有扩展性良好、实现负载均衡等多维度的优势,Oracle RAC提供了相应的集群软件和存储管理软件,今天我...

2664
来自专栏北京马哥教育

Linux 服务器性能出问题,排查下这些参数指标

一个基于 Linux 操作系统的服务器运行的同时,也会表征出各种各样参数信息。通常来说运维人员、系统管理员会对这些数据会极为敏感,但是这些参数对于开发者来说也十...

4165
来自专栏玩转JavaEE

MongoDB副本集配置

上篇文章我们搭建了MongoDB副本集的环境,验证了数据已经可以成功的复制,本文我们就来看看MongoDB副本集的其他操作。 ---- 环境准备 三台服务器,地...

2775
来自专栏ZKEASOFT

.Net Core内存回收模式及性能测试对比

Server GC / Workstation GC

24411

扫码关注云+社区