DAY70:阅读API Reference

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

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

本文共计171字,阅读时间12分钟

前情回顾:

DAY61:关于Cooperative Groups

DAY62:阅读Glossary

DAY63:阅读Execution Environment

DAY64:阅读 Memory Model

DAY65:阅读Device-Side Kernel Launch

DAY66:阅读Streams

DAY67:阅读阅读Events、Synchronization和Device Management

DAY68:阅读 Memory Declarations

DAY69:阅读API Errors and Launch Failures

D.3.1.8. API Reference

The portions of the CUDA Runtime API supported in the device runtime are detailed here. Host and device runtime APIs have identical syntax; semantics are the same except where indicated. The table below provides an overview of the API relative to the version available from the host.

Table 5. Supported API Functions

Runtime API Functions

Details

cudaDeviceSynchronize

Synchronizes on work launched from thread's own block only

cudaDeviceGetCacheConfig

cudaDeviceGetLimit

cudaGetLastError

Last error is per-thread state, not per-block state

cudaPeekAtLastError

cudaGetErrorString

cudaGetDeviceCount

cudaDeviceGetAttribute

Will return attributes for any device

cudaGetDevice

Always returns current device ID as would be seen from host

cudaStreamCreateWithFlags

Must pass cudaStreamNonBlocking flag

cudaStreamDestroy

cudaStreamWaitEvent

cudaEventCreateWithFlags

Must pass cudaEventDisableTiming flag

cudaEventRecord

cudaEventDestroy

cudaFuncGetAttributes

cudaMemcpyAsync

Notes about all memcpy/memset functions:Only async memcpy/set functions are supportedOnly device-to-device memcpy is permittedMay not pass in local or shared memory pointers

cudaMemcpy2DAsync

cudaMemcpy3DAsync

cudaMemsetAsync

cudaMemset2DAsync

cudaMemset3DAsync

cudaRuntimeGetVersion

cudaMalloc

May not call cudaFree on the device on a pointer created on the host, and vice-versa

cudaFree

cudaOccupancyMaxActiveBlocksPerMultiprocessor

cudaOccupancyMaxPotentialBlockSize

cudaOccupancyMaxPotentialBlockSizeVariableSMem

本文备注/经验分享:

今天章节列出了所有能再动态并行的时候, 在设备端调用的CUDA Runtime API函数.请注意使用这些函数需要链接cuda的devrt(Device Runtime)库, 包括手册上也这么说. 但实际上从某个版本的CUDA开始(可能是CUDA 6?), 链接变成的自动的. 用户是否手工指定均不妨碍. 手册还没有更改, 但用户很容易发现这点. 这个库就是提供了这个表格里面的所有函数. 在我们实际编译的时候, 因为CUDA C语言写的源kernel代码, 在最终编译成底层GPU汇编的时候(SASS), 中间会有一个PTX的公开通用GPU虚拟机层次.使用这些表格的函数的时候, 你会看到你的PTX里面, 生成了一些占位用的, 用.weak标记的空白同名函数.这些看到的函数在最终生成目标代码的时候, 会和实际的设备端runtime链接.也就是说, 你看不到设备端runtime函数的PTX代码的.如果想学习参考一下它们是怎么实现的, 可以用cuobjdump --dump-sass看一下它们的最终代码方可看到.注意PTX里面的.weak链接标志是允许链接器(nvlink或者其他的)能在最终链接的时候替换掉它们, 很多linker都有这个支持.这是这些函数的存在性问题.也就是使用这些函数均会需要一段NV提供的, 在GPU上的辅助代码的,这些代码最终会和你的kernel链接在一起, 称为你最终运行的代码的一部分. 使用它们有一定的代价.而回到这些函数的使用上来说,则是基本我们之前说的, 它们只是完整对应Host上同名函数的功能子集, 以及, 部分这些函数具有轻微区别.实际上这些区别已经在之前的章节说过了, 本章节是做一次汇总. 例如从功能的子集上来说: 只有一些关于设备管理的含有Get名字的Runtime API, 却缺少对应的Set版本.例如你可以GetLimit, 却不能Set它.例如你可以GetDevice, 却同样不能Set另外一个GPU设备.例如你可以Get缓存设备, 却不能设定它.等等.设备管理之外的, 用来进行kernel启动的逻辑控制和检测的Stream和Event,也同样相比Host上的完整版本缩减了.例如不能有一个全局的(per device的实际)默认能同步所有工作的默认流(动态并行里面的默认流是per block的)。例如Event不用用来计时.例如Event和Stream都不能用来单独同步(只能整体cudaDeviceSynchronize(), 而不能cudaStreamSynchornize(), cudaEventSynchronize()的. 这三者是我们常说的3大阻塞同步函数, 在动态并行的时候, 你只能使用第一个.而第一个也只能同步本block内部的启动的子kernel们, 而管不到其他blocks启动的其他子kernel们)。 类似的, 非阻塞同步里面的cudaStream和EventQuery也都不存在.你会看到功能的确少了很多. 但这些一般足够用了.需要完整的控制能力, 必须从CPU进行控制.

CPU我们之前说过相当于CEO,GPU相当于奴隶, 或者员工,员工里面可以有小组长, 管理部分任务,但显然它不能负责公司的全面运营.所以这些动态并行时候, 设备端Runtime的缩减可以理解.因为动态并行主要是进行kernel的动态启动, 所以这些函数比较专精于这些方面. 其他的基本木有.cudaEventRecord和cudaStreamWaitEvent这个需要说一下.一般这个是分开在2个流中使用的。后者cudaStreamWaitEvent(), 要求一个流中的任务在前者Event被Record之前, 不能运行.但是前者并不能阻塞调用者(父kernel, 也包括从Host上调用此函数时候的CPU)的运行,该函数会立刻返回. 只是对特定的流中的特定位置做了标记而已。以前这函数无论从设备端还是从Host上, 都对使用者造成了迷惑.主要原因是, 手册说到该函数的时候, 语焉不详, 这里特别强调一下.注意, 还有一些函数, (例如Stream和Event的创建)强制需要一些标志, 来限制自己的功能,达成和Host上具有此标志的同名函数同等功能的效果.因为设备端动态并行的时候, 只有这种较少的功能, 例如Event必须用禁用计时功能的标志创建。这种方式维持了和Host上的表面上看起来的统一(语义上的统一)。避免造成代码维护者(你写完代码后走人了, 的其他负责给你打扫尾巴的同事)误以为这些函数还有这些功能.这样明确的要求你写出这些标志才能调用这些设备端函数, 有助于你直接知道不能如何如何. 此外, cudaMalloc/cudaFree和之前的非动态并行的时候,从Fermi 2.0开始引入的设备端的malloc/free一样,都只能用设备端的释放函数释放分配到的指针.不能混合调用的,因为时两个独立的heap,这个之前曾经稍微提到过一点,有一个小heap的. runtime api初始化后(一般在你调用大部分的host端的runtime api, 首次调用的时候有个显著的延迟,也就是说很多人抱怨的为何cuda第一次调用这么慢的原因), 这个小heap就会分配好空间,所以你也不能从设备端分配过大的缓冲区. 会失败的。 注意这样实际上从设备端分配global memory, 具有2个版本.一个是从Fermi 2.0+开始的, malloc/free(没有cuda字样开头);另外一个则是从3.5+开始的, cudaMalloc/cudaFree,这两个实际上是一样的.但后者的使用, 需要你链接设备端的runtime, 只能用在动态并行的时候.而前者可以用在老卡, 例如2.x上和3.0上(第一代的kepler),并不需要具有动态并行支持. 这个表格很短小,建议直接背下来. 注意这个表格里面, 有一些函数完全没有说明.这种则是表明和Host上的对应的同名函数, 完全一样.可以直接按照Host上的方式调用.很方便.CUDA在引入动态并行的时候的一个设计目标就是尽量让你不需要学习新东西.所以这些设备端的runtime api被设计的如此相似.这是一个好事, 也是一个坏事.好事是的确几乎可以0成本入门动态并行(假设你之前有过普通host上的cuda经验),坏处则是很多地方容易造成迷惑.不过用常了也就习惯了.

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

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

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

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

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

发表于

我来说两句

0 条评论
登录 后参与评论

相关文章

来自专栏顶级程序员

干货 | Python 爬虫的工具列表大全

源 / 伯乐头条 这个列表包含与网页抓取和数据处理的Python库。 网络 通用 urllib -网络库(stdlib)。 requests -网络库。 gr...

3726
来自专栏Python中文社区

如何使用爬虫做一个网站

大家如果有兴趣做网站,在买了VPS,部署了wordpress,配置LNMP环境,折腾一番却发现内容提供是一个大问题,往往会在建站的大(da)道(...

2705
来自专栏java学习

学习java需要会哪些知识才能够去应聘工作?

Java基础 | 数据库 | Android | 学习视频 | 学习资料下载 按照我去培训机构的学习经历,给初学还有自学Java 的同学一个基本的学习脉络,希望...

3616
来自专栏java学习

学习java需要会哪些知识才能够去应聘工作?

按照我去培训机构的学习经历,给初学还有自学Java 的同学一个基本的学习脉络,希望对大家有帮助。 不建议找到一本书死啃,没啥用,不要有这一页看不明白我就不往下看...

29410
来自专栏Linux Python 加油站

Python学习干货 史上最全的 Python 爬虫工具列表大全

链接:https://mp.weixin.qq.com/s/UkXT20Oko6oYbeo7zavCNA

1602
来自专栏风火数据

hadoop大数据面试题

以下资料来源于互联网,很多都是面试者们去面试的时候遇到的问题,我对其中有的问题做了稍许的修改了回答了部分空白的问题,其中里面有些考题出的的确不是很好,但是也不乏...

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

关于Java健壮性的一些思考与实践 No.102

上面两种模式都可以实现标准的 response 的封装,那么具体要封装哪些东西呢?其实最主要的就是统一的 try catch,防止出现任何的 500 错误给到调...

1012
来自专栏非著名程序员

推荐几个非常有用的开发工具之Android Studio插件

推荐几个非常有用的开发工具 非著名程序员 我们都知道Eclipse开发Android将在今年年底google不再继续提供相应的开发支持,转而开始强烈发展Andr...

1938
来自专栏CDA数据分析师

Python 爬虫的工具列表

这个列表包含与网页抓取和数据处理的Python库 网络 通用 urllib -网络库(stdlib)。 requests -网络库。 grab – 网络库(基于...

36410
来自专栏架构师小秘圈

你所不知道的库存超限做法

作者:程序诗人,来自:cnblogs.com/scy251147 零,题记 在互联网企业中,限购的做法,多种多样,有的别出心裁,有的因循守旧,但是种种做法皆想达...

3346

扫码关注云+社区