OpenACC编译报acc例程外部变量、嵌套构造错误如何解决

报错明细 #

• 警告 NVC++-W-1056：acc routine中使用的外部变量未纳入#pragma acc create()子句管理，涉及变量及代码位置：
  • 变量i：sim_xy1.c 第985行、第1011行
  • 变量j：sim_xy1.c 第993行、第1001行
• 严重错误 NVC++-S-1065：计算构造或acc routine中存在不支持的嵌套计算构造，对应代码位置：sim_xy1.c 第976行、第1028行、第1042行

触发原因 #

警告NVC++-W-1056 #

#pragma acc routine标记的函数会被编译为设备端可执行代码，函数内部访问的非局部变量（全局变量、上层作用域定义的变量）默认不会自动分配设备显存、也不会自动完成主机端到设备端的数据拷贝。编译器检测到routine内部访问了外层作用域的循环索引i、j，且这两个变量没有被OpenACC数据子句纳入设备内存管理范围，因此抛出警告。该警告如果不修复，运行时会触发设备端非法内存访问、计算结果错乱等问题。

错误NVC++-S-1065 #

OpenACC规范对计算构造（#pragma acc parallel、#pragma acc kernels等会在设备端启动内核的指令）的嵌套有明确限制：

1. 标记为acc routine的设备函数内部，不能再嵌套定义新的顶层计算构造（设备函数由已启动的设备内核调用，本身不具备独立启动内核的能力）；
2. 主机端发起的parallel/kernels计算构造内部，不能直接嵌套另一层独立的parallel/kernels构造，内层循环仅能通过#pragma acc loop标记，交由外层并行构造统一调度。
   编译器检测到上述违规嵌套写法时，就会抛出该严重错误并终止编译。

修复方案 #

修复NVC++-W-1056警告 #

1. 优先调整变量作用域：如果i、j是循环索引变量，直接将两个变量的定义移动到acc routine内部对应循环的初始化位置，作为循环局部变量。这类循环局部变量OpenACC会自动完成设备端内存管理，不需要额外添加数据子句。
2. 如果i、j是全局变量或上层函数传入的共享变量，在调用该acc routine的外层计算构造中补充对应数据子句：仅需设备端创建内存不需要主机初始值时加create(i,j)，需要拷贝主机端初始值到设备时加copyin(i,j)，避免设备端直接访问未被OpenACC管理的主机内存地址。

修复NVC++-S-1065错误 #

1. 逐个定位报错行号，检查嵌套逻辑：
   • 若报错位置在acc routine标记的设备函数内部，直接删除函数内部写的#pragma acc parallel、#pragma acc kernels等启动内核的构造。如果函数内部有循环需要并行调度，仅需在循环前添加#pragma acc loop标记即可，不能在内核调用的设备函数里重复启动新内核。
   • 若报错位置在主机端定义的parallel/kernels块内部，删除内层独立的parallel/kernels构造，将内层需要并行的循环通过#pragma acc loop标记，合并到外层并行构造的调度逻辑中。
2. 非必要不使用嵌套并行：如果业务逻辑确实需要嵌套内核启动，编译时需要添加-acc=gpu参数开启嵌套并行支持，同时给外层计算构造添加nested子句。注意消费级GPU对嵌套并行的支持度极低，会带来非常高的性能开销。

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