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; // 测试单 GPUcudaSetDevice(0);ncclUniqueId id;ncclGetUniqueId(&id); // 获取唯一通信 IDncclCommInitRank(&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作为“入口”。 -
生命周期：
-
创建 → 调
ncclCommInitRank或ncclCommInitAll得到； -
使用 → 调通信函数时传入；
-
销毁 → 用
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, // 通信组的唯一 IDint 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 内部资源。
-
用法：程序结束前调用，避免内存泄漏。
所以上面那段代码的流程就是：
-
生成唯一 ID (
ncclGetUniqueId)； -
用 ID + rank 数量初始化通信器 (
ncclCommInitRank)； -
打印一句话，证明初始化成功；
-
销毁通信器 (
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; // 两个 GPUconst int size = 32; // 每个 GPU 32 个 floatint devs[nDev] = {0, 1}; // GPU 的 IDfloat *sendbuff[nDev], *recvbuff[nDev];ncclComm_t comms[nDev];//每次调用 NCCL 操作时，需要告诉它属于哪个通信器。cudaStream_t streams[nDev];// 为每个GPU分配数据 + 创建 streamfor (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 bufferfloat *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 初始化 NCCLcheckNCCL(ncclCommInitAll(comms, nDev, devs));//NCCL 集团通信开始ncclGroupStart();// 执行 AllReducefor (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);
| 参数 | 类型 | 含义 |
|---|---|---|
comms | ncclComm_t* | 输出数组，每个 GPU 对应一个通信器。初始化后，每个 GPU 的 NCCL 操作都要用它。 |
nDevices | int | GPU 数量，即要参与通信的设备个数。 |
devs | const int* | GPU ID 数组，长度等于 nDevices，指定哪些 GPU 参与通信。 |
用途：
-
单进程场景下，方便地同时初始化多个 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
);
| 参数 | 类型 | 含义 |
|---|---|---|
sendbuff | const void* | 输入缓冲区，存放当前 GPU 的数据。 |
recvbuff | void* | 输出缓冲区，存放 AllReduce 后的结果（每个 GPU 都有同样结果）。 |
count | size_t | 元素数量（例如 float 的个数）。 |
datatype | ncclDataType_t | 数据类型，例如 ncclFloat、ncclDouble、ncclInt 等。 |
op | ncclRedOp_t | 聚合操作类型，例如 ncclSum（求和）、ncclProd（乘积）、ncclMax、ncclMin。 |
comm | ncclComm_t | NCCL 通信器，指定 GPU 所在的通信上下文。 |
stream | cudaStream_t | CUDA 流，通信在这个流上异步执行。 |
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的模式