CUDA Academy · 课时

双缓冲流水线

将数据分块,持续为 GPU 提供任务

第 4 / 4 课13 个步骤

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

GPU 空闲问题

先拷贝一个巨大数组,再运行内核,最后拷贝回来。每次拷贝期间 GPU 都处于空闲状态,而计算期间传输引擎也处于空闲状态。

拆分工作

解决办法是先将一个巨大的数组拆成更小的数据块。每个数据块都可以独立拷贝和处理,从而实现重叠执行。

两个缓冲区,两条通道

双缓冲使用两个设备缓冲区和两个流。当流 A 处理一个数据块时,流 B 可以拷入下一个数据块。

重叠拷贝与计算

关键在于并发:传输和内核可以同时处理不同的数据块。传输时间会被有用的计算隐藏起来。

需要锁页内存

这只有在主机数据位于锁页内存中时才有效。可分页缓冲区会强制执行阻塞拷贝,使整个流水线重新退化为串行执行。

流水线循环

您需要遍历各个数据块,并在每一步中发出异步拷入、一个内核和异步拷出操作,且全部使用同一流。

cudaMemcpyAsync(d_in, h_in + off, csz, H2D, s);
proc<<<g, b, 0, s>>>(d_in, d_out);
cudaMemcpyAsync(h_out + off, d_out, csz, D2H, s);

交替使用各个流

通过交替使用流,将每个数据块分配给不同的流。偶数数据块使用流 0,奇数数据块使用流 1,这样相邻数据块就能顺利重叠执行。

cudaStream_t s = streams[i % 2];

为什么两个就够了

两个流已经可以让拷贝与计算重叠执行。更多缓冲区虽然能增加并行深度,但收益会逐渐减少,还会增加内存开销,因此请从两个开始。

最后统一等待

将所有数据块加入队列后,请等待全部工作完成,再读取结果。简单调用 cudaDeviceSynchronize 就能一次性等待所有流完成。

cudaDeviceSynchronize();

速度提升

当拷贝时间被计算隐藏后,总耗时会接近两者中较长者的耗时,而不是两者之和。这才是真正的收益。🚀

何时收益最大

当传输时间和计算时间大致平衡时,双缓冲的效果最好。如果其中一项占绝大多数时间,请先集中精力缩短那一项。

快速检查

双缓冲流水线的核心优势是什么?

回顾

将数据拆成数据块,使用两个锁页缓冲区和两个流,并让拷贝与计算重叠执行。在最后进行同步,观察传输时间消失。🎉

免费开始

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

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

课程
30
课程
120

常见问题解答

「双缓冲流水线」课时是免费的吗?

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

「双缓冲流水线」这节课中我会学到什么?

将数据分块,持续为 GPU 提供任务 你通过在浏览器中直接运行的动手代码来练习 CUDA Academy,全天候 AI 导师会在你学习这节课的过程中回答你的问题。

学习 CUDA Academy 需要有经验吗?

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

「双缓冲流水线」课时需要多长时间?

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

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

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

此课程中的所有课时

  1. 为什么可分页内存很慢
  2. 使用 cudaMallocHost 分配页锁定内存
  3. 在流中使用 cudaMemcpyAsync
  4. 双缓冲流水线
← 返回 CUDA Academy