【NCCL】transport数据传输(二)
参考NCCL版本 v2.25.1-1
数据传输概要
p2pTransport
P2P数据传输需要launct kernel,由kernel来真正完成数据传递,以zero copy方式为例,其调用流程大致如下:
zero copy方式传递数据
Zero copy方式传递数据有两类,一类是kernel运行中才知道对端buff B的地址,图解如下:
Rank GPU的kernel将buff B的地址放入待ptrExchange变量中
Recvpeer GPU的kernel读取ptrExchange,获得buff B的地址
Recvpeer GPU的kernel将buff A中数据写入buff B中
Recvpeer GPU的kernel更新rank GPU的tail
Rank GPU的kernel检测到tail更新后,更新Recvpeer GPU的head,表示已接收数据
上图中ptrExchange用于交换直接传输的接收buff地址。
另一类是kernel运行前就已经知道了对端buff B的地址,图解如下:
Recvpeer GPU的kernel将buff A中数据写入buff B中
Recvpeer GPU的kernel更新rank GPU的tail
Rank GPU的kernel检测到tail更新后,更新Recvpeer GPU的head,表示已接收数据
注意,上述两种数据交互都有个前提,就是Recvpeer GPU能直接访问buff B,这也意味着buff B导入了Recvpeer GPU中。
copy方式传递数据
图解如下:
Recvpeer GPU的kernel将buff A中数据写入对端PROTO_LL BUFF中
Recvpeer GPU的kernel更新rank GPU的tail
Rank GPU的kernel检测到tail更新后,将数据拷贝到buff B中
Rank GPU的kernel更新Recvpeer GPU的head
上述交换完成一小块数据传递。
上述交换数据使用的PROTO_LL协议,对于PROTO_SIMPLE协议,发送端将数据填入PROTO_SIMPLE buff中,由对端主动读走,图解如下:
Recvpeer GPU的kernel将buff A中数据写入本端PROTO_SIMPLE BUFF中
Recvpeer GPU的kernel更新rank GPU的tail
Rank GPU的kernel检测到tail更新后,主动读取Recvpeer GPU的PROTO_SIMPLE BUFF中数据到buff B
Rank GPU的kernel更新Recvpeer GPU的head
上述交换完成一小块数据传递。
cudaMemcopy方式(CE模式)
非常有意思的一种数据传递方式,即NCCL检测到用户是使能了NCCL_P2P_USE_CUDA_MEMCPY,将强制使用Copy Engine来搬运数据,图解如下所示。
GP0作为send端,kernel将用户数据写入发送buff中;
GP0作为send端,更新host内存中粉色tail;
host代理线程sendProxyProgress检测到红色tail更新;
host代理线程调用cudaMemcpy接口将GPU0上发送buff数据拷贝到GPU1的接收缓存区buff中(这一步会下发DMA命令给GPU0的Copy Engine);
host代理线程更新绿色tail;
GPU1的kernel检测到绿色tail更新;
GPU1的kernel将接收缓存区中数据拷贝到用户buff中;
GPU1的kernel更新host内存中head;
GPU0的kernel检测到host内存中head更新,知道数据已被消耗,如有必要可进行下一笔传输。
由于CE命令需要host下发,因此传输需要在send端创建代理线程。
netTransport
下面以rank0接收数据(recv端),rank1发送数据(send端),使用NCCL_PROTO_SIMPLE传输协议为例。
copy方式传递
第一阶段:信息交换
在transport建立时候,两端会交换控制信息和QP信息,其中QP信息用于两端QP绑定建链,而控制信息则用于协调数据传输。对于recv端来说,就是将sizeFIFO的虚拟地址va2和rkey2传给send端。对于send端来说,就是将FIFO的虚拟地址va1和rkey1传给recv端。
第二阶段:launch kernel及数据接收BUFF信息获取
rank0作为recv端,将本端transport建立时申请的接收缓存区一部分区域信息(va4、rkey4)写到send端FIFO中,这样send端就知道数据应该传输到recv端哪个位置,其实这个操作就是send端获得发送的credit,rank1通过对比预期idx就可以知道有可用的credit。如下图源码:
初始状态rank1的fifo head为0,接收到idx为1的条目时,认为有可用credit。rank0第一个写过来的idx便是1。
这个阶段两侧都已经launch kernel了,kernel也拿到了用户数据buff地址。
第三阶段:数据传输
数据传输又分为两种情况:单笔传输和多笔传输。
单笔传输
对于单笔数据传输,使用的是RDMA write with immediate命令,其中immediate就是单笔传输的数据长度,如下图所示。
send端kernel将用户数据拷贝到发送缓冲区;
send端kernel更新rank1的tail变量;
send端sendProxyProgress线程检查到tail更新;
send端sendProxyProgress线程下发RDMA write with immediate命令将数据写入recv端接收缓存区中;
recv端RDMA网卡上送CQE告诉rank0有数据接收到,这个CQE携带的immediate就是单笔传输数据长度。
上面第三步对应代码如下,不仅检查tail,还会检查size是否写入。
对于单笔传输来说,一笔RDMA write with immediate不仅将数据写入,也通知了rank0有数据到达,还顺带告诉了rank0这笔传输数据长度是多少,一举多得。
多笔传输
send端kernel将用户数据拷贝到发送缓冲区;
send端kernel更新rank1的tail变量;
send端sendProxyProgress线程检查到tail更新;
send端sendProxyProgress线程下发RDMA write命令将数据写入recv端接收缓存区中;
send端sendProxyProgress线程下发RDMA write with immediate命令将每笔数据长度size写入recv端sizeFIFO中;
RDMA网卡上送CQE告诉rank0有数据接收到,这个CQE携带的immediate数据是无效的。
对于多笔传输来说,需要将每笔传输的长度size信息也告诉recv端,所以一个immediate(32bit)无法承载这些信息,需要额外一笔write告诉recv。
第四阶段:确保数据落盘(写入HBM)
send端(rank1)这轮数据传输完成后会更新rank1的head变量;
send端(rank1)GPU中kernel检查到head更新就知道有新的空闲buff可用,如果需要可以进行下一轮数据传输;
recv端(rank0)感知到数据已经到达(通过第三阶段产生的CQE),然后下发本地GPU的read命令读取最后一笔数据传输的地址头部一个byte,这样可以确保数据强制刷新到HBM中;
recv端(rank0)的RDMA网卡产生CQE告诉recv端刷新动作已完成。
第三步对应代码如下,其中last表示最后一步数据传输(如果是多笔传输则是倒数第二笔)。
原因:当通过gdr将数据写到GPU的PCIE端后,其实并不保证数据已经落在HBM上了,如果不在HBM上,GPU的SM是不能访问数据的,所以为了确保GDR写入的数据落HBM,会进行iflush,也就是将数据刷入HBM。这里有两种实现,一种是发起rdma本地读,通过gdr再读一个写入数据的片段,由于物理硬件设计的原因,该操作可以强制数据入hbm,然后再被读出来;另一种是通过gdrcopy,在cpu侧读取该数据的一个片段,也可以达到flush的效果。iflush是nccl目前唯一会使用rdma read的场景。
在两端建链时候,recv端会创建一个QP跟这个QP本身建链,这个QP就是用于做iflush的。
第五阶段:recv端消费数据
recv端(rank0)的recvProxyProgress线程感知到刷新动作完成(通过第四阶段产生的CQE),然后更新rank0上tail变量;
recv端(rank0)的GPU kernel检测到tail更新;
recv端(rank0)的GPU kernel将数据拷贝到用户buff地址中;
recv端(rank0)的GPU kernel更新rank0上head变量,表示数据已消费。
通过五个阶段完成了一轮数据传输,如果需要,可进行下一轮数据传输。
zero copy方式传递数据
Zero copy方式与copy方式主要差异在第二阶段和第三阶段。
第二阶段:数据接收BUFF信息获取及launch kernel
主要差异是两端都会将用户buff注册MR。recv端会将用户buff注册mr产生的rkey6和用户buff va6等信息告诉send端,这样send端可以直接将数据写入recv端用户buff。
第三阶段:数据传输
单笔传输
send端直接将数据写入recv端用户buff中。
多笔传输
send端直接将数据写入recv端用户buff中。
注意,从代码上看,大多数都采用copy方式,少数满足特定条件的可以使用zero copy方式,这里特定条件还未弄清楚,算是留个疑问点待后续查明。
shmTransport
正常方式
正常情况下,数据buff处于host内存中,由GPU直接操作交换数据,如下所示。
send端GPU kernel将数据搬运到recv端创建的共享数据buff中;
send端GPU kernel更新recv端创建的tail变量;
recv端GPU检测到tail变量更新;
recv端GPU从共享数据buff中搬运数据到用户buff;
recv端GPU kernel更新send端创建的head变量;
send端GPU kernel检测到head更新,表示一轮数据传输完成。
useMemcpy方式
在使用Memcpy的方式下,需要在host起代理线程协助数据传输。
send端GPU kernel将数据从用户buff拷贝到发送缓存buff;
send端GPU kernel更新tail(蓝色);
send端的sendProxyProgress线程检测到tail更新;
send端的sendProxyProgress线程执行cudaMemcpy(cudaMemcpyDeviceToHost)将数据从发送buff拷贝到recv端创建的数据共享内存中;
send端的sendProxyProgress线程更新tail(黄色);
recv端的recvProxyProgress线程检测到tail更新;
recv端的recvProxyProgress线程执行cudaMemcpy(cudaMemcpyHostToDevice)将数据从数据共享buff中拷贝到recv缓冲区中;
recv端的recvProxyProgress线程更新tail(绿色);
recv端GPU kernel检测到tail更新;
recv端GPU kernel将数据从recv缓存区拷贝到用户buff中;
recv端GPU kernel更新head(红色);
send端GPU kernel检测到head更新,表示一轮数据传输完成。
p2p与net传输差异
对于P2P来说,其实就是两个GPU提前知道对方接受缓存区地址和发送缓存区地址,分别导入到自己地址空间中,这样就能通过P2P直接访问对端内存,再通过head、tail等变量协调数据收发过程。
对于net来说,其实是两个代理线程交换了控制信息buff,然后通过控制信息buff协调传输数据地址信息以及实现同步功能。两个代理线程与各自rank上kernel协调产生数据或消费数据。最终实现数据夸node传输。
简图如下所示。
参考文章:
https://blog.csdn.net/shanleo1986/article/details/137777787?spm=1001.2014.3001.5502
https://blog.csdn.net/u014443578/article/details/136052751
https://blog.csdn.net/kidgin7439/category_11998768.html
