感谢昨天 here 的回答，我想我现在已经使用 Pascal 1080Ti 对统一内存进行了正确的基本测试。它分配一个 50GB 的单维数组并将其相加。如果我理解正确，它应该是内存限制的，因为这个测试非常简单(添加整数)。然而，需要 24 秒，相当于大约 2GB/s。当我运行 CUDA8 带宽测试时，我看到更高的速率:固定速率为 11.7GB/s，可分页速率为 8.5GB/s。
有什么方法可以让测试运行速度超过 24 秒吗？
这是完整的测试代码:
$ cat firstAcc.c
#include <stdio.h>
#include <openacc.h>
#include <stdlib.h>
#include <time.h>
#define GB 50
static double wallclock()
{
double ans = 0;
struct timespec tp;
if (0==clock_gettime(CLOCK_REALTIME, &tp))
ans = (double) tp.tv_sec + 1e-9 * (double) tp.tv_nsec;
return ans;
}
int main()
{
int *a;
size_t n = (size_t)GB*1024*1024*1024/sizeof(int);
size_t s = n * sizeof(int);
printf("n = %lu, GB = %.3f\n", n, (double)s/(1024*1024*1024));
a = (int *)malloc(s);
if (!a) { printf("Failed to malloc.\n"); return 1; }
setbuf(stdout, NULL);
double t0 = wallclock();
printf("Initializing ... ");
for (long i = 0; i < n; ++i) {
a[i] = i%7-3;
}
double t1 = wallclock();
printf("done in %f (single CPU thread)\n", t1-t0);
t0=t1;
int sum=0.0;
#pragma acc parallel loop reduction (+:sum)
for (long i = 0; i < n; ++i) {
sum+=a[i];
}
t1 = wallclock();
printf("Sum is %d and it took %f\n", sum, t1-t0);
free(a);
return 0;
}
我编译如下:
$ pgcc -fast -acc -ta=tesla:managed:cc60 -Minfo=accel firstAcc.c
main:
40, Accelerator kernel generated
Generating Tesla code
40, Generating reduction(+:sum)
41, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
40, Generating implicit copyin(a[:13421772800])
然后我运行它两次:
$ ./a.out
n = 13421772800, GB = 50.000
Initializing ... done in 36.082607 (single CPU thread)
Sum is -5 and it took 23.902612
$ ./a.out
n = 13421772800, GB = 50.000
Initializing ... done in 36.001578 (single CPU thread)
Sum is -5 and it took 24.180615
结果 (-5) 是正确的，因为我以这种方式设置数据。这些数字是 7 个整数 -3:+3 的重复序列，当相加时，除最后 2 的余数外，所有数字都被抵消 (-3 -2 = -5)。
可分页的带宽测试(CUDA 8 个样本/1_Utilities)结果为:
$ ./bandwidthTest --memory=pageable
[CUDA Bandwidth Test] - Starting...
Running on...
Device 0: GeForce GTX 1080 Ti
Quick Mode
Host to Device Bandwidth, 1 Device(s)
PAGEABLE Memory Transfers
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 8576.7
Device to Host Bandwidth, 1 Device(s)
PAGEABLE Memory Transfers
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 11474.3
Device to Device Bandwidth, 1 Device(s)
PAGEABLE Memory Transfers
Transfer Size (Bytes) Bandwidth(MB/s)
33554432 345412.1
Result = PASS
NOTE: The CUDA Samples are not meant for performance measurements. Results may vary when GPU Boost is enabled.
我看到那张纸条了。但我应该用什么来代替呢？这些测量值看起来是否正确？
有什么办法可以使测试运行时间更像 6 秒(50GB/8.5GB/s)而不是 25 秒？
使用 --mode=shmoo 的结果实际上显示可分页达到了更高的速率:11GB/s。
$ ./bandwidthTest --memory=pageable --mode=shmoo
[CUDA Bandwidth Test] - Starting...
Running on...
Device 0: GeForce GTX 1080 Ti
Shmoo Mode
.................................................................................
Host to Device Bandwidth, 1 Device(s)
PAGEABLE Memory Transfers
Transfer Size (Bytes) Bandwidth(MB/s)
1024 160.3
2048 302.1
3072 439.2
4096 538.4
5120 604.6
6144 765.3
7168 875.0
8192 979.2
9216 1187.3
10240 1270.6
11264 1335.0
12288 1449.3
13312 1579.6
14336 1622.2
15360 1836.0
16384 1995.0
17408 2133.0
18432 2189.8
19456 2289.2
20480 2369.7
22528 2525.8
24576 2625.8
26624 2766.0
28672 2614.4
30720 2895.8
32768 3050.5
34816 3151.1
36864 3263.8
38912 3339.2
40960 3395.6
43008 3488.4
45056 3557.0
47104 3642.1
49152 3658.5
51200 3736.9
61440 4040.4
71680 4076.9
81920 4310.3
92160 4522.6
102400 4668.5
204800 5461.5
307200 5820.7
409600 6003.3
512000 6153.8
614400 6232.5
716800 6285.9
819200 6368.9
921600 6409.3
1024000 6442.5
1126400 6572.3
2174976 8239.3
3223552 9041.6
4272128 9524.2
5320704 9824.5
6369280 10065.2
7417856 10221.2
8466432 10355.7
9515008 10452.8
10563584 10553.9
11612160 10613.1
12660736 10680.3
13709312 10728.1
14757888 10763.8
15806464 10804.4
16855040 10838.1
18952192 10820.9
21049344 10949.4
23146496 10990.7
25243648 11021.6
27340800 11028.8
29437952 11083.2
31535104 11098.9
33632256 10993.3
37826560 10616.5
42020864 10375.5
46215168 10186.1
50409472 10085.4
54603776 10013.9
58798080 10004.8
62992384 9998.6
67186688 10006.4
提前致谢。
$ pgcc -V
pgcc 17.4-0 64-bit target on x86-64 Linux -tp haswell
PGI Compilers and Tools
Copyright (c) 2017, NVIDIA CORPORATION. All rights reserved.
$ cat /usr/local/cuda-8.0/version.txt
CUDA Version 8.0.61
最佳答案
页面错误过程显然比纯粹的数据副本更复杂。因此，当您通过页面错误将数据驱动到 GPU 时，它无法在性能方面与数据的纯副本竞争。
页面错误实质上引入了 GPU 需要处理的另一种延迟。 GPU是一个隐藏延迟的机器，但它需要程序员给它隐藏延迟的机会。这可以粗略地描述为暴露足够的并行工作。
从表面上看，您似乎已经暴露了很多并行工作(数据集中有约 12B 个元素)。但检索的每个字节或元素的工作强度非常小，因此 GPU 隐藏与页面错误相关的延迟的机会仍然有限。换句话说，GPU 具有执行延迟隐藏的瞬时能力，该能力基于该 GPU 上可以运行的线程的最大补充量(上限:2048 * SM 数量)以及每个线程中公开的工作。不幸的是，您的示例中每个线程中公开的工作可能非常小 - 基本上是单个添加。
帮助隐藏 GPU 延迟的方法之一是增加每个线程的工作量，有多种技术可以实现这一点。一个好的起点是选择具有高计算复杂性的算法(如果可能)。矩阵-矩阵乘法是每个数据元素计算复杂性较高的典型示例。
在这种情况下，一些建议是认识到您正在尝试做的事情是相当有序的，因此从编程的角度来看，通过将工作分解为多个部分并自己管理数据传输来管理并不困难。这将使您能够实现数据传输操作的链路的全部带宽，实现主机->设备带宽的大致充分利用，以及(在本例中在很小的程度上)复制和计算的重叠。对于这样一个简单且易于分解的问题，程序员不使用 UM/oversubscription/page-faulting 是有意义的。
例如，这种方法(UM/超额订阅/页面错误)可能发挥作用的地方是程序员很难提前预测访问模式的算法。遍历一个大图(不能同时全部位于 GPU 内存中)可能就是一个例子。如果您遇到图形遍历问题，每个边遍历都需要大量工作，那么在图形中从节点到节点进行页面错误跳跃的成本可能不是什么大问题，并且可以简化编程工作(不是必须显式管理图形数据移动)可能是值得的。
关于预取，即使它可用，它在这里是否有很大用处也是值得怀疑的。预取本质上仍然取决于在预取请求进行时还有其他事情要做。当要处理的每个数据项的工作量如此之低时，尚不清楚巧妙的预取方案是否真的能为该示例带来很多好处。我们可以想象可能聪明、复杂的预取策略，但这样的努力可能最好花在为这样的问题设计一个分区显式数据传输系统上。
关于cuda - Pascal CUDA8 1080Ti统一内存的速度，我们在Stack Overflow上找到一个类似的问题： https://stackoverflow.com/questions/43768717/