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

相关文章

来自专栏IT笔记

SpringBoot开发案例之整合mail队列篇

? 科帮网邮件队列.png 前言 前段时间搞了个SpringBoot开发案例之整合mail发送服务,也是基于目前各项目平台的邮件发送功能做一个抽离和整合。 问...

5147
来自专栏从流域到海域

如何在Mule 4 Beta中实现自动流式传输

原文地址:https://dzone.com/articles/how-automatic-streaming-in-mule-4-beta-works

1855
来自专栏程序员宝库

教你在不使用框架的情况下也能写出现代化 PHP 代码

我为你们准备了一个富有挑战性的事情。接下来你们将以无框架的方式开启一个项目之旅。 首先声明, 这篇并非又臭又长的反框架裹脚布文章。也不是推销非原创(https:...

3195
来自专栏FreeBuf

滥用Edge浏览器的“恶意站点警告”特性,实现地址栏欺骗

前言 在过去的几个月里,我们看到使用这种以技术支撑的骗术日益增多,用户的浏览器会被辣眼睛的红屏以及类似”你的电脑可能存在风险”的提示消息”锁定”。当然,这种情形...

2109
来自专栏Albert陈凯

手把手教你用python抓取网页导入模块 urllib2随便查询一篇文章,比如On random graph。对每一个查询googlescholar都有一个url,这个url形成的规则是要自己分析的。

http://www.1point3acres.com/bbs/thread-83337-1-1.html **前言: ** 数据科学越来越火了,网页是数据...

2967
来自专栏木东居士的专栏

从0写一个爬虫,爬取500w好友关系数据

5586
来自专栏Golang语言社区

Go 语言编写的缓存及缓存过滤库:groupcache

groupcache 是使用 Go 语言编写的缓存及缓存过滤库,作为 memcached 许多场景下的替代版本。 对比原始 memcached ? 首先,gr...

2484
来自专栏生信宝典

上传高通量测序原始文件

在我们发表高通量测序文章之前通常要上传测序数据到GEO数据库,现总结流程如下。 注册账户、填写MetaSheet 在NCBI GEO官网注册一个账号,然后登陆。...

2589
来自专栏Vamei实验室

被解放的姜戈01 初试天涯

Django是Python下的一款网络服务器框架。Python下有许多款不同的框架。Django是重量级选手中最有代表性的一位。许多成功的网站和APP都基于Dj...

1.7K6
来自专栏架构师之旅

关于Socket高并发的原理介绍及使用Apache Mina带来线上的问题分析

上周在线上出现了一个很低级的问题,但是正是这个低级的问题引起了我的兴趣,其实所谓的低级是因为配置文件配置错了,原本线上是为每个客户端设置了一个席位,就说是客户端...

1833

扫码关注云+社区