写点什么

从 Core Dump 中提取 CUDA 的报错信息

作者:OneFlow
  • 2022 年 9 月 01 日
    重庆
  • 本文字数:3223 字

    阅读完需:约 11 分钟

从Core Dump中提取CUDA的报错信息

近期,Meta AI 团队在生产 PyTorch AI 模型时遇到了一个难题。这一问题由 CUDA 非法内存访问引起,号称集结了 Meta 全公司最牛的 AI 工程师才搞定,这篇博客记录了他们使用 CUDA 的 core dump 来确定报错位置所使用的技巧和实践。


作者|Zachary DeVito

翻译|贾川、程浩源


如果 GPU 读取了无效内存,那么 CUDA API 将会开始从发生错误的地方开始,后续所有 API 调用都会返回 cudaErrorIllegalAddress:


设备在无效内存地址上使用了加载或存储指令。这使得进程处于不一致的状态,任何后续的 CUDA 工作都将返回相同的错误。若要继续使用 CUDA,进程必须终止并重新启动。


因为 CUDA kernel 是从 CPU 异步启动,所以在启动异常 kernel 的地方不会报告此错误,而是在 GPU 上实际发生异常并传播到 CPU 之后的任何 CUDA API 调用时报告此错误。

当然,要是使用 CUDA_LAUNCH_BLOCKING=1 环境变量,CUDA 就会在 kernel 启动后运行完成才返回,但这会使得程序运行明显变慢,可能会改变报错时机,以致某些不确定性问题不再被触发。

此外,如果有多个线程使用 CUDA API,cudaErrorIllegalAddress 可能首先在另一个线程上报错,而不是在启动线程上报错。因此,即使在 CUDA_LAUNCH_BLOCKING=1 的情况下,我也不信任堆栈跟踪呈现的信息。

 

相反,对于“非法地址(illegal address)”这一 bug,我们希望能找到更多、更准确的报错原因。类似于其他处理器,当故障发生时,GPU 上的 SM 会记录有关故障指令的信息。

不幸的是,我意识到没有进程内的方法可以获取这类信息。我们只能在运行之前,通过将 cuda-gdb 或 cuda-memcheck 附加到进程中来访问此类信息。但这对于那些发生率很低的 bug 来说,在这种模式下重新运行这个进程来重现 bug 是不切实际的。

幸运的是,通过设置环境变量 CUDA_ENABLE_COREDUMP_ON_EXCEPTION=1,我们可以使 CUDA 在发生异常后生成 core dumps 来呈现 GPU 的状态,然后用 cuda-gdb 来检查该文件。

 

本文讨论了如何从这些 core dumps 中生成提取信息,以便在没有调试信息的情况下,也能恢复诸多信息,比如参数值和出错指令等。

 

1、生成 core dumps

 

在有故障的进程上设置 CUDA_ENABLE_COREDUMP_ON_EXCEPTION=1。如此一来,当故障发生时,它会生成一个 core dumps 文件 cudacoredump.hostname.pid。


2、使用 cuda-gdb 打开 core dumps

$ /usr/local/cuda/bin/cuda-gdb(cuda-gdb) target cudacore /tmp/cudacoredump.hostname.pidOpening GPU coredump: /tmp/cudacoredump.hostname.pid
复制代码


这应该报告一些关于故障发生地点的信息:


CUDA Exception: Warp Illegal AddressThe exception was triggered at PC 0x7ff8b63ce440[Current focus set to CUDA kernel 0, grid 132575338, block (1240,0,0), thread (0,1,0), device 0, sm 1, warp 62, lane 0]#0  0x00007ff8b63ce570 in void (anonymous namespace)::softmax_warp_forward<c10::Half, c10::Half, float, 8, false, true>(c10::Half*, c10::Half const*, int, int, int, bool const*, int, bool)<<<(1824,1,1),(32,4,1)>>> ()
复制代码

相关信息如下:

 

  • 触发 Warp Illegal Address 的指令地址:The exception was triggered at PC 0x7ff8b63ce440

  • 正在运行的 kernel 名称:softmax_warp_forward

  • 执行停止的地址:0x00007ff8b63ce570


请注意,GPU 的停止地址(...570)是在触发地址(...440)之后。因为内存是异步读取,所以 GPU 会继续执行指令,之后才能发现故障。在查看寄存器的值时要注意这一点,因为你从中看到的是执行停止时的状态,而错误发生时指令中所使用寄存器的值可能也已经被覆盖。

 

最后,除非编译生成的代码中包含调试信息,否则将看不到代码行或文件名信息。但通过后续介绍的方法,即使没有如上内容,你也能从转储中恢复大量信息。


3、反汇编 kernel


使用 disas 查看 kernel 的 shader assembly(SASS)列表:

(cuda-gdb) disas...0x00007ff8b63ce420 <+1056>:  IADD3 R8, R6.reuse, 0xc0, RZ0x00007ff8b63ce430 <+1072>:  IADD3 R18, R6, 0xe0, RZ0x00007ff8b63ce440 <+1088>:  LDG.E.U8.SYS R19, [R2+0xe0]0x00007ff8b63ce450 <+1104>:  ISETP.GE.AND P3, PT, R8, R13, PT...
复制代码


要查看错误指令,请找到与之匹配的 PC:


0x00007ff8b63ce440 <+1088>:  LDG.E.U8.SYS R19, [R2+0xe0]
复制代码


在这种情况下,LDG 是“从全局内存加载”,从地址[R2+0xe0]读取 1 字节(“U8”)到寄存器 R19。出错的原因大概是 R2+0xe0 越界(out of bounds)了。

4、检查寄存器


使用 info reg 查看所有 GPU 寄存器的值:


(cuda-gdb) info regR0             0xb8198             754072R1             0xfffc80            16776320R2             0xff800000          -8388608R3             0xff800000          -8388608R4             0xff800000          -8388608R5             0x7ff8              32760R6             0x0                 0R7             0x2                 2R8             0x407ce000          1081925632...
复制代码

虽然这里能看到 R2 的值,但其实 R2 在 PC...440 和...570 之间的值已经被覆盖了,因此我们很难找到故障地址的值。


5、读取 GPU 内存

 

使用 print 从内存中读取值:

# read a void* from CUDA's global memory:(cuda-gdb) print *(void * @global *)0x7ff841000000
# read an int from CUDA's global memory(cuda-gdb) print *(int @global *)0x7ff841000000
复制代码

6、恢复传递给 kernel 的参数

 

kernel 的参数在常量“参数”内存中传递。加载它们的指令包括对常量内存的引用,如 c[0x0][0x174]:


0x00007ff8b63ce080 <+128>:   IMAD R0, R3.reuse, c[0x0][0x174], R6
复制代码


可以使用以下方法读取此内存:


(cuda-gdb) print *(int @parameter *)0x174152
复制代码

 

要真正获取所有 kernel 参数的值,我们需要了解它们在内存中的排列方式。假设 kernel 有参数:


_global__ void softmax_warp_forward(  output_t *dst,  const input_t *src,  int batch_size, int stride,  int element_count,  const bool *mask = nullptr,  const int head_chunk_size = -1, bool is_transformer_mask = false) {...
复制代码


常量内存中参数的布局与将它们放入 struct 中的布局相同:

struct Args {                  // offset    output_t *dst;             // 0    const input_t *src;        // 8    int batch_size;            // 16    int stride;                // 20    int element_count;         // 24    // <4 bytes padding>    const bool *mask;          // 32    const int head_chunk_size; // 40    bool is_transformer_mask;  // 44};
复制代码


这意味着结构体的值通常与其自身大小的下一个倍数对齐(8 字节类型与 8 字节倍数对齐),必要时插入一些填充字节(padding bytes)。

 

kernel 参数的开头不是 0x0(低位的地址包含一些关于 kernel 的额外元数据),你可能需要查看程序集中对 c[0x0][...]的所有引用,根据值的使用方式,查看参数缓冲区可能从何处开始。我自己运行时,参数看起来从 0x160 开始,这是 cuda-gdb 能对常量内存返回一个合理的值的条件下,对该常量内存的最小引用。

 

知道了布局和起始地址后,就可以用 print 来获取值(在 print 中指定正确的类型):

# stride(cuda-gdb) print *(int @parameter *) (0x160 + 20)152
复制代码

 

SASS 文档(https://docs.nvidia.com/cuda/cuda-binary-utilities/index.html)有更多关于正在运行的汇编语言的文档,但目前还不甚完善,且会随着 GPU 的更新换代而有所改变。


(本文经授权后编译发布。原文:https://github.com/zdevito/zdevito.github.io/blob/main/_posts/2022-07-27-cuda-core-dumps.markdown


其他人都在看


欢迎下载体验 OneFlow v0.8.0 最新版本:https://github.com/Oneflow-Inc/oneflow/

发布于: 刚刚阅读数: 3
用户头像

OneFlow

关注

不至于成为世界上最快的深度学习框架。 2022.03.23 加入

★ OneFlow深度学习框架:github.com/Oneflow-Inc/oneflow ★ OF云平台:oneflow.cloud

评论

发布
暂无评论
从Core Dump中提取CUDA的报错信息_深度学习_OneFlow_InfoQ写作社区