OpenACC 管理 2D 数据移动

Question

如何在主机 CPU 和 GPU 之间交换二维元素？ 我尝试将这个 2d 元素用作 w[0:(n_hidden*i)-1]，但编译器反馈告诉我有一些问题

这是RBM算法函数：

double RBM::propdown(int *h, int i, double b) {
#pragma acc data region \
copyin(w[0:(n_hidden*i)-1],h[0:n_hidden],b) create(pre_sigmoid_activation)
double pre_sigmoid_activation = 0.0;
#pragma acc parallel loop reduction(+:pre_sigmoid_activation) 
for(int j=0; j<n_hidden; j++) {
pre_sigmoid_activation += W[j][i] * h[j];
}
pre_sigmoid_activation += b;  
return sigmoid(pre_sigmoid_activation);
#pragma acc exit data \
delete ( pre_sigmoid_activation)
}

Answer 1

无需对数组进行线性化。只需使用多个括号。

#pragma acc data copyin(W[0:n_hidden][0:N])

我还看到了许多其他问题。

数据指令没有“区域”子句。您可能会将其与“输入数据”子句或作为 OpenACC 基础的 PGI 加速器模型混淆。

无需将“b”放在数据子句中，因为它实际上并未在计算区域中使用。同样，通过将它放在数据子句中，您可以使“b”成为设备上的全局引用。最好将只读标量留在数据子句之外，以便将值作为参数传入，而不是需要从全局内存中获取。

再次通过将标量变量“pre_sigmoid_activation”放在数据子句中，您已经创建了一个全局变量。在这里，减少的结果将存储在这个设备变量中，并且不会在主机上自动更新。为此，您需要添加一个“更新”指令。更好的是，只需将其从数据子句中删除，减少的结果就会更新到宿主变量中。

您有一个不匹配的“退出数据”指令（应该有一个相应的“输入数据”指令）。此外，该指令放置在 return 语句之后，因此永远不会执行，从而将数据留在设备上。

最后，由于 C++ 区分大小写，请确保 OpenACC 指令中的变量名与实际变量名匹配。即“W”而不是“w”。

以下是我编写循环的方式。请注意，我不知道“W”第二维的大小，所以只使用“N”。请相应更新。

#pragma acc data copyin(W[0:n_hidden][0:N],h[0:n_hidden])
{
  #pragma acc parallel loop reduction(+:pre_sigmoid_activation)
  for(int j=0; j<n_hidden; j++) {
    pre_sigmoid_activation += W[j][i] * h[j];
  }
}