封面图片来自互联网

本文翻译自NVIDIA官方博客Parallel Forall，内容仅供参考，如有疑问请访问原网站:https://devblogs.nvidia.com/parallelforall/how-implement-performance-metrics-cuda-cc/.

在这个系列的第一篇文章中，我们通过用CUDA C/C++实现SAXPY，学习了CUDA C/C++编程的基本要素。在这篇文章中，我们会学习如何衡量这个程序以及其他CUDAC/C++程序的性能。我们在之后的文章中经常用到这种性能度量技术，因为程序的性能优化将会变得越来越重要。

译者注：这个系列是指原文的系列，并不是笔者的专栏。

CUDA性能度量通常是在主机端进行的，我们既可以使用CPU的计时器也可以使用CUDA专门的计时器。在开始学习性能度量技术之前，我们需要讨论一下如何同步主机和设备之间的操作。

主机-设备同步

让我们来看一下上一篇博客中SAXPY的数据传输和核函数启动的主机端代码：

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

saxpy<<<(N+255)/256, 256>>>(N, 2.0, d_x, d_y);

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

这里使用cudaMemcpy进行数据传输的方式是同步传输(或者是阻塞传输)方式。同步数据传输直到前面所有发布的CUDA调用全部结束之后才会开始，而且同步数据传输结束之后，随后的CUDA调用才会开始。因此上面第三行的saxpy核函数只有到第二行的y到d_y的数据传输结束之后才会启动。而在另一方面，核函数启动却是异步的。一旦核函数被启动，控制权就立刻返回到CPU，并不会等待核函数执行完成。这样的话就会对最后一行的设备到主机数据传输产生竞态条件(race condition)，但是数据传输的阻塞特性会确保核函数执行完成后再开始数据传输。

译者注：这里的竞态条件前面提到过，简单说就是前面的数据操作还未完成，后面的操作却又要使用前面的数据，这样就会导致错误的结果。

使用CPU的计时器来计算核函数的执行时间

现在我们来看一下如何使用CPU的计时器来给核函数计时。

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

t1 = myCPUTimer();
saxpy<<<(N+255)/256, 256>>>(N, 2.0, d_x, d_y);
cudaDeviceSynchronize();
t2 = myCPUTimer();

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

在上面的代码中，我们除了使用一般的主机时间戳函数myCPUTimer()，还用到了显式的同步障碍cudaDeviceSynchronize()来阻塞CPU执行，直到设备上发布的指令全部执行结束为止。如果没有这个同步障碍，这个代码测试的就是核函数的启动时间而不是执行时间。

使用CUDA事件计时

使用类似cudaDeviceSynchronize()函数的主机设备同步点的一个问题就是它会拖延GPU管道(stall GPU pipeline)。基于这个原因，CUDA提供了一个相比CPU计时器更轻量级的选择，那就是使用CUDA事件API。CUDA事件API包括调用事件创建和销毁函数、事件记录函数以及以毫秒为单位计算两个被记录事件的运行时间的函数。

译者注：这里拖延GPU管道(stall GPU pipeline)的直接结果就是造成CPU和GPU轮流执行，而不再是并行执行。于是就使得程序的运行时间等于CPU与GPU时间之和。具体可以参考:https://blogs.msdn.microsoft.com/shawnhar/2008/04/14/stalling-the-pipeline/

CUDA事件使用的是CUDA streams的概念。一个CUDA流只是一系列在设备上顺序执行的操作。不同流中的操作可以交替执行，在某些情况下甚至可以交叠执行，这个特性可以被用在隐藏主机和设备间的数据传输。(我们会在之后的文章中讨论)。到目前为止，我们所有的操作都是在默认的流中进行的，或者0号流(也叫做空流)。

下面的代码中，我们使用了CUDA事件API来对SAXPY代码进行性能度量。

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);
saxpy<<<(N+255)/256, 256>>>(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);

cuda事件是cudaEvent_t类型，通过cudaEventCreate()和cudaEventDestroy()进行事件的创建和销毁。在上面的代码中cudaEventRecord()将事件start和stop放在默认的流中，即0号stream。函数cudaEventSynchronize()用来阻塞CPU执行直到指定的事件被记录。函数cudaEventElapsedTime()的第一个参数返回start和stop两个记录之间消逝的毫秒时间。这个值的精度大约是0.5ms。

内存带宽

既然我们已经可以精确地测量核函数的运行时间，那么我们就可以用它来计算带宽。我们需要使用理论的峰值带宽和有效内存带宽来评估带宽效率。

理论带宽

理论带宽可以通过产品资料中的硬件规格来计算。例如英伟达Tesla M2050 GPU使用的是时钟频率为1546MHz显存位宽为384-bit的DDR(双倍数据速率)RAM。

使用这些数据，我们可以计算出英伟达Tesla M2050的理论峰值带宽是148 GB/sec：

$BW_{Theoretical}=1546 * 106 * (384/8) * 2 / 109 = 148 GB/s$

在这个表达式中，我们将内存的时钟频率的单位转化为Hz，然后乘以显存宽度(除以8之后，单位由比特转化为字节)，又乘以2是因为该显卡的RAM是DDR(双倍数据速率)。最后我们将结果除以10^9得到以GB/s的计算结果。

有效带宽

我们是通过计算特定程序的活动时间和程序如何访问数据来计算机有效带宽的。我们使用下面的公式：

$BW_{Effective} = (R_B + W_B) / (t * 109)$

这里， $BW_{Effective}$ 是以GB/s的有效带宽， $R_B$ 是每个核函数被读取的字节数， $W_B$ 是每个核函数被写入的字节数， $t$ 是以秒为单位的运行时间。我们可以修改SAXPY例子来计算有效带宽，下面是完整的代码：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

__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: %f\n", maxError);
    printf("Effective Bandwidth (GB/s): %f\n", N*4*3/milliseconds/1e6);
}

在上面的带宽计算(译者注：即表达式N*4*3/milliseconds/1e6)中，N*4是每次数组读或写的字节数，因子3的含义是对x的读以及y的读和写共3次读写操作。程序运行时间被存在变量milliseconds中，把它作为分母即可算出单位时间的带宽大小。注意源程序中除了添加了一些计算带宽的功能外，我们也改变了数组的大小和块的大小(译者注：由于该代码来自之前的博客，所以具体的变化可以对比原来的程序，在这里)。编译并执行上面的代码，我们可以得到：

$ ./saxpy

Max error: 0.000000

Effective Bandwidth (GB/s): 110.374872

测定计算吞吐量

我们刚刚只演示了如何测定带宽，也叫做数据吞吐量。另一种非常重要的性能指标叫做计算吞度量。一种比较通用的测量计算吞吐量的方法是计算GFLOP/s(Giga-FLoating-point OPerations per second)，代表“每秒10亿次的浮点运算数”，这里的Giga就是千兆，即10^9。对于我们的SAXPY计算，测量有效的吞吐量是很简单的：每个SAXPY元素都会做一次乘法加法操作，因此是典型的2FLOPS，所以我们可以得到：

$GFLOP/{s_{Effective}} = 2N / (t * 109)$

其中， $N$ 是SAXPY操作的元素个数， $t$ 是以秒为单位的运行时间。就像理论峰值带宽一样，理论峰值 $GFLOP/s$ 也可以从产品资料查到(但是计算它却很难，因为它具有架构依赖性)。例如，Tesla M2050 GPU的理论单精度浮点峰值吞吐量是 $1030GFLOP/s$ ，而双精度浮点峰值吞吐量是 $515GFLOP/s$ 。SAXPY每次计算读取12个字节，但是仅仅只有一条单独的乘法加法指令(2 FLOPs)，所以很明显这(数据吞吐量)就是带宽限制。而且在这种情况(实际上是大部分情况)下，带宽是最重要的衡量和优化指标。在更复杂的计算中，FLOPs级别的性能测定是很困难的。因此更普遍的方法是使用分析工具来分析计算吞吐量是否是一个瓶颈。这些应用测出的的常常是问题依赖的吞吐量(而不是架构依赖的)，这其实对用户会更有用。例如天文学里每秒百万次交互作用的N体问题，或者每天纳秒级的分子动态模拟。

总结

这篇文章主要介绍了如何用CUDA事件API获取核函数的执行时间。CUDA事件使用GPU计时器，因此避免了与主机设备同步相关的问题。我们也介绍了有效带宽和计算吞吐量的性能测定方法，而且也应用这些方法测定了SAXPY例子中核函数的有效带宽。另外我们也得出，它的内存带宽占了很大比例，因此在性能测试中，计算有效吞吐量是首要的一步。在之后的文章中，我们会进一步讨论在带宽、指令、或者延迟这些因素中，哪一个是限制程序性能的因素。

CUDA事件也可以用来计算主机和设备之间数据传输的速率，方法很简单只要将记录事件的函数放到cudaMemcpy()调用的两边就可以了。

如果你在一个很小的GPU上运行文章中的代码，那么如果你没有减小数组的大小，你可能会得到一个关于不充足设备内存的错误消息。实际上，我们的实例代码目前为止还没有特别检查运行时错误。在下一篇文章中，我们会学习如何进行错误处理以及如何访问现有设备来确定已有资源，这样的话我们就可以写出更鲁棒的代码。

[译]在CUDA C/C++中如何测试代码性能