CUDA C / C ++中的性能测量

CUDA性能测量通常是通过主机代码完成的，可以使用CPU计时器或特定于CUDA的计时器实现。在进入这些性能评估技术之前，我们需要讨论如何在主机和设备之间同步执行。

1、主机设备同步

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调用才能开始。

因此，在第二行从y到d_y的传输完成之前，不会在第三行上启动saxpy内核，另一方面，内核启动是异步的。

在第三行启动内核后，控制权立即返回给CPU，并且不等待内核完成。

尽管这似乎为最后一行中的设备到主机数据传输建立了竞争条件，但数据传输的阻塞性质可确保内核在传输开始之前完成。

2、带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执行，直到设备上所有先前发出的命令都完成为止。如果没有阻塞，此代码将测量内核启动时间，而不是内核执行时间。

3、用CUDA事件来测量GPU执行时间

使用主机设备同步函数（如cudaDeviceSynchronize()）的一个问题是它们会阻塞GPU管道。因此，CUDA通过CUDA事件API提供了相对较轻的CPU计时器替代方案。CUDA事件API包括调用以创建和销毁事件，记录事件以及计算两次记录的事件的时间（以毫秒为单位）。

CUDA事件利用了CUDA流的概念。CUDA流只是在设备上按顺序执行的一系列操作。不同流中的操作可以交错，在某些情况下可以重叠—该属性可用于隐藏主机与设备之间的数据传输。到目前为止，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() 函数在第一个参数中返回记录开始和停止之间经过的毫秒数。 该值的分辨率大约为半微秒。

4、内存带宽

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

（1）理论带宽

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

BW(Theoretical) = 1546 * 10^6 * (384 / 8) * 2 / 10^9  = 148GB/s

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

（2）有效带宽

我们通过定时特定的程序活动并知道程序如何访问数据来计算有效带宽。 我们使用以下等式。

BW(Effective)  =  ( Rb + Wb )  /  ( t * 10^9 )

此处，BW(Effective) 是有效带宽，以GB / s为单位，Rb 是每个内核读取的字节数，Wb 是每个内核写入的字节数，t 是经过的时间（以秒为单位）。

我们可以修改我们的SAXPY示例来计算有效带宽。 完整的代码如下。

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

Max error: 0.000000
Deive cost:1.619936ms
Effective Bandwidth (GB/s): 155.350736

测试了好几次，感觉有效带宽并不是很稳定，有时带宽不到100GB/s。

5、测量计算吞吐量

我们刚刚演示了如何测量带宽，这是对数据吞吐量的度量。 对性能非常重要的另一个指标是计算吞吐量。 通用的计算吞吐量度量是GFLOP / s，它表示“每秒千兆浮点运算”，其中Giga是10^9的前缀。对于我们的SAXPY计算，有效吞吐量的测量很简单：每个SAXPY元素都会做一个乘法和加法操作，通常测量为两个FLOP，所以我们有

GFLOP/S（Effective） = 2 N / ( t * 10^9 )

N是SAXPY操作中元素的数量，t 是消耗的时间（以秒为单位）。 像理论峰值带宽一样，可以从产品资料中收集理论峰值GFLOP / s（但计算起来可能有些棘手，因为它非常依赖于体系结构）。 例如，Tesla M2050 GPU的理论峰值单精度浮点吞吐量为1030 GFLOP / s，理论峰值双精度吞吐量为515 GFLOP / s。

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

6、总结

这篇文章描述了如何使用CUDA事件API计时内核执行时间。 CUDA事件使用GPU计时器，因此避免了与主机设备同步相关的问题。 我们介绍了有效带宽和计算吞吐量性能指标，并在SAXPY内核中实现了有效带宽。 很大一部分内核受内存带宽限制，因此有效带宽的计算是性能优化的良好第一步。因此，我们也需要了解哪个因素（带宽，指令或延迟）是提高性能的限制因素。

通过在cudaMemcpy() 调用的任意端记录事件，CUDA事件还可以用于确定主机与设备之间的数据传输速率。

如果您在较小的GPU上运行本文中的代码，除非减小阵列大小，否则可能会收到有关设备内存不足的错误消息。所以接下来我们需要考虑代码异常处理的问题。