对称内存管理：NVSHMEM 的基石

好，咱们正式开始聊 NVSHMEM 的核心——对称内存。

说实话，我第一次接触 NVSHMEM 的时候，最困惑的就是这个「对称」到底是什么意思。你想想看，我们平时写 CUDA 程序，用 cudaMalloc 分配显存，每个 GPU 管自己那一亩三分地，数据要交换就得显式地拷贝。但 NVSHMEM 不一样，它搞了一套「对称内存」，让所有 GPU 看起来像在共享同一块地址空间。

嗯，这里要注意，这个「看起来像」是重点。它不是真的共享，而是通过硬件和软件协同，让你在编程时感觉不到数据在跨 GPU 移动。

对称内存的分配：nvshmem_malloc

分配对称内存，最直接的方式就是 nvshmem_malloc。我个人习惯把它类比成 malloc，但它是为 NVSHMEM 通信量身定做的。

void *ptr = nvshmem_malloc(size);

就这么一行，所有参与通信的 PE（Processing Element，你可以理解为每个 GPU 进程）上，都会分配一块 大小相同、地址相同 的内存。对，你没看错，地址也相同。这就是「对称」的核心。

我在项目中遇到过一个问题：一开始用 cudaMalloc 分配内存，然后试图用 NVSHMEM 做通信，结果程序直接崩溃。后来才意识到，NVSHMEM 的通信操作只认对称内存。你传一个普通指针进去，它根本不知道该怎么处理。

对称内存的释放：nvshmem_free

有分配就有释放，这是铁律。nvshmem_free 就是干这个的。

nvshmem_free(ptr);

调用这个函数时，所有 PE 上的对称内存都会被释放。你可能会问：「我只在一个 PE 上调用了 nvshmem_free，其他 PE 上的内存也会被释放吗？」

答案是：会的。这就是对称内存的集体行为。我曾经犯过一个错，在某个 PE 上提前释放了内存，而其他 PE 还在用它做通信，结果就是各种段错误。所以，请务必在所有 PE 都完成通信操作后，再统一释放。

对称内存的特性

对称内存有几个关键特性，我总结了一下：

• 全局可见性： 所有 PE 都能通过相同的地址访问到对方的数据。说白了，你拿着一个地址，就能直接读写任何 PE 上的数据。
• 一致性保证： NVSHMEM 提供了多种一致性模型，比如 nvshmem_fence 和 nvshmem_quiet，确保你的 Put/Get 操作按预期完成。
• 性能优化： 对称内存通常分配在 GPU 的高带宽显存上，并且 NVSHMEM 会利用 NVLink 等高速互联，实现低延迟的数据传输。
• 生命周期管理： 对称内存的生命周期由 NVSHMEM 管理，你不需要关心底层是哪个 GPU 的显存。

我记得有一次调试一个多 GPU 的并行算法，数据总是不对。查了半天，发现是某个 PE 上的对称内存被意外覆盖了。因为所有 PE 的地址都一样，一个 PE 的越界写操作，可能正好覆盖了另一个 PE 的关键数据。嗯，这种 bug 特别难找，所以写代码时一定要小心边界。

对称内存与普通内存的区别

咱们来做个对比，这样更清楚：

特性	对称内存 (nvshmem_malloc)	普通内存 (cudaMalloc)
分配方式	所有 PE 同时分配，地址相同	每个 GPU 独立分配，地址不同
通信支持	原生支持 NVSHMEM 的 Put/Get	需要显式拷贝 (cudaMemcpy)
地址空间	全局统一的虚拟地址空间	每个 GPU 独立的虚拟地址空间
释放方式	集体释放 (nvshmem_free)	独立释放 (cudaFree)
性能	针对 RDMA 和 NVLink 优化	通用分配，无特殊优化

说白了，对称内存就是为「多 GPU 直接通信」这个场景量身定做的。如果你只是在一个 GPU 上做计算，用 cudaMalloc 完全没问题。但一旦涉及多个 GPU 之间的数据交换，对称内存就是你的不二之选。

知识体系结构图

下面这张图，帮你理清对称内存的核心逻辑：

这张图把对称内存的三大块——分配、释放、特性——串起来了。你从中心出发，沿着分支往下看，就能快速掌握它的全貌。

好了，对称内存这块就聊到这儿。记住一句话：对称内存是 NVSHMEM 通信的根基，搞懂了它，后面的 Put/Get 操作就是水到渠成的事。

---