Direct GPU Communication: Understanding P2P Memory Access

The Default Path: Staging Through Host Memory

Before we explore the direct path, let's clarify the inefficient path that P2P access is designed to avoid. Without P2P, if you need to move a block of data from GPU 0 to GPU 1, the process involves the host (CPU) as an intermediary:

1. Copy Device to Host: You issue a cudaMemcpy to transfer the data from GPU 0's global memory to the CPU's system RAM. This transfer happens over the PCI Express (PCIe) bus.
2. Copy Host to Device: You issue a second cudaMemcpy to transfer the same data from the CPU's RAM to GPU 1's global memory. This is another trip across the PCIe bus.

This "staging" through host memory is slow for two main reasons: it involves two separate copy operations, and data must traverse the PCIe bus twice, often competing with other system traffic. For the high-frequency communication required in AI workloads, this CPU bottleneck is unacceptable.

The Direct Path: GPUDirect Peer-to-Peer Access

NVIDIA's solution is a family of technologies under the umbrella term GPUDirect. The specific technology for direct communication within a single node (i.e., between GPUs in the same computer) is GPUDirect Peer-to-Peer (P2P).

GPUDirect P2P allows the DMA (Direct Memory Access) engine of one GPU to read from or write to the memory of another GPU directly, without involving the CPU or host memory. This creates a high-speed, low-latency data path.

There are two primary ways to leverage P2P access, as illustrated in the diagram below.

To formalize this, let's turn to a chapter from a book on HPC for AI.

Enabling Technology I: The Hardware Interconnects

This direct data path relies on a physical connection between the GPUs.

• PCI Express (PCIe): All GPUs in a system are connected to the CPU and each other via the PCIe bus. P2P transfers can occur over PCIe, provided the GPUs are under the same "PCIe root complex" (roughly, connected to the same CPU controller). While faster than staging through RAM, the bandwidth is still a limiting factor.
• NVLink: To overcome the limitations of PCIe for GPU-centric workloads, NVIDIA developed NVLink. This is a proprietary, point-to-point interconnect that provides a much higher bandwidth and lower latency "super-highway" directly between GPUs. Modern multi-GPU servers, like NVIDIA's DGX systems, use NVLink to create a high-speed fabric connecting all GPUs.

The following diagram illustrates the physical data path. When data moves between two NVLink-connected GPUs, it can bypass the PCIe bus entirely, flowing from one GPU's streaming multiprocessors (SMs) and memory controller directly across the NVLink bridge to the other GPU.

For a high-level overview of how NVLink provides this performance boost, the following video offers a good analogy.

Enabling Technology II: Unified Virtual Addressing (UVA)

A fast physical connection is necessary but not sufficient. A crucial software-level abstraction is also required: Unified Virtual Addressing (UVA).

In early versions of CUDA, the CPU and each GPU had their own separate and isolated address spaces. A memory address like 0xDEADBEEF on GPU 0 was completely unrelated to the same address on GPU 1 or the CPU.

UVA, introduced with CUDA 4.0, changed this by creating a single, unified virtual address space for the entire system. With UVA, any memory allocated on the host or any GPU receives a unique pointer value within this 48-bit (or larger) virtual address space.

This is the key that unlocks P2P access. When a kernel on GPU 0 is given a pointer, the CUDA driver can now determine from the pointer's value alone whether it points to memory on GPU 0, GPU 1, or the host. This allows the hardware to route memory requests to the correct physical device. Without UVA, a pointer to another GPU's memory would be meaningless.

Using P2P Access in CUDA Code

Now that we've covered the underlying principles, let's look at the CUDA API for using P2P. The process involves three main steps: checking for compatibility, enabling access, and then performing the transfer.

Let's summarize the workflow for a P2P memcpy:

int gpu0 = 0;
int gpu1 = 1;

// 1. Check if a P2P link exists
int can_access_peer;
cudaDeviceCanAccessPeer(&can_access_peer, gpu0, gpu1);

if (can_access_peer) {
    // 2. Enable P2P access from the current device (gpu0) to the peer (gpu1)
    cudaSetDevice(gpu0);
    cudaDeviceEnablePeerAccess(gpu1, 0); // Second argument is reserved, must be 0

    // Symmetrically, enable access from gpu1 to gpu0
    cudaSetDevice(gpu1);
    cudaDeviceEnablePeerAccess(gpu0, 0);

    // Set active device back to whichever you're launching from
    cudaSetDevice(gpu0);

    // ... allocate memory d_ptr0 on gpu0 and d_ptr1 on gpu1 ...

    // 3. Perform the peer-to-peer copy
    // This will use the direct hardware path (NVLink or PCIe)
    cudaMemcpyPeerAsync(d_ptr1, gpu1, d_ptr0, gpu0, num_bytes, stream);

} else {
    // Fallback to manual staging through host memory
    // ... cudaMemcpy(h_ptr, d_ptr0, ...) followed by cudaMemcpy(d_ptr1, h_ptr, ...)
}

The performance gains from using this direct path are substantial, often saturating the bandwidth of the underlying interconnect. This is precisely what allows high-level libraries like NCCL and CUDA-aware MPI to achieve near-linear scaling on multi-GPU systems.

Conclusion

This lesson completes our journey from the fundamentals of C++ to the intricacies of multi-GPU communication, providing a complete stack of knowledge for HPC in AI. We've seen that Peer-to-Peer (P2P) access is the low-level mechanism that breaks the CPU bottleneck, enabling direct and efficient data exchange between GPUs.

Key Takeaways:

• Peer-to-Peer (P2P) access allows one GPU to directly read from or write to another GPU's memory, bypassing the slow path of staging through host RAM.
• This is enabled by two key technologies:
  1. Hardware Interconnects: Fast physical links like NVLink or the PCIe bus provide the data highway.
  2. Unified Virtual Addressing (UVA): A single system-wide address space that allows the CUDA driver to identify the physical location of any pointer.
• P2P communication can be used via explicit copies, like cudaMemcpyPeerAsync, or through direct memory access, where a kernel on one GPU can dereference a pointer to another GPU's memory.
• Before using P2P, you must query for compatibility with cudaDeviceCanAccessPeer and explicitly enable it with cudaDeviceEnablePeerAccess.

By understanding P2P, you now grasp the foundational link between high-level distributed training frameworks and the low-level hardware architecture. You have seen how kernels are written, how their performance is tied to the memory hierarchy and warp execution, and finally, how multiple GPUs are orchestrated to work together as a cohesive, high-performance computing engine. This comprehensive view, from bits and pointers to distributed systems, forms the bedrock of a GPU kernel engineer's expertise.