用户缓冲区注册

1. 摘要：

一项功能，NCCL直接通过用户缓冲区发送/接收/操作数据，而无需额外的内部复制（零复制）。**加速集合通信和降低资源的使用（#channel使用率）。**NCCL给了两种注册方式：

• CUDA Graph注册
• 本地注册

在使用 NCCL进行通信时，如果任何一个 rank（即进程或线程）传递了已经注册的缓冲区（registered buffer）给 NCCL 的通信函数（比如 allreduce、sendrecv 等），那么同一个 communicator 中的其他所有 rank 也必须传递已注册的缓冲区。如果有 rank 传递了已注册的缓冲区，而其他 rank 传递了非注册的缓冲区（non-registered buffer），则会导致 未定义的行为。

2. [2.19.x]起，用NVLS

自 2.19.x 起，NCCL 支持 NVLink Sharp (NVLS) 的用户缓冲区注册；任何支持 NVLS 算法的 NCCL 集合（例如 allreduce）都可以利用此功能。

1. 启用CUDA Graph 的NVLS 缓冲区注册，用户必须遵守：
   • 缓冲区是通过**ncclMemAlloc()**或合格的分配器分配的
   • NCCL 操作在每个rank的 CUDA 图捕获的流上启动
   • 对于每个rank，缓冲区头地址的偏移量是相同的。（offset） 当CUDA graph销毁时，已注册的缓冲区将被注销。 https://docs.nvidia.com/deeplearning/nccl/user-guide/docs/usage/bufferreg.html#ib-sharp-buffer-registration
2. 启用基于本地的 NVLS 缓冲区注册，用户必须遵守：
   • 缓冲区是通过**ncclMemAlloc()**或合格的分配器分配的
   • **ncclCommRegister()**在调用每个rank的集体之前注册缓冲区
   • 像往常一样调用 NCCL 集合，但同样保持每个rank的缓冲区头地址的偏移量相同

当用户显式调用 时，已注册的缓冲区将被注销**ncclCommDeregister()**。对于基于本地的注册，用户可以在程序开始时注册一次缓冲区，并多次重复使用缓冲区以利用注册优势。

3. [2.21.x]IB Sharp Buffer 注册

目前，NCCL 仅支持每个节点包含 1 个rank的通信器的 IB Sharp 缓冲区注册，并且注册可以将 NCCL SM 的使用数量减少到 1。

要通过 CUDA 图启用 IB Sharp 缓冲区注册：

• 使用任何 CUDA allcator（例如，cudaMalloc/ncclMemAlloc）分配发送和接收缓冲区
• 使用 CUDA 图形启动 NCCL 集合

要通过本地注册启用 IB Sharp 缓冲区注册：

• 使用任何 CUDA allcator（例如，cudaMalloc/ncclMemAlloc）分配发送和接收缓冲区
• 使用ncclCommRegister为通信器中的每个等级注册发送和接收缓冲区
• 启动 NCCL 集体

4. [2.23.x]通用缓冲区注册(General Buffer Registration)

该功能面向所有点对点节点内通信（例如 Allgather Ring），可减少内存压力，并提升通信和计算重叠性能。无论是在初始阶段使用ncclCommRegister注册缓冲区，还是应用 CUDA 图，都可以为 NCCL 集合和 sendrecv 启用节点内缓冲区注册。

用户缓冲区可以通过 VMM API（例如cuMem*）分配，任何基于 VMM 的分配器（内存分配器）或ncclMemAlloc均可使用。通过旧版 cuda API（例如cudaMalloc ）分配的缓冲区也可用于注册。然而，由于执行期间可能出现挂起，以及失败和中止期间出现分段错误，因此使用旧版缓冲区注册并不安全；目前，旧版缓冲区注册默认处于禁用状态，用户可以设置NCCL_LEGACY_CUDA_REGISTER=1来启用它。

5. 内存分配器(Memory Allocator)

NCCL 提供了ncclMemAlloc函数，帮助用户通过 VMM API 分配缓冲区，以便后续注册 NCCL。该函数仅适用于 NCCL，因此不建议在应用程序中到处使用ncclMemAlloc分配的缓冲区。

对于高级用户，如果要为NVLS UB创建自己的内存分配器，则分配器的分配缓冲区需要满足以下要求：

• 在支持的 GPU 上使用共享标志CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR和CU_MEM_HANDLE_TYPE_FABRIC分配缓冲区。
• 缓冲区物理内存大小是 CUMEM 推荐粒度的倍数（即 cuMemGetAllocationGranularity(…, CU_MEM_ALLOC_GRANULARITY_RECOMMENDED`））
• 缓冲区虚拟头地址至少与 CUMEM 建议粒度对齐，并且大小是 CUMEM 建议粒度的倍数。

对于使用 VMM API 的通用缓冲区注册，分配器需要满足与 NVLS UB 分配器相同的要求。