第1章：基础API使用 — NCCL通信域初始化与点对点通信

大家好，我是你们这门课的主讲。今天咱们来聊聊NCCL最基础的东西——通信域的初始化，以及最常用的点对点通信和集合通信操作。

说实话，很多同学一上来就追着问AllReduce怎么调优，但往往忽略了最基础的东西。我见过不少项目，明明硬件环境没问题，结果通信域初始化就卡住了，排查半天发现是网卡绑定没配好。嗯，咱们先把地基打牢。

1.1 NCCL通信域初始化

NCCL的通信域，说白了就是一组可以互相通信的GPU。你可以把它理解成一个群——每个GPU是一个群成员，初始化就是拉群的过程。

我个人习惯用 ncclCommInitAll 或者 ncclCommInitRank 来创建通信域。前者适合单节点多卡，后者适合多节点场景。

来看一个最简单的单节点初始化示例：

#include <nccl.h>
#include <cuda_runtime.h>

int main() {
    ncclComm_t comm;
    int nDevices = 4;
    int devs[4] = {0, 1, 2, 3};
    
    // 初始化4块GPU的通信域
    NCCLCHECK(ncclCommInitAll(&comm, nDevices, devs));
    
    // 你的业务逻辑...
    
    // 销毁通信域
    NCCLCHECK(ncclCommDestroy(comm));
    return 0;
}

这里有个坑，我曾经遇到过：ncclCommInitAll 要求所有GPU都在同一个节点上，而且必须对CUDA可见。如果你用 CUDA_VISIBLE_DEVICES 屏蔽了某块卡，初始化就会失败。

1.2 点对点通信：Send / Recv

点对点通信是NCCL最基础的通信方式。一个GPU发，一个GPU收，就这么简单。但你别小看它，很多集合通信的底层实现就是靠Send/Recv拼出来的。

核心API就两个：

• ncclSend(sendbuff, count, datatype, peer, comm, stream)
• ncclRecv(recvbuff, count, datatype, peer, comm, stream)

注意，这两个函数都是异步的。什么意思？就是调用后立即返回，实际数据传输在stream上排队。你需要用 cudaStreamSynchronize 或者 ncclGroupEnd 来确保完成。

来看一个实际例子：rank 0 发数据给 rank 1。

// rank 0: 发送
float* sendbuff;
cudaMalloc(&sendbuff, 1024 * sizeof(float));
// 填充数据...
ncclSend(sendbuff, 1024, ncclFloat, 1, comm, stream);

// rank 1: 接收
float* recvbuff;
cudaMalloc(&recvbuff, 1024 * sizeof(float));
ncclRecv(recvbuff, 1024, ncclFloat, 0, comm, stream);

// 等待完成
cudaStreamSynchronize(stream);

这里有个细节：Send和Recv必须配对。如果rank 0发了但rank 1没收，程序就会死锁。我刚开始学的时候犯过这个错，debug了一下午才发现是rank号写反了。

1.3 集合通信操作

集合通信是NCCL的精髓。你想想看，分布式训练里最耗时的就是梯度同步，而AllReduce就是干这个的。

常用的集合操作有这些：

操作	说明	典型场景
AllReduce	所有GPU的数据归约后，结果分发给所有GPU	梯度同步
AllGather	每个GPU收集所有GPU的数据	参数同步、模型并行
ReduceScatter	先归约，再分散到各GPU	Ring AllReduce的中间步骤
Broadcast	一个GPU的数据广播到所有GPU	初始化参数

咱们重点看AllReduce。它的API长这样：

ncclAllReduce(sendbuff, recvbuff, count, datatype, op, comm, stream);

其中 op 可以是 ncclSum、ncclProd、ncclMax 等。训练中最常用的是 ncclSum，因为梯度累加就是求和。

举个例子，4块GPU做AllReduce求和：

// 每块GPU上都有1024个float
float* sendbuff;
float* recvbuff;
cudaMalloc(&sendbuff, 1024 * sizeof(float));
cudaMalloc(&recvbuff, 1024 * sizeof(float));

// 填充sendbuff...

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

// 现在recvbuff里是4块GPU数据的和

为什么会这样？因为NCCL内部会自动选择最优的通信算法。可能是Ring AllReduce，也可能是Tree AllReduce，取决于你的拓扑和卡数。

AllGather也类似，不过它不做归约，只是收集：

ncclAllGather(sendbuff, recvbuff, sendcount, datatype, comm, stream);

注意 sendcount 是每块GPU发送的元素数，而 recvbuff 的大小需要是 sendcount * nranks。

1.4 知识体系总览

为了让你更直观地理解本章内容，我画了一张图：

这张图把本章的三个核心模块串起来了。从初始化开始，到点对点通信，再到集合通信，每一步都是下一步的基础。

1.5 避坑总结

最后，我把自己踩过的坑总结一下，你写代码时多留个心眼：

• 初始化失败：检查CUDA_VISIBLE_DEVICES，确保设备号连续且可见
• Send/Recv死锁：确保发送方和接收方的rank号、数据类型、数据量完全一致
• 异步操作未同步：别忘了 cudaStreamSynchronize，否则数据可能还没传完
• 集合操作的数据类型：AllReduce的sendbuff和recvbuff类型必须一致，否则结果会乱

嗯，基础API就讲到这里。这些内容虽然简单，但却是后续所有优化的根基。你先把这些API跑通，后面咱们再聊怎么调优。