如何在CUDA C/C++中实现主机和设备同步执行

如何在CUDA C/C++中实现主机和设备同步执行

主机设备同步

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

用 CPU 计时器计时内核执行

现在让我们来看看如何使用 CPU 计时器为内核执行计时。

使用 CUDA 事件计时

在下面的清单中，我们将 CUDA 事件应用于 SAXPY 代码。

内存带宽

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

理论带宽

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是以秒为单位的运行时间。下面是完整的代码。

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

$ ./saxpyMax error: 0.000000Effective 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 。

总结

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

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

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

版权声明：本文内容由网络用户投稿，版权归原作者所有，本站不拥有其著作权，亦不承担相应法律责任。如果您发现本站中有涉嫌抄袭或描述失实的内容，请联系我们jiasou666@gmail.com 处理，核实后本网站将在24小时内删除侵权内容。