NCCL 集合通信入门实战：通信器创建与 Group 语义的六种标准模式全解析

高性能计算通信深度学习分布式训练【免费下载链接】ncclOptimized primitives for collective multi-GPU communication项目地址https://gitcode.com/gh_mirrors/nc/nccl点击查看免费下载本文基于 NCCL 官方用户指南的 Examples 章节examples.rst系统讲解在多进程、多线程、多 GPU 等不同部署环境下如何创建与销毁 NCCL 通信器、如何以 Group 语义发起集合通信并结合 docs/examples 示例目录与 init.cc 源码帮助读者掌握从单进程多卡到 MPI 多节点通信的完整可运行代码与底层实现原理。NCCLNVIDIA Collective Communications Library是面向多 GPU 集合通信优化的通信库。在编写任何 NCCL 应用之前需要先理解一个核心问题你的并行运行时环境parallel runtime environment是如何组织 GPU 的。官方示例覆盖了三类典型组合每个线程/进程使用多个 GPUmultiple GPUs per thread/process使用多个线程multiple threads使用多个进程multiple processes。多进程示例以 MPI 作为并行运行时但任何具备进程间通信能力的多进程系统都可以类似工作。一个贯穿所有示例的重要原则始终检查 NCCL 函数的返回值。为了篇幅清晰下文代码大多省略了错误检查实际生产代码中请通过NCCLCHECK/CUDACHECK等宏检查每一个 NCCL/CUDA 调用。一、通信器的创建与销毁NCCL 的通信器ncclComm_t是所有集合通信操作的载体其创建方式取决于进程的拓扑结构。官方示例给出四种典型场景。1. 单进程、单线程、多设备ncclCommInitAll最简单的场景一个进程、一个线程管理所有 GPU。此时可以直接使用ncclCommInitAll它为devlist中的每一张卡创建一个通信器。例如为 4 张卡创建通信器会产生 4 个通信器对象ncclComm_t comms[4]; int devs[4] { 0, 1, 2, 3 }; ncclCommInitAll(comms, 4, devs);之后即可用单线程 Group 调用的方式或多个线程各自持有一个 comm 对象发起集合操作。程序结束前销毁所有通信器for (int i0; i4; i) ncclCommDestroy(comms[i]);下面是单进程管理多设备的完整可运行示例#include stdlib.h #include stdio.h #include cuda_runtime.h #include nccl.h #define CUDACHECK(cmd) do { \ cudaError_t err cmd; \ if (err ! cudaSuccess) { \ printf(Failed: Cuda error %s:%d %s\n, \ __FILE__,__LINE__,cudaGetErrorString(err)); \ exit(EXIT_FAILURE); \ } \ } while(0) #define NCCLCHECK(cmd) do { \ ncclResult_t res cmd; \ if (res ! ncclSuccess) { \ printf(Failed, NCCL error %s:%d %s\n, \ __FILE__,__LINE__,ncclGetErrorString(res)); \ exit(EXIT_FAILURE); \ } \ } while(0) int main(int argc, char* argv[]) { ncclComm_t comms[4]; //managing 4 devices int nDev 4; int size 32*1024*1024; int devs[4] { 0, 1, 2, 3 }; //allocating and initializing device buffers float** sendbuff (float**)malloc(nDev * sizeof(float*)); float** recvbuff (float**)malloc(nDev * sizeof(float*)); cudaStream_t* s (cudaStream_t*)malloc(sizeof(cudaStream_t)*nDev); for (int i 0; i nDev; i) { CUDACHECK(cudaSetDevice(i)); CUDACHECK(cudaMalloc((void**)sendbuff i, size * sizeof(float))); CUDACHECK(cudaMalloc((void**)recvbuff i, size * sizeof(float))); CUDACHECK(cudaMemset(sendbuff[i], 1, size * sizeof(float))); CUDACHECK(cudaMemset(recvbuff[i], 0, size * sizeof(float))); CUDACHECK(cudaStreamCreate(si)); } //initializing NCCL NCCLCHECK(ncclCommInitAll(comms, nDev, devs)); //calling NCCL communication API. Group API is required when using //multiple devices per thread NCCLCHECK(ncclGroupStart()); for (int i 0; i nDev; i) NCCLCHECK(ncclAllReduce((const void*)sendbuff[i], (void*)recvbuff[i], size, ncclFloat, ncclSum, comms[i], s[i])); NCCLCHECK(ncclGroupEnd()); //synchronizing on CUDA streams to wait for completion of NCCL operation for (int i 0; i nDev; i) { CUDACHECK(cudaSetDevice(i)); CUDACHECK(cudaStreamSynchronize(s[i])); } //free device buffers for (int i 0; i nDev; i) { CUDACHECK(cudaSetDevice(i)); CUDACHECK(cudaFree(sendbuff[i])); CUDACHECK(cudaFree(recvbuff[i])); } //finalizing NCCL for(int i 0; i nDev; i) ncclCommDestroy(comms[i]); printf(Success \n); return 0; }源码视角ncclCommInitAll在 init.cc 中的实现揭示了它简单背后的完整流程——内部先调用ncclGetUniqueId生成交互标识随后进入ncclGroupStartInternal/ncclGroupEndInternal分组语义在组内对devlist中的每张卡依次执行ncclCommInitRankDevrank 即为设备在列表中的下标最后恢复原有 CUDA 设备上下文。同时它对devlist做了合法性校验设备号越界会返回ncclInvalidArgument且默认拒绝同一设备出现两次除非MultiRankGpu相关参数开启。这与头文件 nccl.h.in 中的注释一致devlist为 NULL 时使用前ndev个 CUDA 设备devlist的顺序即定义了通信器内处理器的用户排序。仓库中对应的可直接编译运行的示例是 01_multiple_devices_single_process它在ncclCommInitAll之外还演示了通过ncclCommUserRank/ncclCommCount/ncclCommCuDevice查询并校验通信器属性的方法以及先ncclCommFinalize再ncclCommDestroy、先销毁通信器再销毁 CUDA stream的资源清理顺序。2. 每进程/每线程一个设备ncclCommInitRank当一个进程或宿主线程最多负责一张 GPU 时多机多卡训练最常见的模式应使用ncclCommInitRank这一集合式调用来创建通信器每个线程/进程各自拿到一个通信器对象。由于需要跨进程共享一个 bootstrap 标识标准流程分三步以 MPI 为例第一步获取 MPI 进程信息int myRank, nRanks; MPI_Comm_rank(MPI_COMM_WORLD, myRank); MPI_Comm_size(MPI_COMM_WORLD, nRanks);第二步由某一个 rank 生成唯一 ID再广播给所有其他 rankncclUniqueId id; if (myRank 0) ncclGetUniqueId(id); MPI_Bcast(id, sizeof(id), MPI_BYTE, 0, MPI_COMM_WORLD);第三步各 rank 用同一 ID 创建通信器ncclComm_t comm; ncclCommInitRank(comm, nRanks, id, myRank);此后即可用该通信器发起集合操作如ncclAllReduce( ... , comm);最终调用ncclCommDestroy(comm);销毁。头文件 nccl.h.in 对该 API 的关键约束值得注意rank必须在0到nranks-1之间且在通信器内唯一调用前必须已经cudaSetDevice选定该 rank 绑定的 CUDA 设备ncclCommInitRank会隐式与其他 rank 同步因此必须由不同的线程/进程并发调用或者用ncclGroupStart/ncclGroupEnd包起来这正是第 3 种场景的写法。下面是多 MPI 进程、每进程一个设备的完整可运行示例。注意其中用主机名 /proc/sys/kernel/random/boot_id的哈希来区分同一物理节点从而计算用于选卡的localRank#include stdio.h #include cuda_runtime.h #include nccl.h #include mpi.h #include unistd.h #include stdint.h #include stdlib.h #define MPICHECK(cmd) do { \ int e cmd; \ if( e ! MPI_SUCCESS ) { \ printf(Failed: MPI error %s:%d %d\n, \ __FILE__,__LINE__, e); \ exit(EXIT_FAILURE); \ } \ } while(0) #define CUDACHECK(cmd) do { \ cudaError_t e cmd; \ if( e ! cudaSuccess ) { \ printf(Failed: Cuda error %s:%d %s\n, \ __FILE__,__LINE__,cudaGetErrorString(e)); \ exit(EXIT_FAILURE); \ } \ } while(0) #define NCCLCHECK(cmd) do { \ ncclResult_t r cmd; \ if (r! ncclSuccess) { \ printf(Failed, NCCL error %s:%d %s\n, \ __FILE__,__LINE__,ncclGetErrorString(r)); \ exit(EXIT_FAILURE); \ } \ } while(0) static uint64_t getHash(const char* string, size_t n) { // Based on DJB2a, result result * 33 ^ char uint64_t result 5381; for (size_t c 0; c n; c){ result ((result 5) result) ^ string[c]; } return result; } /* Generate a hash of the unique identifying string for this host * that will be unique for both bare-metal and container instances * Equivalent of a hash of; * * $(hostname)$(cat /proc/sys/kernel/random/boot_id) */ #define HOSTID_FILE /proc/sys/kernel/random/boot_id static uint64_t getHostHash(const char* hostname) { char hostHash[1024]; // Fall back is the hostname if something fails (void) strncpy(hostHash, hostname, sizeof(hostHash)); int offset strlen(hostHash); FILE *file fopen(HOSTID_FILE, r); if (file ! NULL) { char *p; if (fscanf(file, %ms, p) 1) { strncpy(hostHashoffset, p, sizeof(hostHash)-offset-1); free(p); } } fclose(file); // Make sure the string is terminated hostHash[sizeof(hostHash)-1]\0; return getHash(hostHash, strlen(hostHash)); } static void getHostName(char* hostname, int maxlen) { gethostname(hostname, maxlen); for (int i0; i maxlen; i) { if (hostname[i] .) { hostname[i] \0; return; } } } int main(int argc, char* argv[]) { int size 32*1024*1024; int myRank, nRanks, localRank 0; //initializing MPI MPICHECK(MPI_Init(argc, argv)); MPICHECK(MPI_Comm_rank(MPI_COMM_WORLD, myRank)); MPICHECK(MPI_Comm_size(MPI_COMM_WORLD, nRanks)); //calculating localRank based on hostname which is used in selecting a GPU uint64_t hostHashs[nRanks]; char hostname[1024]; getHostName(hostname, 1024); hostHashs[myRank] getHostHash(hostname); MPICHECK(MPI_Allgather(MPI_IN_PLACE, 0, MPI_DATATYPE_NULL, hostHashs, sizeof(uint64_t), MPI_BYTE, MPI_COMM_WORLD)); for (int p0; pnRanks; p) { if (p myRank) break; if (hostHashs[p] hostHashs[myRank]) localRank; } ncclUniqueId id; ncclComm_t comm; float *sendbuff, *recvbuff; cudaStream_t s; //get NCCL unique ID at rank 0 and broadcast it to all others if (myRank 0) ncclGetUniqueId(id); MPICHECK(MPI_Bcast((void *)id, sizeof(id), MPI_BYTE, 0, MPI_COMM_WORLD)); //picking a GPU based on localRank, allocate device buffers CUDACHECK(cudaSetDevice(localRank)); CUDACHECK(cudaMalloc(sendbuff, size * sizeof(float))); CUDACHECK(cudaMalloc(recvbuff, size * sizeof(float))); CUDACHECK(cudaStreamCreate(s)); //initializing NCCL NCCLCHECK(ncclCommInitRank(comm, nRanks, id, myRank)); //communicating using NCCL NCCLCHECK(ncclAllReduce((const void*)sendbuff, (void*)recvbuff, size, ncclFloat, ncclSum, comm, s)); //completing NCCL operation by synchronizing on the CUDA stream CUDACHECK(cudaStreamSynchronize(s)); //free device buffers CUDACHECK(cudaFree(sendbuff)); CUDACHECK(cudaFree(recvbuff)); //finalizing NCCL ncclCommDestroy(comm); //finalizing MPI MPICHECK(MPI_Finalize()); printf([MPI Rank %d] Success \n, myRank); return 0; }源码视角init.cc 中ncclCommInitRank的实现非常薄——它先初始化环境与 CUDA 驱动句柄通过cudaGetDevice(cudaDev)获取当前线程已绑定的设备这就是为什么必须先cudaSetDevice再以NCCL_CONFIG_INITIALIZER默认配置转发给内部函数ncclCommInitRankDev。也就是说rank 绑定的 GPU 由调用时的当前 CUDA 设备隐式决定ncclCommInitRank参数里并不传设备号。仓库中的 03_one_device_per_process_mpi 提供了这个模式的生产级实现其getLocalRank()用MPI_Comm_split_type(comm, MPI_COMM_TYPE_SHARED, ...)按共享内存即同一节点拆分通信组来获得节点内 rank比上文示例中主机名 boot_id 哈希的自实现方案更稳健。该目录还提供 Python 变体基于nccl4pympi4py运行命令为mpirun -np num_processes ./03_one_device_per_process_mpi/c/one_device_per_process_mpi mpirun -np num_processes python ./03_one_device_per_process_mpi/python/one_device_per_process_mpi.py3. 每线程多设备Group 语义 ncclCommInitRank可以把多进程和每进程多设备叠加起来每个 MPI rank 管理ngpus张卡所有 rank 的 GPU 组成一个更大的ngpus*nRanks规模通信器。此时必须使用 Group 语义包裹ncclCommInitRank因为同一个线程要为多张卡并发创建通信器。流程仍是先取 MPI 信息、再由一个 rank 生成并广播 IDint myRank, nRanks; MPI_Comm_rank(MPI_COMM_WORLD, myRank); MPI_Comm_size(MPI_COMM_WORLD, nRanks); ncclUniqueId id; if (myRank 0) ncclGetUniqueId(id); MPI_Bcast(id, sizeof(id), MPI_BYTE, 0, 0, MPI_COMM_WORLD);然后用 Group 包裹ncclCommInitRank本进程创建ngpus个通信器对象它们共同属于ngpus*nRanks的组全局 rank 为myRank*ngpusincclComm_t comms[ngpus]; ncclGroupStart(); for (int i0; ingpus; i) { cudaSetDevice(devs[i]); ncclCommInitRank(commsi, ngpus*nRanks, id, myRank*ngpusi); } ncclGroupEnd();之后调用集合操作单线程 Group 调用或多线程各持一个 comm程序结束前销毁所有通信器for (int i0; ingpus; i) ncclCommDestroy(comms[i]);完整可运行示例多 MPI 进程、每进程 2 张 GPU#include stdio.h #include cuda_runtime.h #include nccl.h #include mpi.h #include unistd.h #include stdint.h #define MPICHECK(cmd) do { \ int e cmd; \ if( e ! MPI_SUCCESS ) { \ printf(Failed: MPI error %s:%d %d\n, \ __FILE__,__LINE__, e); \ exit(EXIT_FAILURE); \ } \ } while(0) #define CUDACHECK(cmd) do { \ cudaError_t e cmd; \ if( e ! cudaSuccess ) { \ printf(Failed: Cuda error %s:%d %s\n, \ __FILE__,__LINE__,cudaGetErrorString(e)); \ exit(EXIT_FAILURE); \ } \ } while(0) #define NCCLCHECK(cmd) do { \ ncclResult_t r cmd; \ if (r! ncclSuccess) { \ printf(Failed, NCCL error %s:%d %s\n, \ __FILE__,__LINE__,ncclGetErrorString(r)); \ exit(EXIT_FAILURE); \ } \ } while(0) static uint64_t getHash(const char* string) { // Based on DJB2a, result result * 33 ^ char uint64_t result 5381; for (int c 0; string[c] ! \0; c){ result ((result 5) result) ^ string[c]; } return result; } /* Generate a hash of the unique identifying string for this host * that will be unique for both bare-metal and container instances * Equivalent of a hash of; * * $(hostname)$(cat /proc/sys/kernel/random/boot_id) */ #define HOSTID_FILE /proc/sys/kernel/random/boot_id static uint64_t getHostHash(const char* hostname) { char hostHash[1024]; // Fall back is the hostname if something fails (void) strncpy(hostHash, hostname, sizeof(hostHash)); int offset strlen(hostHash); FILE *file fopen(HOSTID_FILE, r); if (file ! NULL) { char *p; if (fscanf(file, %ms, p) 1) { strncpy(hostHashoffset, p, sizeof(hostHash)-offset-1); free(p); } } fclose(file); // Make sure the string is terminated hostHash[sizeof(hostHash)-1]\0; return getHash(hostHash, strlen(hostHash)); } static void getHostName(char* hostname, int maxlen) { gethostname(hostname, maxlen); for (int i0; i maxlen; i) { if (hostname[i] .) { hostname[i] \0; return; } } } int main(int argc, char* argv[]) { int size 32*1024*1024; int myRank, nRanks, localRank 0; //initializing MPI MPICHECK(MPI_Init(argc, argv)); MPICHECK(MPI_Comm_rank(MPI_COMM_WORLD, myRank)); MPICHECK(MPI_Comm_size(MPI_COMM_WORLD, nRanks)); //calculating localRank which is used in selecting a GPU uint64_t hostHashs[nRanks]; char hostname[1024]; getHostName(hostname, 1024); hostHashs[myRank] getHostHash(hostname); MPICHECK(MPI_Allgather(MPI_IN_PLACE, 0, MPI_DATATYPE_NULL, hostHashs, sizeof(uint64_t), MPI_BYTE, MPI_COMM_WORLD)); for (int p0; pnRanks; p) { if (p myRank) break; if (hostHashs[p] hostHashs[myRank]) localRank; } //each process is using two GPUs int nDev 2; float** sendbuff (float**)malloc(nDev * sizeof(float*)); float** recvbuff (float**)malloc(nDev * sizeof(float*)); cudaStream_t* s (cudaStream_t*)malloc(sizeof(cudaStream_t)*nDev); //picking GPUs based on localRank for (int i 0; i nDev; i) { CUDACHECK(cudaSetDevice(localRank*nDev i)); CUDACHECK(cudaMalloc(sendbuff i, size * sizeof(float))); CUDACHECK(cudaMalloc(recvbuff i, size * sizeof(float))); CUDACHECK(cudaMemset(sendbuff[i], 1, size * sizeof(float))); CUDACHECK(cudaMemset(recvbuff[i], 0, size * sizeof(float))); CUDACHECK(cudaStreamCreate(si)); } ncclUniqueId id; ncclComm_t comms[nDev]; //generating NCCL unique ID at one process and broadcasting it to all if (myRank 0) ncclGetUniqueId(id); MPICHECK(MPI_Bcast((void *)id, sizeof(id), MPI_BYTE, 0, MPI_COMM_WORLD)); //initializing NCCL, group API is required around ncclCommInitRank as it is //called across multiple GPUs in each thread/process NCCLCHECK(ncclGroupStart()); for (int i0; inDev; i) { CUDACHECK(cudaSetDevice(localRank*nDev i)); NCCLCHECK(ncclCommInitRank(commsi, nRanks*nDev, id, myRank*nDev i)); } NCCLCHECK(ncclGroupEnd()); //calling NCCL communication API. Group API is required when using //multiple devices per thread/process NCCLCHECK(ncclGroupStart()); for (int i0; inDev; i) NCCLCHECK(ncclAllReduce((const void*)sendbuff[i], (void*)recvbuff[i], size, ncclFloat, ncclSum, comms[i], s[i])); NCCLCHECK(ncclGroupEnd()); //synchronizing on CUDA stream to complete NCCL communication for (int i0; inDev; i) CUDACHECK(cudaStreamSynchronize(s[i])); //freeing device memory for (int i0; inDev; i) { CUDACHECK(cudaFree(sendbuff[i])); CUDACHECK(cudaFree(recvbuff[i])); } //finalizing NCCL for (int i0; inDev; i) { ncclCommDestroy(comms[i]); } //finalizing MPI MPICHECK(MPI_Finalize()); printf([MPI Rank %d] Success \n, myRank); return 0; }源码视角从 init.cc 可以看到ncclCommInitAll内部正是Group 包裹多个ncclCommInitRankDev——这与示例 3 中手动用ncclGroupStart/ncclGroupEnd包裹ncclCommInitRank的写法在语义上是同一套机制。Group 的本质是把若干必须并发推进才能完成的调用融合fuse成一个整体在 nccl.h.in 中对ncclGroupEnd的注释明确指出它等待组内所有调用完成但对集合通信而言ncclGroupEnd只保证操作已入队enqueued真正完成仍需等待 CUDA stream。4. 每设备多个通信器阻塞与非阻塞通信器NCCL 允许在同一张设备上创建多个通信器。上面前三个示例创建的都是阻塞式通信器blocking communicatorsncclCommInitRank会阻塞直到通信器在所有 rank 上全局就绪。若希望初始化异步完成例如避免不同通信器的初始化相互等待、实现容错重启时的优雅中止可以设置ncclConfig_t中的blocking 0创建非阻塞通信器non-blocking communicators// blocking communicators CUDACHECK(cudaSetDevice(localRank)); for (int i 0; i commNum; i) { if (myRank 0) ncclGetUniqueId(id); MPICHECK(MPI_Bcast((void *)id, sizeof(id), MPI_BYTE, 0, MPI_COMM_WORLD)); NCCLCHECK(ncclCommInitRank(blockingComms[i], nRanks, id, myRank)); } // non-blocking communicators CUDACHECK(cudaSetDevice(localRank)); ncclConfig_t config NCCL_CONFIG_INITIALIZER; config.blocking 0; for (int i 0; i commNum; i) { if (myRank 0) ncclGetUniqueId(id); MPICHECK(MPI_Bcast((void *)id, sizeof(id), MPI_BYTE, 0, MPI_COMM_WORLD)); NCCLCHECK(ncclCommInitRankConfig(nonblockingComms[i], nRanks, id, myRank, config)); do { NCCLCHECK(ncclCommGetAsyncError(nonblockingComms[i], state)); } while(state ncclInProgress checkTimeout() ! true); }其中checkTimeout()是用户自定义的超时判断函数。几点关键约束均可在源码中验证config 必须先用NCCL_CONFIG_INITIALIZER赋值未初始化的 config 会直接报错——nccl.h.in 中NCCL_CONFIG_INITIALIZER展开了size、magic、version三个字段并给出blocking、cgaClusterSize、minCTAs、maxCTAs、netName、splitShare、trafficClass、commName、nvlsCTAs等全部用户可定制属性非阻塞模式下初始化不会等待全局就绪应用需轮询ncclCommGetAsyncError直到不再返回ncclInProgress从 init.cc 可以看到ncclCommInitRankConfig在ncclGroupEndInternal之后若config.blocking为假就会立刻调用ncclCommGetAsyncError探测当前异步状态——这正是上面轮询模式的官方实现依据除直接创建新通信器外也可以用 ncclCommSplit 从已有通信器按color/key分裂出子通信器color相同的 rank 落入同一新通信器NCCL_SPLIT_NOCOLOR表示该 rank 退出返回 NULL 通信器。关于非阻塞通信器的完整容错用法ncclCommAbort中止、失败处理官方用户指南在 communicators.rst 的 Failure Handling 一节原文:ref:ft指向该处有详细说明。二、通信操作示例集合通信如何入队与完成通信器建好之后如何发起集合通信同样取决于每线程管理几张卡。1. 每进程/每线程一个设备直接调用当每个线程/进程对应一张卡时每个线程只需为自己的设备直接调用集合操作例如 AllReducencclAllReduce(sendbuff, recvbuff, count, datatype, op, comm, stream);调用返回后操作只是被入队到了 CUDA stream并不表示已完成。若希望等待完成再调用cudaStreamSynchronize(stream);完整可运行版本即第一节的示例 2MPI 单设备每进程。2. 每线程多设备Group 语义批量入队当一个线程管理多张设备时必须用 Group 语义一次性在多个设备上启动操作ncclGroupStart(); for (int i0; ingpus; i) ncclAllReduce(sendbuffs[i], recvbuff[i], count, datatype, op, comms[i], streams[i]); ncclGroupEnd();ncclGroupEnd返回后组内所有操作均已入队到各自的 stream。随后对每张卡的 stream 逐一同步即可等待全部完成for (int i0; ingpus; i) cudaStreamSynchronize(streams[i]);为什么同一线程必须用 Group从 nccl.h.in 对ncclSend/ncclRecv的注释可以看到同样的原则NCCL 操作对 GPU 是阻塞式的如果多个操作需要并发推进才能完成它们必须被融合进一个ncclGroupStart/ncclGroupEnd区间。集合操作同理——同一线程对多个通信器分别发起的调用会相互等待只有 Group 融合后才被视为一批并发请求。三、构建、运行与调试上述模式在仓库 docs/examples 中有可直接构建的 C/CUDA 示例部分带nccl4pyPython 变体其组织与通信器模式的对应关系如下来自 01_communicators/README.md特性ncclCommInitAll单进程多卡pthread每线程一卡MPI每进程一卡多节点支持不支持不支持支持进程隔离无无有搭建复杂度低中高最佳适用简单测试、单节点应用单节点线程并行集群生产训练构建示例可以随 NCCL 源码一起构建也可以独立构建。NCCL 不自带启动器仓库示例内置 pthread 与 MPI 两种 bootstrap 机制需要 MPI 时加MPI1非标准位置的 MPI 库可用MPI_HOME指定# 随 NCCL 源码构建 make -j examples [MPI1] # 或 NCCL 已构建后独立构建示例 cd docs/examples make NCCL_HOMEpath-to-nccl [MPI1]构建期可选环境变量NCCL_HOMENCCL 安装根目录、MPI0/1、MPI_HOME、CUDA_HOME。运行单进程多卡示例直接执行即可MPI 示例用mpirun -np n启动线程示例可通过NTHREADSn控制线程数默认等于可见 GPU 数CUDA_VISIBLE_DEVICES控制可见 GPU其他所有 NCCL 环境变量如NCCL_DEBUG同样生效。调试遇到运行时问题可设置NCCL_DEBUGINFO打开详细日志。各示例目录的c/README.md还包含各自的 Common Issues and Solutions 常见问题小节。四、小结选对初始化 API单进程多卡用ncclCommInitAll内部就是 Group 包裹的ncclCommInitRank见 init.cc每进程/线程一卡用ncclGetUniqueId 广播 ncclCommInitRank单线程管多卡时所有跨卡调用包括初始化本身都要用ncclGroupStart/ncclGroupEnd融合。设备绑定是隐式的ncclCommInitRank通过cudaGetDevice获取当前线程绑定的设备因此每次调用前必须cudaSetDevice多卡时全局 rank 需按myRank*nDev i之类的规则唯一编排。需要精细控制时切换到ncclCommInitRankConfig未初始化的ncclConfig_t必须先用NCCL_CONFIG_INITIALIZER赋值blocking 0创建非阻塞通信器后用ncclCommGetAsyncError轮询直到就绪从而实现可中止、可并发的初始化配合ncclCommAbort/ncclCommSplit做容错。NCCL 调用返回 ≠ 通信完成所有集合操作都只是入队到 CUDA stream必须cudaStreamSynchronize后才能确认完成同时每个 NCCL 调用都应检查ncclResult_t返回值。赞分享高性能计算通信深度学习分布式训练【免费下载链接】ncclOptimized primitives for collective multi-GPU communication项目地址https://gitcode.com/gh_mirrors/nc/nccl点击查看免费下载相关推荐NCCL 集体通信CollectiveAPI 全面解析八大集合操作的签名、语义与实现原理NCCL 集体通信CollectiveAPI 全面解析八大集合操作的签名、语义与实现原理 本指南以 NCCL 用户指南的 Collective Commu高性能计算通信深度学习分布式训练MLX 分布式通信mlx.core.distributed完全指南Group、集合通信与四大后端MLX 分布式通信mlx.core.distributed完全指南Group、集合通信与四大后端 导读 MLXArray framework for A人工智能深度学习机器学习本地部署NCCL C API 完全指南通信器管理、集合通信、Device API 与运行时参数系统的深度解析NCCL C API 完全指南通信器管理、集合通信、Device API 与运行时参数系统的深度解析 本文以 NCCL 用户指南的 API 参考 docs/高性能计算通信深度学习分布式训练上一篇Mac Mouse Fix深度配置指南如何实现专业级鼠标定制与平滑滚动优化下一篇VisiData 菜单系统完全指南从鼠标操作到菜单定制创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考