如何正确使用OpenMP Target Enter/Exit/Update实现非结构化异步设备计算？

目标 #

我希望在C++中以当前使用CUDA的方式使用OpenMP target：

• 分配任意数量的设备端数组（可与主机端数组镜像），按需初始化。
• 使用这些数组执行任意系列设备端操作，无需复制到主机端。
• 完成后，将所需数组（如输出数组）复制回主机端。

最终我希望将此代码作为库通过PyBind11等工具使用，但目前这不是重点。我的实现思路是：
创建一个包含数组及所需元数据、工具方法的类，在构造函数、析构函数、to_host、set_data等方法中添加#pragma omp target enter/exit/update语句。
然后编写接收这些数组对象作为输入的函数，执行如#pragma omp target parallel for的操作。
之后按需调用这些函数，完成后调用to_host方法在主机端使用数据。

问题 #

我在实现过程中遇到了一些问题：首先我不确定使用enter后target指令的正确用法，但IBM文档表明可正常使用map语句，社区文档也支持在方法中使用enter和exit。

我为数组异步操作编写的代码示例如下：

void async_op(AsyncArray2D result, AsyncArray2D a, AsyncArray2D b) {
    #pragma omp target map(to: a.data_ptr[:a.size], b.data_ptr[:b.size], result.data_ptr[:result.size])
    #pragma omp teams distribute parallel for
    for (int i = 0; i < result.size; i++) {
        result.data_ptr[i] = a.data_ptr[i] * b.data_ptr[i];
        result.data_ptr[i] -= 1 / b.data_ptr[i];
        result.data_ptr[i] += b.data_ptr[i] / (a.data_ptr[i] + i);
    }
}

注：我知道需要使用nowait实现异步操作，但这不是当前问题的重点。

问题在于，只有在上述函数中添加#pragma omp target update from(result.data_ptr[:result.size])，才能将数组正确复制回主机端。如果在to_host方法或main()中调用该语句，复制无法正常工作（仅得到全零输出）。我也可以使用to/from或隐式映射，但这只是自动完成复制，update仍然无效。

因此我甚至不确定enter语句是否生效，因为这几乎和不使用该语句的效果一样（除了内存传输可能更高效）。

我猜测根本问题是OpenMP存在某种上下文限制，但我需要了解具体问题和可选方案，以设计符合需求的解决方案。是否只能在同一类的方法中执行异步操作？

代码示例 #

以下是一个非最小示例，我使用g++-12 -O2 -fopenmp -fno-stack-protector -fcf-protection=none -foffload=nvptx-none example.cpp -o example编译，结果可能因系统而异。

#include <cstring>
#include <iostream>
#include <omp.h>


class AsyncArray2D {
public:
    float* data_ptr;
    size_t size;
    bool is_on_device;
    int shape[2], strides[2];

    AsyncArray2D(size_t rows, size_t cols) {
        size = rows * cols;
        shape[0] = rows;
        shape[1] = cols;
        strides[0] = cols;
        strides[1] = 0;
        data_ptr = new float[size];
        #pragma omp target enter data map(alloc: data_ptr[:size])
    }

    void to_host() {
        // 此方法似乎无效
        #pragma omp target update from(data_ptr[0:size])
    }

    void set_data(float * buf_ptr) {
        std::memcpy(data_ptr, buf_ptr, size * sizeof(float));
        // 不确定此语句是否生效
        #pragma omp target update to(data_ptr[:size])
    }

    ~AsyncArray2D() {
        #pragma omp target exit data map(delete: data_ptr[:size])            
    }
};

void async_op(AsyncArray2D result, AsyncArray2D a, AsyncArray2D b) {
    #pragma omp target map(to: a.data_ptr[:a.size], b.data_ptr[:b.size], result.data_ptr[:result.size])
    #pragma omp teams distribute parallel for
    for (int i = 0; i < result.size; i++) {
        result.data_ptr[i] = a.data_ptr[i] * b.data_ptr[i];
        result.data_ptr[i] -= 1 / b.data_ptr[i];
        result.data_ptr[i] += b.data_ptr[i] / (a.data_ptr[i] + i);
    }
   // 取消注释下方语句可得到正确结果
   // #pragma omp target update from(result.data_ptr[:result.size])
}

int main() {
    size_t dim_size = 8;
    
    float * buffer_a = new float[dim_size * dim_size];
    float * buffer_b = new float[dim_size * dim_size];
    float * buffer_result = new float[dim_size * dim_size];
    std::fill_n(buffer_a, dim_size * dim_size, 1.2);
    std::fill_n(buffer_b, dim_size * dim_size, 2.7);
    std::fill_n(buffer_result, dim_size * dim_size, 0);
    std::cout << "Created buffer" << std::endl;
    AsyncArray2D a(dim_size, dim_size);
    AsyncArray2D b(dim_size, dim_size);
    AsyncArray2D result(dim_size, dim_size);
    std::cout << "Created arrays" << std::endl;
    a.set_data(buffer_a);
    b.set_data(buffer_b);
    result.set_data(buffer_result);
    std::cout << "Set data" << std::endl;
    // 循环执行以观察GPU是否工作
    for (int i = 0; i < 100; i++) {
        async_op(result, a, b);
    }
    std::cout << "Added results" << std::endl;
    result.to_host();
    std::cout << "Moved to host" << std::endl;
    for (int j = 0; j < dim_size; j++) {
        for (int i = 0; i < dim_size; i++) {
            std::cout << result.data_ptr[j * dim_size + i] << ' ' ;
        }
        std::cout << std::endl;
    }
    return 0;
}

运行此代码得到以下结果：

0 0 0 0 0 0 0 0 
0 0 0 0 0 0 0 0 
0 0 0 0 0 0 0 0 
0 0 0 0 0 0 0 0 
0 0 0 0 0 0 0 0 
0 0 0 0 0 0 0 0 
0 0 0 0 0 0 0 0 
0 0 0 0 0 0 0 0

但取消注释update语句，或将result的映射改为from，可得到预期结果：

5.11963 4.0969 3.71338 3.51249 3.38886 3.30511 3.24463 3.1989 
3.16311 3.13434 3.1107 3.09094 3.07418 3.05977 3.04726 3.0363 
3.02661 3.01798 3.01025 3.00329 2.99699 2.99125 2.98601 2.9812 
2.97677 2.97268 2.96889 2.96537 2.9621 2.95903 2.95617 2.95348 
2.95096 2.94858 2.94633 2.94422 2.94221 2.94031 2.93851 2.93679 
2.93516 2.93361 2.93213 2.93072 2.92936 2.92807 2.92683 2.92565 
2.92451 2.92341 2.92236 2.92135 2.92038 2.91945 2.91854 2.91767 
2.91683 2.91602 2.91524 2.91448 2.91375 2.91304 2.91235 2.91169

据我观察，to_host方法完全无效，该pragma语句未产生任何作用。

我不想切换到其他框架，因为我需要尽可能通用的方案，也不想局限于英特尔生态或使用其他会增加代码分发复杂度的工具。

编辑补充：我刚测试将async_op改为AsyncArray2D的方法，通过result数组调用，此时to_host可正常工作。但将所有操作改为类方法对我来说有局限性（会导致代码臃肿，方法列表过长），因此这不是我的理想选择。

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