当前位置: 首页 > news >正文

【NCCL】transport数据传输(二)

参考NCCL版本 v2.25.1-1

数据传输概要

p2pTransport

P2P数据传输需要launct kernel,由kernel来真正完成数据传递,以zero copy方式为例,其调用流程大致如下:

zero copy方式传递数据

Zero copy方式传递数据有两类,一类是kernel运行中才知道对端buff B的地址,图解如下:

  1. Rank GPU的kernel将buff B的地址放入待ptrExchange变量中

  2. Recvpeer GPU的kernel读取ptrExchange,获得buff B的地址

  3. Recvpeer GPU的kernel将buff A中数据写入buff B中

  4. Recvpeer GPU的kernel更新rank GPU的tail

  5. Rank GPU的kernel检测到tail更新后,更新Recvpeer GPU的head,表示已接收数据

上图中ptrExchange用于交换直接传输的接收buff地址。

另一类是kernel运行前就已经知道了对端buff B的地址,图解如下:

  1. Recvpeer GPU的kernel将buff A中数据写入buff B中

  2. Recvpeer GPU的kernel更新rank GPU的tail

  3. Rank GPU的kernel检测到tail更新后,更新Recvpeer GPU的head,表示已接收数据

注意,上述两种数据交互都有个前提,就是Recvpeer GPU能直接访问buff B,这也意味着buff B导入了Recvpeer GPU中。

copy方式传递数据

图解如下:

  1. Recvpeer GPU的kernel将buff A中数据写入对端PROTO_LL BUFF中

  2. Recvpeer GPU的kernel更新rank GPU的tail

  3. Rank GPU的kernel检测到tail更新后,将数据拷贝到buff B中

  4. Rank GPU的kernel更新Recvpeer GPU的head

上述交换完成一小块数据传递。

上述交换数据使用的PROTO_LL协议,对于PROTO_SIMPLE协议,发送端将数据填入PROTO_SIMPLE buff中,由对端主动读走,图解如下:

  1. Recvpeer GPU的kernel将buff A中数据写入本端PROTO_SIMPLE BUFF中

  2. Recvpeer GPU的kernel更新rank GPU的tail

  3. Rank GPU的kernel检测到tail更新后,主动读取Recvpeer GPU的PROTO_SIMPLE BUFF中数据到buff B

  4. Rank GPU的kernel更新Recvpeer GPU的head

上述交换完成一小块数据传递。

cudaMemcopy方式(CE模式)

非常有意思的一种数据传递方式,即NCCL检测到用户是使能了NCCL_P2P_USE_CUDA_MEMCPY,将强制使用Copy Engine来搬运数据,图解如下所示。

  1. GP0作为send端,kernel将用户数据写入发送buff中;

  2. GP0作为send端,更新host内存中粉色tail;

  3. host代理线程sendProxyProgress检测到红色tail更新;

  4. host代理线程调用cudaMemcpy接口将GPU0上发送buff数据拷贝到GPU1的接收缓存区buff中(这一步会下发DMA命令给GPU0的Copy Engine);

  5. host代理线程更新绿色tail;

  6. GPU1的kernel检测到绿色tail更新;

  7. GPU1的kernel将接收缓存区中数据拷贝到用户buff中;

  8. GPU1的kernel更新host内存中head;

  9. 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就是单笔传输的数据长度,如下图所示。

  1. send端kernel将用户数据拷贝到发送缓冲区;

  2. send端kernel更新rank1的tail变量;

  3. send端sendProxyProgress线程检查到tail更新;

  4. send端sendProxyProgress线程下发RDMA write with immediate命令将数据写入recv端接收缓存区中;

  5. recv端RDMA网卡上送CQE告诉rank0有数据接收到,这个CQE携带的immediate就是单笔传输数据长度。

上面第三步对应代码如下,不仅检查tail,还会检查size是否写入。

对于单笔传输来说,一笔RDMA write with immediate不仅将数据写入,也通知了rank0有数据到达,还顺带告诉了rank0这笔传输数据长度是多少,一举多得。

多笔传输

  1. send端kernel将用户数据拷贝到发送缓冲区;

  2. send端kernel更新rank1的tail变量;

  3. send端sendProxyProgress线程检查到tail更新;

  4. send端sendProxyProgress线程下发RDMA write命令将数据写入recv端接收缓存区中;

  5. send端sendProxyProgress线程下发RDMA write with immediate命令将每笔数据长度size写入recv端sizeFIFO中;

  6. RDMA网卡上送CQE告诉rank0有数据接收到,这个CQE携带的immediate数据是无效的。

对于多笔传输来说,需要将每笔传输的长度size信息也告诉recv端,所以一个immediate(32bit)无法承载这些信息,需要额外一笔write告诉recv。

第四阶段:确保数据落盘(写入HBM)

  1. send端(rank1)这轮数据传输完成后会更新rank1的head变量;

  2. send端(rank1)GPU中kernel检查到head更新就知道有新的空闲buff可用,如果需要可以进行下一轮数据传输;

  3. recv端(rank0)感知到数据已经到达(通过第三阶段产生的CQE),然后下发本地GPU的read命令读取最后一笔数据传输的地址头部一个byte,这样可以确保数据强制刷新到HBM中;

  4. 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端消费数据

  1. recv端(rank0)的recvProxyProgress线程感知到刷新动作完成(通过第四阶段产生的CQE),然后更新rank0上tail变量;

  2. recv端(rank0)的GPU kernel检测到tail更新;

  3. recv端(rank0)的GPU kernel将数据拷贝到用户buff地址中;

  4. 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直接操作交换数据,如下所示。

  1. send端GPU kernel将数据搬运到recv端创建的共享数据buff中;

  2. send端GPU kernel更新recv端创建的tail变量;

  3. recv端GPU检测到tail变量更新;

  4. recv端GPU从共享数据buff中搬运数据到用户buff;

  5. recv端GPU kernel更新send端创建的head变量;

  6. send端GPU kernel检测到head更新,表示一轮数据传输完成。

useMemcpy方式

在使用Memcpy的方式下,需要在host起代理线程协助数据传输。

  1. send端GPU kernel将数据从用户buff拷贝到发送缓存buff;

  2. send端GPU kernel更新tail(蓝色);

  3. send端的sendProxyProgress线程检测到tail更新;

  4. send端的sendProxyProgress线程执行cudaMemcpy(cudaMemcpyDeviceToHost)将数据从发送buff拷贝到recv端创建的数据共享内存中;

  5. send端的sendProxyProgress线程更新tail(黄色);

  6. recv端的recvProxyProgress线程检测到tail更新;

  7. recv端的recvProxyProgress线程执行cudaMemcpy(cudaMemcpyHostToDevice)将数据从数据共享buff中拷贝到recv缓冲区中;

  8. recv端的recvProxyProgress线程更新tail(绿色);

  9. recv端GPU kernel检测到tail更新;

  10. recv端GPU kernel将数据从recv缓存区拷贝到用户buff中;

  11. recv端GPU kernel更新head(红色);

  12. 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

http://www.cnnetsun.cn/news/2632404.html

相关文章:

  • MLIR与CGRA编译优化技术解析
  • Cloudflare AI Labyrinth:用数字迷宫反制AI爬虫,保护原创内容
  • ELK日志平台实战
  • 告别手动操作:用Python脚本批量调用SAP BAPI,自动化FICO凭证与MM物料创建
  • 搞定7nm DRC收敛:一份来自Innovus和ICC2实战的避坑清单(附脚本)
  • 多软件互通避坑:模型互导不碎面、不丢材质
  • 智能戒指技术解析:从多模态传感到开源生态
  • AI与机器学习驱动的智能运营:从数据到决策的自动化闭环
  • Claude Code + GLM-5 深度赋能测试:开发 8 大 Skill 构建 AI 测试助手集群
  • 自动语音识别技术原理与实战:从MFCC到端到端模型
  • 神仙免费云服务器 - 阿贝云
  • GEO(生成式引擎优化)完全指南:让你的技术内容被AI看见
  • AI搜索优化值不值?价格与效果真实解析
  • 软件设计师备考 第0章 题型分布、示例、学习路线
  • 为什么92%的Gemini正则失败源于上下文锚定错误?——6个生产环境真实Case逆向拆解
  • iPaaS集成平台选型参考:五款热门产品能力介绍
  • FPGA如何精准控制三片ADS1282同步采样?SPI时序与同步逻辑的保姆级解析
  • 聊天机器人数据分析实战:从黑盒到白盒的优化闭环
  • Linux dd命令实战:手把手教你用/dev/zero和seek参数精准擦除eMMC分区
  • 从CTF实战看LFSR与BM算法:如何破解流密码与伪随机生成器
  • Windows 10/11系统下,用YOLOv8改进YOLOv5的C3模块:一份给CV新手的保姆级数据集训练指南
  • 告别同步烦恼:手把手教你用AD9680+LMK04828搭建多板卡JESD204B采集系统(附Vivado调试技巧)
  • 你的STM32循迹小车跑不直?可能是编码器测速的‘坑’没避开
  • 保姆级教程:用CarSim 2020和Simulink手把手搭建平行泊车仿真(附MPC控制器模型)
  • Cadence Allegro铺铜实战:从动态避让到静态优化,手把手教你高效处理PCB电源层
  • 终极热键侦探:3分钟快速定位Windows快捷键占用程序
  • AI系统审计:如何识别数据投毒与对抗性攻击的微观威胁
  • 房地产AI应用:从自动化到价值创造的务实路径与案例解析
  • 单片机RTC实验
  • 从VOC到YOLO v5/v8:手把手教你构建标准目标检测数据集(含数据划分脚本)