CUDA - 如何在CUDA C/C++中实现性能度量

原文链接：How to Implement Performance Metrics in CUDA C/C++

---

在本系列的第一篇文章中，我们通过研究SAXPY的CUDA C/C++实现来了解CUDA C/C++的基本元素。在本文中，我们将讨论如何分析CUDA C/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()进行的数据传输是同步（synchronous）（或阻塞（blocking））传输。同步数据传输不会在所有先前发出的CUDA调用完成之前开始，并且在同步传输完成之前，后续CUDA调用也无法开始。因此，在第二行从y到d_y的传输完成之前，第三行的saxpy内核启动不会发出。另一方面，内核启动是异步的。一旦内核在第三行启动，控制就会立即返回到CPU，不会等待内核完成。虽然这似乎为最后一行的设备到主机的数据传输设置了一个竞速条件，但数据传输的阻塞性质确保了内核在传输开始之前完成。

用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事件（event）计时

使用主机设备同步点（如cudaDeviceSynchronize()）的一个问题是它们会暂停GPU pipeline。因此，CUDA通过CUDA事件API提供了一种相对轻量级的CPU计时器替代品。CUDA事件API包括以下调用：创建和销毁事件、记录事件以及计算两个记录事件之间所用时间（以毫秒为单位）。

CUDA事件利用了CUDA流（stream） 的概念。CUDA流只是在设备上按顺序执行的一系列操作。不同流中的操作可以交错，在某些情况下可以重叠——这是一个可用于隐藏主机和设备之间的数据传输的属性（我们稍后将对此进行详细讨论）。到目前为止，GPU上的所有操作都发生在默认流或流0（也称为“空流（null stream）”）中。

在下面的代码中，我们将CUDA事件应用于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()将开始和停止事件放入默认流中。当设备到达流中的事件时，它将记录该事件的时间戳。函数cudaEventSynchronize()会阻止CPU执行，直到指定的事件被记录为止。cudaEventElapsedTime()函数在第一个参数中返回记录start和stop之间经过的时间（毫秒）。该值的精度约为半微秒。

内存带宽

现在我们有了一种精确计时内核执行的方法，我们将使用它来计算带宽。在评估带宽效率时，我们同时使用理论峰值带宽（theoretical peak bandwidth）和观察或有效内存带宽（observed or effective memory bandwidth）。

理论带宽

理论带宽可以使用产品手册中的硬件规格来计算。例如，NVIDIA Tesla M2050 GPU使用DDR（double data rate）RAM，内存时钟频率为1546MHz，内存接口位宽为384-bit。使用这些数据，NVIDIA Tesla M2050的峰值理论内存带宽为148GB/秒，计算如下。$$B{W}_{Theoretical}=1546\ast 10^6\ast (384/8)\ast 2/10^9=148GB/s$$在这个计算中，我们将内存时钟频率转换为Hz，将其乘以接口宽度（除以8，将bit转换为byte），并由于DDR而乘以2。最后，我们除以$10^9$，将结果转换为GB/s。

有效带宽

我们通过计时特定的程序活动和了解程序如何访问数据来计算有效带宽。我们使用以下方程式。$$B{W}_{Effective}=(R_B+W_B)/(t\ast 10^9)$$这里，$B{W}_{Effective}$是以GB/s为单位的有效带宽，$R_B$是每个内核读取的字节数，$W_B$是每个内核写入的字节数。$t$是以秒为单位的经过时间。我们可以修改我们的SAXPY示例来计算有效带宽。下面是完整的代码。

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

在带宽计算中，N*4是每个数组读取或写入传输的字节数，系数3表示x的读取和y的读取和写入。经过的时间存储在变量milliseconds中，以使单位清晰。请注意，除了添加带宽计算所需的功能外，我们还更改了数组大小和线程块大小。在Tesla M2050上编译和运行此代码可以得到：

$ ./saxpy
Max error: 0.000000
Effective Bandwidth (GB/s): 110.374872

测量计算吞吐量（computational throughput）

我们刚刚演示了如何测量带宽，带宽是衡量数据吞吐量的指标。另一个对性能非常重要的指标是计算吞吐量。计算吞吐量的一个常见度量是GFLOP/s，它代表“千兆浮点运算每秒（Giga-FLoating-point OPerations per second）”，其中千兆是$10^9$。对于我们的SAXPY计算，测量有效吞吐量（effective throughput）很简单：每个SAXPY元素都执行一个乘加（multiply-add）运算，通常测量为2FLOPs，因此我们有$$GFLOP/s_{Effective}=2N/(t\ast 10^9)$$$N$是SAXPY操作中的元素数，$t$是以秒为单位的经过时间。与理论峰值带宽一样，理论峰值GFLOP/s可以从产品手册中得到（但计算它可能有点棘手，因为它非常依赖于架构）。例如，Tesla M2050 GPU的理论峰值单精度浮点吞吐量为1030GFLOP/s，理论峰值双精度吞吐量为515GFLOP/s。

SAXPY每个元素计算读取12字节，但只执行一条乘加指令（2FLOPs），因此很明显，它将受到带宽限制，因此在这种情况下（事实上在许多情况下），带宽是衡量和优化的最重要指标。在更复杂的计算中，在FLOPs级别上测量性能可能非常困难。因此，更为常见的做法是使用评测工具来了解计算吞吐量是否是一个瓶颈。应用程序通常提供针对特定问题（而不是特定架构）的吞吐量度量，因此对用户更有用。例如，天文上的n体问题“每秒十亿次交互”，或分子动力学仿真的“纳秒/天”。

总结

这篇文章描述了如何使用CUDA事件API对内核执行计时。CUDA事件使用GPU定时器，因此避免了与主机设备同步相关的问题。我们提出了有效带宽和计算吞吐量性能指标，并在SAXPY内核中实现了有效带宽。很大一部分内核受内存带宽限制，因此计算有效带宽是性能优化的第一步。在未来的文章中，我们将讨论如何确定带宽、指令或延迟三者谁是性能的限制因素。

CUDA事件还可以通过记录cudaMemcpy()调用两侧的事件来确定主机和设备之间的数据传输速率。

如果你在一个较小的GPU上运行本文中的代码，你可能会收到一条关于设备内存不足的错误消息，除非你减少数组大小。事实上，到目前为止，我们的示例代码还没有费心去检查运行时错误。在下一篇文章中，我们将学习如何在CUDA C/C++中进行错误处理，以及如何查询现有设备以确定其可用资源，这样我们就可以编写更健壮的代码。