点对点内存访问
通过 NVLink 直接进行 GPU 到 GPU 的复制
点对点内存访问 是 CoddyKit 上的免费 CUDA Academy 课时。 这是第 3 节课，共 4 节。 你可以在下方免费阅读本课时的完整内容 — 然后在浏览器中使用内置代码编辑器和全天候 AI 导师进行实践。 这是 CUDA Academy 学习路径的一部分，你的进度在网页和 CoddyKit 应用中同步。 CUDA Academy 课程共包含 4 节课。
缓慢的绕行路径
将数据从一个 GPU 移到另一个 GPU 时，通常会先经过 CPU 的内存。这次往返速度很慢，还会浪费主机的带宽。
GPU 直接通信
现代 GPU 可以完全绕过 CPU。点对点访问允许一个 GPU 通过直接链路读写另一个 GPU 的内存。
NVLink 高速公路
卡之间的高速链路通常是 NVLink，速度远快于共享的 PCIe 总线。如果条件允许，P2P 传输会沿这条高速公路进行。
使用前先检查
并非每一对 GPU 都支持 P2P。请先使用 cudaDeviceCanAccessPeer 进行查询，它会报告一个设备是否能够访问另一个设备。
int can;
cudaDeviceCanAccessPeer(&can, 0, 1);启用访问权限
访问权限默认处于关闭状态。当 GPU 0 为当前设备时，请调用 cudaDeviceEnablePeerAccess，允许它访问 GPU 1 的内存。
cudaSetDevice(0);
cudaDeviceEnablePeerAccess(1, 0);访问是单向的
启用点对点访问只会授予您请求的方向。若要让 GPU 1 读取 GPU 0，您还必须从设备 1 启用该方向。
进行点对点复制
要在两张卡之间直接移动缓冲区，请使用 cudaMemcpyPeer。您需要指定目标指针和设备，以及源指针和设备。
cudaMemcpyPeer(dst, 1, src, 0, bytes);跨设备读取的内核
启用访问后，GPU 0 上的内核可以解引用位于 GPU 1 上的指针。硬件会通过点对点链路透明地获取数据。
没有 P2P 时
如果两张卡无法进行点对点访问，运行时会悄悄退回到经由主机内存中转的方式。复制仍然可以完成，只是速度会降到较慢的 PCIe 速率。
异步点对点复制
点对点传输可以与其他工作重叠。将 cudaMemcpyPeerAsync 与一个流配合使用，就能让复制操作在内核继续计算时运行。
cudaMemcpyPeerAsync(dst, 1, src, 0, bytes, stream);完成后关闭访问
点对点访问会占用资源，因此完成后请释放它。cudaDeviceDisablePeerAccess 会拆除您之前建立的链路。
cudaDeviceDisablePeerAccess(1);快速检查
回想一下，两个 GPU 之间的点对点访问有什么好处。
回顾
P2P 允许 GPU 直接共享内存，理想情况下会通过 NVLink 完成。请检查是否支持，按方向启用访问，然后使用 cudaMemcpyPeer。下一步：使用 NCCL 扩展规模。✨
用 AI 导师学习 C++ — 免费
在浏览器中编写并运行真实代码，获得全天候 AI 导师的即时帮助，并在网页或应用中继续学习。
- 课程
- 30
- 课程
- 120
常见问题解答
「点对点内存访问」课时是免费的吗？
是的 — 「点对点内存访问」的完整文本可在网页上免费阅读。要进行交互式练习（内置代码编辑器和全天候 AI 导师）并解锁 CUDA Academy 课程的其余内容，请升级到 CoddyKit PRO。 CUDA Academy 课程共包含 4 节课。
「点对点内存访问」这节课中我会学到什么？
通过 NVLink 直接进行 GPU 到 GPU 的复制 你通过在浏览器中直接运行的动手代码来练习 CUDA Academy，全天候 AI 导师会在你学习这节课的过程中回答你的问题。
学习 CUDA Academy 需要有经验吗？
无需任何先前经验。CoddyKit 上的 CUDA Academy 课程适合初学者到高级学习者，你可以从这里开始或从头开始，按照自己的节奏学习。 这是第 3 节课，共 4 节。
「点对点内存访问」课时需要多长时间？
大多数 CoddyKit 课程大约需要 5–10 分钟。每节课都很精短且互动，所以你能稳步进步，并在网页和应用中从离开的地方继续。
我能在这节 CUDA Academy 课中编写并运行代码吗？
能。每节 CUDA Academy 课都包含内置代码编辑器，你可以在浏览器中直接编写并运行真实代码，并获得即时 AI 反馈 — 无需本地设置。