第4章：点对点通信：Send/Recv原语详解

多卡通信里最基础的东西，就是点对点通信。

说白了，就是一张卡把数据发给另一张卡。你想想看，无论是All-Reduce还是All-Gather，底层拆开来看，全是Send和Recv的组合。我当年刚接触分布式训练时，以为NCCL有多玄乎，结果发现核心原语就这么几个。

4.1 Send/Recv原语：最朴素的通信方式

Send和Recv，顾名思义，一个发一个收。在NCCL里，它们是最底层的通信原语。我习惯把它们比作两个人传话——一个人喊，一个人听。

来看一个最简单的例子：

// GPU 0 发送数据给 GPU 1
if (rank == 0) {
    ncclSend(sendbuff, size, ncclFloat, 1, comm, stream);
} else if (rank == 1) {
    ncclRecv(recvbuff, size, ncclFloat, 0, comm, stream);
}

这段代码里，rank 0 把 sendbuff 里的数据发给 rank 1，rank 1 把数据收进 recvbuff。就这么简单。

但这里有个坑——发送和接收必须配对。如果只有Send没有Recv，或者Recv的rank号写错了，通信就会卡死。我在项目中遇到过这种情况，排查了半天才发现是rank号写反了。

4.2 阻塞与非阻塞通信

NCCL的Send/Recv有两种模式：阻塞和非阻塞。

阻塞模式：调用ncclSend或ncclRecv后，函数会一直等待，直到数据真正发送或接收完成。这时候CPU啥也干不了，只能干等。

非阻塞模式：函数立即返回，数据在后台传输。CPU可以继续做其他事，比如准备下一批数据。

我个人的习惯是：能用非阻塞就别用阻塞。为什么呢？因为GPU计算和通信可以重叠。你想想看，如果每次通信都阻塞，GPU就得停下来等数据，效率太低了。

来看一个非阻塞的例子：

// 非阻塞发送
ncclSend(sendbuff, size, ncclFloat, 1, comm, stream);
// 这里CPU可以继续做其他事
compute_something_on_cpu();
// 等待通信完成
cudaStreamSynchronize(stream);

嗯，这里要注意：非阻塞不代表不等待。你最终还是要调用cudaStreamSynchronize来确保数据传完了。否则你可能会读到不完整的数据。

4.3 同步与异步模式

同步和异步，这两个词经常被混淆。我简单解释一下：

• 同步模式：发送方和接收方必须同时就绪。发送方发出数据后，必须等接收方确认收到，才能继续。
• 异步模式：发送方发出数据后，不管接收方有没有收到，直接继续执行。接收方在需要数据时再去检查。

NCCL的Send/Recv默认是同步的。也就是说，rank 0调用ncclSend后，必须等rank 1调用ncclRecv，数据才会真正传输。如果rank 1迟迟不调用Recv，rank 0就会卡住。

为什么会这样？因为NCCL需要保证数据一致性。你想想看，如果发送方不管接收方就跑了，接收方可能读到垃圾数据。

但同步模式也有缺点——容易死锁。我曾经在项目中遇到过这种情况：两个GPU互相发送数据，结果都先调用了Send，然后都等着对方Recv，结果谁也等不到谁，死锁了。

解决方案很简单：让一个GPU先发后收，另一个先收后发。

// GPU 0：先发后收
ncclSend(sendbuff0, size, ncclFloat, 1, comm, stream);
ncclRecv(recvbuff0, size, ncclFloat, 1, comm, stream);

// GPU 1：先收后发
ncclRecv(recvbuff1, size, ncclFloat, 0, comm, stream);
ncclSend(sendbuff1, size, ncclFloat, 0, comm, stream);

这样就不会死锁了。

4.4 实战：两个GPU之间传数据

光说不练假把式。我们来写一个完整的例子：两个GPU互相发送一个浮点数数组。

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

int main() {
    int nGPUs = 2;
    int size = 1024 * 1024; // 1M个float
    float *sendbuff, *recvbuff;
    ncclComm_t comm;
    cudaStream_t stream;

    // 初始化NCCL
    ncclCommInitAll(&comm, nGPUs, NULL);

    // 为每个GPU分配显存
    cudaSetDevice(0);
    cudaMalloc(&sendbuff, size * sizeof(float));
    cudaMalloc(&recvbuff, size * sizeof(float));
    cudaStreamCreate(&stream);

    // 填充数据（略）

    // GPU 0 发送给 GPU 1
    if (rank == 0) {
        ncclSend(sendbuff, size, ncclFloat, 1, comm, stream);
        ncclRecv(recvbuff, size, ncclFloat, 1, comm, stream);
    } else if (rank == 1) {
        ncclRecv(recvbuff, size, ncclFloat, 0, comm, stream);
        ncclSend(sendbuff, size, ncclFloat, 0, comm, stream);
    }

    // 等待完成
    cudaStreamSynchronize(stream);

    // 清理
    ncclCommDestroy(comm);
    cudaFree(sendbuff);
    cudaFree(recvbuff);
    cudaStreamDestroy(stream);

    return 0;
}

这段代码里，我用了先发后收和先收后发的模式，避免了死锁。你可以在实际项目中直接套用这个模板。

4.5 知识体系图

下面这张图展示了本章的核心逻辑：

这张图把本章的知识点串起来了。从Send/Recv原语出发，分支出阻塞和非阻塞，再分支出同步和异步。底部是实战中最重要的避坑指南。

好了，这一章就到这里。点对点通信是NCCL的基石，搞懂了它，后面的集合通信就好理解了。