时间戳

首页 > 解决方案 > OPEN ACC - 如何处理 Struct 内部例程的数据管理？

loops - OPEN ACC - 如何处理 Struct 内部例程的数据管理？

问题描述

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

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

在 main 函数中，我有一个 while 循环，在其中它被称为一个声明数据的例程，我使用以下 pragma：

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 编译时出现此错误： 致命错误：设备上部分存在数据子句中的变量：名称 = 数据

标签： loopsstructopenaccdata-management

解决方案

这里有几个问题。

#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”。同样，这将执行浅拷贝，因此更新数据将复制主机/设备指针，这将导致运行时错误。

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

推荐阅读