openacc - async - 在使用多个异步队列时获得数据传输的加速
openacc - async - gaining speedup in data transfer when using many async queues
我的问题是关于在主机和设备之间的不同异步队列中传输多个数组的影响。
假设我们有四个数组：
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 （注意异步子句没有编号）：
#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
CopyOut：70.929us
more
复制输入：65.944us
更新设备：62.271us
自我更新：60.592us
CopyOut：59.565us
more_nonumber
CopyIn: 66.018us
更新设备：62.735us
自我更新：70.862us
CopyOut：72.317us
平均 9800 次运行！
与 one 相比，使用 more 方法时，copyout (70.929/59.565) 的速度提高了 19%，而自我更新 (69.502/60.592) 的速度提高了 14%。
我的问题：
这些数字合法吗？我们可以依靠这些数字吗？
为方便起见，我已将我的代码放在 github 上。你可以看看
异步在设备上交织数据移动和计算时最有用。因此，这个练习有点无关紧要，但我会尽力解释发生了什么。我应该注意，这是 PGI 当前 (v16.10) 实现 "async" 的方式，而不一定是其他 OpenACC 实现如何实现 "async".
默认情况下，PGI 使用双缓冲系统来执行数据传输。由于 DMA 传输必须驻留在物理内存中，因此缓冲区是固定的。运行时将虚拟内存复制到固定缓冲区，开始缓冲区的异步传输，然后开始第二个缓冲区的虚拟到固定副本。缓冲区被填满，然后依次传输，直到复制完整数组。请注意，每个异步队列都有自己的缓冲区。
如果您在没有任何异步子句的情况下为代码计时，您会发现将 4 个变量放在一个编译指示中比每个变量都有自己的编译指示要快得多。原因是对于 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
当您将 "async" 添加到 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
---------
我的问题是关于在主机和设备之间的不同异步队列中传输多个数组的影响。
假设我们有四个数组：
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 （注意异步子句没有编号）：
#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
CopyOut：70.929us
more
复制输入：65.944us
更新设备：62.271us
自我更新：60.592us
CopyOut：59.565us
more_nonumber
CopyIn: 66.018us
更新设备：62.735us
自我更新：70.862us
CopyOut：72.317us
平均 9800 次运行！
与 one 相比，使用 more 方法时，copyout (70.929/59.565) 的速度提高了 19%，而自我更新 (69.502/60.592) 的速度提高了 14%。
我的问题： 这些数字合法吗？我们可以依靠这些数字吗？
为方便起见，我已将我的代码放在 github 上。你可以看看
异步在设备上交织数据移动和计算时最有用。因此，这个练习有点无关紧要，但我会尽力解释发生了什么。我应该注意，这是 PGI 当前 (v16.10) 实现 "async" 的方式，而不一定是其他 OpenACC 实现如何实现 "async".
默认情况下，PGI 使用双缓冲系统来执行数据传输。由于 DMA 传输必须驻留在物理内存中，因此缓冲区是固定的。运行时将虚拟内存复制到固定缓冲区，开始缓冲区的异步传输，然后开始第二个缓冲区的虚拟到固定副本。缓冲区被填满，然后依次传输，直到复制完整数组。请注意，每个异步队列都有自己的缓冲区。
如果您在没有任何异步子句的情况下为代码计时，您会发现将 4 个变量放在一个编译指示中比每个变量都有自己的编译指示要快得多。原因是对于 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
当您将 "async" 添加到 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
---------