OpenACC英伟达GPU编译报错：计算区域调用exit无acc routine信息

这是OpenACC异构计算编译的典型适配问题，触发逻辑如下：

• 你编写的exit_tool()存在OpenACC计算区域（即被#pragma acc kernels/#pragma acc parallel标记、会被卸载到GPU上运行的代码块）内的调用。
• 使用-ta=multicore参数编译时，计算区域的目标运行设备是CPU主机端，主机侧C标准库的printf、exit函数默认可直接调用，编译器不会校验设备侧函数兼容性，因此不会报错。
• 使用-ta=nvidia:managed参数编译时，所有要在GPU设备端运行的函数，必须显式通过#pragma acc routine标记告知编译器生成设备端可执行代码，且调用链上的所有函数都要满足这个要求。但C标准库的exit()是主机端专属函数，作用是终止主机侧进程，GPU上的运行线程没有权限直接终止整个进程，也不存在对应的GPU设备端实现，编译器找不到exit()的设备侧可用版本，就会抛出该错误。
• 不要尝试给exit_tool()加#pragma acc routine标记解决问题，这种方式只会把报错层级下沉，编译器依然会因为找不到exit()的设备实现报同样的错。

不要在GPU设备端的计算区域内直接调用进程终止逻辑，采用「设备端打错误标记+主机端校验后退出」的兼容方案，参考实现如下：

#include <stdio.h>
#include <stdlib.h>

// 主机端专属错误处理函数，仅在主机侧调用
void exit_tool()
{
  printf("Simulacao interrompida");
  exit(1);
}

int main()
{
  int err_flag = 0;
  // 把错误标记变量拷贝到设备端，计算完成后拷回主机
  #pragma acc parallel copy(err_flag)
  {
    // 正常GPU计算逻辑
    // ...
    if (/* 匹配错误触发条件 */) {
      // nvidia目标下设备端可直接调用printf打印提示
      printf("Runtime error detected in device code\n");
      err_flag = 1;
    }
  }
  // 执行流回到主机端后，检查错误标记触发退出
  if (err_flag) {
    exit_tool();
  }

  // 其余业务逻辑
  return 0;
}

如果使用的是支持OpenACC 2.7及以上版本的编译器，也可以在设备端调用acc_abort()直接中止当前计算区域的执行，但该接口不会执行你自定义exit_tool()里的打印、资源回收逻辑，不同编译器的实现行为也存在差异，兼容性不如错误标记回传方案，不推荐优先使用。

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