DAY70：阅读API Reference
我们正带领大家开始阅读英文的《CUDA C Programming Guide》,今天是第70天，我们正在讲解CUDA 动态并行，希望在接下来的30天里，您可以学习到原汁原味的CUDA，同时能养成英文阅读的习惯。
关注微信公众号，查看历史信息，可以看到之前的阅读
本文共计171字，阅读时间12分钟
前情回顾：
DAY61：关于Cooperative Groups
DAY65:阅读Device-Side Kernel Launch
DAY67：阅读阅读Events、Synchronization和Device Management
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上发帖