DAY69：阅读API Errors and Launch Failures

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

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

本文共计344字，阅读时间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

D.3.1.7. API Errors and Launch Failures

As usual for the CUDA runtime, any function may return an error code. The last error code returned is recorded and may be retrieved via the cudaGetLastError() call. Errors are recorded per-thread, so that each thread can identify the most recent error that it has generated. The error code is of type cudaError_t.

Similar to a host-side launch, device-side launches may fail for many reasons (invalid arguments, etc). The user must call cudaGetLastError() to determine if a launch generated an error, however lack of an error after launch does not imply the child kernel completed successfully.

For device-side exceptions, e.g., access to an invalid address, an error in a child grid will be returned to the host instead of being returned by the parent's call to cudaDeviceSynchronize().

D.3.1.7.1. Launch Setup APIs

Kernel launch is a system-level mechanism exposed through the device runtime library, and as such is available directly from PTX via the underlying cudaGetParameterBuffer() andcudaLaunchDevice() APIs. It is permitted for a CUDA application to call these APIs itself, with the same requirements as for PTX. In both cases, the user is then responsible for correctly populating all necessary data structures in the correct format according to specification. Backwards compatibility is guaranteed in these data structures.

As with host-side launch, the device-side operator <<<>>> maps to underlying kernel launch APIs. This is so that users targeting PTX will be able to enact a launch, and so that the compiler front-end can translate <<<>>> into these calls.

The APIs for these launch functions are different to those of the CUDA Runtime API, and are defined as follows:

本文备注/经验分享：

今天的章节主要是两个内容, 一个是动态并行时候的子kernel的出错的处理. 另外一个则是不使用CUDA Runtime API的特色的<<<>>>语法, 如何启动子kernel。 先说一下前面这个, 子kernel的错误处理实际上分成两个部分, 一个是启动kernel后(用<<<>>>, 或者本章的后半部分的方法), 需要立刻检测是否能启动; 另外一个部分则是, 如果子kernel能启动, 则在后续的运行中, 发生错误(例如子kernel运行30ms后挂了)如何检测. 这两个部分实际上是在Host上也是需要注意的问题, 特别是很多CUDA老用户都不知道应当怎么正确处理。 我先说一下Host上, Host上需要对1个kernel, 检测这两点, 1点是kernel启动后立刻跟随cudaError_t r = cudaGetLastError()来检测kernel能否启动. 然后如果这个r值不是cudaSuccess, 那么kernel就完全没有启动. 如果是cudaSuccess, 则kernel已经开始启动了, 但这不代表后续的一段时间内, kernel能稳定运行, 需要在下次进行同步调用之类的场合(例如cudaMemcpy和cudaDeviceSynchronize()之类的)检查是否返回了之前的kernel异步执行错误. 根据论坛上的诸多用户的问题, 往往集中在两点: (1)新用户遇到后续的同步调用之类的场合, cudaMemcpy之类的, 返回的之前的kernel执行期间导致的错误, 会感到迷惑. 例如她们会感觉这是cudaMemcpy自己出错了, 而没有意识到这是之前的错误, 因为CUDA有这个特性, 之前的kernel出错后, 后面的所有API都会从检测到出错的时候(往往是后续第一个同步调用), 持续的返回该错误.新人需要注意到CUDA的这个特性, 否则往往找不对是哪里的问题(论坛上屡见不鲜）。 (2) kernel启动后, 必须在<<<>>>的下一行立刻跟随检查, kernel是否启动成功. 也就是说, 实际上一个kernel需要检查两个地方: <<<>>>启动后立刻检测是否能启动. 和能启动后续的异步执行期间(无论对于Host还是动态并行都是异步的)是否能持续运行不出错. 第二点是老手容易忽略的问题. 特别是一个项目离开了开发环境(例如你的开发机器), 往往会可能直接启动都不能启动. 这是两点. 而在动态并行的时候, 其中后一点和Host上是一样的, 都需要立刻跟随检查.而前一点则不同: 动态并行一旦子kernel挂掉后(例如子kernel胡乱访存, 将显存乱写一通), 因为父子kernel都在同一个卡上运行, 父kernel的正常运行环境(例如需要的显存内容), 已经被挂掉的子kernel破坏掉了,此时无法像Host上那样, 通过后续的同步之类的位置, 能返回异步错误检测到子kernel挂了, 这个做不到了.因为此时, 父kernel已经不能正常运行了, 同样的挂了(CPU上能这样检测是因为CPU和GPU分开运行, GPU上的运行环境(CUDA Context)挂掉后, CPU上面的完好).此时唯一的结果是父kernel和所有子kernel一起整体挂掉, 然后返回给Host一个统一的出错返回值(在后续的CPU上的最近同步调用的检测处), 也就是说, 动态并行的子kernel运行期间出错只能从Host上检测到. 而不能在上一级检测到.这点需要注意. 以及, 需要说明的是, 曾经见过一些kernel挂的很快, 导致上面说的两处检测的第一处都会直接挂掉,例如<<<>>>后面的下一行立刻就已经完蛋了. 此时不代表kernel不能启动, 而是代表它速挂了.这种情况可以从<<<<>>>下面的下一行的cudaGetLastError()处检测返回的类型,如果只是说, 启动配置参数无效(例如你的过大的block形状, 在特定的卡上), 那么是不能启动.如果上去得到ULP(unspeified launch failure)或者cudaErrorUnknown或者其他类似错误, 证明kernel快速的挂掉了(快速到刚刚从<<<>>>返回).这点需要注意判断.因为绝大部分kernel总是因为访存挂掉的,我建议在遇到出错的情况的时候, 用打开了cuda memory checker的nsight, 或者cuda-memcheck来检测具体是哪个kernel出现问题.避免找不到地方.

本章节的第二部分比较有特色. 是说如何不使用<<<>>>来启动kernel的.不使用<<<>>>实际上是完全可以的, 例如Host上的CUDA Driver API启动kernel的时候, 就不能使用<<<>>>来启动kernel(因为做不到啊，Driver API被设计成跨语言的, 不要求Host语言也是CUDA C, 例如从其他第三方软件中，或者第三方开发环境中, 例如C#, golang之类，此时这些环境或者语言不能使用<<<>>>). 但是本章节的不使用<<<>>>是特意为动态并行而说明的, 也就是如何从父kernel中这样用, 而不是从Host代码中.实际上, Host CUDA Runtime API也可以不使用<<<>>>, 通过一个叫cudaLaunchKernel()的函数即可启动,但是本手册故意没有提到这点，却提到了为动态并行的时候如何使用它. 这需要涉及到两部分: (1)如何获取一个为kernel启动所准备的参数缓冲区.然后在这个缓冲区中, 按照一种特定的方式填充上参数. (2)用这个缓冲区, 外加特定的启动配置(例如启动形状, shared memory动态大小配置)来启动特定的kernel. 你会看到, 这个实际上就是<<启动配置<>>>(参数);的方式. 只是给拆分成两部分了.实际上这个才是真正的如何在动态并行的时候启动kernel, 而<<<>>>则会自动被CUDA编译器转换成这种真正的调用. 这里面需要有两点注意的: (1) 应当直接使用cudaGetParameterBuffer来获取参数缓冲区, 而不要尝试自己使用动态并行时候cudaMalloc()来分配(或者malloc)，因为后者可能代价更高. 前者可能代价较轻. (2) 参数的传递比较特殊, 需要在缓冲区中具有特殊的对齐格式. 因为手册的后续章节会说, 所以这里暂时不进行说明. 但是提醒一点, 用户总是可以通过查看一个动态并行的kernel的PTX代码, 来观察编译器是如何将<<<>>>改成这两个分布调用的时候, 进行填充的. 往往会有启发. 这个过程(参数填充)实际上对于来自CUDA 3.2之前的老用户, 是非常熟悉的过程, 因为很像当时流行的driver api的参数填充过程.这里就不说了, 手册也已经删除了相关内容。CUDA只所以为CUDA C在动态并行的时候提供这个, 是为了方便PTX用户，例如虽然说, 一些PTX用户的实际代码风格是:

__global__ void your_kernel(....) { asm嵌入的PTX.... .... 离开ptx, 使用<<<>>>>进行动态并行 继续asm嵌入的PTX.... ... }

这样就很尴尬了. 需要反复的离开进入asm模式, 就为了用动态并行.而提供了这两个函数, 则可以直接通过PTX进行调用. 非常方便.而CUDA C的对应版本, 则可以在你尝试从刚才的上面这个kernel外形中, 能快速的继续先保持CUDA C的动态并行启动,先改称从CUDA C里将<<<>>>改成参数缓冲区填充+启动.如果这一步进行顺利, 则可以完全脱离CUDA C,再将CUDA C的参数填充+启动继续改成PTX格式.这样逐步的完全能无障碍的迁移到PTX.而能全体PTX, 则方便了很多软件的二次开发.例如她们需要动态的生成kernel, 却因为授权或者其他原因, 不能随着携带一个巨大的CUDA Toolkit, 也不能要求用户总是具有例如VC之类的环境, 能随时调用nvcc工具链编译出PTX.这样可以完全脱离nvcc, 直接生成PTX即可(文本格式, 很方便).需要说明的是, 内置CUDA二次开发支持还有其他方式, 但是已经超过了本章节相关的内容了. 以及, 需要说明的是, cooperative group启动的时候有类似的风格, 但是本系列之前越过了它, 用户感兴趣可以回头看一下.

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

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