OpenACC：深度复制和统一内存

Question

我想清楚地了解我经常遇到的使用 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 副本，它指向主机内存上的一个地址，但不是它所指向的数据的副本。设备无法读取主机地址，这会产生错误。

深拷贝问题是我错误的正确解释吗？
此外，如果我使用的是统一内存，这确实是一个深拷贝问题，设备不应该能够读取地址吗，至少实际上是 A，位于统一内存和地址 space?

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

#pragma acc enter data(data)

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

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

Answer 1

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

OpenACC：深度复制和统一内存

OpenACC: Deep Copy and Unified Memory

deep-copy

openacc

unified-memory