OpenACC vs C++:致命错误:设备上部分存在变量
OpenACC vs C++: FATAL ERROR: variable is partially present on the device
我正在尝试使用 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();
}
- 为什么我的变量存在于其范围之外
- "new"创建的实例的所有成员变量是否都存在于堆上而不是堆栈上?
- 如何从文件中存在的路径中检索环境变量
- 如何通过结构中的变量找出结构向量中是否存在结构
- VS 中的 C++ 17 会导致 C++14 中不存在的变量(重新)评估错误
- OpenACC vs C++:致命错误:设备上部分存在变量
- 变量在临时存储中存在多长时间
- 对条件表达式结果的赋值(其中第二个和第三个操作数是相同类型和值类别的变量)是否仍然存在?
- 根据成员变量的类型是否存在,有条件地定义该变量
- SFINAE:检测成员变量的存在在 g++ 上不起作用
- 返回成员变量(如果存在)
- 为什么堆栈中的函数局部变量之间存在内存空间
- 原子线程围栏:为什么在这个非原子变量上存在数据竞争?这有关系吗
- 当变量仍然存在时调用C++析构函数
- c++11 lambda 真的支持闭包吗?函数变量中存在语义冲突
- 是否存在32位变量无法正确对齐的情况
- 防止变量脱离范围,因此它们仍然存在另一个线程
- 当只有一个线程写入 c++ 中的布尔变量时,是否存在争用条件
- 保存在 16 位的双变量样本中
- 如何检查 std 中是否存在变量或函数?(在 C++ 中)