Peer-to-Peer Memory Access
Direct GPU-to-GPU copies over NVLink.
Peer-to-Peer Memory Access is a free CUDA Academy lesson on CoddyKit — lesson 3 of 4. You can read the complete lesson below for free — then practise it hands-on in the browser with a built-in code editor and a 24/7 AI tutor. It is part of the CUDA Academy learning path, one of 4 lessons in the course, and your progress syncs across the web and the CoddyKit app.
The Slow Detour
Moving data from one GPU to another usually bounces through the CPU's memory first. That round trip is slow and wastes the host's bandwidth.
GPUs Talking Directly
Modern GPUs can skip the CPU entirely. Peer-to-peer access lets one GPU read and write another's memory over a direct link.
The NVLink Highway
The fast link between cards is often NVLink, far quicker than the shared PCIe bus. P2P transfers ride this highway when it is available.
Check Before You Trust
Not every pair of GPUs can do P2P. Ask first with cudaDeviceCanAccessPeer, which reports whether one device may reach another.
int can;
cudaDeviceCanAccessPeer(&can, 0, 1);Turn On Access
Permission is off by default. While GPU 0 is current, call cudaDeviceEnablePeerAccess to let it reach into GPU 1's memory.
cudaSetDevice(0);
cudaDeviceEnablePeerAccess(1, 0);Access Is One-Way
Enabling peer access grants only the direction you ask for. For GPU 1 to read GPU 0, you must enable that direction too, from device 1.
Copying Peer to Peer
To move a buffer directly between cards, use cudaMemcpyPeer. You name the destination pointer and device plus the source pointer and device.
cudaMemcpyPeer(dst, 1, src, 0, bytes);Kernels Reading Across
With access enabled, a kernel on GPU 0 can dereference a pointer that lives on GPU 1. The hardware fetches it over the peer link transparently.
When P2P Is Not There
If two cards cannot peer, the runtime quietly falls back to staging through host memory. You still get a copy, just at the slower PCIe rate.
Async Peer Copies
Peer transfers can overlap other work. Pair cudaMemcpyPeerAsync with a stream so the copy runs while kernels keep computing.
cudaMemcpyPeerAsync(dst, 1, src, 0, bytes, stream);Turn It Off When Done
Peer access uses resources, so release it when finished. cudaDeviceDisablePeerAccess tears down the link you opened earlier.
cudaDeviceDisablePeerAccess(1);Quick Check
Recall the benefit of peer-to-peer access between two GPUs.
Recap
P2P lets GPUs share memory directly, ideally over NVLink. Check support, enable access per direction, then use cudaMemcpyPeer. Next: scaling with NCCL. ✨
Frequently asked questions
Is the “Peer-to-Peer Memory Access” lesson free?
Yes — the full text of “Peer-to-Peer Memory Access” is free to read here on the web, and the CUDA Academy course includes 4 lessons in total. To practise it interactively (a built-in code editor and a 24/7 AI tutor) and unlock the rest of the CUDA Academy course, upgrade to CoddyKit PRO.
What will I learn in “Peer-to-Peer Memory Access”?
Direct GPU-to-GPU copies over NVLink. You practise CUDA Academy with hands-on code you run directly in the browser, and a 24/7 AI tutor answers your questions as you work through the lesson.
Do I need any experience to start CUDA Academy?
No prior experience is required. CUDA Academy on CoddyKit is structured for beginners through advanced learners; this is — lesson 3 of 4, so you can start here or from the beginning and move at your own pace.
How long does the “Peer-to-Peer Memory Access” lesson take?
Most CoddyKit lessons take about 5–10 minutes. Each one is bite-sized and interactive, so you make steady progress and pick up exactly where you left off across the web and the app.
Can I write and run code in this CUDA Academy lesson?
Yes. Every CUDA Academy lesson includes a built-in code editor, so you write and run real code right in your browser and get instant AI feedback — no local setup required.
All lessons in this course
- Enumerating and Selecting Devices
- Partitioning Work Across GPUs
- Peer-to-Peer Memory Access
- Multi-GPU with NCCL