cuda-NCCL笔记（1）-- 初步了解使用NCCL

NCCL是英伟达提供的多GPU之间通信的库。NCCL只支持Linux平台；所以学习NCCL的前提是有一个Linux环境，并且还有多个GPU。环境只能自己想办法去获取了。

默认读者已经自己配好NCCL的环境了；网上有很多用apt下载的教程，没有root权限可以看我的教程：如何在没有权限的服务器上下载NCCL-CSDN博客

配置头文件和库目录就不在我的内容范围了，自己去用cmake或者vscode或者设置环境变量。

和cuda一样，先来看最简单的NCCL代码，然后再一一解释新出现的内容

#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <nccl.h>
#include <iostream>
#include<cstdio>


int main() {
    ncclComm_t comm;
    int nDev = 1;  // 测试单 GPU
    
    cudaSetDevice(0);

    ncclUniqueId id;
    ncclGetUniqueId(&id); // 获取唯一通信 ID

    ncclCommInitRank(&comm, nDev, id, 0); // 初始化 NCCL 通信

    std::cout << "NCCL initialized on device 0" << std::endl;

    ncclCommDestroy(comm);
    return 0;
}

ncclComm_t

• 类型：typedef struct ncclComm* ncclComm_t;

• 作用：这是 NCCL 的 通信器对象，用来表示一组 GPU 的通信环境。

• 你要做任何通信（比如 AllReduce、Broadcast），都需要一个 ncclComm_t 作为“入口”。

• 生命周期：

  1. 创建 → 调 ncclCommInitRank 或 ncclCommInitAll 得到；

  2. 使用 → 调通信函数时传入；

  3. 销毁 → 用 ncclCommDestroy 释放。

ncclUniqueId

作用：唯一标识一个 NCCL 通信域（communicator）。

• NCCL 要在多进程之间建立通信，每个进程都需要知道自己属于哪个通信组，这就是 ncclUniqueId 的作用。

ncclUniqueId id;
ncclGetUniqueId(&id);//生成方式

• 通常在 rank 0 进程里生成一次，然后通过 进程间通信机制（比如 socket、MPI、或者文件共享）传给所有其它进程。

• 使用方式：
  所有进程拿到相同的 id 后，就可以调用 ncclCommInitRank 加入同一个通信组。

ncclCommInitRank

ncclResult_t ncclCommInitRank(
    ncclComm_t* comm,     // 输出：通信器对象
    int nranks,           // 通信组里总共有多少个参与者
    ncclUniqueId commId,  // 通信组的唯一 ID
    int rank              // 当前参与者的编号（0 ~ nranks-1）
);

• 作用：初始化一个通信器，把自己加入到某个通信组里。

• 参数说明：

  • comm：输出参数，返回的通信器对象。

  • nranks：组里总共有多少个参与者。

  • commId：组的唯一 ID（大家必须一致）。

  • rank：当前参与者的编号。

ncclResult_t

• 作用：这是 NCCL 的返回码类型，用来表示函数调用是否成功。

• 常见返回值：

  • ncclSuccess：成功

  • ncclUnhandledCudaError：底层 CUDA 出错

  • ncclSystemError：系统调用失败（比如 socket 出错）

  • ncclInvalidArgument：传入的参数非法

  • ncclInternalError：NCCL 内部错误

  • ncclInProgress：异步操作还没完成（较少见）

官方一般建议用一个宏来统一错误检查：

#define NCCLCHECK(cmd) do {                         \
  ncclResult_t r = cmd;                             \
  if (r != ncclSuccess) {                           \
    printf("NCCL failure %s:%d '%s'\n",             \
           __FILE__,__LINE__,ncclGetErrorString(r));\
    exit(EXIT_FAILURE);                             \
  }                                                 \
} while(0)

ncclCommDestroy

ncclResult_t ncclCommDestroy(ncclComm_t comm);

• 作用：销毁通信器，释放 NCCL 内部资源。

• 用法：程序结束前调用，避免内存泄漏。

所以上面那段代码的流程就是：

1. 生成唯一 ID (ncclGetUniqueId)；

2. 用 ID + rank 数量初始化通信器 (ncclCommInitRank)；

3. 打印一句话，证明初始化成功；

4. 销毁通信器 (ncclCommDestroy)。

多GPU通信

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

void checkCuda(cudaError_t res) {
    if (res != cudaSuccess) {
        std::cerr << "CUDA Error: " << cudaGetErrorString(res) << std::endl;
        exit(EXIT_FAILURE);
    }
}

void checkNCCL(ncclResult_t res) {
    if (res != ncclSuccess) {
        std::cerr << "NCCL Error: " << ncclGetErrorString(res) << std::endl;
        exit(EXIT_FAILURE);
    }
}

int main() {
    const int nDev = 2;        // 两个 GPU
    const int size = 32;       // 每个 GPU 32 个 float
    int devs[nDev] = {0, 1};    // GPU 的 ID
    
    float *sendbuff[nDev], *recvbuff[nDev];
    ncclComm_t comms[nDev];//每次调用 NCCL 操作时，需要告诉它属于哪个通信器。
    cudaStream_t streams[nDev];

    // 为每个GPU分配数据 + 创建 stream
    for (int i = 0; i < nDev; i++) {
        checkCuda(cudaSetDevice(devs[i]));
        checkCuda(cudaMalloc(&sendbuff[i], size * sizeof(float)));
        checkCuda(cudaMalloc(&recvbuff[i], size * sizeof(float)));
        checkCuda(cudaStreamCreate(&streams[i]));

        // 初始化 send buffer
        float *h_data = new float[size];
        for (int j = 0; j < size; j++) h_data[j] = float(i + 1);
        checkCuda(cudaMemcpy(sendbuff[i], h_data, size * sizeof(float), cudaMemcpyHostToDevice));
        delete[] h_data;
    }

    // 单进程多 GPU 初始化 NCCL
    checkNCCL(ncclCommInitAll(comms, nDev, devs));

    //NCCL 集团通信开始
    ncclGroupStart();
    // 执行 AllReduce
    for (int i = 0; i < nDev; i++) {
        checkCuda(cudaSetDevice(devs[i]));
        checkNCCL(ncclAllReduce(sendbuff[i], recvbuff[i], size,
                                 ncclFloat, ncclSum, comms[i], streams[i]));
    }
    ncclGroupEnd();
    // 等待完成
    for (int i = 0; i < nDev; i++) {
        checkCuda(cudaSetDevice(devs[i]));
        checkCuda(cudaStreamSynchronize(streams[i]));
    }

    // 拷贝结果并打印
    for (int i = 0; i < nDev; i++) {
        checkCuda(cudaSetDevice(devs[i]));
        float h_result[size];
        checkCuda(cudaMemcpy(h_result, recvbuff[i], size * sizeof(float), cudaMemcpyDeviceToHost));
        std::cout << "GPU " << i << " result[0] = " << h_result[0] << std::endl;
    }

    // 清理
    for (int i = 0; i < nDev; i++) {
        checkNCCL(ncclCommDestroy(comms[i]));
        checkCuda(cudaFree(sendbuff[i]));
        checkCuda(cudaFree(recvbuff[i]));
        checkCuda(cudaStreamDestroy(streams[i]));
    }

    return 0;
}

ncclCommInitAll

作用：在 单进程多 GPU 情况下初始化通信器（ncclComm_t）。

ncclResult_t ncclCommInitAll(ncclComm_t* comms, int nDevices, const int* devs);

用途：

• 单进程场景下，方便地同时初始化多个 GPU 的通信器。

• 不需要手动生成 ncclUniqueId，NCCL 内部会自己处理。

ncclGroupStart 与 ncclGroupEnd

NCCL 支持 批量操作，可以把多次通信操作 打包提交，减少同步开销。

• 在 ncclGroupStart 到 ncclGroupEnd 之间的所有 NCCL 调用会 批量执行。

• 对多 GPU 或多操作的场景，可以显著提高性能。

用途：

• 减少多 GPU 并发操作中的同步延迟。

• 推荐在 同时执行多个 AllReduce、Reduce、Broadcast 等操作 时使用。

ncclAllReduce

作用：执行 AllReduce 通信，把每个 GPU 的数据按指定操作汇总到所有 GPU。

ncclResult_t ncclAllReduce(
    const void* sendbuff,
    void* recvbuff,
    size_t count,
    ncclDataType_t datatype,
    ncclRedOp_t op,
    ncclComm_t comm,
    cudaStream_t stream
);

AllReduce 的特点：所有 GPU 都能得到相同的结果。

为什么操作之前都要setDevice？

CUDA 的多 GPU 上下文

• CUDA 每个 GPU 有一个独立的 设备上下文。

• 当前线程只能访问它“当前设置”的 GPU 上下文。

• cudaSetDevice(int dev) 就是告诉 CUDA 后续的所有操作都在这个 GPU 上执行（分配内存、启动 kernel、创建 stream 等）。

为什么 cudaMalloc 前需要 cudaSetDevice

• 如果不 cudaSetDevice，默认 GPU 是 0。

• 所以即使你想给 GPU 1 分配内存，不设置 device，内存也会分配到 GPU 0 上。

• 多 GPU 时，每个 GPU 的 buffer 都必须在它自己的上下文中分配。

为什么 ncclAllReduce 前也要 cudaSetDevice

• NCCL 是 基于 CUDA 流（cudaStream_t）执行的。

• 每个 stream 属于某个 GPU 上下文。

• 如果当前线程上下文不是 stream 所在 GPU，会出现：

  • 内存访问错误

  • NCCL 操作 hang（卡住）

• 因此在执行 NCCL 操作前，确保线程当前上下文对应正确 GPU 是安全做法。

什么时候可以省略

• 如果你用 ncclCommInitAll + 每个 GPU 独立线程：

  • 每个线程只操作自己的 GPU

  • 那么每个线程固定上下文，可以在线程初始化时只设置一次 device

• 但在 单线程控制多 GPU 的情况下：

  • 每次访问不同 GPU，都要 cudaSetDevice

ncclGroupStart / ncclGroupEnd 的真正作用

NCCL 的操作默认是异步的

• ncclAllReduce 等 NCCL 函数本质上 不会阻塞 CPU，它们只是把操作 提交到 GPU 流。

• NCCL 会把命令发送给 底层的 NCCL 通信库，然后 GPU 执行。

• 如果你在 单线程同时对多个 GPU 调用 NCCL，每个 GPU 的操作会立即尝试启动通信，但 NCCL 需要保证所有 GPU 的操作被正确“排列”，否则可能出现死锁。

如果示例代码不用ncclGroupStart 会卡住

• 如果你没有用 ncclGroupStart/ncclGroupEnd：

  • CPU 会顺序调用 GPU0 的 ncclAllReduce → GPU1 的 ncclAllReduce。

  • NCCL 可能在 GPU0 等待 GPU1 发来的数据，但 GPU1 的操作还没提交 → 死锁。

• 使用 Group API 后：

  • NCCL 会把所有 GPU 的操作收集起来，再一次性启动 → 避免等待死锁。

  • ncclGroupStart()：告诉 NCCL 接下来的一系列操作属于同一组，不要立即启动通信。

  • ncclGroupEnd()：提交整个组的操作，NCCL 会对所有 GPU 的操作 一次性调度，保证不会出现 GPU 等待其它 GPU 的情况。

简单总结

• 单 GPU → 不用 group API 也没问题。

• 多 GPU 同线程 → 建议使用 ncclGroupStart/End，尤其是 AllReduce、Broadcast 这种涉及所有 GPU 的 collective 操作。

• 多线程每线程一个 GPU → 每线程只处理自己的 GPU，一般不需要 group API。

下一节，会尝试多线程每个线程一个GPU的模式