浅谈HTTPS传输协议原理

本文深入解析HTTPS通信中SSL加密的基本实现方法,包括对称加密、非对称加密算法特点,以及HTTPS通信过程的详细步骤,确保客户端与服务器间通信的安全性。

我们常常在使用网上银行时看到的连接都是以“https”开始的,那么这个https是什么呢?这其实是表示目前连接使用了SSL进行加密,能保证客户端到服务器端的通信都在被保护起来,那么浏览器是如果实现的呢?下面让我们来介绍一下SSL基本的实现方法。

首先我们有两种基本的加解密算法类型:对称加密,非对称加密(公私钥加密),现在介绍一下这两种加密算法的特点:

对称加密:密钥只有一个,加密解密为同一个密码,且加解密速度快,典型的对称加密算法有DESAES等,示意图如下:


1 对称加密

非对称加密:密钥成对出现(且根据公钥无法推知私钥,根据私钥也无法推知公钥),加密解密使用不同密钥(公钥加密需要私钥解密,私钥加密需要公钥解密),相对对称加密速度较慢,典型的非对称加密算法有RSADSA等,示意图如下:


2 非对称加密

根据上面的两种加密方法,现在我们就可以设计一种无法让他人在互联网上知道你的通讯信息的加密方法:

  1. 在服务器端存在一个公钥及私钥

  2. 客户端从服务器取得这个公钥

  3. 客户端产生一个随机的密钥

  4. 客户端通过公钥对密钥加密(非对称加密)

  5. 客户端发送到服务器端

  6. 服务器端接受这个密钥并且以后的服务器端和客户端的数据全部通过这个密钥加密(对称加密)

    HTTPS通信过程的时序图如下:


    3 HTTPS通信时序图

正如上图所示,我们能保证下面几点:

  1. 客户端产生的密钥只有客户端和服务器端能得到

  2. 加密的数据只有客户端和服务器端才能得到明文

  3. 客户端到服务端的通信是安全的

当然实际的SSL实现算法复杂的多,并有数据校验、身份验证等功能,如果需要更多了角请参看RFC2246RFC4346文档。

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
最新发布
06-27
评论
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

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

抵扣说明:

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

余额充值