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,GPU与GPU之间建立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
最新发布