openacc - async - 在使用许多异步队列时加快数据传输速度

Question

我的问题是关于在主机和设备之间的不同异步队列中传输多个数组的效果。

假设我们有四个数组：

double *a, *b, *c, *d;

而且，每个都分配了N 的大小。

a = (double*) malloc(N * sizeof(double));
b = (double*) malloc(N * sizeof(double));
c = (double*) malloc(N * sizeof(double));
d = (double*) malloc(N * sizeof(double));

现在，我们可以在一个子句中在设备和主机之间传输它们：

#pragma acc enter data copyin(a[0:N], b[0:N], c[0:N], d[0:N]) async
#pragma acc wait

或者，我们可以使用许多异步子句并将它们转移到不同的队列中：

#pragma acc enter data copyin(a[0:N]) async(1)
#pragma acc enter data copyin(b[0:N]) async(2)
#pragma acc enter data copyin(c[0:N]) async(3)
#pragma acc enter data copyin(d[0:N]) async(4)
#pragma acc wait

上述两种方法的结果是相同的。然而，就性能而言，在某些情况下，第二个似乎更好。

我做了一些测量，发现对于copyouting 和更新主机似乎使用more 比一个队列在性能方面优于one。

我们称第一种方法为one，第二种方法为more，下面的方法为more_nonumber（注意async 子句没有编号）：

#pragma acc enter data copyin(a[0:N]) async
#pragma acc enter data copyin(b[0:N]) async
#pragma acc enter data copyin(c[0:N]) async
#pragma acc enter data copyin(d[0:N]) async
#pragma acc wait

然后，这里是 10,000 次迭代的测量值（不包括 100 次第一次和 100 次最后一次，导致其间的平均迭代次数为 9,800 次）：

one

复制：64.273us

更新设备：60.928us

更新自我：69.502us

复制输出：70.929us

---

more

复制：65.944us

更新设备：62.271us

更新自我：60.592us

复制输出：59.565us

---

more_nonumber

复制：66.018us

更新设备：62.735us

更新自我：70.862us

复制输出：72.317us

---

平均 9800 次运行！

与one 相比，使用more 方法时，copyout (70.929/59.565) 的速度提高了 19%，或者更新自我 (69.502/60.592) 的速度提高了 14%。

我的问题： 这些数字合法吗？我们可以依赖这些数字吗？

为方便起见，我已将我的代码放在github 上。你可以看看。

Answer 1

在设备上交错数据移动和计算时，异步最有用。因此，这个练习有点无关紧要，但我会尽力解释发生了什么。我应该注意，这就是 PGI 当前（v16.10）实现“异步”的方式，而不一定是其他 OpenACC 实现方式实现“异步”的方式。

默认情况下，PGI 使用双缓冲系统来执行数据传输。由于 DMA 传输必须驻留在物理内存中，因此缓冲区是固定的。运行时将虚拟内存复制到固定缓冲区，开始缓冲区的异步传输，然后开始第二个缓冲区的虚拟到固定副本。缓冲区被填充，然后依次传输，直到复制完整的数组。请注意，每个异步队列都有自己的缓冲区。

如果您在没有任何 async 子句的情况下对代码进行计时，您会发现将 4 个变量放在一个 pragma 中比分别拥有自己的变量要快得多。原因是使用 4 个 pragma 时，主机会等到最后一个缓冲区被发送后才移动到下一个变量。当它们都在同一个 pragma 中时，一旦一个数组的最后一个缓冲区开始传输，运行时就可以开始用下一个数组的数据填充另一个缓冲区。

当您将“async”添加到单个 pragma 中，然后是等待时，您应该会发现与根本不使用“async”或“wait”相比没有性能差异。换句话说，它们是相同的：

#pragma acc update device(a,b,c,d) 

#pragma acc update device(a,b,c,d) async
#pragma acc wait

当您将“异步”添加到 4 个单独的 pragma 时，您将获得与将它们全部放在同一个 pragma 中相同的性能，因为 CPU 不会等待开始缓冲下一个数组。

我无法解释为什么当每个阵列都在自己的异步队列上时，从设备复制回来（更新自身、复制输出）会更快。在我看来，这并不重要。

对于此测试，最佳方案是不使用双缓冲区，而是将整个数组固定在物理内存中。这样您就可以节省从虚拟内存复制到固定内存的成本。默认情况下编译器不能这样做，因为物理内存是有限的，它不能保证程序的内存适合。对于大部分时间都花在使用固定内存在设备上进行计算的大多数代码而言，如果根本没有帮助的话。最后，释放固定内存有一些性能开销（设备需要同步以保证不会传输到固定内存）。因此，如果数组在单个数据区域中但有很多更新，这是最有利的。

以下是使用固定内存的时间：

% pgcc -fast -acc -ta=tesla:cc35,pinned transfer.c -o tpinned.out
% numactl -C 2 ./tpinned.out
one
CopyIn: 50.161us
Update device: 49.679us
Update self: 47.595us
CopyOut: 47.631us
---------
more
CopyIn: 52.448us
Update device: 52.135us
Update self: 49.904us
CopyOut: 47.926us
---------
more_nonumber
CopyIn: 52.172us
Update device: 51.712us
Update self: 49.363us
CopyOut: 49.430us
---------