DAY54：阅读Assertion

B.18. Assertion

Assertion is only supported by devices of compute capability 2.x and higher. It is not supported on MacOS, regardless of the device, and loading a module that references the assert function on Mac OS will fail.

stops the kernel execution if expression is equal to zero. If the program is run within a debugger, this triggers a breakpoint and the debugger can be used to inspect the current state of the device. Otherwise, each thread for which expression is equal to zero prints a message to stderr after synchronization with the host via cudaDeviceSynchronize(), cudaStreamSynchronize(), orcudaEventSynchronize(). The format of this message is as follows:

Any subsequent host-side synchronization calls made for the same device will return cudaErrorAssert. No more commands can be sent to this device until cudaDeviceReset() is called to reinitialize the device.

If expression is different from zero, the kernel execution is unaffected.

For example, the following program from source file test.cu

will output:

test.cu:19: void testAssert(): block: [0,0,0], thread: [0,0,0] 
Assertion `should_be_one` failed.

Assertions are for debugging purposes. They can affect performance and it is therefore recommended to disable them in production code. They can be disabled at compile time by defining the NDEBUG preprocessor macro before including assert.h. Note that expression should not be an expression with side effects (something like (++i > 0), for example), otherwise disabling the assertion will affect the functionality of the code.

本文备注/经验分享：

今天是assert()，该函数即是所谓的"断言"函数. 此函数接受一个条件做为断言(assertion), 当条件不成立的时候, 异常的终止你的kernel。此时即是所谓的"断言失败"(assertion failed) 这个函数实际上有两个版本, 一个是CPU上用的版本, 一个是device(GPU)上用的版本.无论哪个版本都需要#include <assert.h> 这点比较特别.因为assert.h是host compiler提供的, 但是你在GPU上使用, 却依然要include它.否则无法通过编译.类似的还有printf。 后者也是GPU上使用的, 但同样如果你的代码不包含host端的stdio.h, 同样无法使用.这是使用的时候需要注意的. 其次, 如果你在Host上曾经用过此函数, 则知道此函数主要是用来辅助调试的.我们耳熟能详的某个程序突然崩溃了, 提示在某某地方断言失败, 然后可选的问你是否发送错误报告之类的,就是CPU版本的常见表现(这个对话框有多种形式的, 具体形式和系统, 以及系统级别的调试标志, 以及你安装的调试器--例如VS, WinDBG之类的有关)，而GPU版本的同样类似, 但不会弹出对话框, 只是会提示你, 某代码文件, 某行处, 断言失败云云.此时根据情况, 如果此CUDA软件是独立运行的, 直接就会彻底的挂掉, context被损坏, 完全无法继续, 必须重置设备或者重建context(后者对于driver api)，如果是在调试器下运行(例如nsight下, 或者cuda-gdb下),则会给你一次检查死亡现场的机会. 例如当前发生assertion失败时刻的局部变量, global memory, shared memory内容之类的信息.所以这是一个很好的调试工具. 精确的说, assert()它和printf()一起, 是在没有调试器(例如NSight)的时候, 能够现场调试的很称手的两个工具.很多人喜欢使用这种方式, 例如刚才说的后者的printf()大法.甚至有极端的客户, 当我们以前在QQ群进行现场调试指导的时候,被拒绝接受nsight, 而是坚持使用printf和assert,同时对我们的工程师破口谩骂. 以及, 这也是NV的OpenCL之类的, 完全无调试器的情况下, 很好的调试工具.(NV的NSight可以调试DX的shader, OpenGL的Shader, C++ AMP, CUDA, 但就是不能调试OpenCL.） assert很大程度上等价于, 你使用NSight的时候, 改变断点为条件断点，例如assert(condition)，可以直接很多情况下, 将condition设定为断点的条件.这两种方式各有利弊。 assert()因为是做为一个函数提供, 你使用它需要敲入它, 从而代码发生改变, 从而需要重新编译.有的大项目编译起来是个时间上的灾难(好在CUDA 9进一步提速了编译速度)，好处是你不需要有单独的调试器.而后者NSight之类则你无需改变代码, 直接设定条件断点即可, 但坏处是并非所有时候你可能手头都有NSight的. 此外, 需要说明的是,很多来自CPU的用户, 习惯大量对一些罕见事件, 大量的添加assert()，因为在CPU上的编程中, 该函数非常轻量, 几乎可以认为是无代价.但是在GPU上, CUDA中使用它, 却不同.这个函数(assert)和printf, 以及设备端的malloc, free一样,都是所谓的设备端系统调用(device-side system call)，如果你用cuobjdump观察代码, 会直接看到有类似call NULL(范意, 实际略有不同)的特殊调用.这种特别的系统调用代价很昂贵.所以以前如果有来自CPU上的大量使用assert(), 甚至是在生产发布的情况下的习惯的客户,不建议将此习惯保留到GPU上, 这是一个需要注意的地方.类似的printf也不应当使用, 同样代价昂贵(但比AMD的printf好, AMD以前的很多驱动版本, 使用printf等于自动禁用优化, 性能降低甚至1个数量级).同理论坛上说的malloc, 和free(设备端调用)同样应当避免. 都是一个道理. 这些是对断言函数的基本作用说明和注意事项.无太多需要说明的东西. 此外, 需要补充的是:手册上提供了一种方式(在#include你的assert.h之前), 定义NDEBUG宏, 从而能自动移除你所有代码中存在的assert()效果.(例如你可以在你的代码文件最前面定义,也可以在编译的时候通过-D方式指定, 都可以)，此外, 除了辅助调试, 这也是能快速中途异常终止你的kernel的一个很好的方式.但是需要注意的是, 一旦用这种方式终止了你的kernel, 下次就必须cudaDeviceReset()，而cudaDeviceReset()后, 再下次的任意Runtime API将引入隐式初始化(implicit initialization) 而这个是一个耗时的操作, 一般情况下请不要这样做, 除非必要.这某种方式和你用*(int *)NULL = 0;来中途快速终止(比block每个线程都return掉, 然后再下一个block被切换上SM, 再同样return掉要快)有点类似, 都是使用后, 当前cuda context被破坏了.此外, 考虑到很多时候, 需要一种能从host上通知device的方式, 好中途快速的终止kernel, 同时还能保留继续下次启动kernel的能力.我建议进行映射的内存上的标志轮询. 方式也很简单, 分配一段较小的page-locked host memory, 然后映射到device上,然后kernel做为一个volatile 参数接受(例如volatile int *p)， 然后host可以随时对该端内存设定标志, 例如1，然后GPU的kernel端可以随时读取p[0]判断是否为1, 如果是, 可以直接提前终止.该方法适用性较广.以及, 代价也较低, 当L2不是非常忙碌的时候,论询(system-memory polling)的一次延迟只有大约100-200ns，volatile关键字保证了跨PCI-E的传输.这是一种好点的方式. assert()其实人人都会用的.我只是说了一下可能的坑或者要点.

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

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