OPEN ACC - 如何在例程中处理 Struct 的数据管理？

Question

我有一个这样定义的结构：

typedef struct Data_{
  double **v;   
  .
  .
  .
  double *press; 
}Data;

在 main 函数中，我有一个 while 循环，在其中称为例程，其中声明了数据，我使用以下编译指示：

static Data data;
#pragma acc enter data copyin(data[:7])

在这个例程中，它被称为另一个例程 (RHS (&data, ...))。同样，在后者内部我调用了另一个例程（RHS1(data,...)），其中有我想要加速的循环：

#pragma acc parallel loop present(data[:7])
  for (i = beg; i <= end; i++) {
    rhs[i][MX1] += dt*data->src[i][MX1];
    . += .
    . += .
    . += .
    . += .
    . += .
    rhs[i][ENG] += dt*sweep->src[i][ENG];
  }

我在使用 -managed 编译时遇到此错误： 致命错误：数据子句中的变量部分存在于设备上：名称=数据

Answer 1

这里有几个问题。

#pragma acc enter data copyin(data[:7])

这是说您想复制一个包含 7 个“数据”结构的数组，而不是具有 7 个成员的结构。将其更改为：

#pragma acc enter data copyin(data)

编译器知道结构的大小，并将在设备副本中分配适当的大小。

但是，复制子句只执行结构的浅表复制。所以在这里你将数据成员的主机指针复制到设备并访问“src”会给你一个非法地址错误。因此，您需要采取额外步骤，对将在设备上使用的所有成员进行手动深层复制。

所以在程序的后面，在你分配了“src”之后，添加一个额外的指令：

#pragma acc enter data copyin(data.src[0:nx][0:ny])

(将nx和ny改为维度的实际大小变量)

这将首先在设备上创建“src”数组，然后将设备 src“附加”到数据的设备副本。 “attach”基本上就是把设备“data.src”的值填到设备src数组的地址上。

然后在你的并行循环中，只使用“present(data)”。

注意，如果使用update指令同步数据，一定要只更新“src”，不要更新“data”。同样，这将执行浅拷贝，因此更新数据将复制 host/device 指针，这会导致运行时错误。

我在 Parallel Programming with OpenACC 一书第 5 章中写了一个关于如何执行此操作的基本示例，您可以在以下位置找到示例：https://github.com/rmfarber/ParallelProgrammingWithOpenACC/blob/master/Chapter05/array_of_structs.c

OPEN ACC - 如何在例程中处理 Struct 的数据管理？

OPEN ACC - how to handle data management of Struct inside routines?

struct

loops

data-management

openacc