1

我想清楚地了解我经常遇到的使用 OpenACC 加速应用程序的情况。假设我有这个循环:

#pragma acc parallel loop collapse(4)
for (k = KBEG; k <= KEND; k++){
for (j = JBEG; j <= JEND; j++){
for (i = IBEG; i <= IEND; i++){
  for (nv = 0; nv < NVAR; nv++) A0[k][j][i][nv] =
                                data->A[k][j][i][nv];
}}}

作为数据结构化类型变量:

typedef struct Data_{
  double ****A;    
  double ****B;    
} Data;

我注意到无论是否使用统一内存(-ta=tesla:managed),执行时都会出现错误:error 700: Illegal address during kernel execution. 我发现了我在文献中读到的深拷贝问题的问题:编译器完成的隐式拷贝对 A 做了一个简单的拷贝,它指向主机内存上的一个地址,但不是它所指向的数据的拷贝。设备无法读取主机地址,这会产生错误。

  1. 深拷贝问题是对我的错误的正确解释吗?

  2. 此外,如果我正在使用统一内存并且它确实是一个深拷贝问题,那么设备是否应该能够读取地址,至少虚拟地位于统一内存和地址空间上的 A?

我可以轻松解决添加指令的错误:

#pragma acc enter data(data)

并添加present(data)到并行编译指示。请注意,我不需要手动复制 A 和 B。

我想了解问题的原因和解决方案。

4

1 回答 1

1

统一内存仅适用于已分配(堆)内存。我假设“数据”本身没有分配?在这种情况下,您确实需要将它包含在数据区域中,并且应该添加“present”子句,这样编译器就不会尝试隐式复制它。

于 2020-11-16T18:57:51.897 回答