8. 性能指标

博客补充：CUDA C++ 最佳实践指南-CSDN博客https://blog.csdn.net/qq_62704693/article/details/141267262?spm=1001.2014.3001.5502

在尝试优化 CUDA 代码时，了解如何准确测量性能并了解带宽在性能测量中的作用是值得的。本章讨论如何使用 CPU 计时器和 CUDA 事件正确测量性能。然后，本文探讨了带宽如何影响性能指标，以及如何缓解它带来的一些挑战。

8.1. 时序

CUDA 调用和内核执行可以使用 CPU 或 GPU 计时器进行计时。本节将介绍这两种方法的功能、优点和缺陷。

8.1.1. 使用 CPU 计时器

任何 CPU 计时器都可用于测量 CUDA 调用或内核执行的运行时间。各种 CPU timing 方法的详细信息不在本文档的讨论范围之内，但开发人员应始终了解其 timing calls 提供的分辨率。

使用 CPU 计时器时，重要的是要记住许多 CUDA API 函数是异步的;也就是说，它们在完成工作之前将控制权返回给调用 CPU 线程。所有内核启动都是异步的，名称上带有后缀的 memory-copy 函数也是如此。因此，要准确测量特定调用或 CUDA 调用序列的运行时间，有必要在启动和停止 CPU 计时器之前立即调用，使 CPU 线程与 GPU 同步。阻塞调用 CPU 线程，直到该线程之前发出的所有 CUDA 调用都完成。Async cudaDeviceSynchronize() cudaDeviceSynchronize()

尽管也可以将 CPU 线程与 GPU 上的特定流或事件同步，但这些同步函数不适用于默认流以外的流中的计时代码。阻塞 CPU 线程，直到之前向给定流发出的所有 CUDA 调用都已完成。阻塞，直到 GPU 记录了特定流中的给定事件。由于驱动程序可能会交错执行来自其他非默认流的 CUDA 调用，因此其他流中的调用可能包含在计时中。cudaStreamSynchronize() cudaEventSynchronize()

由于默认流（流 0）表现出设备上工作的序列化行为（默认流中的操作只能在任何流中的所有先前调用完成后开始;并且任何流中的后续操作都不能开始，直到它完成），因此这些函数可以可靠地用于默认流中的计时。

请注意，本节中提到的 CPU 到 GPU 同步点意味着 GPU 的处理管道停顿，因此应谨慎使用，以尽量减少其性能影响。

8.1.2. 使用 CUDA GPU 计时器

CUDA 事件 API 提供了创建和销毁事件、记录事件（包括时间戳）以及将时间戳差异转换为浮点值（以毫秒为单位）的调用。如何使用 CUDA 事件进行计时编码说明了它们的用途。

如何使用 CUDA 事件对 CUDA 事件进行时间编码

cudaEvent_t start, stop;
float time;cudaEventCreate(&start);
cudaEventCreate(&stop);cudaEventRecord( start, 0 );
kernel<<<grid,threads>>> ( d_odata, d_idata, size_x, size_y,NUM_REPS);
cudaEventRecord( stop, 0 );
cudaEventSynchronize( stop );cudaEventElapsedTime( &time, start, stop );
cudaEventDestroy( start );
cudaEventDestroy( stop );

此处用于将 and 事件放入默认流 stream 0 中。当设备到达流中的事件时，设备将记录事件的时间戳。该函数返回记录 and 事件之间经过的时间。此值以毫秒为单位，分辨率约为半微秒。与此列表中的其他调用一样，CUDA 工具包参考手册中描述了它们的具体操作、参数和返回值。请注意，计时是在 GPU clock 上测量的，因此 timing resolution 与操作系统无关。cudaEventRecord() start stop cudaEventElapsedTime() startstop

8.2. 带宽

带宽 - 数据传输的速率 - 是影响性能的最重要因素之一。几乎所有对代码的更改都应该在它们如何影响带宽的上下文中进行。如本指南的内存优化中所述，带宽可能会受到存储数据的内存选择、数据的布局方式和访问顺序以及其他因素的显著影响。

为了准确测量性能，计算理论带宽和有效带宽非常有用。当后者远低于前者时，设计或实现细节可能会减少带宽，增加带宽应该是后续优化工作的主要目标。

注意

高优先级：在衡量性能和优化优势时，使用计算的有效带宽作为指标。

8.2.1. 理论带宽计算

理论带宽可以使用产品资料中提供的硬件规格来计算。例如，NVIDIA Tesla V100 使用内存时钟速率为 877 MHz 和 4096 位宽内存接口的 HBM2（双倍数据速率）RAM。

使用这些数据项，NVIDIA Tesla V100 的峰值理论内存带宽为 898 GB/s：

$(0.877\times 10^{9}\times(4096/8)\times2)\div 10^{9}=898GB/s$

在此计算中，内存时钟速率转换为 Hz，乘以接口宽度（除以 8，将位转换为字节），然后由于数据速率翻倍而乘以 2。最后，这个乘积除以 $10^{9}$ 将结果转换为 GB/s。

注意

某些计算使用 $1024^{3}$ 而不是 $10^{9}$ 进行最终计算。在这种情况下，带宽将为 836.4 GiB/s。在计算理论带宽和有效带宽时，使用相同的除数非常重要，这样比较才有效。

注意

在启用了 ECC 的 GDDR 内存的 GPU 上，可用 DRAM 减少了 6.25%，以允许存储 ECC 位。与禁用 ECC 的同一 GPU 相比，为每个内存事务获取 ECC 位也会将有效带宽减少约 20%，但 ECC 对带宽的确切影响可能更高，并且取决于内存访问模式。另一方面，HBM2 存储器提供专用的 ECC 资源，允许无开销的 ECC 保护。2

8.2.2. 有效带宽计算

有效带宽是通过对特定程序活动进行计时以及了解程序如何访问数据来计算的。为此，请使用以下公式：

$Effective\setminus bandwidth=((B_{r}+B_{w})\div 10^{9})\div time$

此处，有效带宽以 GB/s、B 为单位r是每个内核读取的字节数 $B_{w}$ 是每个内核写入的字节数，时间以秒为单位。

例如，要计算 2048 x 2048 矩阵副本的有效带宽，可以使用以下公式：

$Effective\setminus bandwidth=((2048^{2}\times 4\times 2)\div 10^{9})\div time$

元素数乘以每个元素的大小（float 为 4 字节），再乘以 2（由于读取和写入），再除以 $10^{9}$ （或 $1024^{3}$ ）获取传输的内存（GB）。此数字除以获得 GB/s 的时间（以秒为单位）。

8.2.3. Visual Profiler 报告的吞吐量

对于计算能力为 2.0 或更高的设备，Visual Profiler 可用于收集多种不同的内存吞吐量度量。以下吞吐量指标可以显示在 Details （详细信息）或 Detail Graphs （详细信息图表）视图中：

请求的全局负载吞吐量
请求的 Global Store 吞吐量
全局负载吞吐量
全球商店吞吐量
DRAM 读取吞吐量
DRAM 写入吞吐量

Requested Global Load Throughput 和 Requested Global Store Throughput 值表示内核请求的全局内存吞吐量，因此对应于 Effective Bandwidth Calculation 下显示的计算获得的有效带宽。

由于最小内存事务大小大于大多数字大小，因此内核所需的实际内存吞吐量可以包括内核未使用的数据传输。对于全局内存访问，此实际吞吐量由 Global Load Throughput （全局负载吞吐量）和 Global Store Throughput （全局存储吞吐量）值报告。

请务必注意，这两个数字都很有用。实际内存吞吐量显示代码与硬件限制的接近程度，并且将有效带宽或请求带宽与实际带宽进行比较，可以很好地估计内存访问的次优合并所浪费的带宽量（请参阅对全局内存的合并访问）。对于全局内存访问，请求的内存带宽与实际内存带宽的比较由 Global Memory Load Efficiency （全局内存负载效率）和 Global Memory Store Efficiency （全局内存存储效率）指标报告。

本文来自互联网用户投稿，该文观点仅代表作者本人，不代表本站立场。本站仅提供信息存储空间服务，不拥有所有权，不承担相关法律责任。如若转载，请注明出处：http://www.rhkb.cn/news/460251.html

如若内容造成侵权/违法违规/事实不符，请联系长河编程网进行投诉反馈email:809451989@qq.com，一经查实，立即删除！