CUDA Academy · 课时

测量有效带宽

比较实际达到的吞吐量与峰值吞吐量

第 4 / 4 课13 个步骤

测量有效带宽 是 CoddyKit 上的免费 CUDA Academy 课时。 这是第 4 节课,共 4 节。 你可以在下方免费阅读本课时的完整内容 — 然后在浏览器中使用内置代码编辑器和全天候 AI 导师进行实践。 这是 CUDA Academy 学习路径的一部分,你的进度在网页和 CoddyKit 应用中同步。 CUDA Academy 课程共包含 4 节课。

从理论走向数字

合并访问很有帮助,但您还应通过测量加以验证。最能反映实际情况的指标,是内核以每秒千兆字节计的有效带宽。

有效带宽的含义

有效带宽统计内核实际需要的字节数,再除以内核耗时。它反映的是有用工作,而不是浪费的内存行。

简单公式

将读取的字节数与写入的字节数相加,再除以经过的秒数。这个比值就是实际达到的吞吐量。

double gbps = (bytesRead + bytesWritten) / seconds / 1e9;

使用事件计时

使用 CUDA 事件包围内核。这样可以精确测量 GPU 工作,避免 CPU 端时钟带来的干扰。

cudaEventRecord(start);
myKernel<<<g, b>>>(...);
cudaEventRecord(stop);

读取经过的时间

在停止事件上完成同步后,请让 CUDA 查询时间间隔。它返回的是毫秒,因此代入公式前要先除以 1000。

float ms;
cudaEventElapsedTime(&ms, start, stop);

确认字节数

对于包含 n 个浮点数的 C = A + B,需要读取 2n 个浮点数并写入 n 个,因此总共移动 3n 个浮点数。乘以 4 即可得到移动的字节数。

size_t bytes = 3 * n * sizeof(float);

找到峰值

每个 GPU 都会列出理论峰值带宽。只有将有效带宽看作该上限的一部分时,这个数值才有意义。

效率比率

用有效带宽除以峰值带宽即可得到百分比。合并的受内存带宽限制的内核通常能达到峰值的 70% 到 90%;使用跨步访问的内核则会低得多。

先预热

第一次启动需要支付一次性的初始化代价。先运行一次内核进行预热,然后对多次计时运行取平均值,以获得稳定的结果。

让性能分析器确认

Nsight Compute 会直接报告内存吞吐量,因此您可以核对计算结果。它还会指出由糟糕的合并访问造成的字节浪费。🔧

用它来做判断

比率较低意味着内存是瓶颈,因此应优化合并访问和布局。比率较高则意味着应从其他方面寻找速度。

快速检查

计算有效带宽。

回顾

现在,您可以使用事件测量有效带宽,将其与峰值进行比较,并根据比率判断内存是否是瓶颈。🎉

免费开始

用 AI 导师学习 C++ — 免费

在浏览器中编写并运行真实代码,获得全天候 AI 导师的即时帮助,并在网页或应用中继续学习。

课程
30
课程
120

常见问题解答

「测量有效带宽」课时是免费的吗?

是的 — 「测量有效带宽」的完整文本可在网页上免费阅读。要进行交互式练习(内置代码编辑器和全天候 AI 导师)并解锁 CUDA Academy 课程的其余内容,请升级到 CoddyKit PRO。 CUDA Academy 课程共包含 4 节课。

「测量有效带宽」这节课中我会学到什么?

比较实际达到的吞吐量与峰值吞吐量 你通过在浏览器中直接运行的动手代码来练习 CUDA Academy,全天候 AI 导师会在你学习这节课的过程中回答你的问题。

学习 CUDA Academy 需要有经验吗?

无需任何先前经验。CoddyKit 上的 CUDA Academy 课程适合初学者到高级学习者,你可以从这里开始或从头开始,按照自己的节奏学习。 这是第 4 节课,共 4 节。

「测量有效带宽」课时需要多长时间?

大多数 CoddyKit 课程大约需要 5–10 分钟。每节课都很精短且互动,所以你能稳步进步,并在网页和应用中从离开的地方继续。

我能在这节 CUDA Academy 课中编写并运行代码吗?

能。每节 CUDA Academy 课都包含内置代码编辑器,你可以在浏览器中直接编写并运行真实代码,并获得即时 AI 反馈 — 无需本地设置。

此课程中的所有课时

  1. 什么是内存事务
  2. 合并读取与跨步读取
  3. 数组结构与结构体数组
  4. 测量有效带宽
← 返回 CUDA Academy