Fast Multi-GPU collectives with NCCL-翻译

NCCL是一个针对多GPU集群通信的轻量级库,能够利用拓扑结构提高通信效率。它支持多种集体通信模式,如all-reduce、all-gather、broadcast等,并提供了接近最优的带宽性能。
部署运行你感兴趣的模型镜像

本文是对NCCL官方英文博客的翻译,方便大家学习交流。因水平有限,如有问题欢迎更正。
PS:我只翻译了介绍,且翻译的不严谨。nv官方博客给的介绍比较简单,大家有个简单认知就可以了。博客里对NCCL的使用介绍的也非常简单,随后我会翻译并丰富NCCL的使用,单独成一篇博客。


目前许多服务器包含8个或更多的GPU。 理论上,将应用程序从一个GPU扩展到多个GPU应该提供巨大的性能提升。然而在实践中,性能提升往往很难达到理论值。
这个问题往往是由两个原因导致。 第一是分配任务时没有充分利用的并行性,从而导致处理器不饱和。 第二个原因是任务分配不合理,导致GPU间需要交换大量数据,花费更多的时间进行通信而不是计算。 解决通信瓶颈的最重要的方式是充分利用可用的GPU间带宽,这是NCCL所关心的。
NCCL(发音为“Nickel”)是一个多GPU集群通信原语库,它具有拓扑感知能力,可以轻松集成到您的应用程序中。 NCCL最初作为一个开源研究项目开发,因此NCCL的特点为轻量级,仅依赖于通常的C ++和CUDA库。 NCCL可以部署在单进程或多进程应用程序中,透明地处理进程间通信。 最后,NCCL的API格式与MPI collective的API格式基本相同,如果你有MPI collective经验,那可以很快上手NCCL。

Collective Communication

Collective communication routines是许多处理器之间的数据传输的常见模式。 如果你有MPI的经验,那么你应该已经很熟悉常见的集体操作。例如,图1是一个all-reduce的例子。共四个GPU(0,1,2,3),每个GPU上有一个大小为N的vector(A,B,C,D),计算为求vectorAdd,S[k]=A[k]+B[K]+C[k]+D[k]( 0<=k<N ).

Figure 1: Illustration of the All-Reduce collective.

Figure 1: Illustration of the All-Reduce collective.

另一个常见的collective操作是all-gather,其中K个GPU中有一个N大小的数组,每个GPU将自己的数组传播到其他所有GPU中,最后所有GPU都会得到N * K大小的结果数据,如图2所示。

Figure 2: Illustration of the All-Reduce collective.

Figure 1: Illustration of the All-Reduce collective.

broadcast是第三例子。broadcast即将一个GPU内的N大小数据传输到其他所有GPU中。如图3所示。

Figure 2: Illustration of the All-Reduce collective.

Figure 3: Illustration of the Broadcast collective.

以上所有提到的collectives实现都有“out-of-place”和“in-place”版本。“out-of-place”的意思是输入数据和输出数据是互相独立的数据空间。“in-place”的意思是输出数据覆盖输入数据。

Figure 2: Illustration of the All-Reduce collective.

Figure 4: A common PCIe topology for 4 GPUs attached to a single CPU. Red arrows represent PCIe x16 connections.

有许多方法可以有效地实现collectives。然而,与众不同的是,我们的实现考虑了处理器之间的互连拓扑。例如,在如图4所示的PCIe树拓扑中,考虑将数据从GPU0广播所有其它GPU。

在这种情况下,two-step树算法是常见的选择。第一步,数据先从GPU0传输到一个GPU中。第二步,数据从这两个GPU中同时传输到剩余的其他GPU。
在本例子中,two-step树算法有两种选择。(1)第一步中将数据从GPU0传输到GPU1,然后第二步GPU0 to GPU2,GPU1 to GPU3。(2)第一步中将数据从GPU0传输到GPU2,然后第二步GPU0 to GPU1,GPU2 to GPU3。根据拓扑结果,很明显能够看出第二选项是较优的,因为从GPU0到GPU2和GPU1到GPU3同时发送数据将导致在PCIe链路上的争用,从而导致在该步骤中,有效带宽减半。一般来说,为了实现良好的collectives性能,需要仔细注意互连拓扑。

Figure 5: Ring order of GPUs in PCIe tree.

Figure 5: Ring order of GPUs in PCIe tree.

一个环形(ring)拓扑结构是一个很好的方法来优化broadcast的带宽。
然后通过将环周围的输入的小块从GPU0中继到GPU3来执行广播。 有趣的是,即使当应用于“树形”PCIe拓扑时,环算法也为几乎所有的标准集合操作提供接近最优的带宽。 但请注意,选择正确的振铃顺序仍然很重要。

GPU Collectives with NCCL

NCCL内collectives的实现方式为ring-style,以提供最大带宽。NCCL能够自动对GPU进行编号并实现最佳ring顺序。因此程序员不需要对硬件进行配置即可时应用程序在获得很好的性能。

许多collectives实现时需要一个缓冲区来存储中间结果。 为了最小化每个GPU上的几MB的内存开销,NCCL将大型集合分成许多小块。 对于集合算法的每个步骤和块。
启动单独的内核和cudaMemcpy调用是非常低效的。对于一个collectives算法,将每一个步骤实现为一个kernel,每一个chunk数据都执行cudaMemcpy,这种做法是非常低效的。因此,NCCL内的每一个collective操作都由一个大kernel实现。
NCCL广泛使用GPUDirect peer to peer直接访问方式在处理器之间传输数据。
在两个GPU P2P 访问不可用的情况下(比如两个GPU分别挂载在两个CPU下的PCI-E switch下,此时这两个GPU是不能P2P的),数据要先传输到pinned memory暂存,再传输到目标GPU。类似地,通过轮询device memory 或 pinned memory中volatile变量来实现同步。

在NCCL中,NCCL使用三个primitives来实现每个collectives:Copy, Reduce, and ReduceAndCopy。每一个步骤都被优化,目的是能够有效地在GPU间传输细颗粒度的数据片段(4-16KB)。NVCCL中的kernels经过优化后,实现了在低占用率下实现最大带宽。因此,NCCL可以达到使用单个CUDA block就可以使PCIe 3.0*16互联带宽饱和。这使得小部分线程通信、大部分线程计算同时进行。

NCCL目前支持all-gather,all-reduce,broadcast,reduce和reduce-scatter collectives操作。NCCL只适用于单节点环境,节点内可以由任意数量的GPU。

Figure 2: Illustration of the All-Reduce collective.

您可能感兴趣的与本文相关的镜像

Wan2.2-I2V-A14B

Wan2.2-I2V-A14B

图生视频
Wan2.2

Wan2.2是由通义万相开源高效文本到视频生成模型,是有​50亿参数的轻量级视频生成模型,专为快速内容创作优化。支持480P视频生成,具备优秀的时序连贯性和运动推理能力

NCCL:chunk,slice,step的关系 原创 已于 2025-03-01 17:01:41 修改 · 1k 阅读 · 31 · 29 · CC 4.0 BY-SA版权 文章标签: #gpu算力 #算法 关于Simple协议的4MB缓存的使用: https://github.com/NVIDIA/nccl/issues/544 For collectives we have a 4MB buffer cut into 8x 512K chunks indeed. AllReduce/Ring, Allgather, and ReduceScatter will send 2 chunks at a time (hence 1MB), while Broadcast, Reduce and AllReduce/Tree will use the 8 chunks independently. 运行环境 单机4卡,4个channel,每个channel的ring顺序如下: GPU0--->GPU1--->GPU2-GPU3 提示: runRing函数里的:chunkSize,loopSize,size,minChunkSize,realChunkSize,gridOffset,offset,nelem等变量,全都是数量,而不是空间大小!!!! 一、Broadcast集合通信(NCCL_ALGO_TREE) 1、Simple协议分配的共享显存大小是DEFAULT_BUFFSIZE=4MB,那如果Broadcast发送256MB数据的话(数据类型为unit8或者int32都是一样的),是怎么划分的?chunk,step和slice间的关系。 设计case发送256MB的数据(在调用ncclBroadcast的时候传入的size为:long long size = 64*1024*1024,数据类型为 int,broadcast会把数据类型统一转换为char类型:ncclInfoSetDerived(),总字节数不变),打印runRing中的数据如下: 命令为:all_reduce_perf -b 256M -e 256M -f 2 -g 1 -n 1 -d int32 leos: broadcast.h runRing nthreads = 256 leos: broadcast.h runRing args->nWarps = 4 leos: broadcast.h runRing bid = 0 1 2 3 //bid是使用的channel的编号,这里4个channel全部用上 leos: broadcast.h runRing nChannels = 4 //nChannels是使用的channel的个数,用了全部4个channel leos: broadcast.h runRing chunkSize = 524288 //一个buff是4MB,分为了NCCL_STEPS个(8个)chunk,所以broadcast里的chunkSize就是512K(这里是数量),一个chunk占用的空间大小就是512K*1B=512KB(这里是空间) Broadcast的runRing中计算chunkSize const ssize_t chunkSize = int(Proto::calcBytePerStep()/sizeof(T) * (Proto::Id == NCCL_PROTO_SIMPLE ? BROADCAST_CHUNKSTEPS : 1)); leos: broadcast.h runRing loopSize = 2097152 //loopSize是nChannels*chunkSize = 4 * 524288 = 2M(这个是数量),每个循环处理的数据量大小为:2M*1B=2MB(这里是空间),这256MB的数据,每次要处理2MB,这2MB数据是由4个channel共同处理的,每个channel也就是每个block处理512KB数据 leos: broadcast.h runRing size = 268435456 //真正的数据总量(这里是数量),大小为:256M*1B=256MB(这里是空间) leos: broadcast.h nelem= 524288 //nelem是每个Primitives函数处理的数据长度,是512K(这里是数量),占用空间大小为512K*1B=512KB(这里是空间),这一个chunk的大小,每个block处理512KB数据,总共256MB数据,需要分为512次处理完(4个channel),分到每个channel上就是64次 所以总结一下就是:总共256MB数据,每次处理2MB数据(loopSize是2M(这个是数量)),总共需要128次发送完256MB数据。 而每次处理的这2MB数据,是4个channel共同完成的,每个channel每次发送就是512KB数据。而每个channel有256个线程,就是这256个线程共同处理512KB数据。 共享缓存是4MB,就是最多可以存8个512KB数据,就是8个chunk。 2、broadcast集合通信Simple协议buffer的划分: 几个概念的解释,chunk,step和slice chunk:通信原语Primitives的函数比如copySend,一次传输的就是一个chunk的数据。对应上图一个loopSize分为4个nelem,一个nelem就对应一个chunk,即一次Primitives函数调用所传递的数据大小。 step:是用来记录数据传递的计数的,broadcast里设置了一个chunk对应一个step,每发送一个chunk,step就要加1,也就是说step记录了总共发送的chunk大小。因此broadcast的stepSize就是bufferSize / (sizeof(T) * NCCL_STEPS)=512KB。(broadcast里sizeof(T)是1) 一个chunk可以有多个step,broadcast里是即BROADCAST_CHUNKSTEPS=1,所以chunkSize就是stepSize * BROADCAST_CHUNKSTEPS =512KB; slice:一次Primitives函数内部所传递的数据大小是一个chunk,genericOp里这个chunk可能分为多个slice分次传递,在broadcast里设置了一个slice就对应了一个chunk。每个Simple协议的buffer大小是4MB,broadcast里设置了一个slice就是一个chunk,所以每个slice大小是512KB, 每个step和slice对应的,每次消耗一个slice,就给step加1。 在broadcast通信里,可以只考虑chunk和step,忽略slice的概念(不应该忽略,因为节点间通信的单位也是slice) 每次调用Primitives的函数发送一个chunk,就把step++,用来做同步使用 3、使用LL协议的情况为 LL协议分配的共享显存大小是DEFAULT_LL_BUFFSIZE = 256KB 设计case发送256MB的数据,打印runRing中的数据如下: leos: broadcast.h runRing nthreads = 256 leos: broadcast.h runRing args->nWarps = 4 leos: broadcast.h runRing bid = 0 1 2 3 //bid是使用的channel的编号,这里4个channel全部用上 leos: broadcast.h runRing nChannels = 4 //nChannels是使用的channel的个数,用了全部4个channel leos: broadcast.h runRing chunkSize = 16384 //chunkSize = DEFAULT_LL_BUFFSIZE / 8 / 2 = 16K,一个chunk是16K*1B=16KB leos: broadcast.h runRing loopSize = 65536 //loopSize是nChannels*chunkSize = 4 * 16384 = 64K,每个循环处理的数据量大小为:64K*1B=64KB,这256MB的数据,每次要处理64KB,这64KB数据是由4个channel共同处理的,每个channel也就是每个block处理16KB数据 leos: broadcast.h runRing size = 268435456 //真正的数据总量,256MB leos: broadcast.h nelem= 16384 //nelem是每个Primitives函数处理的数据长度,是16KB,这一个chunk的大小,每个block处理16KB数据,总共256MB数据,需要分为16,384次处理完(4个channel),分到每个channel上就是4096次 所以总结一下就是:总共256MB数据,每次处理64KB数据(loopSize是64KB),总共需要4,096次发送完256M数据。 而每次处理的这64KB数据,是4个channel共同完成的,每个channel每次发送就是16KB数据。而每个channel有256个线程,就是这256个线程共同处理16KB数据。 三种协议上计算chunkSize的方式有不同,每种协议的缓存大小为: DEFAULT_LL_BUFFSIZE = 256KB 262144 DEFAULT_LL128_BUFFSIZE = 448KB 458752 DEFAULT_BUFFSIZE = 4MB 4194304 Simple:chunkSize = DEFAULT_BUFFSIZE / 8 = 512K LL: chunkSize = DEFAULT_LL_BUFFSIZE / 8 / 2 = 16K LL128: chunkSize = (DEFAULT_LL128_BUFFSIZE / 8) * 7 / 8 = 49K #define NCCL_STEPS 8 #define BROADCAST_SLICESTEPS 1 #define BROADCAST_CHUNKSTEPS 1 二、Allreduce集合通信(NCCL_ALGO_RING) 1、Simple协议分配的共享显存大小是DEFAULT_BUFFSIZE=4MB,那如果allreduce发送256MB数据,int32类型的话,是怎么划分的?chunk,step和slice间的关系。 设计case发送256MB(int32类型)的数据,打印runRing中的数据如下: leos: all_reduce.h runRing nthreads = 256 //一个block中有256个线程 leos: all_reduce.h runRing args->nWarps = 4 //256个线程划分为4个warp leos: all_reduce.h runRing bid = 0 1 2 3 //bid是使用的channel的编号,这里4个channel全部用上 leos: all_reduce.h runRing nChannels = 4 //nChannels是使用的channel的个数,用了全部4个channel leos: all_reduce.h runRing chunkSize = 524288 //一个buff是4MB,分为了NCCL_STEPS个(8个)chunk,所以allreduce里的chunkSize就是512K(4MB/8/4B*4,这里是数量),一个chunk占用的空间大小(并不是chunkSize的值)就是512K*4B=2MB(这里是空间) Allreduce的runRing中计算chunkSize __device__ static int calcBytePerStep() { return ncclShmem.comm.buffSizes[NCCL_PROTO_SIMPLE]/NCCL_STEPS; } const ssize_t chunkSize = int(Proto::calcBytePerStep()/sizeof(T) * (Proto::Id == NCCL_PROTO_SIMPLE ? ALLREDUCE_CHUNKSTEPS : 1)); leos: all_reduce.h runRing loopSize = 8388608 //loopSize是nChannels*nranks*chunkSize = 4 * 4 * 524288 = 8M(这里是数量),每个循环处理的数据量大小为:8M*4B=32MB(这里是空间),这256MB的数据,每次要处理32MB,这32MB数据是由4个channel共同处理的,每个channel也就是每个block处理8MB数据 leos: all_reduce.h runRing size = 67108864 //真正的数据总量(这里是数量),大小为:64M * 4B = 256MB(这里是空间) StepPerSlice = 2 SlicePerChunk = 2 #define ALLREDUCE_CHUNKSTEPS (NCCL_STEPS/2) 4 #define ALLREDUCE_SLICESTEPS (NCCL_STEPS/4) 2 下面的变量都是数量 leos: step 0 all_reduce.h realChunkSize= 524288 leos: step 0 all_reduce.h nelem= 524288 leos: k-2 steps reduce all_reduce.h realChunkSize= 524288 leos: k-2 steps reduce all_reduce.h nelem= 524288 leos: k-1 step all_reduce.h realChunkSize= 524288 leos: k-1 step all_reduce.h nelem= 524288 leos: k-2 steps copy to next GPU reduce all_reduce.h realChunkSize= 524288 leos: k-2 steps copy to next GPU reduce all_reduce.h nelem= 524288 leos: step final copy all_reduce.h realChunkSize= 524288 leos: step final copy all_reduce.h nelem= 524288 2、Simple协议分配的共享显存大小是DEFAULT_BUFFSIZE=4MB,那如果allreduce发送256MB数据,int8类型的话,是怎么划分的?chunk,step和slice间的关系。 设计case发送256MB(int32类型)的数据,打印runRing中的数据如下: leos: all_reduce.h runRing nthreads = 256 //一个block中有256个线程 leos: all_reduce.h runRing args->nWarps = 4 //256个线程划分为4个warp leos: all_reduce.h runRing bid = 0 1 2 3 //bid是使用的channel的编号,这里4个channel全部用上 leos: all_reduce.h runRing nChannels = 4 //nChannels是使用的channel的个数,用了全部4个channel leos: all_reduce.h runRing chunkSize = 2097152 //一个buff是4MB,分为了NCCL_STEPS个(8个)chunk,所以allreduce里的chunkSize就是2M(4MB/8/1B*4,这里是数量),一个chunk占用的空间大小(并不是chunkSize的值)就是2M*1B=2MB(这里是空间) leos: all_reduce.h runRing loopSize = 33554432 //loopSize是nChannels*nranks*chunkSize = 4 * 4 * 2097152 = 32M(这里是数量),每个循环处理的数据量大小为:32M*1B=32MB(这里是空间),这256MB的数据,每次要处理32MB,这32MB数据是由4个channel共同处理的,每个channel也就是每个block处理8MB数据 leos: all_reduce.h runRing size = 268435456 //真正的数据总量(这里是数量),大小为:256M * 1B = 256MB(这里是空间) 3、allreduce集合通信Simple协议buffer的划分: 几个概念的解释,chunk,step和slice chunk:通信原语Primitives的api比如copySend,一次传输的就是一个chunk的数据。对应上图一个loopSize分为4个nelem,一个nelem就对应一个chunk,即一次Primitives函数调用所传递的数据大小。 step:是用来记录数据传递的个数的,Allreduce里设置了1个chunk对应2个slice,1个slice对应2个step,每发送一个chunk,就要对应发送2个slice,每发送1个slice,step就要加2,也就是说step记录了总共发送的slice的个数。Allreduce中Ring的stepSize就是bufferSize / (sizeof(T) * NCCL_STEPS)=128K。(Allreduce里sizeof(T)是4),一个step代表空间大小为:128K*4B=512KB 一个chunk可以有多个step,Allreduce里是即ALLREDUCE_CHUNKSTEPS=4,所以chunkSize就是stepSize * ALLREDUCE_CHUNKSTEPS=512K;(Simple协议Ring算法Allreduce集合通信, chunk大小为512K*4B=2MB) slice:一次Primitives函数调用所传递的数据大小是一个chunk,genericOp里这个chunk可能分为多个slice分次传递。每个Simple协议的buffer大小是4MB,Allreduce里设置了两个slice就是一个chunk,所以每个slice大小是1MB, Allreduce里每两个step是一个slice,每次消耗一个slice,就给step加2。 每次调用Primitives的函数发送一个chunk,就等于发送两个slice,就等于发送四个step,就把step+4,用来做同步使用 Ring算法,Simple协议,int32类型,4个channel,Allreduce集合通信buff的划分: Ring算法,Simple协议,int8类型,4个channel,Allreduce集合通信buff的划分: Allreduce,数据类型是int32,选择Ring算法,Simple协议的话,chunksize大小是多少: // Data bytes (no flags etc) in one step of the fifo queue. __device__ static int calcBytePerStep() { return ncclShmem.comm.buffSizes[NCCL_PROTO_SIMPLE]/NCCL_STEPS; } const ssize_t chunkSize = int(Proto::calcBytePerStep()/sizeof(T) * (Proto::Id == NCCL_PROTO_SIMPLE ? ALLREDUCE_CHUNKSTEPS : 1)); =512K/4*4 = 512K(个)(这个是数量) 1G数据,Allreduce Ring算法跟rank个数有关系 场景1: Ring算法,Simple协议,int8类型,4个channel,Allreduce中关键数据结构打印(两个节点,每个节点4个GPU卡): opCount 0 slicesteps 2 spl 14 cpl 8 nbytes 1073741824 -> algorithm 1 protocol 2 nchannels 4 nthreads 256, nloops 16 nsteps 896 chunksize 2097152 comm 0x560d1a939d00 Ring算法,Simple协议,int8类型,4个channel,Allreduce中关键数据结构打印(一个节点,4个GPU卡): opCount 0 slicesteps 2 spl 6 cpl 4 nbytes 1073741824 -> algorithm 1 protocol 2 nchannels 4 nthreads 256, nloops 32 nsteps 768 chunksize 2097152 comm 0x55f13c275550 场景2: Ring算法,Simple协议,int32类型,4个channel,Allreduce中关键数据结构打印(两个节点,每个节点4个GPU卡): opCount 0 slicesteps 2 spl 14 cpl 8 nbytes 1073741824 -> algorithm 1 protocol 2 nchannels 4 nthreads 256, nloops 16 nsteps 896 chunksize 2097152 comm 0x5558c097ac80 Ring算法,Simple协议,int32类型,4个channel,Allreduce中关键数据结构打印(一个节点,4个GPU卡): opCount 0 slicesteps 2 spl 6 cpl 4 nbytes 1073741824 -> algorithm 1 protocol 2 nchannels 4 nthreads 256, nloops 32 nsteps 768 chunksize 2097152 comm 0x55f9599dcf10 Tree算法Allreduce,不管是一个节点还是两个节点,nloops nsteps 都一样,Tree算法跟rank个数没有关系 场景3: Tree算法,Simple协议,int8类型,4个channel,Allreduce中关键数据结构打印(两个节点,每个节点4个GPU卡): opCount 0 slicesteps 1 spl 1 cpl 1 nbytes 1073741824 -> algorithm 0 protocol 2 nchannels 4 nthreads 256, nloops 512 nsteps 512 chunksize 524288 comm 0x55d39ec0bfe0 Tree算法,Simple协议,int8类型,4个channel,Allreduce中关键数据结构打印(一个节点,4个GPU卡): opCount 0 slicesteps 1 spl 1 cpl 1 nbytes 1073741824 -> algorithm 0 protocol 2 nchannels 4 nthreads 256, nloops 512 nsteps 512 chunksize 524288 comm 0x556cfa4b2930 场景4: Tree算法,Simple协议,int32类型,4个channel,Allreduce中关键数据结构打印(两个节点,每个节点4个GPU卡): opCount 0 slicesteps 1 spl 1 cpl 1 nbytes 1073741824 -> algorithm 0 protocol 2 nchannels 4 nthreads 256, nloops 512 nsteps 512 chunksize 524288 comm 0x55b4f5817d60 Tree算法,Simple协议,int32类型,4个channel,Allreduce中关键数据结构打印(一个节点,4个GPU卡): opCount 0 slicesteps 1 spl 1 cpl 1 nbytes 1073741824 -> algorithm 0 protocol 2 nchannels 4 nthreads 256, nloops 512 nsteps 512 chunksize 524288 comm 0x5588838e5b10 Allreduce Tree算法里chunk: slice: step = 1:1:1 问题: all_reduce.h 的 runRing 里的 chunkSize,是描述了类型T的数据的数量,还是类型T的数量占用的空间呢? 当然在T==char的时候,两个值是相等的没有什么差别。 但是如果T==int32,那如果chunkSize描述的是数量还是比较合理的,因为chunkSize的计算公式是要/T的 const ssize_t chunkSize = int(Proto::calcBytePerStep()/sizeof(T) * (Proto::Id == NCCL_PROTO_SIMPLE ? ALLREDUCE_CHUNKSTEPS : 1)); 解答: all_reduce.h 里的 runRing chunkSize的意思是有多少个类型为T的数据(这个是数量),而不是总的数据量大小bytes。比如T为int32的时候,在Simple协议下,打印的chunkSize就是524288=512K,而T为char的时候,同样是Simple协议,chunkSize就是2097152=2M 但是两个数据类型下,每次处理的数据量应该都是2MB,只不过在int类型的时候,chunkSize小,在char类型的时候,chunkSize大,这也能说明chunkSize其实是个数量,不是个空间大小Bytes 用T==int和T==char做了对比测试: int32: leos: all_reduce.h runRing chunkSize = 524288 //数据类型为int32,allreduce里的一个chunk就是512K个int32类型的数据,那一个chunk的空间大小就是512K*4B=2MB leos: all_reduce.h runRing loopSize = 8388608 //loopSize是nChannels*nranks*chunkSize = 4 * 4 * 524288 = 8M,那数据量为:8M*4B=32 MB,这256MB的数据,每次要处理32MB,这32MB数据是由4个channel共同处理的,每个channe也就是每个block处理8MB数据 leos: all_reduce.h runRing size = 67108864 //真正的数据总量,64M * 4B = 256MB char: leos: all_reduce.h runRing chunkSize = 2097152 //数据类型为char,allreduce里的一个chunk就是2M个char类型的数据,那一个chunk的空间大小就直接是2MB leos: all_reduce.h runRing loopSize = 33554432 //loopSize是nChannels*nranks*chunkSize = 4 * 4 * 2097152 = 32M,那数据量为:32M*1B=32 MB,这256MB的数据,每次要处理32MB,这32MB数据是由4个channel共同处理的,每个channe也就是每个block处理8MB数据 leos: all_reduce.h runRing size = 268435456 //真正的数据总量,268435456 = 256MB 总结一下: 对于broadcast来说,在核函数处理之前,NCCL已经把数据统一转换为char类型了,就不需要考虑数量和空间大小的关系了。 对于Allreduce来说,核函数收到的数据就是原始的数据,如果数据类型是char,那么数量和空间大小是对应的。如果数据类型是int32,那么空间大小==数量*4B 对于chunkSize,loopSize和nelem这几个代码中的变量,可以理解为是描述了数据类型为T的数据的数量,要是计算占用空间的话,需要根据数据类型T做相应的转换。 无论数据类型是int还是char,broadcast和Allreduce内部的一个step代表的空间大小都是512KB。broadcast的一个chunk和一个slice大小都是512KB,Allreduce一个chunk大小是2MB,一个slice大小是1MB 在Allreduce过程中,每个Rank接收和发送的数据量都为:2K(N - 1)/N,其中K是每个Rank的数据总量,N是Rank个数。 3、NCCL来说 Simple协议 leos: genericOp stepSize = 131072(128K), StepPerSlice = 2, sliceSize = 262144(512K) leos: genericOp sliceSize = 262144(512K) leos: waitPeer nelts = 262144(512K) leos: waitPeer sizeof(T) = 4 森林好小子999 关注 31 29 0 分享 TP、PP、DP、集群GPU卡数、grad_accum_steps之间关系【DP=集群中GPU卡总数world_size/(TP*PP)】【grad_accum_steps=gbz/(mbz*DP)】 u013250861的博客 986 【代码】TP、PP、DP、grad_accum_steps之间关系【DP=集群中GPU卡总数world_size/(TP*PP)】 5. PyTorch+NCCL源码编译_pytorch nccl 6-23 # 新建conda虚拟环境,取名为nccl2conda create -n nccl2python=3.11conda activate nccl2#下载v2.2.1 源码gitclone --branch v2.2.1 --recursive https://github.com/pytorch/pytorchcdpytorch# v2.2.1# 安装依赖包pipinstall-r requirements.txt#以开发模式安装torch,不使用系统nccl,而是torch自带的,位于third ... NVIDIA NCCL 源码学习(十)- 多机间ncclSend和ncclRecv的过程_nccl编 ... 6-22 intpeerrecv = (info->comm->nRanks+info->comm->rank-info->delta)%info->comm->nRanks; args.nsteps =DIVUP(info->recvbytes, info->comm->buffSizes[NCCL_PROTO_SIMPLE]/NCCL_STEPS/SENDRECV_SLICEFACTOR); if(args.nsteps ==0) args.nsteps =1; NCCLCHECK(SaveProxy<proxyRecv>(peerrecv, &ar... LLM学习笔记 weixin_40103562的博客 855 常见术语。 理解NCCL源码必看:详细论述NCCL源码做了哪些事 这个地瓜真是甜啊 7622 详解NCCL 源码的实现的功能以及如何实现的 NVIDIA NCCL 源码学习(十一)- ring allreduce_nccl源码分析 6-25 ncclDataType_t datatype,ncclRedOp_t op,ncclComm*comm,cudaStream_t stream){structncclInfoinfo={ncclCollAllReduce,"AllReduce",sendbuff,recvbuff,count,datatype,op,0,comm,stream,/* Args */ALLREDUCE_CHUNKSTEPS,ALLREDUCE_SLICESTEPS};returnncclEnqueueCheck(&info);}... 浅谈NCCL Proxy线程 6-23 本文档针对NCCL 2.19版本。 在intra-node,GPUGPU之间建立P2Ptransport的时候,和inter-node中,通过NET建立NET transport的时候,都需要proxy线程的参与,其实总共有两个proxy线程,一个叫做proxyService线程,是每个NODE中每个GPU对应的一个,主要维护连接建立,建立transport的setup和connect阶段。另一个叫做proxyProgress线程,... NCCL 原理 最新发布 哦豁灬 2726 NCCL是Nvidia Collective multi-GPU Communication Library的简称,它是一个实现多GPU的collective communication通信(all-gather, reduce, broadcast)库,Nvidia做了很多优化,以在PCIe、Nvlink、InfiniBand上实现较高的通信速度。 NVIDIA NCCL 源码学习(十)- 多机间ncclSend和ncclRecv的过程 热门推荐 KIDGIN7439的专栏 1万+ 通信由kernel和proxy线程协调完成,send端kernel负责将数据从input搬运到buf,proxy线程负责将buf中数据通过网络发送给recv端,kernel和proxy间通过队列实现生产者消费者模式,send端通过rdma send发送数据,和recv端通过队列实现生产者消费者模式,队列位于send端,recv端每次下发一个wr到rq之后会执行rdma write通知send端 NVIDIA GPGPU的通信架构解析 6-25 sendbuff, recvbuff, count, datatype, op, root, comm, stream, /* Args */ REDUCE_CHUNKSTEPS, REDUCE_SLICESTEPS }; return ncclEnqueueCheck(&info);}前面的代码主要是参数设置,核心在于ncclEnqueueCheck。其基本逻辑如下:主机通过异步提交任务(包括通信命令),经过排队和组调度,利用cudalaunchkernel发送至GPU... 【分布式】小白看Ring算法 - 03 6-25 NCCL(NVIDIACollective Communications Library)是由NVIDIA开发的一种用于多GPU间通信的库。NCCL的RING算法是NCCL库中的一种通信算法,用于在多个GPU之间进行环形通信。 RING算法的基本思想是将多个GPU连接成一个环形结构,每个GPU与相邻的两个GPU进行通信。数据沿着环形结构传递,直到到达发送方的位置。这样的环形结构可以有效... NVIDIA NCCL 源码学习(六)- channel搜索 KIDGIN7439的专栏 9550 nccl中channel的概念表示一个通信路径,为了更好的利用带宽和网卡,以及同一块数据可以通过多个channel并发通信,另外后续可以看到一个channel对应了一个GPU SM,所以基于这些原因,nccl会使用多channel,搜索的过程就是搜索出来一组channel。 NVIDIA NCCL 源码学习(十一)- ring allreduce KIDGIN7439的专栏 1万+ 之前的章节里我们看到了nccl send/recv通信的过程,本节我们以ring allreduce为例看下集合通信的过程。整体执行流程和send/recv很像,所以对于相似的流程只做简单介绍,主要介绍ring allreduce自己特有内容。 NVIDIA NCCL 源码学习(九)- 单机内ncclSend和ncclRecv的过程_ncclenq... 6-24 NCCLCHECK(ncclRecv((void*)(recvbuff[i] + j * chunk), chunk, ncclFloat, j, comms[i], s[i])); } } NCCLCHECK(ncclGroupEnd()); //synchronizing on CUDA stream to complete NCCL communication for(inti=0; i<nDev; i++) CUDACHECK(cudaStreamSynchronize(s[i])); ... NVIDIA NCCL 源码学习(十二)- double binary tree 6-8 上节我们以ring allreduce为例看到了集合通信的过程,但是随着训练任务中使用的gpu个数的扩展,ring allreduce的延迟会线性增长,为了解决这个问题,NCCL引入了tree算法,即double binary tree。 double binary tree 朴素的tree算法将所有机器节点构造成一棵二叉树,支持broadcast,reduce,前缀和。假设root节点要broadcast一个消... 【乱写的】收集一些和GPU以及NCCL相关的定义(持续更新) 参谋谋的博客 1839 SHArP 论文,其实没必要细看。简单来说,SHArP是一个软硬结合的通信协议,实现在了NVIDIA Quantum HDR Switch的ASIC里。它可以把从各个node收到的数据进行求和,并发送回去。再说的通俗一点,通过使用SHArP,我们把求和(聚合/Reduce,随便怎么叫)的操作交由交换机完成了。这种做法,业界叫做In-network Computing(在网计算)。用术语展开来讲,就是将计算卸载到网络中进行。更多相关的知识可以看这个英伟达的汇报。 NCCL中QP和Channel是什么关系(来自deepseek) rjc_lihui的专栏 707 是 RDMA 的底层队列,负责“干活”(执行数据传输)。 NVIDIA NCCL 源码学习(十四)- NVLink SHARP 6-5 由于type为-1,因此ncclTopoFollowPath直接返回gpu0,从gpu0开始搜索。 ncclResult_tncclTopoSearchRecGpu(structncclTopoSystem*system,structncclTopoGraph*graph,structncclTopoGraph*saveGraph,structncclTopoNode*gpu,intstep,intbackToNet,intbackToFirstRank,intforcedOrder,int*time){if((*time)<=0)returnncclSucc... 【PyTorch中的分布式训练入门】:提升模型训练效率的秘密武器 ![torch_cluster-1.6.2+pt... # 摘要 本文深入探讨了PyTorch分布式训练的技术细节和实践应用,涵盖了分布式训练的基础理论、PyTorch实现、调试优化技巧,以及未来发展趋势。文章首先介绍了并行计算的模型、分布式训练的 NCCL_IB_SPLIT_DATA_ON_QPS效果分析 aashuii的博客 257 DMA Length是61440(0xf000),基本等于NCCL_P2P_NET_CHUNKSIZE=65535可以看出内存中的数据以NCCL_P2P_NET_CHUNKS ———————————————— 版权声明:本文为优快云博主「森林好小子999」的原创文章,遵循CC 4.0 BY-SA版权协议,转载请附上原文出处链接及本声明。 原文链接:https://blog.youkuaiyun.com/shanleo1986/article/details/145949742 what does loopshzie here
06-27
### NCCL 测试工具及使用方法 NCCL(NVIDIA Collective Communications Library)是 NVIDIA 提供的一个用于高性能 GPU 集群通信的库。为了验证 NCCL 的正确性和性能,可以使用 NCCL 自带的测试工具。这些工具可以帮助用户检测多 GPU 和多节点环境下的通信性能和稳定性。 以下是一些常用的 NCCL 测试工具及其使用方法: #### 1. **nccl-tests** `nccl-tests` 是 NVIDIA 提供的一组测试程序,用于验证 NCCL 的功能和性能。可以通过以下步骤安装和运行 `nccl-tests`: - 安装依赖项并编译 `nccl-tests`: ```bash git clone https://github.com/NVIDIA/nccl-tests.git cd nccl-tests make MPI=1 CUDA_HOME=/usr/local/cuda # 根据系统配置调整 CUDA_HOME 路径 ``` - 运行单机多 GPU 测试: ```bash mpirun -np <NUM_GPUS> -bind-to none -map-by slot \ -H localhost:<NUM_GPUS> -x NCCL_DEBUG=INFO -x LD_LIBRARY_PATH \ ./build/all_reduce_perf -b 8 -e 1G -f 2 -g <NUM_GPUS> ``` 其中: - `<NUM_GPUS>`:指定 GPU 的数量。 - `-b` 和 `-e`:分别表示测试数据大小的起始值和结束值。 - `-f`:指定每次测试的数据大小倍增因子。 - `-g`:指定每个进程使用的 GPU 数量。 - 运行多节点测试: ```bash mpirun -np <TOTAL_NUM_GPUS> -H <NODE1>:<GPUS_ON_NODE1>,<NODE2>:<GPUS_ON_NODE2> \ -x NCCL_DEBUG=INFO -x LD_LIBRARY_PATH ./build/all_reduce_perf -b 8 -e 1G -f 2 -g <NUM_GPUS_PER_NODE> ``` #### 2. **NCCL 环境变量调试** 在运行 NCCL 测试时,可以通过设置环境变量来调试和优化性能。例如: - `NCCL_DEBUG=INFO`:启用调试日志,帮助排查问题[^1]。 - `NCCL_IB_HCA=<IB_DEVICE>`:指定 InfiniBand 设备。 - `NCCL_SOCKET_IFNAME=<NETWORK_INTERFACE>`:指定网络接口。 #### 3. **常见错误排查** 如果在运行 NCCL 测试时遇到类似以下错误: ``` RuntimeError: NCCL error in: /opt/conda/conda-bld/pytorch_1640811805959/work/torch/csrc/distributed/c10d/ProcessGroupNCCL.cpp:957, invalid usage, NCCL version 21.0.3 ncclInvalidUsage: This usually reflects invalid usage of NCCL library (such as too many async ops, too many collectives at once, mixing streams in a group, etc). ``` 可以尝试以下方法解决: - 检查是否正确配置了 NCCL 环境变量。 - 确保所有节点上的 NCCL 版本一致。 - 减少并发操作的数量或调整批次大小[^1]。 #### 4. **性能分析工具** 除了 `nccl-tests`,还可以使用以下工具进行性能分析: - **NVIDIA Nsight Systems**:提供系统级性能分析,帮助识别瓶颈。 - **NVIDIA Nsight Compute**:专注于 GPU 内核性能分析。 ### 示例代码 以下是一个简单的 Python 脚本,演示如何使用 PyTorch 的分布式训练功能进行 NCCL 测试: ```python import torch import torch.distributed as dist import torch.multiprocessing as mp def run(rank, world_size): dist.init_process_group(backend='nccl', init_method='env://', rank=rank, world_size=world_size) tensor = torch.ones(1).cuda(rank) dist.all_reduce(tensor) print(f"Rank {rank}: {tensor.item()}") def main(): world_size = torch.cuda.device_count() mp.spawn(run, args=(world_size,), nprocs=world_size, join=True) if __name__ == "__main__": main() ```
评论
成就一亿技术人!
拼手气红包6.0元
还能输入1000个字符
 
红包 添加红包
表情包 插入表情
 条评论被折叠 查看
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

当前余额3.43前往充值 >
需支付:10.00
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明:

1.余额是钱包充值的虚拟货币,按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载,可以购买VIP、付费专栏及课程。

余额充值