GPU 通信流程

1. 核心修正：谁是路，谁是车？

RDMA 不是这些硬件之间的“交换渠道”，而是一种交通规则（机制）。

我们可以把数据传输比作快递运输：

数据 (Payload)：你要寄的货物（KV Cache）。
RDMA：“VIP 直送服务”（一种不经过中转站/CPU 处理的规则）。
InfiniBand (IB) / RoCE：高速公路（连接两台机器的网线/光纤协议）。
网卡 (NIC/HCA)：货车（负责在公路上跑的硬件）。
PCIe：公司大门到仓库的水泥路（网卡连接到主板/GPU 的内部通道）。
NVLink：任意门/虫洞（NVIDIA 专用的、比 PCIe 快得多的 GPU 之间互联通道）。

在一次**跨机 GPU 通信（GPUDirect RDMA）**中是这样串联的：

完整路径是：

$$\text{GPU显存 (A)} \xrightarrow{\text{PCIe/NVLink}} \text{网卡 (A)} \xrightarrow{\text{InfiniBand/Ethernet}} \text{网卡 (B)} \xrightarrow{\text{PCIe/NVLink}} \text{GPU显存 (B)}$$

在这个过程中：

CPU (中央处理器)：完全被旁路（Bypass）。它只负责最开始发个号令说“开始传吧”，然后就去睡觉或干别的了，不参与搬运数据。
GPU (图形处理器)：是数据的源头或终点，但 GPU 的计算核心 (CUDA Cores) 不参与通信，只是显存被读取/写入。
PCIe vs NVLink：
- 如果网卡插在 PCIe 插槽上，数据就走 PCIe。
- 如果用的是 NVIDIA 的超级节点（如 DGX GH200），网卡可能直接通过 NVLink Switch 互联，那就不走 PCIe，速度更快。

3. 深入辨析：NVLink 的特殊地位

NVLink 和 PCIe/InfiniBand 放在一起提，这是一个非常好的直觉，因为它们都在竞争“数据通道”的角色。

PCIe：是通用的标准（大家都用）。目前的 PCIe 5.0 x16 带宽大概是 64 GB/s。
InfiniBand (NDR)：目前主流的高性能网络，带宽单口 50 GB/s (400Gbps) 左右。
NVLink (4.0)：NVIDIA 自己的私有协议，单条通道带宽看似不高，但通常它是 18 条甚至更多一起用，总带宽能达到 900 GB/s！

关键区别：

InfiniBand/RDMA 主要解决 “节点与节点之间” (Server to Server) 的通信。
NVLink 主要解决 “卡与卡之间” (GPU to GPU) 的通信。

但是在 NVIDIA 的野心中，界限正在模糊：
现在的 GB200 NVL72 架构，实际上是想用 NVLink 技术把 72 张 GPU 连成一张超级大卡，原本需要走 InfiniBand RDMA 的跨机通信，现在在机柜内部直接变成了走 NVLink。

4. 总结

修正后的准确表述应该是：

RDMA 是一种通信机制，它利用网卡 (NIC) 的硬件能力，通过 PCIe 或 NVLink 通道直接读取 GPU/内存中的数据，并通过 InfiniBand 或 Ethernet 网络发送到远端，整个过程不经过 CPU 的干预和内存拷贝。

GPU 显存和显存之间不仅可以用 NVLink 传输，而且在 CUDA 编程模型中，这种传输甚至可以不表现为“传输（Copy）”，而表现为直接的内存访问（Load/Store）。

这里有三个层面的解释，帮你把理解从“物理连接”提升到“编程模型”：

1. 物理层：不仅是传输，更是“合并”

如果两张显卡（GPU A 和 GPU B）之间用 NVLink 连上了：

没有 NVLink 时（走 PCIe）：
GPU A 要把数据给 GPU B，通常路径是：
GPU A HBM -> PCIe -> CPU RAM (中转) -> PCIe -> GPU B HBM
（虽然有 PCIe P2P 技术可以绕过 CPU RAM，但还要走 PCIe 总线，带宽受限且延迟高）。
有 NVLink 时：
GPU A 和 GPU B 之间就像搭了一座私有的高速立交桥。
GPU A HBM <===> NVLink <===> GPU B HBM
这个过程完全不经过 CPU，也不经过 PCIe。

2. 逻辑层：Peer-to-Peer (P2P) Access

这是 HPC 开发中最关键的概念。

NVLink 的强大之处在于它支持 P2P Direct Access。这意味着 GPU A 的 CUDA Core 可以直接读取或写入 GPU B 的显存地址，就像读写自己的显存一样。

CUDA 代码体现：
你只需要在代码里调用 cudaDeviceEnablePeerAccess(target_gpu_id, 0)。
之后，你在 GPU A 上运行的 Kernel，可以直接解引用一个指向 GPU B 显存的指针。
```
// GPU A 的 Kernel
__global__ void kernel(float* ptr_to_gpu_b) {
    // 直接读取另一张卡的显存，数据通过 NVLink 瞬间拉过来
    float val = *ptr_to_gpu_b; 
}
```
NUMA 架构：
在操作系统眼里，这实际上构成了一个 NUMA (Non-Uniform Memory Access) 架构。显存实际上变成了一个统一的地址空间，只是访问“隔壁显卡”的显存比访问自己的稍微慢那么一点点（延迟极低），但带宽极高（600-900 GB/s）。

3. 架构层：NVSwitch 的作用

如果只有 2 张卡，它们可以用 NVLink 直连。
但如果你有 8 张卡（比如一台 HGX H800 服务器），它们怎么两两互联？

这时候就需要 NVSwitch 芯片。它就像一个路口的“环岛”或“交换机”。
8 张 GPU 全部连到 NVSwitch 上，任意两张 GPU 都可以通过 NVSwitch 进行 NVLink 通信。
效果：这 8 张 GPU 在逻辑上可以被视为一张拥有超大显存的巨型 GPU。

总结与对比

特性	PCIe (传统通道)	NVLink (高速通道)
主要用途	CPU与GPU通信，网卡通信	GPU与GPU之间互联 (甚至 CPU-GPU，如 Grace Hopper)
带宽 (双向)	~64 GB/s (Gen5 x16)	900 GB/s (H100 NVLink)
延迟	较高 (微秒级)	极低 (纳秒级)
语义	主要是 Copy (搬运)	支持 Load/Store (直接读写地址)
在 AI 中的地位	瓶颈	生命线 (用于 Tensor Parallelism)

回到你的应用场景（Chunked Prefill / PD 分离）：

如果你的 Prefill 和 Decode 任务是在同一台服务器的通过 NVLink 连接的两张卡上，那么 KV Cache 的传输是瞬间完成的（甚至不需要传输，只需要传个指针）。
只有当它们在不同的物理机上时，才需要用到我们刚才聊的 RDMA/InfiniBand。