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

在 本系列文章的第一篇 中，我们通过检查 CUDA C/C++ SAXPY 来研究 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 调用在同步传输完成之前无法开始。因此，第三行的 saxpy 内核启动在第二行从 y 到 d_y 的传输完成后才会发出。另一方面，内核启动是异步的。一旦内核在第三行启动，控制权立即返回到 CPU ，而不是等待内核完成。而 MIG ht 似乎为设备在最后一行主机数据传输设置了一个竞争条件，数据传输的阻塞性质确保了内核在传输开始之前完成。

用 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 管道。因此， CUDA 通过 CUDA 事件 API 为 CPU 定时器提供了一个相对轻量级的替代方案。 CUDA 事件 API 包括在两个记录的事件之间调用 create 和 破坏 事件、 record 事件和 以毫秒为单位计算已用时间 。

CUDA 事件利用 CUDA streams . CUDA 流只是按顺序在设备上执行的操作序列。在某些情况下[vx3 . 4 可以交叉使用 vx3 . 4]的流。到目前为止， GPU 上的所有操作都发生在默认流或流 0 （也称为“空流”）中。

在下面的清单中，我们将 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() 将启动和停止事件放入默认流 stream 0 。当事件到达流中的事件时，设备将记录事件的时间戳。函数 cudaEventSynchronize() 会阻止 CPU 的执行，直到记录指定的事件为止。 cudaEventElapsedTime() 函数在第一个参数中返回录制 start 和 stop 之间经过的毫秒数。该值的分辨率约为半微秒。

内存带宽

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

理论带宽

理论带宽可以使用产品文献中提供的硬件规格计算。例如， NVIDIA Tesla M2050 GPU 使用内存时钟速率为 1546 MHz 的 DDR （双数据速率） RAM 和 384 位宽的内存接口。使用这些数据项， NVIDIA Tesla M2050 的峰值理论内存带宽为 148 GB / s ，如下所示。

BW _Theoretical= 1546 * 10⁶* (384 / 8) * 2 / 10⁹= 148 GB / s

在这个计算中，我们将内存时钟速率转换为赫兹，乘以接口宽度（除以 8 ，将位转换为字节），再乘以 2 ，这是由于数据速率加倍。最后，我们除以 10⁹将结果转换为 GB / s 。

有效带宽

我们通过计时特定的程序活动和了解程序如何访问数据来计算有效带宽。我们用下面的等式。

BW _Effective=（ R _B+ W _B（ VZX50]* 10⁹)

这里， BW _Effective有效带宽，单位为 GB / s ， R _B是每个内核读取的字节数， W _B是每个内核写入的字节数， 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);

}

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

$ ./saxpy

Max error: 0.000000

Effective Bandwidth (GB/s): 110.374872

测量计算吞吐量

我们刚刚演示了如何测量带宽，带宽是数据吞吐量的度量。另一个对性能非常重要的指标是计算吞吐量。计算吞吐量的常用度量是 GFLOP / s ，它代表“每秒千兆浮点运算”，其中 Giga 是 10 的前缀⁹. 我们通常测量 SAXPY 的吞吐量，因为每一个 SAXPY 运算都是有效的

_{GFLOP/s  Effective}== 2 N /（ t ：<* 10⁹)

N  是 SAXPY 操作中的元素数， t 是以秒为单位的运行时间。与理论峰值带宽一样，理论峰值 GFLOP / s 可以从产品文献中获得（但是计算它可能有点棘手，因为它与体系结构非常相关）。例如， Tesla M2050 GPU 的单精度浮点吞吐量理论峰值为 1030 GFLOP / s ，双倍精度的理论峰值吞吐量为 515 GFLOP / s 。

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

总结

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

CUDA 事件还可以用于确定主机和设备之间的数据传输速率，方法是在 cudaMemcpy() 调用的任一侧记录事件。

如果你在这个设备上运行一个关于内存不足的错误[ZC9]，你可能会得到一个更小的错误。实际上，到目前为止，我们的示例代码还没有费心检查运行时错误。在[VZX337]中，我们将学习如何在 CUDA C / C ++中执行错误处理以及如何查询当前设备以确定它们可用的资源，以便我们可以编写更健壮的代码。