速度Pascal CUDA8 1080Ti统一内存

Question

由于答案在这里昨天,我想我现在已经使用帕斯卡1080Ti统一存储的正确的基本测试.它分配一个50GB的单维数组并将其加起来.如果我理解正确,它应该是内存绑定,因为这个测试很简单(添加整数).但是,它需要24秒,相当于大约2GB/s.当我运行CUDA8 bandwidthTest时,我看到更高的速率: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).

可分页的bandwidthTest(CUDA 8 samples/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

Answer 1

页面错误处理显然比纯数据副本更复杂.因此,当您通过页面错误将数据驱动到GPU时,它无法在性能方面与纯数据副本竞争.

页面错误实质上引入了GPU处理的另一种延迟.GPU是一种延迟隐藏的机器,但它需要程序员给它隐藏延迟的机会.这可以粗略地描述为暴露足够的并行工作.

从表面上看,您似乎已经暴露了大量的并行工作(数据集中的~12B元素).但是每个字节或检索到的元素的工作强度非常小,因此GPU仍然有机会隐藏与页面错误相关的延迟.换句话说,GPU具有基于可在GPU上飞行的最大线程补充(上限:2048*SM的SM)以及在每个线程中暴露的工作来执行延迟隐藏的瞬时容量.不幸的是,在你的例子中每个线程中暴露的工作可能很小 - 基本上是一个单独添加.

帮助GPU延迟隐藏的一种方法是增加每个线程的工作,并且有各种技术可以做到这一点.一个好的起点是选择具有高计算复杂度的算法(如果可能).矩阵 - 矩阵乘法是每个数据元素的大计算复杂度的经典示例.

在这种情况下,一些建议是认识到你要做的事情是非常有序的,因此从编程的角度来看,通过将工作分解成碎片并自己管理数据传输并不困难.这将允许您实现数据传输操作的链路的全部带宽,实现主机 - >设备带宽的近似完全利用,以及(在该示例中非常小的程度)复制和计算的重叠.对于这样一个简单且易于分解的问题,程序员不使用UM /超额订阅/页面错误是有意义的.

例如,这种方法(UM /超额订阅/页面错误)可能会发光的地方将是一种算法,其中程序员很难提前预测访问模式.遍历大图(可能不会同时在GPU内存中)可能就是一个例子.如果你有一个图遍历问题,每个边遍历有大量的工作,那么图中的页面故障跳跃节点到节点的成本可能不是什么大问题,并且简化了编程工作(不是必须明确管理图数据移动可能是值得的.

关于预取,这是值得怀疑的,即使它可用,它是否在这里有用.预取仍然主要取决于在预取请求正在进行时还有其他事情要做.如果要处理的每个数据项的工作量如此之低,则不清楚一个聪明的预取方案是否真的会为此示例带来很多好处.我们可以想象可能是聪明,复杂的预取策略,但是这样的努力可能更好地用于为这样的问题制作分区显式数据传输系统.