3. RDMA基本元素

本文旨在讲解RDMA中最基本的元素及其含义。介绍了WQ、WQE、QP、SQ、RQ、SRQ、CQ、CQE等概念,阐述了它们在RDMA通信中的作用和关系,还说明了WR和WC是WQE和CQE在用户层的映射,最后总结了用户态与硬件间的任务流转过程。

摘要生成于 C知道 ,由 DeepSeek-R1 满血版支持, 前往体验 >

        RDMA技术中经常使用缩略语,很容易让刚接触的人一头雾水,本篇的目的是讲解RDMA中最基本的元素及其含义。

        我将常见的缩略语对照表写在前面,阅读的时候如果忘记了可以翻到前面查阅。

WQ

        Work Queue简称WQ,是RDMA技术中最重要的概念之一。WQ是一个储存工作请求的队列,为了讲清楚WQ是什么,我们先介绍这个队列中的元素WQE(Work Queue Element,工作队列元素)

WQE

        WQE可以认为是一种“任务说明”,这个工作请求是软件下发给硬件的,这份说明中包含了软件所希望硬件去做的任务以及有关这个任务的详细信息。比如,某一份任务是这样的:“我想把位于地址0x12345678的长度为10字节的数据发送给对面的节点”,硬件接到任务之后,就会通过DMA去内存中取数据,组装数据包,然后发送。

        WQE的含义应该比较明确了,那么我们最开始提到的WQ是什么呢?它就是用来存放“任务书”的“文件夹”,WQ里面可以容纳很多WQE。有数据结构基础的读者应该都了解,队列是一种先进先出的数据结构,在计算机系统中非常常见,我们可以用下图表示上文中描述的WQ和WQE的关系:

        WQ这个队列总是由软件向其中增加WQE(入队),硬件从中取出WQE,这就是软件给硬件“下发任务”的过程。为什么用队列而不是栈?因为进行“存”和“取“操作的分别是软件和硬件,并且需要保证用户的请求按照顺序被处理,在RDMA技术中,所有的通信请求都要按照上图这种方式告知硬件,这种方式常被称为“Post”

       

QP

        Queue Pair简称QP,就是“一对”WQ的意思。

SQ和RQ

        任何通信过程都要有收发两端,QP就是一个发送工作队列和一个接受工作队列的组合,这两个队列分别称为SQ(Send Queue)和RQ(Receive Queue)。我们再把上面的图丰富一下,左边是发送端,右边是接收端:

        WQ怎么不见了?SQ和RQ都是WQ,WQ只是表示一种可以存储WQE的单元,SQ和RQ才是实例。

        SQ专门用来存放发送任务,RQ专门用来存放接收任务。在一次SEND-RECV流程中,发送端需要把表示一次发送任务的WQE放到SQ里面。同样的,接收端软件需要给硬件下发一个表示接收任务的WQE,这样硬件才知道收到数据之后放到内存中的哪个位置。上文我们提到的Post操作,对于SQ来说称为Post Send,对于RQ来说称为Post Receive

        需要注意的是,在RDMA技术中通信的基本单元是QP,而不是节点。如下图所示,对于每个节点来说,每个进程都可以使用若干个QP,而每个本地QP可以“关联”一个远端的QP。我们用“节点A给节点B发送数据”并不足以完整的描述一次RDMA通信,而应该是类似于“节点A上的QP3给节点C上的QP4发送数据”。

        每个节点的每个QP都有一个唯一的编号,称为QPN(Queue Pair Number),通过QPN可以唯一确定一个节点上的QP。

SRQ

        Shared Receive Queue简称SRQ,意为共享接收队列。概念很好理解,就是一种几个QP共享同一个RQ时,我们称其为SRQ。以后我们会了解到,使用RQ的情况要远远小于使用SQ,而每个队列都是要消耗内存资源的。当我们需要使用大量的QP时,可以通过SRQ来节省内存。如下图所示,QP2~QP4一起使用同一个RQ:

CQ

        Completion Queue简称CQ,意为完成队列。跟WQ一样,我们先介绍CQ这个队列当中的元素——CQE(Completion Queue Element)。可以认为CQE跟WQE是相反的概念,如果WQE是软件下发给硬件的“任务书”的话,那么CQE就是硬件完成任务之后返回给软件的“任务报告”。CQE中描述了某个任务是被正确无误的执行,还是遇到了错误,如果遇到了错误,那么错误的原因是什么。

        而CQ就是承载CQE的容器——一个先进先出的队列。我们把表示WQ和WQE关系的图倒过来画,就得到了CQ和CQE的关系:

        每个CQE都包含某个WQE的完成信息,他们的关系如下图所示:

        下面我们把CQ和WQ(QP)放在一起,看一下一次SEND-RECV操作中,软硬件的互动(图中序号顺序不表示实际时序):

2022/5/23:下图及后面的列表顺序有修改,将原来第2条的“接收端硬件从RQ中拿到任务书,准备接收数据”移动到“接收端收到数据,进行校验后回复ACK报文给发送端”之后,并且修改了描述,现在为第6条。
这里我犯了错误的点是RQ和SQ不同,是一个“被动接收”的过程,只有收到Send报文(或者带立即数的Write报文)时硬件才会消耗RQ WQE。感谢 

@连接改变世界

 的指正。

  1. 接收端APP以WQE的形式下发一次RECV任务到RQ。
  2. 发送端APP以WQE的形式下发一次SEND任务到SQ。
  3. 发送端硬件从SQ中拿到任务书,从内存中拿到待发送数据,组装数据包。
  4. 发送端网卡将数据包通过物理链路发送给接收端网卡。
  5. 接收端收到数据,进行校验后回复ACK报文给发送端。
  6. 接收端硬件从RQ中取出一个任务书(WQE)。
  7. 接收端硬件将数据放到WQE中指定的位置,然后生成“任务报告”CQE,放置到CQ中。
  8. 接收端APP取得任务完成信息。
  9. 发送端网卡收到ACK后,生成CQE,放置到CQ中。
  10. 发送端APP取得任务完成信息。

NOTE: 需要注意的一点是,上图中的例子是可靠服务类型的交互流程,如果是不可靠服务,那么不会有步骤6的ACK回复,而且步骤9以及之后的步骤会在步骤5之后立即触发。关于服务类型以及可靠与不可靠,我们将在《RDMA基本服务类型》一文中讲解。

至此,通过WQ和CQ这两种媒介,两端软硬件共同完成了一次收发过程。

WR和WC

        说完了几个Queue之后,其实还有两个文章开头提到的概念没有解释,那就是WR和WC(不是Water Closet的缩写)。

        WR全称为Work Request,意为工作请求;WC全称Work Completion,意为工作完成。这两者其实是WQE和CQE在用户层的“映射”。因为APP是通过调用协议栈接口来完成RDMA通信的,WQE和CQE本身并不对用户可见,是驱动中的概念。用户真正通过API下发的是WR,收到的是WC

        WR/WC和WQE/CQE是相同的概念在不同层次的实体,他们都是“任务书”和“任务报告”。于是我们把前文的两个图又加了点内容:

 

总结

        好了,我们用IB协议[1]3.2.1中的Figure 11这张图总结一下本篇文章的内容:

        用户态的WR,由驱动转化成了WQE填写到了WQ中,WQ可以是负责发送的SQ,也可以是负责接收的RQ。硬件会从各个WQ中取出WQE,并根据WQE中的要求完成发送或者接收任务。任务完成后,会给这个任务生成一个CQE填写到CQ中。驱动会从CQ中取出CQE,并转换成WC返回给用户。

基础概念就介绍到这里,下一篇将介绍RDMA的几种常见操作类型。

 

<think>嗯,我现在需要了解NCCL的Reduce通信的执行步骤,并且要深入到源码层面。首先,我应该回忆一下NCCL是什么。NCCL是Nvidia Collective Communication Library,主要用于多GPU之间的高效通信,比如在深度学习训练中进行AllReduce、Broadcast等操作。Reduce操作是其中一种,它可能将多个GPU上的数据通过某种操作(比如求和)汇总到一个GPU上。 接下来,我需要理解Reduce的基本概念。比如,在MPI中,Reduce操作会将所有进程的数据通过某种操作合并到根进程中。NCCL的Reduce可能类似,但专门针对GPU优化。那么,NCCL的Reduce具体是如何实现的呢?执行步骤是怎样的? 首先,可能需要考虑数据在GPU间的传输方式。比如,使用PCIe或者NVLink这样的高速互联。然后,NCCL可能利用这些硬件特性来优化数据传输。但具体的执行步骤可能涉及到多个阶段,比如内存准备、通信协议的确定、数据传输、同步等。 接下来,我需要深入到源码层面。NCCL的源码是开源的,可以在GitHub上找到。比如,在源码中,Reduce相关的函数可能位于某个文件中,比如nccl/src/collectives/目录下可能有reduce的实现。例如,ncclReduce函数可能负责处理这个操作。 执行步骤可能包括:确定参与通信的GPU(即通信域),确定根节点,每个GPU准备发送缓冲区,根节点准备接收缓冲区。然后,根据不同的算法(比如环状算法、树状算法)来组织数据传输。比如,树状算法可能将数据逐层汇总到根节点,而环状算法可能通过循环的方式传递数据,逐步累积结果。 在源码中,可能看到对不同算法的选择,比如在某个函数中根据条件选择使用Ring还是Tree算法。例如,在ncclReduce函数中,可能会调用ncclTree或ncclRing相关的函数。每个算法都有不同的步骤,比如Ring算法可能需要将数据分成多个块,依次传递,每次传递时进行本地Reduce操作,直到所有数据被处理完。 另外,数据的分段传输可能也是一个优化点,比如将大的数据分成多个块并行传输,减少单次传输的开销。在源码中,可能会有循环处理每个数据块的代码,调用CUDA核函数来执行实际的Reduce操作(比如加法)。 还需要考虑同步机制,确保所有GPU在正确的时间点进行数据传输和计算。比如,使用CUDA事件或者流来同步各个GPU的操作,避免数据竞争或未完成的操作。 举个例子,假设使用Ring算法进行Reduce操作。步骤可能是这样的:每个GPU将数据分成若干块,然后依次向右的GPU发送一个块,同时接收左边GPU的块。每次接收到数据后,执行本地的Reduce操作(比如累加),然后将结果继续传递。经过若干次循环后,所有数据块都会被汇总到根GPU上。 在源码中,可能会看到类似这样的逻辑:确定环的拓扑结构,为每个GPU分配左右邻居,然后循环发送和接收数据块,调用CUDA核函数进行Reduce。每个数据块的处理可能通过一个循环,遍历所有参与Reduce的GPU,逐步累积结果。 此外,NCCL可能针对不同的硬件架构进行优化,比如针对Volta、Ampere架构的不同特性调整传输策略。例如,使用GPU的RDMA能力来加速数据传输,或者利用共享内存减少通信延迟。 在代码层面,可能需要查看具体的设备代码,比如在src/devices/目录下,不同GPU架构的实现可能有不同的优化。例如,ncclPrimitives结构体可能封装了发送和接收数据的函数,而针对Reduce操作可能会有专门的Kernel函数,使用CUDA C编写,高效地在GPU上执行数据合并。 总结起来,NCCL的Reduce执行步骤可能包括:初始化通信域和拓扑结构,选择最优算法(如Ring或Tree),分段数据传输,执行本地Reduce操作,同步等待所有传输完成。源码中的实现会涉及多个层次,从顶层的API调用到底层的CUDA核函数,以及不同硬件架构的优化代码。 不过,我可能需要具体查看源码来验证这些推测。例如,查找ncclReduce函数的实现,看看它调用了哪些内部函数,如何选择算法,以及数据传输的具体流程。此外,了解NCCL中的数据结构,如ncclComm结构体,可能存储了通信域的详细信息,如rank、节点数、拓扑结构等,这些信息在Reduce操作中会被使用到。 可能还需要注意,NCCL中的collectives操作通常由多个步骤组成,比如在Reduce中,可能需要将数据从各个rank发送到root,这可能涉及到多个步骤的协调和数据传输。在源码中,可能会有状态机或者类似的机制来管理这些步骤的执行顺序和同步。 最后,考虑到NCCL的高性能设计,可能还有诸如流水线传输、带宽优化、延迟隐藏等技术应用在Reduce操作中,这些都可能体现在源码的具体实现细节里。</think>NCCL(NVIDIA Collective Communication Library)中的 **Reduce** 操作用于将多个 GPU 上的数据按指定操作(如求和)汇总到目标 GPU(根节点)。其实现结合了硬件特性(如 NVLink、RDMA)和高效算法(如 Ring、Tree)。以下是其执行步骤及源码层面的关键分析: --- ### **1. 初始化通信域** - **通信域(ncclComm)**:存储 GPU 拓扑、Rank 信息、算法选择等。 - **源码关键结构**:`ncclComm`(定义于 `include/nccl.h`),包含成员如 `rank`(当前 GPU 的 ID)、`nRanks`(总 GPU 数)、`bootstrap`(通信初始化信息)等。 --- ### **2. 算法选择** NCCL 根据数据大小、GPU 拓扑动态选择算法: - **Ring 算法**:适合中等数据,通过环形通信逐段聚合。 - **Tree 算法**:适合大数据,分层聚合减少跳数。 **源码示例**(`src/collectives/collectives.cc`): ```cpp ncclResult_t ncclReduce(...) { // 选择算法 if (algorithm == NCCL_ALGO_AUTO) { if (bytes < 256 * 1024) algorithm = NCCL_ALGO_RING; else algorithm = NCCL_ALGO_TREE; } // 调用具体算法实现 if (algorithm == NCCL_ALGO_RING) { return ncclReduceRing(...); } else { return ncclReduceTree(...); } } ``` --- ### **3. Ring 算法实现** **步骤**: 1. **分块**:将数据分为 `chunkSize` 的块。 2. **环形传输**:每个 GPU 向右邻居发送数据块,同时接收左邻居数据块。 3. **本地 Reduce**:接收数据后执行本地操作(如累加)。 4. **循环迭代**:经过 `nRanks-1` 次迭代后,所有数据聚合到根节点。 **源码关键函数**: - `ncclReduceRing`(`src/collectives/ring.cc`): ```cpp ncclResult_t ncclReduceRing(...) { for (int step=0; step < nsteps; step++) { // 计算当前块的发送/接收地址 char* sendBuf = sendbuff + step * chunkSize; char* recvBuf = recvbuff + step * chunkSize; // 异步发送/接收 NCCL_CHECK(ncclTransportSend(...)); NCCL_CHECK(ncclTransportRecv(...)); // 本地 Reduce(CUDA 核函数) reduceKernel<<<blocks, threads>>>(recvBuf, sendBuf, count, op); } } ``` --- ### **4. Tree 算法实现** **步骤**: 1. **分层聚合**:构建二叉树,叶子节点向上层节点发送数据。 2. **递归 Reduce**:中间节点聚合子节点数据,直至根节点。 3. **带宽优化**:减少网络跳数,适合大规模数据。 **源码关键函数**: - `ncclReduceTree`(`src/collectives/tree.cc`): ```cpp ncclResult_t ncclReduceTree(...) { int parent = (rank - 1) / 2; // 父节点 int child1 = 2*rank + 1; // 子节点1 int child2 = 2*rank + 2; // 子节点2 // 接收子节点数据并聚合 if (child1 < nRanks) NCCL_CHECK(recvReduce(child1, ...)); if (child2 < nRanks) NCCL_CHECK(recvReduce(child2, ...)); // 发送到父节点 if (rank != root) NCCL_CHECK(send(parent, ...)); } ``` --- ### **5. CUDA 核函数优化** - **数据并行处理**:每个 CUDA 线程处理多个元素,隐藏内存延迟。 - **向量化访问**:通过 `uint4` 类型(128 位)合并内存操作。 **核函数示例**(`src/collectives/device/common_kernel.h`): ```cpp template<typename T, int UNROLL> __device__ void reduceLoop(...) { #pragma unroll for (int i=0; i < UNROLL; i++) { // 执行操作(如加法) if (op == ncclSum) res[i] += vals[i]; } } ``` --- ### **6. 硬件加速** - **RDMA 传输**:通过 GPU Direct 直接访问远程 GPU 内存。 - **NVLink 拓扑**:自动检测高速链路,优化数据传输路径。 **源码示例**(`src/devices/net.cc`): ```cpp ncclResult_t ncclTransportSend(...) { // 使用 RDMA 或 IPC 通道 if (use_rdma) { NCCL_CHECK(rdmaSend(...)); } else { NCCL_CHECK(ipcSend(...)); } } ``` --- ### **总结** NCCL 的 Reduce 操作通过算法优化(Ring/Tree)、硬件加速(NVLink/RDMA)和 CUDA 核函数高效执行数据聚合。其源码实现涵盖通信域管理、算法选择、数据传输和本地计算,充分适配现代多 GPU 系统的特性。
评论 1
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

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

抵扣说明:

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

余额充值