如何调试代码 700“非法内存访问”，又名“CUDA_EXCEPTION_14，Warp 非法地址”？

Question

我的代码正在显示

CUDA error calling "cudaStreamSynchronize(0)", code is 700 an illegal memory access was encountered on 489
 0# my_func(signed char const*, unsigned char const*, int*, int*, int, int) in libthing.so

我该如何调试这个？

Answer 1

调试非法内存访问/扭曲非法地址

这里有一份全面的指南，可帮助您发现可能犯下的愚蠢错误。前两个步骤有点多余，因为我们最终将使用调试器，但它们是帮助隔离各种问题的好主意。

第 1 步。确保检查主机上的所有 CUDA API 调用是否有错误，包括内核启动。

我的输出看起来像这样

CUDA error calling "cudaStreamSynchronize(0)", code is 700 an illegal memory access was encountered on 489

第 489 行是第一次我检查了 CUDA API 调用是否有错误。这与错误发生的位置几乎无关。

由于如果我们检查它们，我们也可以看到来自 GPU 的 CUDA 错误，因此我们需要确保我们的代码在任何与 GPU 交互时都检查它们。

为此，我使用以下头文件来定义错误检查函数：

#pragma once

#include <boost/stacktrace.hpp>
#include <iostream>
#include <stdexcept>

#define STRINGIZE_DETAIL(x) #x
#define STRINGIZE(x) STRINGIZE_DETAIL(x)

#define CUDA_CHECK(call)                                                                                               \
  do {                                                                                                                 \
    if ((call) != cudaSuccess) {                                                                                       \
      const cudaError_t err = cudaGetLastError();                                                                      \
      std::cerr << "CUDA error calling \"" #call "\", code is " << err << " " << cudaGetErrorString(err) << " on "     \
                << __LINE__ << "\n"                                                                                    \
                << boost::stacktrace::stacktrace() << std::endl;                                                       \
      throw std::runtime_error("Problem.");                                                                            \
    }                                                                                                                  \
  } while (0)

#define CUDA_KERNEL_LAUNCH_CHECK()                                                                                     \
  do {                                                                                                                 \
    const auto cuda_err = cudaGetLastError();                                                                          \
    if (cuda_err != cudaSuccess) {                                                                                     \
      throw std::runtime_error(std::string("CUDA kernel launch failed! ") + cudaGetErrorString(cuda_err));             \
    }                                                                                                                  \
  } while (0)

如果你的代码之前看起来像这样：

cudaMemcpy(dst, src, 20*sizeof(float), cudaMemcpyHostToDevice);
my_kernel<<<blocks, threads, shared_mem, stream>>>(dst, result);

添加错误检查后，它看起来像这样

CUDA_CHECK(cudaMemcpy(dst, src, 20*sizeof(float), cudaMemcpyHostToDevice));
my_kernel<<<blocks, threads, shared_mem, stream>>>(dst, result);
CUDA_KERNEL_LAUNCH_CHECK();

请注意，您希望在每次内核启动后都有一个

CUDA_KERNEL_LAUNCH_CHECK()

立即。

现在您的代码正在检查错误，您将看到更接近问题发生的位置。

2.使用 CUDA_LAUNCH_BLOCKING

运行你的程序

CUDA_LAUNCH_BLOCKING=1 ./my_program.exe

CUDA_LAUNCH_BLOCKING

将导致每个内核在进入下一行之前完成运行。由于每次内核启动后都会出现

CUDA_KERNEL_LAUNCH_CHECK()

，现在这将准确告诉您是哪个内核导致了问题。

3.在调试模式下编译

像这样编译你的代码

nvcc -g G my.cu source.cu files.cu

这确保您可以访问调试器内的 CUDA 源代码并能够单步调试它。

4.使用cuda-gdb

像这样运行你的代码。

cuda-gdb --silent --ex run --args ./my_program.exe

在没有调试器的情况下运行程序给出了 C++ 主机源中的地址：

CUDA error calling "cudaStreamSynchronize(0)", code is 700 an illegal memory access was encountered on 489

在

cuda-gdb

中运行程序给出：

CUDA Exception: Warp Illegal Address
The exception was triggered at PC 0x10000320778 (cuda.cu:414)

Thread 1 "benchmark.exe" received signal CUDA_EXCEPTION_14, Warp Illegal Address.
[Switching focus to CUDA kernel 0, grid 15, block (0,0,0), thread (0,0,0), device 0, sm 0, warp 0, lane 0]
0x000001000031f450 in my_kernel<<<(1,1,1),(128,1,1)>>> (rec=0x7fffc7c00000 '\377' <repeats 200 times>, arg1=0x7fffce800000, arg2=0x7fffc7de8a00, 
    arg3=0x555555a7a9a0, arg4=0x7fffc7cf4400 "", arg5=1000, arg6=1000) at my_source_file.cu:382
382       const auto ci = order[i];

调试器已确定第 414 行是问题所在 - 确实如此！

那条线看起来像这样

array[i] = val;

我们怀疑该分配导致了问题，因为它涉及地址。但为什么呢？

请注意，除了参数 5 为

0x7fffc

之外，所有参数都具有以

0x555555a7a9a0

开头的指针。这不是十六进制垃圾——它实际上是问题的根源。 CUDA 没有在距离所有其他地址很远的 0x555555a7a9a0 处分配单个内存块。相反，主机指针已传递到设备，这导致了故障。这就是为什么指针的值如此不同的原因。

请注意，尽管调试器似乎已停止在

382       const auto ci = order[i];

这一行与问题无关。我不知道为什么编译器在这里停止。