共计 3945 个字符，预计需要花费 10 分钟才能阅读完成。
近期，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 *)0x7ff8410000006
复原传递给 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)
152SASS 文档（https://docs.nvidia.com/cuda/…）有更多对于正在运行的汇编语言的文档，但目前还不甚欠缺，且会随着 GPU 的更新换代而有所扭转。
（本文经受权后编译公布。原文：
https://github.com/zdevito/zd…）
欢送下载体验 OneFlow v0.8.0 最新版本：https://github.com/Oneflow-In…
