OpenACC实现同例程分CPU/GPU执行及pragma不生效问题排查

你添加的两条OpenACC pragma未生效，完全是用法不符合OpenACC规范导致的，不存在编译器识别问题：

• #pragma acc routine(...)作用域错误
  该指令的作用是告知编译器为当前正在定义的函数/例程生成对应架构的目标代码，必须紧贴例程的定义位置（函数声明/实现的最前方）放置。你把它放在调用语句前没有任何语义，编译器会直接忽略这条无效指令。
• #pragma acc loop collapse(3)缺少上层计算区域上下文
  loop属于附属指令，无法独立生效，必须嵌套在parallel/kernels/serial等设备计算区域指令内部，用来指导编译器对区域内循环做并行划分、线程映射。你单独在循环前放置这条指令，编译器找不到关联的设备offload区域，根本不会生成任何GPU加速代码，因此-Minfo=accel不会输出任何加速编译提示，例程始终运行在CPU上。

另外需要明确：OpenACC没有提供自动根据调用场景切换CPU/GPU执行的语义，主机端（CPU）和设备端（GPU）的指令集完全不同，编译器只会为不同架构生成对应版本的二进制，具体执行哪个版本必须由代码显式控制。

按以下步骤配置即可实现首次调用走CPU、后续调用走GPU的需求：

1. 为例程生成CPU、GPU双版本二进制 #

在目标例程的实现定义最前方，添加routine指令，告知编译器同时生成两个架构的可执行代码：

// 紧贴例程实现放置
#pragma acc routine seq  // 生成CPU端顺序执行版本，供主机直接调用
#pragma acc routine seq  // 生成GPU端顺序执行版本，供GPU线程调用
void your_calc_routine(int kb, int ke, int jb, int je, int ib, int ie /* 其余形参自行补充 */) {
    // 这里只保留单层迭代对应的计算逻辑，不要直接放三层循环
    // 即原三层循环k/j/i单次迭代的计算逻辑放在这里即可
}

注意：routine的并行级别必须和调用场景匹配：单个GPU线程内顺序执行的逻辑用seq级别，warp内向量调度的逻辑用vector级别，错配会导致线程映射错误甚至编译失败。

2. 调用处通过显式分支切换执行路径 #

维护一个首次调用的标记位，根据标记选择CPU执行路径或GPU offload路径，GPU路径下必须显式声明并行计算区域，搭配collapse(3)完成三层循环的线程映射：

// 静态标记位，记录是否为首次执行
static bool is_first_call = true;

if (is_first_call) {
    // 首次执行走CPU路径，直接三层循环调用CPU版例程
    for (int k = kb; k <= ke; k++){
        for (int j = jb; j <= je; j++){
            for (int i = ib; i <= ie; i++){
                your_calc_routine(k, ke, j, je, i, ie);
            }
        }
    }
    // 若后续GPU计算需要复用CPU初始化的数据，可在此处添加acc data copyin指令将数据预拷贝到显存，避免后续重复传输
    is_first_call = false;
} else {
    // 非首次执行走GPU路径，显式声明并行计算区域
    #pragma acc parallel loop collapse(3)
    for (int k = kb; k <= ke; k++){
        for (int j = jb; j <= je; j++){
            for (int i = ib; i <= ie; i++){
                your_calc_routine(k, ke, j, je, i, ie);
            }
        }
    }
}

如果你不想把循环内逻辑拆成独立例程，也可以直接把三层循环的计算逻辑同时写在CPU分支和GPU的并行区域内，逻辑完全一致即可，不需要额外封装。

3. 编译与验证 #

编译时保留-Minfo=accel参数，配置正确时会输出两类明确提示：

• 编译器已为your_calc_routine生成主机seq版本、设备seq版本的代码
• 编译器已对GPU路径下的parallel loop collapse(3)区域完成循环合并、线程块/线程束映射
  此时再通过Nsight-System抓取性能剖面，即可看到非首次调用时存在明确的GPU Kernel调度记录，例程正确运行在GPU上。

常见坑点：如果例程计算涉及主机和设备间的数据传输，一定要通过#pragma acc data区域显式管理数据生命周期，避免每次GPU调用都做冗余数据拷贝，同时防止出现主机、设备端数据不一致的问题。

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