如何在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 ，如下所示。

BWTheoretical= 1546 * 106* (384 / 8) * 2 / 109= 148 GB / s

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

有效带宽

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

BWEffective=（RB+WB（ VZX50]* 109)

这里，BWEffective有效带宽，单位为 GB / s ，RB是每个内核读取的字节数，WB是每个内核写入的字节数，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 的前缀9. 我们通常测量 SAXPY 的吞吐量，因为每一个 SAXPY 运算都是有效的

GFLOP/s Effective== 2 N /（ t ：《* 109）

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 ++中执行错误处理以及如何查询当前设备以确定它们可用的资源，以便我们可以编写更健壮的代码。

关于作者

Mark Harris 是 NVIDIA 杰出的工程师，致力于 RAPIDS 。 Mark 拥有超过 20 年的 GPUs 软件开发经验，从图形和游戏到基于物理的模拟，到并行算法和高性能计算。当他还是北卡罗来纳大学的博士生时，他意识到了一种新生的趋势，并为此创造了一个名字： GPGPU （图形处理单元上的通用计算）。

审核编辑：郭婷