CUDA kernel触发断言失败时能否获取断言的文件、行号信息？

CUDA设备端assert触发属于异步错误，kernel启动调用（<<<...>>>语法）本身不会返回这类运行时错误，只有后续执行cudaDeviceSynchronize()或其他CUDA Runtime API调用时，才会收到cudaErrorAssert类型的错误码。你可以通过以下两种方案拿到断言对应的文件名、行号：

方案1：直接解析CUDA自动输出的标准错误流（零侵入） #

设备端原生assert触发时，CUDA驱动会自动将格式化的断言失败信息写入进程的stderr流，输出格式固定为：

<源文件名>:<行号>: <kernel名>: block: [块ID], thread: [线程ID] Assertion `<断言表达式>` failed.

你只需要在检测到cudaDeviceSynchronize()返回断言错误时，捕获进程的stderr输出，做简单字符串解析就能提取到需要的文件名、行号信息。
注意事项：

• 设备端原生assert仅在Debug编译模式下生效，编译时加-DNDEBUG（Release模式默认配置）会直接移除所有断言逻辑
• 设备端断言触发后，当前CUDA上下文会进入不可恢复的错误状态，所有后续CUDA API调用都会返回错误，必须调用cudaDeviceReset()销毁上下文后才能重新初始化CUDA
• 该方案不需要修改任何现有kernel代码，适配成本最低。

方案2：自定义设备端断言宏（结构化获取信息，无需解析日志） #

如果不想依赖日志解析，可以自定义断言逻辑，触发时将文件名、行号写入预分配的设备端全局内存，主机端同步后直接读取即可拿到结构化的错误信息，参考实现如下：

// 设备端全局错误信息存储
__device__ int g_assert_line = -1;
__device__ char g_assert_file[256] = {0};

// 自定义设备端断言宏
#define CUDA_KERNEL_ASSERT(cond) \
  do { \
    if (!(cond)) { \
      g_assert_line = __LINE__; \
      const char* cur_file = __FILE__; \
      int idx = 0; \
      while (cur_file[idx] != '\0' && idx < 255) { \
        g_assert_file[idx] = cur_file[idx]; \
        idx++; \
      } \
      g_assert_file[idx] = '\0'; \
      __trap(); // 触发设备端硬件陷阱，终止kernel执行 \
    } \
  } while(0)

// 带断言信息抓取的错误检查宏
#define CHECK_WITH_ASSERTION_FETCH(call) \
  do { \
    cudaError_t err = (call); \
    if (err != cudaSuccess) { \
      int fail_line = -1; \
      char fail_file[256] = {0}; \
      cudaMemcpyFromSymbol(&fail_line, g_assert_line, sizeof(int), 0, cudaMemcpyDeviceToHost); \
      cudaMemcpyFromSymbol(fail_file, g_assert_file, 256, 0, cudaMemcpyDeviceToHost); \
      if (fail_line != -1) { \
        printf("Kernel assertion failed at %s:%d, error msg: %s\n", \
               fail_file, fail_line, cudaGetErrorString(err)); \
      } else { \
        printf("CUDA runtime error: %s\n", cudaGetErrorString(err)); \
      } \
      cudaDeviceReset(); // 重置损坏的CUDA上下文 \
      exit(EXIT_FAILURE); \
    } \
  } while(0)

// 测试用例
__global__ void my_kernel(int x){
  CUDA_KERNEL_ASSERT(x != 0);
}

int main(){
  my_kernel<<<1,1>>>(0);
  CHECK_WITH_ASSERTION_FETCH(cudaDeviceSynchronize());
  return 0;
}

该方案的优势是不需要捕获、解析标准输出，拿到的信息直接可用，适合集成到工程的基础错误处理框架中。

内容的提问来源于stack exchange，提问作者Richard