随着高性能计算需求的日益增长，数据传输效率成为了制约计算速度的关键因素。英伟达（NVIDIA）公司推出的 NCCL (NVIDIA Collective Communications Library) 库，为 GPU 集群提供了高效的通信支持。近期，NCCL 2.27 版本引入了一项名为 "symmetric memory" 的全新特性，旨在进一步提升通信性能，尤其是在小数据量通信场景下。新媒网了解到，这项技术一经推出，便受到了业界的广泛关注。

据悉，symmetric memory 的设计理念与 NVSHMEM 有着异曲同工之妙，但目前仅限于单节点（intra-node）通信。不过，英伟达公司计划在未来将其与 IBGDA 技术相结合，以实现跨节点（inter-node）通信的支持。

Symmetric Memory：低延迟内核，性能飞跃

从相关测试数据来看，symmetric memory 在低延迟内核模式下，能够将小数据量通信的性能发挥到极致。通过利用 NVLink 技术，symmetric memory 可以在单个 NVLink Domain 内实现高速通信。尤其是在 NVL72 架构上，节点内通信的带宽和延迟表现出了显著优势。即使是在 DGX-H100 等 NVL8 架构上，也能带来可观的性能提升。

用户可以通过 nccl-tests 工具进行验证，只需在运行时添加参数 -R 2 即可开启 Symmetric Memory 模式。

Symmetric Memory 使用示例

以下是一个 Symmetric Memory 的使用示例：

1. 使用 ncclMemMalloc 函数，通过 VMM (Virtual Memory Management) API 分配缓冲区。
2. 使用 ncclCommWindowRegister 接口注册源/目标缓冲区。
3. 调用集合通信算子，例如 ncclAllgather。
4. 使用 ncclCommWindowDeRegister 释放 Symmetric Memory 窗口。
5. 调用 ncclMemFree 释放 VMM API 分配的缓冲区。

要深入理解 Symmetric Memory，需要对 VMM 的关键机制有所了解。

需要注意的是，在上述示例中，Allgather 的输入和输出缓冲区都需要添加偏移量。

CUDA 的虚拟内存管理 (Virtual Memory Management, VMM)

CUDA 10.2 版本引入了虚拟内存管理 (VMM)，提供了一系列新的 CUDA Driver API。VMM 的核心思想是将显存的虚拟地址和物理地址解耦，允许开发者独立管理它们。

CUDA 10.2 之前的内存分配方式

在 CUDA 10.2 之前，开发者主要通过以下 API 分配 GPU 内存：

// cuda api
cudaError_t cudaMalloc(void** devPtr, size_t size);
cudaError_t cudaFree(void* devPtr);
CUresult cuMemAlloc(CUdeviceptr* dptr, size_t bytesize);
CUresult cuMemFree(CUdeviceptr dptr);

这些 API 直接返回可用的虚拟地址。

CUDA 10.2 引入的 VMM 接口

新的 VMM API 允许开发者先保留一段虚拟地址空间，再单独分配物理内存，并将二者手动映射起来：

CUresult cuMemGetAllocationGranularity(size_t* granularity, const CUmemAllocationProp* prop, CUmemAllocationGranularity_flags option);

VMM API 要求分配粒度 (granularity) 对齐，而 cudaMalloc 和 cuMemAlloc 没有这个约束。实际上，cudaMalloc 也会与分配粒度对齐，只不过用户无需感知。

CUDA 实现 Vector 面临的问题

英伟达官方提供了一个基于 VMM API 实现 vector 的完整案例。通过对比传统的 cudaMalloc 实现和基于 VMM API 的实现，可以更清晰地理解它们的差异。

首先，我们来看一下 vector 类的定义：它使用一个 reserve 接口用于提前保留一块内存，grow 接口用于在 reserve size 不够用时，增长分配空间的大小。

cudaMalloc 实现方案

在 cudaMalloc 版本的 reserve 实现中：

1. 当需要的容量 new_sz 大于当前已保留容量 reserve_sz 时，必须重新分配一块更大的内存（new_sz）。
2. 然后把原有数据拷贝到新分配的内存中。
3. 最后释放旧内存，用新指针 new_ptr 作为存储区域。

这种方式的缺点是显而易见的：有额外的数据拷贝开销，而且对内存需求更大。例如，当 reserve_sz 为 1GB、new_sz 为 2GB 时，GPU 上需要同时有 3GB 的空闲显存才能完成扩容，这会造成空间浪费，甚至可能导致分配失败。

基于 VMM API 实现方案

VMM 提供了虚拟地址与物理内存解耦的能力，因此需要在 vector 中记录更多的信息。

• chunk_sz：通过 cuMemGetAllocationGranularity 获取的最小分配粒度。
• handles：存储物理内存句柄（由 cuMemCreate 分配）。
• handle_sizes：记录每个句柄的分配大小。

在 VMM 版本的 reserve 函数中，如果需要的空间比剩余 VA 空间大，则调用 cuMemAddressReserve 试图在 d_p + reserve_sz 之后继续保留一段连续 VA 空间。但是如果连续空间被其他应用占用，则会返回一段不连续的 VA（new_ptr ≠ d_p + reserve_sz），此时就需要走到 slow path，重新分配一段更大的 VA 空间，并且重新 cuMemMap 已有的物理内存到新 VA 空间。当调用到 grow 函数时，新的物理内存通过 cuMemCreate 分配，并映射到 VA 中剩余的空闲位置中。

当需要分配 block3 的空间，但是 d_p 空间不够时，会重新分配一块更大的 VA 空间 d_p_1。新分配的空间并没有和 vector 之间的物理内存建立映射关系，需要重新建立，例如 d_p_1 VA 块 0 → handles[0] 等。原来 d_p 废弃的映射也需要通过 unmap 函数释放，d_p reserve 的 VA 在此之后释放。在 grow 函数中，分配了新的物理空间（handles[3])，并将 d_p_1 VA 的 block3 映射到这块物理空间上。

P2P 场景下使用 VMM

在传统的实现中，如果两个 GPU 之间需要进行 Peer-to-Peer (P2P) 访问，必须显式调用 cudaDeviceEnablePeerAccess API。在使用 VMM 之后，P2P 访问的控制更加精细化：只需调用一次 cuMemSetAccess，即可为目标内存设置特定的 peer access 权限。可以针对单个内存区域设置访问权限，而不是全局启用。这种方式支持更细粒度的安全和性能优化。

CUmemAccessDesc accessDesc = {};

实现 Symmetric Memory - Host 端准备工作

回到第二节的示例，步骤 1 中调用的 ncclMemMalloc 本质上是通过 VMM API 分配空间并完成映射，返回 VA 地址。而步骤 2 中调用的 ncclCommWindowRegister 则注册 src/dst buffer，该 API 是 NCCL 2.27 新增的用户接口，负责完成 symmetric memory 相关注册工作。symmetric memory 的主要机制在该函数中实现。新媒网跨境获悉，该技术能够显著提升数据传输效率。

/* Register memory window  */
ncclResult_t ncclCommWindowRegister(ncclComm_t comm, void* ptr, size_t size);

ncclCommWindowRegister 函数执行流程概览

1. 定位输入 buffer 对应的物理内存。

   调用 cuMemRetainAllocationHandle，根据输入 buffer 的 VA 地址找到对应物理内存的 handle（PA handle）。

2. 构建注册任务。

   根据 PA 信息及 size 相关信息，构建 ncclSymRegTask，将其加入 comm->symRegTaskQueue 任务处理队列。

3. 异步任务处理。

   异步执行 ncclCommGroupRegisterSymmetric 处理队列中的任务。步骤 4 和步骤 5 是本步骤的具体执行过程。

4. 初始化 symmetric memory 以及在 symmetric memory 上分配内部使用结构（首次调用时）。

   分配 symmetric VA space 以及在 symmetric memory 上分配内部管理结构 struct ncclSymDevBase；调用 ncclCommSymmetricRegisterInternal，将 ncclSymDevBase 结构物理内存映射到所有本地 rank 的对称虚拟地址，实现跨进程共享。

5. 注册内存窗口（处理队列中的任务）。

   处理 symRegTaskQueue 队列中的任务。调用 ncclCommSymmetricRegisterInternal 接口，将物理内存映射到所有本地 rank 的对称虚拟地址位置，实现跨进程的内存共享。

异步处理依赖 NCCL 的 group 机制，可同时处理多个通信器任务。

NCCL 中的 Symmetric Memory 管理结构

struct ncclComm {
    // ...
    void* baseUCSymPtr; // 通过 cuMemAddressReserve 分配
    void* baseMCSymPtr; // 通过 cuMemAddressReserve 分配，用于 NVLS 多播
    size_t baseStride;   // 每个 rank 的 VA 窗口大小
    size_t symAllocHead; // 记录当前分配位置
    struct ncclSymRegTaskQueue symRegTaskQueue; // symmetric memory 注册任务队列
    // ...
};

• baseUCSymPtr：通过 cuMemAddressReserve 分配，大小为 comm->baseStride * comm->localRanks。
• baseMCSymPtr：通过 cuMemAddressReserve 分配，大小为 comm->baseStride，用于 NVLS 多播。
• baseStride：每个 rank 的 VA 窗口大小。如果设置了 WIN_STRIDE 环境变量，则直接使用；否则取所有 rank 中最大的显存大小作为窗口大小，例如 96GB。
• symAllocHead：记录当前分配位置。
• symRegTaskQueue：symmetric memory 注册任务的队列。

映射流程示例（3 个 rank）

以 UC 对称内存为例（MC 类似），ncclSymDevBase 是 NCCL Symmetric 内存系统的设备端核心结构，初始化时预分配。在初始化的时候，每个 rank reserve 的 VA 空间大小是一样的。默认 VA 空间大小为 comm->baseStride * comm->localRanks。comm->symAllocHead 初值为 0。

步骤：

1. 分配物理内存：每个 rank 调用 cuMemCreate 为 ncclSymDevBase 分配物理内存。
2. 导出 shared handle：多进程场景下，每个 rank 调用 cuMemExportToShareableHandle 导出 handle。
3. 交换 handle：所有 rank 通过 allgather 获得全部 handle；

Rank 0: [memHandle0, memHandle1, memHandle2]

4. 导入 shared handle → impHandle：每个 rank 调用 cuMemImportFromShareableHandle 导入其他 rank 的 handle。
5. 映射到本地 symmetric memory VA 空间：每个 impHandle 被映射到 baseUCSymPtr + targetRank * baseStride + offset。

ncclSymDevBase 结构，这里 offset=0。例如，对于 rank 1 上分配的 ncclSymDevBase 结构，映射到 rank 0 的 comm->baseUCSymPtr + 1 * comm->baseStride 位置；映射到 rank 1 的 comm->baseUCSymPtr + 1 * comm->baseStride；映射到 rank 2 相同偏移位置。VA 和 PA 绑定后，每个 rank 都能访问其他 rank 的物理空间。

6. 更新 comm->symAllocHead 指针。

新媒网跨境认为，随着 symmetric memory 技术的不断发展和完善，未来的 GPU 集群通信将更加高效、便捷。

新媒网（公号: 新媒网跨境发布），是一个专业的跨境电商、游戏、支付、贸易和广告社区平台，为百万跨境人传递最新的海外淘金精准资讯情报。

本文来源：新媒网 https://nmedialink.com/posts/13998.html