nccl & mpi 跨节点通信

假如有两个节点， 每个节点两个GPU

查看系统与机器架构：

uname -m
lsb_release -a

在每台机器上操作

安装nccl

wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2204/x86_64/cuda-keyring_1.1-1_all.deb
dpkg -i cuda-keyring_1.1-1_all.deb
apt-get update
apt install libnccl2=2.26.5-1+cuda12.9 libnccl-dev=2.26.5-1+cuda12.9

安装MPI

apt-get update
apt-get install openmpi-bin openmpi-common libopenmpi-dev

在主节点上操作

新建mpi_hosts文件，输入：

node1 slots=2  # 节点1，2个GPU, 如果是infiniband网络， 通过ibnetdiscover 查看网络拓扑，。
node2 slots=2  # 节点2，2个GPU

如果是infiniband网络：
(以下两种方式（不建议的配置内容的部分），虽然可以通讯，但是进程无法启动，因为mpirun是通过ssh启动orted（MPI 的远程守护进程）远程节点的进程，所以，即使在infiniband网络中，使用IPoIB（IP over InfiniBand）地址，还是通过hosts配置ip的方式最佳，默认仍会通过 ​​SSH​​（ TCP/IP 协议栈） 在远程节点上启动进程，但实际数据传输会直接走 InfiniBand 网络​​（绕过 TCP/IP 栈）)

# 查看 IPoIB 地址（在所有节点执行）
ip addr show ib0
# 输出示例：
# ib0: <BROADCAST,MULTICAST,UP> mtu 4092 qdisc mq state UP 
#     inet 192.168.1.100/24 scope global ib0

不建议的配置内容

# mpi_hosts 内容：
0x0002c90300002778 slots=2  # node1 
0x0002c90300002f78 slots=2  # node2

或许

# mpi_hosts 内容：
node1.ib slots=2  # node1 
node2.ib slots=2  # node2

InfiniBand 网络通常由 ​​Subnet Manager​​ 自动分配主机名（如 node1.ib、node2.ib），这些名称已内置解析，无需手动修改 /etc/hosts

slots=2 表示每个节点可以运行 2 个 MPI 进程（每个 GPU 一个进程）。

配置/etc/hosts,

192.168.1.100 node1
192.168.1.101 node2

设置免密码登录ssh， OpenMPI默认使用rsh或ssh在远程节点上启动进程。如果未配置免密登录，每次连接都会提示输入密码，导致mpirun卡住。

ssh-keygen -t rsa   # 一直按回车
ssh-copy-id node1   # 输入密码
ssh-copy-id node2

所有节点同步二进制程序

scp ./nccl_allreduce node2:/data/coding

在主节点上启动（使用 NCCL_DEBUG=INFO 可以看到更多信息）

NCCL_DEBUG=WARN mpirun --hostfile mpi_hosts --allow-run-as-root  -np 4 /data/coding/nccl_allreduce 

程序执行输出：

Rank 3 (Node 1) -> GPU 1
Rank 2 (Node 1) -> GPU 0
Rank 0 (Node 0) -> GPU 0
Rank 1 (Node 0) -> GPU 1
NCCL version 2.26.5+cuda12.9
Rank 0 result: 10.0 10.0 ...
Rank 1 result: 10.0 10.0 ...
Rank 2 result: 10.0 10.0 ...
Rank 3 result: 10.0 10.0 ...

程序代码：

/**
 * 
 * 编译： mpic++ -o nccl_allreduce nccl_allreduce.cpp -lnccl -lcudart -lmpi -I/usr/local/cuda-12.3/include -L/usr/local/cuda-12.3/lib64
 */
#include <iostream>
#include <vector>
#include <nccl.h>
#include <cuda_runtime.h>
#include <mpi/mpi.h>

#define NCCL_CHECK(cmd) { ncclResult_t r = cmd; if (r != ncclSuccess) { printf("NCCL error %d at line %d\n", r, __LINE__); exit(EXIT_FAILURE); } }
#define CUDA_CHECK(cmd) { cudaError_t r = cmd; if (r != cudaSuccess) { printf("CUDA error %d at line %d\n", r, __LINE__); exit(EXIT_FAILURE); } }

int main(int argc, char* argv[]) {
    // 初始化MPI
    MPI_Init(&argc, &argv);
    int rank, size;
    MPI_Comm_rank(MPI_COMM_WORLD, &rank);
    MPI_Comm_size(MPI_COMM_WORLD, &size);

 
    // 显式设置设备（必须在所有CUDA/NCCL调用前执行）
    int local_rank = rank % 2;
    CUDA_CHECK(cudaSetDevice(local_rank));
    printf("Rank %d (Node %d) -> GPU %d\n", rank, rank/2, local_rank);

    // 初始化NCCL
    ncclUniqueId id;
    if (rank == 0) ncclGetUniqueId(&id);
    MPI_Bcast(&id, sizeof(id), MPI_BYTE, 0, MPI_COMM_WORLD);

    ncclComm_t comm;
    NCCL_CHECK(ncclCommInitRank(&comm, size, id, rank));

    // 创建CUDA流
    cudaStream_t stream;
    CUDA_CHECK(cudaStreamCreate(&stream));

    // 分配和初始化数据
    const int count = 10;
    float *sendbuff, *recvbuff;
    CUDA_CHECK(cudaMalloc(&sendbuff, count * sizeof(float)));
    CUDA_CHECK(cudaMalloc(&recvbuff, count * sizeof(float)));

    std::vector<float> hostData(count, rank + 1.0f);
    CUDA_CHECK(cudaMemcpy(sendbuff, hostData.data(), count * sizeof(float), cudaMemcpyHostToDevice));

    // 执行AllReduce
    NCCL_CHECK(ncclAllReduce(sendbuff, recvbuff, count, ncclFloat, ncclSum, comm, stream));
    CUDA_CHECK(cudaStreamSynchronize(stream));

    // 验证结果
    std::vector<float> result(count);
    CUDA_CHECK(cudaMemcpy(result.data(), recvbuff, count * sizeof(float), cudaMemcpyDeviceToHost));
    printf("Rank %d result: %.1f %.1f ...\n", rank, result[0], result[1]);

    // 清理资源
    CUDA_CHECK(cudaFree(sendbuff));
    CUDA_CHECK(cudaFree(recvbuff));
    NCCL_CHECK(ncclCommDestroy(comm));
    CUDA_CHECK(cudaStreamDestroy(stream));
    MPI_Finalize();
    return 0;
}

可以看见， 主节点上启动了四个进程， 一个 mpirun 进程， 两个nccl_allreduce进程，orted进程， 从节点 二个nccl_allreduce进程，一个orted进程。（待验证）

问题： 已经有了nccl，为什么还要mpi？

1. mpi在cpu之间的集合通讯， nccl在gpu之间的集合通讯， 因为最初的mpi实现就是在cpu之间的通讯。代码中的MPI_Bcast即实现主机内存中的数据通讯，而ncclAllReduce实现gpu内存中的数据通讯。
2. mpi负责在不同节点（或许同节点）的启动、维护、管理、编号多进程，这不属于nccl的内容， 所以nccl并未实现