OpenACC：深度复制和统一内存答案

【问题标题】：OpenACC: Deep Copy and Unified MemoryOpenACC：深度复制和统一内存
【发布时间】：2021-02-25 06:51:01
【问题描述】：

我想清楚地了解我经常遇到的使用 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 吗？

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

#pragma acc enter data(data)

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

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

【问题讨论】：

标签： deep-copy openacc unified-memory

【解决方案1】：

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

【讨论】：

你的意思是分配在主机内存还是设备内存？ 'data.A' 和 'data.B' 使用 malloc 分配。我不明白统一内存何时可用，因为我猜数据在堆内存上。这确实与深拷贝问题有关吗？（问题 1）
在主机上分配，但在使用 CUDA 统一内存时，它分配在主机或设备上可访问的托管内存池中。因此，当通过 malloc 分配“data.A”和“data.B”时（尽管编译器将 malloc 替换为对“cudaMallocManaged”的调用），它们使用托管指针。由于“数据”不是 malloc 的，它仍然需要通过数据指令手动管理。在某些时候，我们希望能够在堆栈和静态变量上使用统一内存，但这需要对 Linux 操作系统本身进行更新，目前尚未采用。
我想查看编译器反馈消息，但我最好的猜测是它试图隐式复制“data->A”以及“data”。它不知道“A”是被管理的，所以仍然必须尝试复制它。具有动态数据成员的聚合类型的隐式副本对于编译器来说很困难，因此最好手动管理它们并使用“present”子句，这样就不会尝试隐式复制。
是的，这很奇怪，但可能与编译器不知道如何处理标量指针的已知问题有关。默认情况下，它将指针视为数组，因此编译器尝试使用循环边界作为数组维度。我已经和我们的编译器工程师讨论过这个问题，但这是一个棘手的问题，因为编译器很难判断指针是标量还是指向数组，所以他们在进行隐式复制时选择了更常见的数组用法.
作为一般规则，我会查看编译器的反馈，看看它在哪里进行隐式复制，并用显式副本替换它们或使用“present”子句（或更常见的是添加“default(present) "。虽然编译器在绝大多数情况下都能正确获取隐式副本，但这些问题案例太多了，我不喜欢编译器隐式执行它。