# CUDA C BEST PRACTICE 笔记翻译 第三8章

中间的几章都是一些开发的一些基础，个人感觉用处不大，后面有时间再说，先搞一些我自己感兴趣的

8.性能指标

当尝试优化CUDA代码时，了解如何准确地测量性能并理解带宽在性能测量中所起的作用是重要的。本章讨论了如何正确的用 CPU 计时器和 CUDA 时间度量性能,以及探讨贷款如何影响只能指标以及如何应对这些挑战。

8.1 时间

CUDA 调用和核函数运行可以使用 CPU 和 GPU 计时器进行计时。本节将讨论这两种方法的功能、优点和缺陷。

8.1.1 使用CPU 计时器

任何 CPU 计时器都可以被用来测量 CUDA 调用和核函数的耗时。不同 CPU 计时器的细节不在本文的讨论范围之内。但是开发者需要了解所使用的计时器能提供的功能。

当使用 CPU 计时器时，记住许多 CUDA API函数都是异步的是十分重要的。也就是说，这些函数会在执行完自己的任务之前将控制权返回调用的 CPU 线程。所有的核函数启动都是异步的，有 Async 后缀的内存拷贝函数也是异步的。因此，为了准确的度量一系列 cuda 调用的时间，在开始和结束计时器之前调用cudaDeviceSynchronize()函数来同步 CPU 和 GPU 线程是十分重要的。cudaDeviceSynchronize()函数会阻塞 CPU 线程直到之前的 CUDA 调用都完成。

尽管可以同步 CPU 线程和 GPU 上的一个特定流(stream)或事件(event),但是这些同步函数并不适用与除了默认流以外的流上的计时代码。cudaStreamSynchronize()会阻塞 CPU 线程直到给定流中之前的 CUDA 调用执行完毕。cudaEventSynchronize()会阻塞CPU 线程直到特定流中给定的事件执行完成。由于驱动会交替执行其他非默认流中的CUDA 调用，其他流中的调用可能会被包含在计时中。

由于默认流(stream 0)会串行的展示设备上的工作（默认流上的任务只有等其他流上的正在处理的任务完成后才能开始，其他流上的任务只有等默认流上的任务完成后才能执行）。所以在默认流上的计时是可靠的。

本章节提到的CPU-GPU 之间的同步的函数会导致 GPU 执行流水线的停顿，所以使用起来需要谨慎以最小化性能影响。

注：简单说某个流内的操作是顺序的，非默认流之间是异步的，默认流有阻塞作用。也就是说如果默认流上有任务，其他流都必须等默认流执行完成才能继续执行，默认流上也会等其他流执行完成再执行，也就是有排他性。非默认流之间是可以同时执行的，无法保证执行的先后顺序。

8.1.2 使用 GPU 计时器

注：事件是 cuda 的一个常见概念，下文的事件等价于 event，后面可能会 event 和事件混用，在这里说明，避免给读者造成困扰。

CUDA 事件API 提供了函数用来创建和销毁event,记录 event(包含时间戳),将时间戳之间的差值转换为以毫秒为单位的浮点数（就是帮你做个差算耗时）。

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 );

这里 cudaEventRecord()用来将 start 和 end 事件放置在默认流中。当流执行到这个事件时，会记录下这个事件的时间戳。cudaEventElapsedTime()函数返回 start 和 end 这两个 event 之间的时间间隔。单位是 ms，分辨率约为半毫秒。类似于其他表中函数调用，这些操作、参数和返回值都可以在CUDA Toolkit Reference Manual中找到介绍。需要注意的是这个时间是用 GPU 时钟度量的，所以这个时间测量是操作系统无关的。

8.2 带宽(Bandwidth)

带宽——数据的传输速率，是影响性能的最重要的因素之一。几乎所有对代码的改动都需要考虑他们对带宽的影响。带宽会受到数据存储的内存选择，数据的排布、数据访问的顺序和其他因素剧烈的影响。

为了准确地测量性能，需要计算出理论带宽和有效带宽。当后者明显小于前者，说明设计和实现的细节可能会降低带宽，所以增加带宽应该是后续优化的主要目标。

8.2.1 理论带宽计算

理论带宽可以用产品手册提供的硬件参数来计算。例如，NVIDIA Tesla V100 使用HBM2 RAM ，内存时钟频率为877 MHz以及4096位宽的内存接口。

使用这些数据，V100的理论内存带宽峰值是 898GB/s

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

在计算中，内存时钟频率被转换成 Hz，乘以位宽(除以8转换为字节)，乘以2是由于 HBM2的双倍数据速率，最后结果除以$10^9$转换为 GB/s。

Note:
在开启了 ECC 的 GDDR 内存的 GPU 上，可用的 DRAM 会减少6.25%以提供 ECC bits的存储空间。每个内存事务中获取 ECC bits 相比于关闭 ECC 的 GPU 来说也会减少约20%的有效的有效内存带宽。ECC 对内存带宽的影响取决于内存访问模式且可能会更高。另一方面，HBM2内存会提供额外的 ECC 资源，并且允许无额外开销的 ECC 保护。

8.2.2 有效带宽计算

有效带宽通过对特定程序计时以及了解多少数据被该程序访问来计算的。公式为

$bandwidth=((B_r+B_w)\div 10^9 )\div time$

这里的有效带宽单位是 GB/s $B_r$是单个 kernel 读取的字节数，$B_w$是单个 kernel 写入的字节数，time 的单位是 s。

举例来说，为了计算 2048*2048 矩阵拷贝的有效带宽，可以用下列公式

$bandwidth=((2048^2 \times 4 \times 2)\div 10^9 )\div time$

元素数目乘以每个元素的字节数（float 是四字节）再乘2（由于既有读取又有写入），再除$10^9$将单位转换为 GB，再除以以秒为单位的时间得到 GB/s。

8.2.3 由可视化 Profiler 报告的 吞吐量。

对于计算能力大于等于2.0的设备来说，可视化 Profiler 可以用来收集一些不同的内存吞吐指标。下列的吞吐指标可以在Details 或者 Detail Graphs 视图下展示

注：这里的吞吐量和上文说的带宽可以理解为都是对内存的单位时间内处理的元素的大小的一个度量，可以理解为同一种类型的指标。

• Requested Global Load Throughput

• Requested Global Store Throughput

• Global Load Throughput

• Global Store Throughput

• DRAM Read Throughput

• DRAM Write Throughput

Requested Global Load Throughput和Requested Global Store Throughput这两个值只带了核函数请求全局内存的吞吐，是通过有效带宽计算得到的。

由于最小的内存事务大小比大多数的字大小要大，实际的核函数请求的内存吞吐可能包含一些不是核函数需要的数据的传输。对于全局内存访问，实际的吞吐通过Global Load Throughput和Global Store Throughput这两个值反应。

需要注意的是，两个只都是有用的，实际的内存吞吐反应了代码与硬件的极限的接近程度。将有效带宽或请求带宽与实际带宽进行比较，可以很好地估计由于内存的次优的合并访存而浪费了多少带宽。对于全局内存访问,请求内存带宽与实际内存带宽的比较通过Global Memory Load Efficiency和 Global Memory Store Efficiency 这两个指标反映。