RDMA操作类型(三)

Send操作

引用一下,IB协议第9.4.1章节原文:

The SEND Operation is sometimes referred to as a Push operation or as having channel semantics. Both terms refer to how the SW client of the
transport service views the movement of data. With a SEND operation the initiator of the data transfer pushes data to the remote QP. The initiator
doesn’t know where the data is going on the remote node.

Send操作有时被称为“PUSH”操作或具有通道语义。这两个术语都指传输服务的软件客户端(SW client)看数据移动的视角。通过Send操作,数据传输的发起者将数据推送到远端QP。发起程序不知道数据在远程节点上的去向。

也就是说,Send操作就是发送端将数据扔给接收端,至于接收端将数据存放到哪,是接收端来决定的,发送端不需要知道。

那么,接收端又是怎么知道数据放哪呢?

The remote node’s Channel Adapter places the data into the next available receive buffer for that QP. On an HCA, the receive buffer is pointed to by the WQE at the head of the QP’s receive queue.

远程节点的通道适配器(HCA)将数据放入该QP的下一个可用接收缓冲区,接收缓冲区的位置由QP接收队列头的WQE(RWQE)指向。

接收端在接收到对端Send过来的数据后,会从RQ中获取一个可用的RWQE,将接收数据存放到RWQE所描述的内存区域中,这个RWQE是由上层app用户下发的,也就是说,数据存放区域是由接收端上次app用户指定的。

用一张图来描述一下Send操作。
在这里插入图片描述

①接收端app下发RWQE指定下一个接受数据存放的内存区域
②发送端app下发WQE指定需要发送的数据所在内存
③发送端HCA从SQ中将WQE取下解析
④发送端根据WQE指向地址获取数据组包
⑤发送端发送数据到接收端
⑥接收端接收到数据,然后从RQ中取下RWQE解析
⑦接收端将数据写入到RWQE指向的内存中

由上可以知道send操作需要两边都参与,所以send操作属于双边操作。

RDMA write操作

RDMA write操作是单边操作,数据传输过程中,不需要对端CPU参与,相比于send双边操作,单边操作能极大降低对端CPU在数据传输中的负载,将对端CPU从数据传输中解放出来干其他事,这也是IB协议的初衷。
在这里插入图片描述

①在本端执行write操作前,对端需要将对端的内存地址(目的地址)以及密钥告诉本端,至于对端怎么告诉本端,IB协议未做规定,可以通过send操作,也可以TCP或其他传输协议。
②本端将数据源地址、目的地址、密钥填入WQE中下发。
③本端HCA从SQ中取下WQE解析。
④本端HCA根据WQE中源地址拿到数据。
⑤本端HCA将数据组包发送到对端。
⑥对端HCA将数据接收,然后校验合法性无误后将数据写入内存。

为了保证write操作写入的安全性,IB协议定义了PD、MR、Key等元素,并规定了write操作前的准备动作以及对端接收到write数据时的校验动作,详细内容将在后续文章中慢慢展开。

RDMA read操作

RDMA read操作与RDMA write操作类似。
在这里插入图片描述

①在本端执行read操作前,对端需要将对端的内存地址(目的地址)以及密钥告诉本端,至于对端怎么告诉本端,IB协议未做规定,可以通过send操作,也可以TCP或其他传输协议。
②本端将数据源地址、目的地址、密钥填入WQE中下发。
③本端HCA从SQ中取下WQE解析。
④本端HCA将数据组包发送到对端。
⑤对端HCA接收read请求,校验无误后从内存中将数据读取下来。
⑥对端HCA将数据组包发送给本端。
⑦本端HCA接收到数据后放入WQE指定的目的内存中。

注意,read操作的数据源地址在对端内存上,目的地址为本端内存,这与write操作相反。

ATOMIC操作

引用一下,IB协议第9.4.5章节原文:

ATOMIC Operations execute a 64-bit operation at a specified address on a remote node. The operations atomically read, modify and write the destination address and guarantee that operations on this address by other QPs on the same CA do not occur between the read and the write.

原子操作在远程节点上对指定地址执行64位操作。这些操作以原子方式读取、修改并写入目标地址,并确保在读取和写入之间,同一通信适配器上的其他QP不会对该地址进行操作。

简单来说,就是本端指定对端某个地址发起Atomic操作请求给对端,对端在完成响应这个Atomic操作前,不能响应同一个地址的其他Atomic操作(原子性)。

Atomic操作的要求:
1)原子操作仅由可靠的连接和可靠的数据报传输服务支持。
2)强烈建议严格在硬件中提供原子操作支持。
3)原子命令请求包中的虚拟地址应自然地与8字节边界对齐。响应的CA会对此进行检查,如果请求不自然对齐,则返回一个无效的请求NAK。

IB协议定义的Atomic操作类型:

1)FetchAdd (Fetch and Add)

​ FetchAdd原子操作告诉响应端在响应端(对端)内存中自然对齐的虚拟地址读取64位内存值,使用AtomicETH中的64位Add数据字段执行无符号加法,并将结果(必须与请求方的内存类型匹配)写回相同的虚拟地址。

2)CmpSwap (Compare and Swap)

​ CmpSwap原子操作告诉响应端读取响应端内存中自然对齐的虚拟地址处的64位值,将其与Atomicheth头中的比较数据字段进行比较,如果相等,则将AtomicETH头中的交换数据字段写入同一虚拟地址。如果它们不相等,响应端内存的内容不会改变。

ATOMIC操作要求响应端(对端)处理Atomic操作的整个路径都具有原子性。对于HCA来说,其内部的原子性需要内部总线保证,比如AXI5。对于PCIe链路上来说,其原子性有PCIe协议来保证,即如下图。对于HCA与Host侧CPU访问统一内存的原子性,需要RC来保证。
在这里插入图片描述

Resync操作

只有可靠的数据报传输服务才支持RESYNC操作。RESYNC请求本质上与零长度可靠数据报只发送(Reliable Datagram Send-Only)请求相同,但有几个独特的属性:

1)请求端使用RESYNC强制接收端将其预期PSN重置为请求端定义的值;

2)RESYNC携带长度为零的数据负;

3)要求响应端接受重新同步请求,即使当前执行的请求尚未完成;

4)重新同步请求本身不会直接使用请求端的send WQE,也不会直接使用响应端的receive WQE(RWQE)。

Resync操作主要用于重置响应端(对端)预期接收的PSN,顾名思义,强制同步成指定的PSN。

  • 0
    点赞
  • 1
    收藏
    觉得还不错? 一键收藏
  • 2
    评论
MVAPICH2-GDR为GPU RDMA提供了一些接口,可以方便地进行GPU RDMA操作。以下是使用MVAPICH2-GDR进行GPU RDMA操作的基本步骤: 1. 初始化MVAPICH2-GDR库:开发人员需要调用相应的函数来初始化MVAPICH2-GDR库。 2. 创建GPU缓冲区:开发人员需要为GPU分配内存,并创建相应的GPU缓冲区。 3. 注册GPU缓冲区:开发人员需要将GPU缓冲区注册到RDMA适配器的内存区域中,以便进行RDMA操作。 ```c void *gpu_addr; cudaMalloc(&gpu_addr, size); register_gpu(gpu_addr, size); ``` 4. 创建RDMA请求:开发人员可以使用MVAPICH2-GDR提供的接口创建RDMA请求。 ```c MPIDI_CH3I_RDMA_put_gpu(void *buf, int len, int dest, uint64_t addr, int rkey, int tag, MPID_Comm *comm, int *req_complete); ``` 5. 执行GPU RDMA操作:开发人员可以使用MVAPICH2-GDR提供的接口进行GPU RDMA操作。 ```c MPIDI_CH3I_RDMA_post_gpu(int type, void *buf, int len, int dest, uint64_t addr, int rkey, int tag, MPID_Comm *comm, int *req_complete); ``` 6. 等待GPU RDMA操作完成:开发人员需要等待GPU RDMA操作完成后,才能继续进行下一步操作。 ```c MPIDI_CH3I_RDMA_wait(void *request); ``` 7. 释放GPU缓冲区:当GPU RDMA操作完成后,开发人员需要释放相应的GPU缓冲区。 ```c unregister_gpu(gpu_addr); cudaFree(gpu_addr); ``` 需要注意的是,MVAPICH2-GDR提供的GPU RDMA接口是比较底层的接口,需要开发人员具备一定的RDMA编程和GPU编程经验。同时,开发人员需要考虑到一些问题,例如数据同步、内存管理等,以实现高效的GPU RDMA操作
评论 2
添加红包

请填写红包祝福语或标题

红包个数最小为10个

红包金额最低5元

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

抵扣说明:

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

余额充值