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 条评论
登录 后参与评论

相关文章

来自专栏LEo的网络日志

go获取机器的mac地址和ip

3796
来自专栏大数据智能实战

pytorch实践中module 'torch' has no attribute 'form_numpy'问题的解决

最近开始仔细玩了一下pytorch,发现里面有个BUG之前都没有发现。 在测试torch最基本的示例的情况下,居然碰到了个pytorch无法转化numpy为Te...

4179
来自专栏FreeBuf

Windows Server 2012 R2的提权过程解析

近期,我在进行一项安全评估的过程中遇到了一个麻烦。这是某个组织的一台远程桌面服务器,安装的是Windows Server 2012 R2系统,但是我手中的用户账...

35510
来自专栏java一日一条

大量参数与信息丢失之间不可不说的故事

代码越少就越好?对象越少就越好?这些都是真的吗?由绝大多数情况来看,这还真的都不一定。

371
来自专栏智能大石头

微软的软件授权及保护服务(SLPS)试用分析

这些天都在绞尽脑汁地想怎么样设计一个授权方式来保护我的组件,今天看了一下同事从广州带回来的Tech2007的讲稿,里面提到了 微软的软件授权及保护服务(SLPS...

1988
来自专栏QQ会员技术团队的专栏

Android 动态库压缩壳的实现

计算机软件领域所说的壳实际上是一种软件加密技术。壳主要分为两大类:加密壳和压缩壳,加密壳侧重于防止软件被篡改,而压缩壳则侧重于减小软件体积。其实,在Window...

1.4K1
来自专栏区块链技术专栏

EOSIO 智能合约数据库演练

对于数据库,最重要的功能就是查询,如果没有查询功能,数据库里的数据就不能呈现,也就没有意义。查询数据库主要分为两方面,一方面是主键查询,一方面是通过二级索引查询...

4507
来自专栏轮子工厂

1. C语言的第一个程序

(。・∀・)ノ゙嗨!大家好,我是呆博~很开心可以在这里给大家分享我的 C 语言学习笔记~

934
来自专栏小李刀刀的专栏

容易被误解的overflow:hidden

为了页面的健壮性,我们常常需要使用overflow:hidden。有时候是为了防止布局被撑开,有时候是为了配合其它规则实现文字截断,还有时候纯粹是为了创建块级上...

35511
来自专栏about云

日志分析实战之清洗日志小实例7:查看样本数据,保存统计数据到文件

问题导读 1.如何从所有数据中,抽取样本查看? 2.如何保存结果到hdfs? 3.saveAsTextFile的作用是什么? 上一篇 日志分析实战之清洗...

2785

扫码关注云+社区