关于深度学习:从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.pid
Opening GPU coredump: /tmp/cudacoredump.hostname.pid

这应该报告一些对于故障产生地点的信息：

CUDA Exception: Warp Illegal Address
The 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, RZ
0x00007ff8b63ce430 <+1072>:  IADD3 R18, R6, 0xe0, RZ
0x00007ff8b63ce440 <+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 reg
R0             0xb8198             754072
R1             0xfffc80            16776320
R2             0xff800000          -8388608
R3             0xff800000          -8388608
R4             0xff800000          -8388608
R5             0x7ff8              32760
R6             0x0                 0
R7             0x2                 2
R8             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的参数在常量“参数”内存中传递。加载它们的指令包含对常量内存的援用，如c0x0：

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

能够应用以下办法读取此内存：

(cuda-gdb) print *(int @parameter *)0x174
152

要真正获取所有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的额定元数据），你可能须要查看程序集中对c0x0的所有援用，依据值的应用形式，查看参数缓冲区可能从何处开始。我本人运行时，参数看起来从0x160开始，这是cuda-gdb能对常量内存返回一个正当的值的条件下，对该常量内存的最小援用。

晓得了布局和起始地址后，就能够用print来获取值（在print中指定正确的类型）：

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

SASS文档（https://docs.nvidia.com/cuda/…）有更多对于正在运行的汇编语言的文档，但目前还不甚欠缺，且会随着GPU的更新换代而有所扭转。

（本文经受权后编译公布。原文：
https://github.com/zdevito/zd…）

欢送下载体验 OneFlow v0.8.0 最新版本：https://github.com/Oneflow-In…