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

OpenACC vs C++: FATAL ERROR: variable is partially present on the device

本文关键字:存在 变量 vs C++ 致命错误 OpenACC      更新时间:2023-10-16

我正在尝试使用 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 编译指示以将必要的数据发送到设备并最终 使用如下代码:

#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

我不知道代码有什么问题。我将不胜感激任何关于如何修复代码的想法或如何进一步调查问题的建议。谢谢!

这里发生的事情是"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();
}