OpenACC 和面向对象的 C++

Question

我正在尝试编写与 OpenACC 并行化的面向对象的 C++ 代码。我能够在 OpenACC 上找到一些 Whosebug 问题和 GTC 讨论，但我找不到面向对象代码的一些真实世界示例。

在 an example for a OpenACCArray was shown that does some memory management in the background (code available at http://www.pgroup.com/lit/samples/gtc15_S5233.tar)。但是，我想知道是否可以创建一个 class 来管理更高级别的数组。例如

struct Data
{

//    OpenACCArray<float> a;

    OpenACCArray<Vector3<float>> a3;

    Data(size_t len) {
#pragma acc enter data copyin(this)
//        a.resize(len);
        a3.resize(len);
    }
    ~Data() {
#pragma acc exit data delete(this)
    }
    void update_device() {
//        a.update_device();
        a3.update_device();
    }
    void update_host() {
//        a.update_host();
        a3.update_host();
    }
};

int main(int argc, char *argv[])
{
    const size_t len = 32*128;
    Data d(len);

    d.update_device();
 #pragma acc kernels loop independent present(d)
    for (int i=0; i < len; ++i) {
     float val = (float)i/(float)len;

     d.a3[i].x = val;
     d.a3[i].y = i;
     d.a3[i].z = d.a3[i].x / d.a3[i].y;
    }
    d.update_host();
    for (int i=0; i < len/128; ++i) {
       cout << i << ": " << d.a3[i].x << "," << d.a3[i].y << "," << d.a3[i].z << endl;
    }
    cout << endl;
    return 0;
}

有趣的是，这个程序可以运行，但是一旦我取消注释 OpenACCArray<float> a;，即向该数据结构添加另一个成员，我就会收到内存错误。 FATAL ERROR: variable in data clause is partially present on the device.

由于 OpenACCArray 结构是一个平面结构，它自己处理指针间接寻址，因此将其复制为成员应该可行吗？还是需要指向结构的指针并且指针必须与指令硬连接？然后我担心我必须使用别名指针的问题，正如 jeff larkin 在所建议的那样。我不介意完成此工作运行，但我找不到任何参考如何做到这一点。使用编译器指令 keepgpu,keepptx 有助于理解编译器在做什么，但我更喜欢使用逆向工程生成的 ptx 代码的替代方法。

非常感谢任何指向有用的参考项目或文档的指针。

Answer 1

在 OpenACCArray1.h header 中，删除两个“#pragma acc enter data create(this)”编译指示。发生的事情是 "Data" 构造函数正在设备上创建 "a" 和 "a3" objects。因此，当在 OpenACCArray 构造函数中遇到第二个进入数据区域时，设备 this 指针已经存在。

只有一个数据成员时有效，因为 "a3" 和 "Data" 共享 this 指针的相同地址。因此，当遇到第二个 enter data pragma 时，当前检查发现它已经在设备上，因此不会再次创建它。当添加 "a" 时，"Data" 的大小是 "a" 的两倍，因此目前的检查发现 this 指针已经存在，但大小与以前不同。这就是 "partially present" 错误的含义。数据存在，但大小与预期不同。

只有 parent class/struct 应该在设备上创建 this 指针。

希望这对您有所帮助，垫子

OpenACC 和面向对象的 C++

OpenACC and object oriented C++

c++

openacc