4. 点对点通信：ncclSend与ncclRecv的API详解，同步与异步模式
点对点通信，说白了就是两个GPU之间直接传数据。你想想看，在分布式训练里，很多时候并不需要所有GPU都互相通信——比如流水线并行，前一层的GPU算完了，把梯度传给下一层就行。这时候，ncclSend和ncclRecv就是你的主力工具。
我个人习惯把这两个API叫做「快递员」和「收件人」。一个负责把数据打包发出去，一个负责接收并拆包。嗯，听起来简单，但实际用起来坑不少。我刚开始用NCCL时，就因为在同步异步模式上栽过跟头，今天咱们好好捋一捋。
4.1 ncclSend与ncclRecv的API签名
先看函数原型，这是最基础的东西：
ncclResult_t ncclSend(
const void* sendbuff, // 发送缓冲区
size_t count, // 元素个数
ncclDataType_t datatype, // 数据类型
int peer, // 目标rank
ncclComm_t comm, // 通信器
cudaStream_t stream // CUDA流
);
ncclResult_t ncclRecv(
void* recvbuff, // 接收缓冲区
size_t count, // 元素个数
ncclDataType_t datatype, // 数据类型
int peer, // 源rank
ncclComm_t comm, // 通信器
cudaStream_t stream // CUDA流
);参数看着多，其实核心就几个：数据在哪、发给谁、用什么流。我遇到过不少新手，把peer参数填成自己的rank，结果死等——自己给自己发数据，NCCL可不干这事儿。
4.2 同步模式 vs 异步模式
这里我要重点讲一下。NCCL的Send和Recv默认是阻塞式的，但阻塞的是CUDA流，不是CPU线程。很多人搞混了这一点。
| 模式 | 行为 | 适用场景 |
|---|---|---|
| 同步（默认） | ncclSend/ncclRecv提交到流后，CPU立即返回。但GPU端会等待操作完成才执行后续kernel。 | 小数据量、强依赖场景 |
| 异步（配合流） | 使用不同的CUDA流，让Send/Recv与其他计算重叠。 | 大数据量、需要计算通信重叠 |
说白了，同步模式就是「发完再算」，异步模式就是「边发边算」。我当年在训练一个百亿参数模型时，就是因为没用好异步模式，GPU利用率一直上不去。后来改成异步，吞吐量直接提升了30%。
4.3 代码示例：同步点对点通信
先来个最简单的例子，两个GPU之间传一个张量：
// rank 0 发送，rank 1 接收
if (rank == 0) {
float* sendbuf;
cudaMalloc(&sendbuf, size * sizeof(float));
// 填充数据...
NCCL_CHECK(ncclSend(sendbuf, size, ncclFloat, 1, comm, stream));
} else if (rank == 1) {
float* recvbuf;
cudaMalloc(&recvbuf, size * sizeof(float));
NCCL_CHECK(ncclRecv(recvbuf, size, ncclFloat, 0, comm, stream));
}
cudaStreamSynchronize(stream);注意看，这里我用了cudaStreamSynchronize来确保GPU操作完成。如果不加这行，CPU可能会在数据还没传完时就往下执行——嗯，后果你懂的。
4.4 代码示例：异步模式实现计算通信重叠
这才是实战中的重头戏。假设你要做矩阵乘法，同时还要传数据：
// 创建两个流：一个计算流，一个通信流
cudaStream_t computeStream, commStream;
cudaStreamCreate(&computeStream);
cudaStreamCreate(&commStream);
// 在通信流上执行Send/Recv
if (rank == 0) {
NCCL_CHECK(ncclSend(sendbuf, size, ncclFloat, 1, comm, commStream));
// 同时在计算流上做矩阵乘法
matmul_kernel<<<grid, block, 0, computeStream>>>(A, B, C);
} else if (rank == 1) {
NCCL_CHECK(ncclRecv(recvbuf, size, ncclFloat, 0, comm, commStream));
matmul_kernel<<<grid, block, 0, computeStream>>>(A, B, C);
}
// 等待所有流完成
cudaStreamSynchronize(computeStream);
cudaStreamSynchronize(commStream);这里有个坑：两个流是并行的，但如果你在计算流里用到了sendbuf或recvbuf的数据，那就出事了。我建议的做法是：通信流只负责传输，计算流只负责计算，两者不要共享缓冲区。
4.5 同步与异步的选择策略
什么时候用同步，什么时候用异步？我总结了一个简单的判断标准：
- 数据量小于1MB：用同步模式。异步的开销（流切换、事件管理）可能比传输本身还大。
- 数据量在1MB到100MB之间：看你的计算密度。如果计算很轻，同步就行；如果计算很重，异步更划算。
- 数据量大于100MB：强烈建议用异步。这时候通信时间占比很大，不重叠就是浪费。
说白了，异步不是银弹。我见过有人为了炫技，小数据量也硬上异步，结果性能反而下降了。嗯，合适才是最好的。
4.6 知识体系图：点对点通信的核心逻辑
这张图把点对点通信的脉络理清楚了。从左到右是数据流，从上到下是模式选择。你写代码的时候，脑子里先过一遍这个图，基本不会跑偏。
4.7 常见问题与排查思路
最后分享几个我实际踩过的坑：
- 死锁问题：如果两个GPU互相等待对方的Send/Recv，就会死锁。比如rank0先Send再Recv，rank1先Recv再Send——这就死锁了。解决办法是保证调用顺序一致。
- 缓冲区大小不匹配：发送端count=1000，接收端count=500。NCCL不会报错，但你会读到垃圾数据。我建议在调试阶段打印一下count值。
- 流同步遗漏：异步模式下，如果忘了同步流，CPU可能读到未完成的数据。我的习惯是在每个通信操作后加一个cudaStreamSynchronize，等调通了再优化掉。
好了，点对点通信就讲到这里。记住：同步保安全，异步提性能。具体用哪个，看你的数据量和计算密度。多写几遍代码，自然就熟练了。