OpenACC vs C++：致命错误：变量部分存在于设备上

Question

我正在尝试使用 OpenACC 将一些 C++ 应用程序移植到 GPU。尽可能 期望，C++代码有很多封装和抽象。记忆是 在一些类似矢量的类中分配，然后这个类在许多其他类中被重用 围绕应用程序的类。而且我在尝试正确操作时遇到了麻烦 将 OpenACC 编译指示插入代码中。这是我的代码的简化示例 工作：

#define DATASIZE 16

class Data {
  float *arr;
public:
  Data() {arr = new float[DATASIZE];}
  ~Data() { delete [] arr; }
  float &get(int i) { return arr[i]; }
};

class DataKeeper {
  Data a, b, c;
public:
  void init() {
    for (int i = 0; i < DATASIZE; ++i)
      a.get(i) = 0.0;
  }
};

int main() {
  DataKeeper DK;
  DK.init();
}

我插入了一些 OpenACC pragma 以将必要的数据发送到设备并结束 用这样的代码：

#define DATASIZE 16

class Data {
  float *arr;

public:
  Data() {
    arr = new float[DATASIZE];
#pragma acc enter data copyin(this)
#pragma acc enter data create(arr[:DATASIZE])
  }

  ~Data() {
#pragma acc exit data delete(arr)
#pragma acc exit data delete(this)
    delete [] arr;
  }

  float &get(int i) { return arr[i]; }
};

class DataKeeper {
  Data a, b, c;

public:
  DataKeeper() {
#pragma acc enter data copyin(this)
  }

  ~DataKeeper() {
#pragma acc exit data delete(this)
  }

  void init() {
#pragma acc parallel loop
    for (int i = 0; i < DATASIZE; ++i) {
      a.get(i) = 0.0;
    }
  }
};

int main() {
  DataKeeper DK;
  DK.init();
}

但是编译运行后出现如下错误：

$ pgc++ test.cc -acc -g

$ ./a.out 
_T24395416_101 lives at 0x7fff49e03070 size 24 partially present
Present table dump for device[1]: NVIDIA Tesla GPU 0, compute capability 3.5, threadid=1
host:0x1ae6eb0 device:0xc05ca0200 size:64 presentcount:0+1 line:11 name:(null)
host:0x1f33620 device:0xc05ca0600 size:64 presentcount:0+1 line:11 name:(null)
host:0x1f33d10 device:0xc05ca0a00 size:64 presentcount:0+1 line:11 name:(null)
host:0x7fff49e03070 device:0xc05ca0000 size:8 presentcount:0+1 line:11 name:_T24395600_98
host:0x7fff49e03078 device:0xc05ca0400 size:8 presentcount:0+1 line:11 name:_T24395600_98
host:0x7fff49e03080 device:0xc05ca0800 size:8 presentcount:0+1 line:11 name:_T24395600_98
allocated block device:0xc05ca0000 size:512 thread:1
allocated block device:0xc05ca0200 size:512 thread:1
allocated block device:0xc05ca0400 size:512 thread:1
allocated block device:0xc05ca0600 size:512 thread:1
allocated block device:0xc05ca0800 size:512 thread:1
allocated block device:0xc05ca0a00 size:512 thread:1

FATAL ERROR: variable in data clause is partially present on the device: name=_T24395416_101
 file:/home/bozhenovn/tst/test.cc _ZN10DataKeeperC1Ev line:27

我不知道代码有什么问题。对于如何修复代码或如何进一步调查问题的建议，我将不胜感激。谢谢！

Answer 1

这里发生的事情是“a”的主机地址与“DK”的起始地址相同。因此，当编译器在当前表中查找主机地址时，它使用该表将变量的主机地址映射到设备地址，它发现大小不同。 "a" 是 8 号，而 "DK" 是 24 号。

我将在下面展示修复程序，但让我们回溯并了解这里发生了什么。在主机上创建“DK”时，它首先为其每个数据成员创建存储，然后调用每个数据成员的类构造函数。然后它执行它自己的构造函数。因此，对于每个数据成员，您的代码将在设备上创建类 this 指针，然后在设备上分配“arr”数组。完成此操作后，将在设备上创建“DK”，并为每个数据成员提供空间。但是，由于“DK”的设备副本是在数据成员之后创建的，因此编译器无法自动将两者关联起来。

下面，我发布了两个可能的修复。

首先，您可以让“数据”类管理它自己的数据，但您需要动态分配类数据成员。这样，Data 构造函数将出现在 DataKeeper 构造函数之后，因此编译器可以关联设备数据（也称为“附加”）。

其次，您可以让 DataKeeper 类管理 Data 类的数据。但是，这将要求 Data 的数据是公开的。

请注意，我编写了“使用 OpenACC 进行并行编程”一书的第 5 章“高级数据管理”，其中包含有关 C++ 类数据管理的部分。您可以在以下位置找到我的示例代码：https://github.com/rmfarber/ParallelProgrammingWithOpenACC/tree/master/Chapter05 特别是，看看我是如何做通用容器类“accList”的。

修复 #1：

#define DATASIZE 16
#include <iostream>
#ifdef _OPENACC
#include <openacc.h>
#endif

class Data {
  float *arr;

public:
  Data() {
    arr = new float[DATASIZE];
#pragma acc enter data copyin(this)
#pragma acc enter data create(arr[:DATASIZE])
  }

  ~Data() {
#pragma acc exit data delete(arr)
#pragma acc exit data delete(this)
    delete [] arr;
  }

  float &get(int i) { return arr[i]; }
  void updatehost() {
   #pragma acc update host(arr[0:DATASIZE])
  }

};

class DataKeeper {
  Data *a, *b, *c;

public:
  DataKeeper() {
#pragma acc enter data copyin(this)
  a = new Data;
  b = new Data;
  c = new Data;
  }

  ~DataKeeper() {
#pragma acc exit data delete(this)
  delete a;
  delete b;
  delete c;
  }

  void init() {
#pragma acc parallel loop present(a,b,c)
    for (int i = 0; i < DATASIZE; ++i) {
      a->get(i) = i;
    }
    a->updatehost();
    std::cout << "a.arr[0]=" << a->get(0) << std::endl;
    std::cout << "a.arr[end]=" << a->get(DATASIZE-1) << std::endl;
  }
};

int main() {
  DataKeeper DK;
  DK.init();
}

修复 #2

#define DATASIZE 16
#include <iostream>
#ifdef _OPENACC
#include <openacc.h>
#endif

class Data {
public:
  float *arr;

  Data() {
    arr = new float[DATASIZE];
  }

  ~Data() {
    delete [] arr;
  }

  float &get(int i) { return arr[i]; }
};

class DataKeeper {
  Data a, b, c;

public:
  DataKeeper() {
#pragma acc enter data copyin(this)
#pragma acc enter data create(a.arr[0:DATASIZE])
#pragma acc enter data create(b.arr[0:DATASIZE])
#pragma acc enter data create(c.arr[0:DATASIZE])
  }

  ~DataKeeper() {
#pragma acc exit data delete(this)
#pragma acc exit data delete(a.arr)
#pragma acc exit data delete(b.arr)
#pragma acc exit data delete(c.arr)
  }

  void init() {
#pragma acc parallel loop present(a,b,c)
    for (int i = 0; i < DATASIZE; ++i) {
      a.get(i) = i;
    }
#pragma acc update host(a.arr[0:DATASIZE])
    std::cout << "a.arr[0]=" << a.arr[0] << std::endl;
    std::cout << "a.arr[end]=" << a.arr[DATASIZE-1] << std::endl;
  }
};

int main() {
  DataKeeper DK;
  DK.init();
}