用户缓冲区注册¶

用户缓冲区注册是一项功能，允许 NCCL 直接通过用户缓冲区发送/接收/操作数据，而无需额外的内部复制（零复制）。它可以加速集体操作并大大减少资源使用量（例如，#通道使用量）。NCCL 提供了两种注册用户缓冲区的方法；一种是CUDA Graph 注册，另一种是本地注册。NCCL 要求，对于所有 NCCL 通信函数调用（例如，allreduce、sendrecv 等），如果通信器中的任何 rank 将注册缓冲区传递给 NCCL 通信函数，则同一通信器中的所有其他 rank 都必须传递其注册缓冲区；否则，混合使用注册缓冲区和非注册缓冲区可能会导致未定义的行为。

NVLink Sharp 缓冲区注册¶

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

要为 NVLS 启用基于 CUDA Graph 的缓冲区注册，用户必须遵守以下几个要求

缓冲区通过 ncclMemAlloc() 或合格的分配器分配（参见内存分配器）。

NCCL 操作在 CUDA graph 捕获的流上为每个 rank 启动。

对于每个 rank 的集体操作，到缓冲区头地址的偏移量是相同的。

注册的缓冲区将在 CUDA graph 销毁时被注销。这是一个基于 CUDA graph 的缓冲区注册示例

void* sendbuff;
void* recvbuff;
size_t count = 1 << 25;
CHECK(ncclMemAlloc(&sendbuff, count * sizeof(float)));
CHECK(ncclMemAlloc(&recvbuff, count * sizeof(float)));

cudaGraph_t graph;
CHECK(cudaStreamBeginCapture(stream, cudaStreamCaptureModeThreadLocal));
CHECK(ncclAllReduce(sendbuff, recvbuff, 1024, ncclFloat, ncclSum, comm, stream));
// Same offset to the sendbuff and recvbuff head address for each rank
CHECK(ncclAllReduce((void*)((float*)sendbuff + 1024), (void*)((float*)recvbuff + 2048), 1024, ncclFloat, ncclSum, comm, stream));
CHECK(cudaStreamEndCapture(stream, &graph));

cudaGraphExec_t instance;
CHECK(cudaGraphInstantiate(&instance, graph, NULL, NULL, 0));
CHECK(cudaGraphLaunch(instance, stream));
CHECK(cudaStreamSynchronize(stream));
CHECK(cudaGraphExecDestroy(instance));
CHECK(cudaGraphDestroy(graph));

CHECK(ncclMemFree(sendbuff));
CHECK(ncclMemFree(recvbuff));

另一方面，要为 NVLS 启用基于本地的缓冲区注册，用户必须遵守以下要求

缓冲区通过 ncclMemAlloc() 或合格的分配器分配（参见内存分配器）。

在为每个 rank 调用集体操作之前，使用 ncclCommRegister() 注册缓冲区。

像往常一样调用 NCCL 集体操作，但同样保持每个 rank 到缓冲区头地址的偏移量相同。

注册的缓冲区将在用户显式调用 ncclCommDeregister() 时被注销。这是一个基于本地的缓冲区注册示例

void* sendbuff;
void* recvbuff;
size_t count = 1 << 25;
void* sendRegHandle;
void* recvRegHandle;
CHECK(ncclMemAlloc(&sendbuff, count * sizeof(float)));
CHECK(ncclMemAlloc(&recvbuff, count * sizeof(float)));

CHECK(ncclCommRegister(comm, sendbuff, count * sizeof(float), &sendRegHandle));
CHECK(ncclCommRegister(comm, recvbuff, count * sizeof(float), &recvRegHandle));

CHECK(ncclAllReduce(sendbuff, recvbuff, 1024, ncclFloat, ncclSum, comm, stream));
CHECK(ncclAllReduce((void*)((float*)sendbuff + 1024), (void*)((float*)recvbuff + 2048), 1024, ncclFloat, ncclSum, comm, stream));
CHECK(cudaStreamSynchronize(stream));

CHECK(ncclCommDeregister(comm, sendRegHandle));
CHECK(ncclCommDeregister(comm, recvRegHandle));

CHECK(ncclMemFree(sendbuff));
CHECK(ncclMemFree(recvbuff));

对于基于本地的注册，用户可以在程序开始时注册一次缓冲区，并多次重复使用该缓冲区以利用注册优势。

为了节省内存，分配一大块缓冲区并注册一次也是有效的。 sendbuff 和 recvbuff 可以通过大块进一步分配，用于零复制 NCCL 操作，只要 sendbuff 和 recvbuff 满足偏移量要求。以下示例显示了一个用例

void* buffer;
void* handle;
void* sendbuff;
void* recvbuff;
size_t size = 1 << 29;

CHECK(ncclMemAlloc(&buffer, size));
CHECK(ncclCommRegister(comm, buffer, size, &handle));

// assign buffer chunk to sendbuff and recvbuff
sendbuff = buffer;
recvbuff = (void*)((uint8_t*)buffer + (1 << 20));

CHECK(ncclAllReduce(sendbuff, recvbuff, 1024, ncclFloat, ncclSum, comm, stream));
CHECK(cudaStreamSynchronize(stream));

CHECK(ncclCommDeregister(comm, handle));

CHECK(ncclMemFree(sendbuff));

IB Sharp 缓冲区注册¶

NCCL 2.21.x 支持 IB Sharp 缓冲区注册，任何支持 IB Sharp 算法的 NCCL 集体操作都可以从中受益，例如 allreduce、reducescatter 和 allgather。目前，NCCL 仅支持每个节点包含 1 个 rank 的通信器的 IB Sharp 缓冲区注册，并且注册可以将 NCCL SM 使用量减少到 1。

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

使用任何 CUDA 分配器（例如，cudaMalloc/ncclMemAlloc）分配发送和接收缓冲区

使用 CUDA graph 启动 NCCL 集体操作

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

使用任何 CUDA 分配器（例如，cudaMalloc/ncclMemAlloc）分配发送和接收缓冲区

使用 ncclCommRegister 为通信器中的每个 rank 注册发送和接收缓冲区

启动 NCCL 集体操作

通用缓冲区注册¶

自 2.23.x 起，NCCL 支持节点内缓冲区注册，它针对所有点对点节点内通信，并带来更少的内存访问、更少的 SM 使用量和性能提升。在开始时通过 ncclCommRegister 注册缓冲区或应用 CUDA graph 都可以为 NCCL 集体操作和 sendrecv 启用节点内缓冲区注册。注册的缓冲区可以通过旧版 cuda API（例如，cudaMalloc）以及 VMM API（例如，cuMem* 或 ncclMemAlloc）分配。但是，强烈建议使用 VMM 分配的缓冲区，因为它比失败和中止期间的旧版缓冲区更安全。

内存分配器¶

为方便起见，NCCL 提供了 ncclMemAlloc 函数，以帮助用户通过 VMM API 分配缓冲区，该缓冲区稍后可用于 NCCL 注册。它仅为 NCCL 设计，因此不建议在应用程序中的任何地方使用 ncclMemAlloc 分配的缓冲区。对于高级用户，如果您想为 NVLS 缓冲区注册创建自己的内存分配器，则该分配器需要满足以下要求

在支持的 GPU 上，使用共享标志 CU_MEM_HANDLE_TYPE_POSIX_FILE_DESCRIPTOR 和 CU_MEM_HANDLE_TYPE_FABRIC 分配缓冲区。

缓冲区大小是多播建议粒度的倍数 (即 cuMulticastGetGranularity(…, CU_MULTICAST_GRANULARITY_RECOMMENDED))

缓冲区头地址至少与多播最小粒度对齐 (即 cuMulticastGetGranularity(…, CU_MULTICAST_GRANULARITY_MINIMUM))