GPU核心性能指标_gpu性能指标

内存科普

CPU缓存分为三类，一级缓存(L1)、二级缓存(L2)和三级缓存(L3)。CPU在实际数据读取中重要的却是一级缓存，因为一级缓存速度最快，二级缓存其次，三级缓存最慢，只是三级缓存的容量最大。GPU中也有缓存。

速度最快的是寄存器，他能和cpu同步的配合，接着是缓存，在CPU片上，然后是主存储器，现在常见的就是内存条，显卡上也有内存芯片，然后是硬盘，这些内存设备的速度和容量相反，越快的越小，越慢的越大。

CPU和GPU的主存都是采用DRAM——动态随机存取存储器，而低延迟的内存，比如一级缓存，则采用SRAM——静态随机存取存储器。虽然底层的存储器延迟高，容量大，但是其中有数据被频繁使用的时候，就会向更高一级的层次传输，比如我们运行程序处理数据的时候，程序第一步就是把硬盘里的数据传输到主存里面。

GPU和CPU的内存设计有相似的准则和模型。但他们的区别是：CUDA编程模型将内存层次结构更好的呈献给开发者，让我们显示的控制其行为。

GPU主要性能指标

GPU的性能主要由以下几个参数构成：

1. 计算能力（吞吐量）：通常关心的是32位浮点计算能力。16位浮点训练也开始流行，如果只做预测的话也可以用8位整数。
2. 显存大小：当模型越大，或者训练时的批量越大时，所需要的GPU内存就越多。
3. 显存位宽：位数越大则瞬间所能传输的数据量越大
4. 显存带宽：只有当内存带宽足够时才能充分发挥计算能力。

对于大部分用户来说，只要考虑计算能力就可以了。GPU内存尽量不小于4GB。但如果GPU要同时显示图形界面，那么推荐的内存大小至少为6GB。内存带宽通常相对固定，选择空间较小。

计算能力（吞吐量）

一个非常重要的性能指标就是计算吞吐量，单位为GFLOP/s，算力指标
Giga-FLoating-point OPerations per second
表示每秒的浮点操作数量。
每秒浮点运算量，是衡量GPU硬件计算能力的指标。

例如：现在intel purley platform的旗舰skylake 8180是28Core@2.5GHZ，支持AVX512，其理论双精度浮点性能是：28Core2.5GHZ32FLOPs/Cycle=2240GFLPs=2.24TFLOPs

显存大小

而GPU中的我们叫做显存。基本上GPU的内存带宽要比CPU多一个数量级。

显存容量：其主要功能就是暂时储存GPU要处理的数据和处理完毕的数据**。显存容量大小决定了GPU能够加载的数据量大小**。（在显存已经可以满足客户业务的情况下，提升显存不会对业务性能带来大的提升。在深度学习、机器学习的训练场景，显存的大小决定了一次能够加载训练数据的量，在大规模训练时，显存会显得比较重要。

显存位宽

显存位宽：显存在一个时钟周期内所能传送数据的位数，位数越大则瞬间所能传输的数据量越大。
显存在一个时钟周期内所能传送数据的位数，位数越大则瞬间所能传输的数据量越大，这是显存的重要参数之一。目前市场上的显存位宽有64位、128位、256位和512位几种，人们习惯上叫的64位显卡、128位显卡和256位显卡就是指其相应的显存位宽。显存位宽越高，性能越好价格也就越高，因此512位宽的显存更多应用于高端显卡，而主流显卡基本都采用128和256位显存。

内存带宽计算

CUDA GPU的计算速度可以达到gigaflops（每秒10亿次浮点操作），是Core i7/Nethalem速度的十倍。为充分利用强大的计算能力，我们需要从存储器中（全局显存或计算机内存）尽量快地给他们提供数据。所以内存带宽就是衡量GPU的一个重要指标。

显存带宽： 指每秒能传输的数据量。
显存带宽=显存频率X显存位宽/8/1000 (GB/s) eg: GeForce GTX 1060 显存频率8008MHz，显示位宽192bit
则：显存带宽=8008X192/8/1000(GB/s)=192.2 GB/s

峰值理论带宽Theoretical Bandwidth

举例
一个内存始终频率1546MHz，384位宽的内存接口。则他的峰值理论带宽为
NVIDIA Tesla M2050 GPU uses DDR (double data rate) RAM with a memory clock rate of 1,546 MHz and a 384-bit wide memory interface. Using these data items, the peak theoretical memory bandwidth of the NVIDIA Tesla M2050 is 148 GB/sec, as computed in the following.

有效带宽Effective Bandwidth

有效带宽单位为GB/s
RB是每个kernel 读取bytes数量
WB是每个kernel 写入bytes的数量
t是消耗时间

#include 

__global__
void saxpy(int n, float a, float *x, float *y)
{
  int i = blockIdx.x*blockDim.x + threadIdx.x;
  if (i < n) 
  	y[i] = a*x[i] + y[i];
}

int main(void)
{
  int N = 20 * (1 << 20);
  float *x, *y, *d_x, *d_y;
  x = (float*)malloc(N*sizeof(float));
  y = (float*)malloc(N*sizeof(float));

  cudaMalloc(&d_x, N*sizeof(float)); 
  cudaMalloc(&d_y, N*sizeof(float));

  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }

  cudaEvent_t start, stop;
  cudaEventCreate(&start);
  cudaEventCreate(&stop);

  cudaMemcpy(d_x, x, N*sizeof(float), cudaMemcpyHostToDevice);
  cudaMemcpy(d_y, y, N*sizeof(float), cudaMemcpyHostToDevice);

  cudaEventRecord(start);

  // Perform SAXPY on 1M elements
  saxpy<<<(N+511)/512, 512>>>(N, 2.0f, d_x, d_y);

  cudaEventRecord(stop);

  cudaMemcpy(y, d_y, N*sizeof(float), cudaMemcpyDeviceToHost);

  cudaEventSynchronize(stop);
  float milliseconds = 0;
  cudaEventElapsedTime(&milliseconds, start, stop);

  float maxError = 0.0f;
  for (int i = 0; i < N; i++) {
    maxError = max(maxError, abs(y[i]-4.0f));
  }

  printf("Max error: %fn", maxError);
  printf("Effective Bandwidth (GB/s): %fn", N*4*3/milliseconds/1e6);
}
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53

In the bandwidth calculation, N*4 is the number of bytes transferred per array read or write, and the factor of three represents the reading of x and the reading and writing of y.
结果
$ ./saxpy
Max error: 0.000000
Effective Bandwidth (GB/s): 110.374872