企业官网建设流程全解析

• 高性能计算
• 通信
• 深度学习
• 分布式训练

本篇技术指南围绕 NVIDIA Collective Communications Library（NCCL）的点对点（Point-to-Point, P2P）通信原语展开，系统讲解两类操作：以ncclSend/ncclRecv为代表的**双边（two-sided）通信，以及以ncclPutSignal/ncclSignal/ncclWaitSignal为代表的单边（one-sided）远程内存访问（RMA）**通信。文中完整覆盖每个 API 的函数签名、参数语义、完成时序与分组（Group）用法，并结合当前仓库的源码实现与官方示例，给出 sendrecv、scatter/gather、all-to-all、邻居交换、ping-pong、barrier 等可直接落地的编程模式。读完本文，你将能够使用 NCCL 点对点接口构建任意自定义通信拓扑，并掌握对称内存窗口注册与信号同步的完整流程。

一、概述：NCCL 的两类点对点原语

NCCL 提供两类点对点通信原语：

• 双边操作（Two-sided operations）：自 NCCL 2.7 引入。当各 rank 之间需要互发互收彼此不同的数据、且这种交换无法用 broadcast 或 allgather 表达时（即每个 rank 发送和接收的数据内容各不相同），就必须使用双边点对点通信。发送方与接收方都必须显式参与：一方调用ncclSend，另一方调用ncclRecv。
• 单边操作（One-sided operations, RMA）：自 NCCL 2.29 引入。允许一个 rank 直接写入远端内存，无需目标进程显式发起匹配操作。其前提是目标内存已通过ncclCommWindowRegister在对称内存窗口（symmetric memory window）中预注册；同步则通过目标 rank 调用ncclWaitSignal等待信号完成。

这两类原语在 src/include/nccl_common.h 的ncclFunc_t枚举中对应不同的函数类型：ncclFuncSend（6）、ncclFuncRecv（7）以及ncclFuncPutSignal（12）、ncclFuncSignal（13）、ncclFuncWaitSignal（14），与 collective 函数类型并列，共同经统一的入队机制调度。

二、双边点对点操作：ncclSend 与 ncclRecv

2.1 适用场景

双边操作适用于任意 rank 之间数据各不相同的通信模式。每次点对点通信需要两个 NCCL 调用：一个 rank 调用ncclSend，另一个 rank 调用ncclRecv，且双方必须使用相同的count和相同的datatype。两个 API 的完整签名在 API 参考文档 中定义，声明位于 src/nccl.h.in。

2.2 ncclSend：发送数据到指定 peer

ncclResult_t ncclSend(const void* sendbuff, size_t count, ncclDataType_t datatype, int peer, ncclComm_t comm, cudaStream_t stream);

语义要点（出自 API 文档）：

• 将数据从sendbuff发送给 rankpeer；
• rankpeer需要以与本 rank相同的datatype和相同的count调用ncclRecv；
• 该操作对 GPU 是阻塞的（blocking for the GPU）。若多个ncclSend/ncclRecv操作需要并发推进才能完成，则必须将它们融合在ncclGroupStart/ncclGroupEnd分组段内执行。

2.3 ncclRecv：从指定 peer 接收数据

ncclResult_t ncclRecv(void* recvbuff, size_t count, ncclDataType_t datatype, int peer, ncclComm_t comm, cudaStream_t stream);

语义要点：从 rankpeer接收数据到recvbuff；对端需以相同datatype与count调用ncclSend；同样对 GPU 阻塞，多操作并发时需用 Group 融合。

ncclDataType_t支持的类型包括ncclInt8/ncclChar、ncclUint8、ncclInt32/ncclInt、ncclUint32、ncclInt64、ncclUint64、ncclFloat16/ncclHalf、ncclFloat32/ncclFloat、ncclFloat64/ncclDouble、ncclBfloat16、ncclFloat8e4m3、ncclFloat8e5m2，定义见 src/nccl.h.in。

2.4 Group 语义：并发推进与死锁规避

点对点调用本身对 GPU 阻塞，因此构建多跳通信模式时必须使用分组：

• 组内的点对点调用在整个分组完成之前阻塞，但组内各调用可视为独立并发推进，彼此不应相互阻塞；
• 因此，必须把需要并发推进的调用合并进同一个分组，以避免死锁；
• 唯一例外：同一分组内指向同一个 peer的点对点调用会按程序顺序串行执行。

以ncclCommInitAll创建的多 GPU 单进程场景为例，若所有 GPU 都先发送再接收，不分组就会互相等待对方先接收而陷入死锁。仓库示例 c/main.cc 明确注释了这一点：ncclGroupStart/ncclGroupEnd在环形模式中必不可少，分组让 NCCL 以最优顺序执行操作。

2.5 经典双边通信模式（可直接复制运行）

以下示例均来自 点对点使用指南。NCCL 语义允许每个 rank 使用不同的 size、datatype 和 buffer 组合。

Sendrecv（双方同时收发）——用 MPI 术语说，sendrecv 是两个 rank 同时发送并接收数据，只需把两个调用合并到一个分组：

ncclGroupStart(); ncclSend(sendbuff, sendcount, sendtype, peer, comm, stream); ncclRecv(recvbuff, recvcount, recvtype, peer, comm, stream); ncclGroupEnd();

One-to-all（scatter）——从rootrank 向所有其他 rank 分发数据：

ncclGroupStart(); if (rank == root) { for (int r=0; r<nranks; r++) ncclSend(sendbuff[r], size, type, r, comm, stream); } ncclRecv(recvbuff, size, type, root, comm, stream); ncclGroupEnd();

All-to-one（gather）——所有 rank 向rootrank 汇聚数据：

ncclGroupStart(); if (rank == root) { for (int r=0; r<nranks; r++) ncclRecv(recvbuff[r], size, type, r, comm, stream); } ncclSend(sendbuff, size, type, root, comm, stream); ncclGroupEnd();

All-to-all——向所有 peer 发送并向所有 peer 接收：

ncclGroupStart(); for (int r=0; r<nranks; r++) { ncclSend(sendbuff[r], sendcount, sendtype, r, comm, stream); ncclRecv(recvbuff[r], recvcount, recvtype, r, comm, stream); } ncclGroupEnd();

Neighbor exchange（N 维空间邻居交换）——沿各维度与前后邻居交换数据：

ncclGroupStart(); for (int d=0; d<ndims; d++) { ncclSend(sendbuff[d], sendcount, sendtype, next[d], comm, stream); ncclRecv(recvbuff[d], recvcount, recvtype, prev[d], comm, stream); } ncclGroupEnd();

2.6 源码实现与设备端内核

在 src/collectives.cc 中，ncclSend与ncclRecv的宿主端实现将参数封装为struct ncclInfo（函数类型ncclFuncSend/ncclFuncRecv、缓冲区、count、datatype、peer、comm、stream 等），随后调用统一的ncclEnqueueCheck完成入队，并带有 NVTX 性能标记。

设备端执行路径见 src/device/sendrecv.h：RunWorkBatch<ncclFuncSendRecv, ...>针对 P2P 的特化通过runSend/runRecv将工作切分到各 channel，以prims.directSend/prims.directRecv直接读写对端地址完成传输。ncclFuncSend与ncclFuncRecv在底层复用同一套SendRecv内核机制。

三、单边 RMA 操作：PutSignal、Signal 与 WaitSignal

3.1 概念与前置条件

单边远程内存访问操作使 rank 无需目标进程显式参与即可直接访问远端内存。关键前提：

• 目标内存必须预先通过ncclCommWindowRegister注册在对称内存窗口（symmetric memory window）中；
• 点对点同步可通过目标 rank 调用ncclWaitSignal等待信号来实现。

对称内存通常通过ncclMemAlloc分配。这两个辅助 API 的声明位于 src/nccl.h.in 与 src/nccl.h.in，实现分别位于 src/allocator.cc 与 src/dev_runtime.cc。

窗口注册 flag（src/nccl.h.in）：

此外NCCL_WIN_REQUIRED_ALIGNMENT（4096）规定了窗口注册所需的对齐要求。

3.2 ncclPutSignal：带信号的数据写入

ncclResult_t ncclPutSignal(const void* localbuff, size_t count, ncclDataType_t datatype, int peer, ncclWindow_t peerWin, size_t peerWinOffset, int sigIdx, int ctx, unsigned int flags, ncclComm_t comm, cudaStream_t stream);

完成语义（出自 API 文档）：

• 将数据从localbuff写入 rankpeer注册窗口peerWin的peerWinOffset偏移处，随后更新一个远端信号；
• ncclPutSignal返回到 CPU 线程，仅表示操作已成功入队到 CUDA stream；
• 当ncclPutSignal在 CUDA stream 上执行完成时，localbuff即可安全复用或修改；
• 当远端 peer 的信号被更新时，保证对应的ncclPutSignal数据已送达远端内存，并且所有先前发往同一 peer 与同一 ctx 的ncclPutSignal/ncclSignal操作的信号更新也已完成。

3.3 ncclSignal：纯信号，不传输数据

ncclResult_t ncclSignal(int peer, int sigIdx, int ctx, unsigned int flags, ncclComm_t comm, cudaStream_t stream);

向 rankpeer发送一个不携带数据的信号。sigIdx、ctx、flags的约束与ncclPutSignal相同（当前必须为 0）。当远端 peer 的信号被更新时，所有先前发往同一 peer 与同一 ctx 的ncclPutSignal/ncclSignal操作的信号更新也都已完成——因此ncclSignal可用于在已由ncclPutSignal写入的数据之后追加一个同步点。

3.4 ncclWaitSignal：等待来自特定 rank 的信号

等待逻辑通过描述符数组表达。描述符类型ncclWaitSignalDesc_t定义于 src/nccl.h.in：

typedef struct { int opCnt; // 需要等待的信号操作个数 int peer; // 等待哪个 peer 的信号 int sigIdx; // 信号索引标识（当前必须为 0） int ctx; // 上下文标识（当前必须为 0） } ncclWaitSignalDesc_t;

ncclResult_t ncclWaitSignal(int nDesc, ncclWaitSignalDesc_t* signalDescs, ncclComm_t comm, cudaStream_t stream);

• nDesc：signalDescs数组中描述符的个数；
• 每个描述符指明：期望从特定peer、在特定信号索引sigIdx与上下文ctx上收到多少个信号（opCnt）；
• ncclWaitSignal返回 CPU 线程仅表示操作已成功入队；
• 当ncclWaitSignal在 CUDA stream 上完成时，所有指定信号均已被接收，对应的数据在本地内存中可见。

3.5 单边操作的分组语义

多个ncclPutSignal调用可以用ncclGroupStart/ncclGroupEnd分组：

• 组内发往不同 peer 或不同 ctx的操作可并发执行、以任意顺序完成；
• ncclGroupEnd返回即保证组内所有操作均已达成完成；
• 发往同一 peer 与同一 ctx的操作按程序顺序执行，且数据交付与信号更新都遵循程序顺序。

在 src/rma/rma.cc 中可以看到 RMA 的调度实现：ncclFuncPutSignal/ncclFuncSignal走ncclRmaPut路径，ncclFuncWaitSignal走ncclRmaWaitSignal路径，并且 CE（Collective Engine）与 Proxy 两类任务可并行执行；连续的 put/signal 操作会尝试批处理进同一个 plan（canBatchRmaTasks允许 Put 与 Signal 任务合并批处理），而ncclWaitSignal因 API 本身已支持一次等待多个 peer 的信号而不再做进一步批处理。

3.6 经典单边通信模式（可直接复制运行）

以下示例均来自 点对点使用指南。

PutSignal + WaitSignal（ping-pong）——完整展示从内存分配、窗口注册到收尾清理的全流程：

// Allocate symmetric memory for RMA operations void *sendbuff, *recvbuff; NCCLCHECK(ncclMemAlloc((void**)&sendbuff, size)); NCCLCHECK(ncclMemAlloc((void**)&recvbuff, size)); // Register buffers as symmetric windows ncclWindow_t sendWindow, recvWindow; NCCLCHECK(ncclCommWindowRegister(comm, sendbuff, size, &sendWindow, NCCL_WIN_COLL_SYMMETRIC)); NCCLCHECK(ncclCommWindowRegister(comm, recvbuff, size, &recvWindow, NCCL_WIN_COLL_SYMMETRIC)); int peer = (rank == 0) ? 1 : 0; ncclWaitSignalDesc_t waitDesc = {.opCnt = 1, .peer = peer, .sigIdx = 0, .ctx = ctx}; if (rank == 0) { // Rank 0: wait then put NCCLCHECK(ncclWaitSignal(1, &waitDesc, comm, stream)); NCCLCHECK(ncclPutSignal(sendbuff, count, datatype, peer, recvWindow, 0, 0, 0, 0, comm, stream)); } else { // Rank 1: put then wait NCCLCHECK(ncclPutSignal(sendbuff, count, datatype, peer, recvWindow, 0, 0, 0, 0, comm, stream)); NCCLCHECK(ncclWaitSignal(1, &waitDesc, comm, stream)); } CUDACHECK(cudaStreamSynchronize(stream)); // Cleanup NCCLCHECK(ncclCommWindowDeregister(comm, sendWindow)); NCCLCHECK(ncclCommWindowDeregister(comm, recvWindow)); NCCLCHECK(ncclMemFree(sendbuff)); NCCLCHECK(ncclMemFree(recvbuff));

Barrier（屏障）——每个 rank 向所有其他 rank 发信号，并等待来自所有 rank 的信号：

ncclWaitSignalDesc_t *waitDescs = malloc(nranks * sizeof(ncclWaitSignalDesc_t)); for (int r = 0; r < nranks; r++) { waitDescs[r].opCnt = 1; waitDescs[r].peer = r; waitDescs[r].sigIdx = 0; waitDescs[r].ctx = 0; } ncclGroupStart(); for (int r = 0; r < nranks; r++) { ncclSignal(r, 0, 0, 0, comm, stream); } ncclGroupEnd(); ncclWaitSignal(nranks, waitDescs, comm, stream);

All-to-all（单边版本）——每个 rank 向所有其他 rank 写入数据，并等待来自所有 rank 的信号：

size_t offset[nranks]; ncclWaitSignalDesc_t *waitDescs = malloc(nranks * sizeof(ncclWaitSignalDesc_t)); for (int r = 0; r < nranks; r++) { offset[r] = r * count * wordSize(datatype); waitDescs[r].opCnt = 1; waitDescs[r].peer = r; waitDescs[r].sigIdx = 0; waitDescs[r].ctx = 0; } ncclGroupStart(); for (int r = 0; r < nranks; r++) { ncclPutSignal(sendbuff[r], count, datatype, r, window, offset[r], 0, 0, 0, comm, stream); } ncclGroupEnd(); ncclWaitSignal(nranks, waitDescs, comm, stream);

使用注意：需要提前为每个 peer 用ncclCommWindowRegister注册内存窗口；调用ncclPutSignal前用户需保证各缓冲区已就绪（可通过上面的 barrier 保证）。

3.7 源码级参数校验与约束

在 src/enqueue.cc 中，入队路径对 RMA 参数做了严格校验，这正是文档中"当前必须为 0"约束的落地实现：

• sigIdx != 0或flags != 0均返回ncclInvalidArgument；
• ncclPutSignal：peerWin不得为空；源缓冲区srcBuff必须位于一个带NCCL_WIN_COLL_SYMMETRIC标志的对称窗口内，否则报错；当前不支持由多个物理 cuMem 段支撑的 VA，也不支持 host 内存段支撑的 VA；
• ncclSignal：count必须为 0；
• ncclWaitSignal：signalDescs非空且nDesc > 0，每个描述符的opCnt > 0、sigIdx == 0、ctx == 0。

若 RMA CE 尚未初始化，入队时会自动挂起一个 CE 初始化任务，保证信号机制就绪。

四、完整可运行示例：环形通信（Ring Pattern）

仓库在 docs/examples/02_point_to_point/01_ring_pattern/ 提供了完整的双边点对点教学示例，展示如何在单进程多 GPU 下用ncclSend/ncclRecv构建环形通信：

• C 版本（c/main.cc）：流程为cudaGetDeviceCount探测 GPU →ncclCommInitAll建立每 GPU 一个通信子 → 创建 stream、分配 1GB 设备缓冲区 →在ncclGroupStart/ncclGroupEnd内让每个 GPUncclSend给下一跳、ncclRecv自上一跳 →cudaStreamSynchronize同步 → 拷贝回主机端校验数据正确性 →ncclCommFinalize/ncclCommDestroy清理；
• Python 版本（python/ring_pattern.py）：基于 nccl4py 绑定（nccl.core.Communicator.init_all()）与 CuPy，演示同样拓扑的send()/recv()用法。

环形拓扑是理解 NCCL 许多 collective 算法的基础；该示例同时演示了死锁规避（分组）与结果校验两个关键实践。

五、实践要点总结

1. 匹配规则：双边操作中ncclSend与ncclRecv必须成对出现，且count与datatype完全一致；
2. 分组防死锁：任何需要并发推进的多跳模式都必须用ncclGroupStart/ncclGroupEnd包裹；同 peer 的调用在组内保序执行；
3. 单边前置条件：RMA 写入前，目标窗口必须已通过ncclCommWindowRegister（带NCCL_WIN_COLL_SYMMETRIC）注册，内存建议用ncclMemAlloc分配；
4. 信号约束：当前版本sigIdx、ctx、flags均须为 0，ncclWaitSignalDesc_t.opCnt须大于 0，违反约束会在入队时被 src/enqueue.cc 校验拒绝；
5. 完成语义：RMA API 的 CPU 返回只代表入队成功；localbuff的安全复用以 stream 上操作完成为准，远端数据可见性以信号更新为准；
6. 顺序保证：同一 peer 与同一 ctx 的 put/signal 操作按程序顺序执行（数据交付与信号更新均如此），不同 peer/ctx 之间可并发完成。

如需进一步查阅，可参考 点对点 API 参考、点对点使用指南 以及 环形通信示例。

• 高性能计算
• 通信
• 深度学习
• 分布式训练