通过使用CUDA和RDMA库，在GPU和远程内存之间启用GPUDirect RDMA功能，可以实现RDMA支持GPUDirect特性。

下面是使用GPUDirect RDMA的示例代码，本文将对该代码进行逐段解析：

// 1. 初始化CUDA和RDMA
cudaSetDevice(0);
cudaMalloc(&dptr, size);

ibv_device **dev_list = ibv_get_device_list(&num_devices);
context = ibv_open_device(*dev_list);
pd = ibv_alloc_pd(context);
cq = ibv_create_cq(context, 10, NULL, NULL, 0);

// 2. 创建GPUDirect RDMA使用的MR（内存区域）
mr = ibv_reg_mr(pd, dptr, size, IBV_ACCESS_LOCAL_WRITE | IBV_ACCESS_REMOTE_WRITE | IBV_ACCESS_REMOTE_READ);

// 3. 获取目标Infiniband RDMA网卡的IP和端口号
struct sockaddr_in addr;
addr.sin_family = AF_INET;
addr.sin_port = htons(12345);
inet_pton(AF_INET, "192.168.1.2", &addr.sin_addr);

// 4. 获取远程内存区域（MR）的RKEY（指向内存区域的关键字）
ibv_qp_init_attr init_attr;
memset(&init_attr, 0, sizeof(init_attr));
init_attr.send_cq = cq;
init_attr.recv_cq = cq;
init_attr.qp_type = IBV_QPT_RC;
init_attr.cap.max_send_wr = 1;
init_attr.cap.max_recv_wr = 1;
init_attr.cap.max_send_sge = 1;
init_attr.cap.max_recv_sge = 1;
qp = ibv_create_qp(pd, &init_attr);
ibv_connect_qp(qp, &addr, NULL);
ibv_post_recv(qp, &r_wr, &bad_wr);
while (!r_wr.done);

// 5. 在目标主机上写入数据
struct ibv_send_wr s_wr;
struct ibv_sge s_sge;
struct ibv_send_wr *bad_wr = NULL;
int rc;
memset(&s_wr, 0, sizeof(s_wr));
memset(&s_sge, 0, sizeof(s_sge));
void *msg = malloc(size);
memcpy(msg, buf, size);
s_sge.addr = (uintptr_t)msg;
s_sge.length = size;
s_sge.lkey = mr->lkey;
s_wr.wr_id = (uintptr_t)msg;
s_wr.sg_list = &s_sge;
s_wr.num_sge = 1;
s_wr.opcode = IBV_WR_RDMA_WRITE_WITH_IMM;
s_wr.send_flags = IBV_SEND_SIGNALED;
s_wr.wr.rdma.remote_addr = raddr
rc = ibv_post_send(res->qp, &s_wr, &bad_wr);
if (rc)
	fprintf(stderr, "failed to post SR\n");
return rc;

1. 初始化CUDA和RDMA

cudaSetDevice(0);
cudaMalloc(&dptr, size);
ibv_device **dev_list = ibv_get_device_list(&num_devices);
context = ibv_open_device(*dev_list);
pd = ibv_alloc_pd(context);
cq = ibv_create_cq(context, 10, NULL, NULL, 0);

这段代码是在使用CUDA和InfiniBand（IB）进行GPU编程和网络通信时的一些初始化操作。

首先，cudaSetDevice(0)用于将当前设备设置为第一个GPU设备。这个函数用于选择要在其中执行CUDA操作的设备。参数0表示选择第一个设备，如果存在多个设备，则可以选择其他设备。

接下来，cudaMalloc(&dptr, size)用于在GPU上分配一块内存。dptr是一个指向指针的指针，用于存储分配的GPU内存的地址。size表示要分配的内存大小。

然后，ibv_device **dev_list = ibv_get_device_list(&num_devices)用于获取所有可用的InfiniBand设备列表并存储在dev_list中。num_devices用于存储找到的设备数量。

context = ibv_open_device(*dev_list)用于打开第一个InfiniBand设备并返回一个用于后续IB操作的上下文。*dev_list表示从设备列表中选择第一个设备。

pd = ibv_alloc_pd(context)用于为IB上下文分配一个保护域（Protection Domain，PD），PD是IB中进行内存保护和访问控制的基本单位。

最后，cq = ibv_create_cq(context, 10, NULL, NULL, 0)用于创建一个完成队列（Completion Queue，CQ），CQ用于存储IB操作完成的通知和相关信息。参数10表示CQ的容量，NULL表示不指定通知回调函数和用户数据，最后一个参数0表示创建CQ时的附加标志。

2. 创建GPUDirect RDMA使用的MR（内存区域）

mr = ibv_reg_mr(pd, dptr, size, IBV_ACCESS_LOCAL_WRITE | IBV_ACCESS_REMOTE_WRITE | IBV_ACCESS_REMOTE_READ);

这段代码是在使用InfiniBand进行内存注册的操作。

mr = ibv_reg_mr(pd, dptr, size, IBV_ACCESS_LOCAL_WRITE | IBV_ACCESS_REMOTE_WRITE | IBV_ACCESS_REMOTE_READ)用于将GPU内存注册为可被远程读写的内存区域。

pd是之前通过ibv_alloc_pd函数分配的保护域（PD）。

dptr是之前通过cudaMalloc函数在GPU上分配的内存。

size表示要注册的内存大小。

IBV_ACCESS_LOCAL_WRITE | IBV_ACCESS_REMOTE_WRITE | IBV_ACCESS_REMOTE_READ是注册内存时指定的访问权限。其中，IBV_ACCESS_LOCAL_WRITE表示允许本地写操作，IBV_ACCESS_REMOTE_WRITE表示允许远程写操作，IBV_ACCESS_REMOTE_READ表示允许远程读操作。通过指定这些权限，可以控制对内存的访问和操作方式。

ibv_reg_mr函数将返回一个指向注册内存区域的内存区域描述符（Memory Region Descriptor，MRD）的指针。这个描述符可以用于后续的IB操作，如发送和接收数据。

这段代码用于将GPU上分配的内存注册为可供IB进行远程读写的内存区域，以便在后续的网络通信中使用。

3. 获取目标Infiniband RDMA网卡的IP和端口号

struct sockaddr_in addr;
addr.sin_family = AF_INET;
addr.sin_port = htons(12345);
inet_pton(AF_INET, "192.168.1.2", &addr.sin_addr);

这段代码是用于设置目标Infiniband RDMA网卡的IP和端口号。

首先，addr.sin_family = AF_INET 将地址结构的地址族设置为 IPv4，表示使用 IPv4 地址。

接下来，addr.sin_port = htons(12345) 将端口号设置为 12345。htons 函数用于将主机字节序转换为网络字节序，因为网络通信中使用的是网络字节序。

然后，inet_pton(AF_INET, "192.168.1.2", &addr.sin_addr) 将 IP 地址设置为 "192.168.1.2"。inet_pton 函数用于将点分十进制的 IP 地址转换为二进制形式，并将结果存储在 &addr.sin_addr 中。

通过这些设置，struct sockaddr_in 结构体的成员已经包含了目标主机的 IPv4 地址、端口号等信息，可以用于后续的网络通信操作。

4. 获取远程内存区域（MR）的RKEY（指向内存区域的关键字）

ibv_qp_init_attr init_attr;
memset(&init_attr, 0, sizeof(init_attr));
init_attr.send_cq = cq;
init_attr.recv_cq = cq;
init_attr.qp_type = IBV_QPT_RC;
init_attr.cap.max_send_wr = 1;
init_attr.cap.max_recv_wr = 1;
init_attr.cap.max_send_sge = 1;
init_attr.cap.max_recv_sge = 1;
qp = ibv_create_qp(pd, &init_attr);
ibv_connect_qp(qp, &addr, NULL);
ibv_post_recv(qp, &r_wr, &bad_wr);
while (!r_wr.done);

这段代码是使用InfiniBand创建并配置一个队列对（Queue Pair，QP）并进行通信的过程。

首先，定义了 ibv_qp_init_attr 结构体 init_attr 来存储QP的初始化属性。

接下来，通过 memset(&init_attr, 0, sizeof(init_attr)) 将 init_attr 结构体的内存清零，以确保所有成员的初始值为0。

然后，设置 init_attr.send_cq 和 init_attr.recv_cq 分别为之前创建的完成队列 cq，用于处理发送和接收完成的通知。

init_attr.qp_type 设置为 IBV_QPT_RC，表示创建的QP类型为 Reliable Connection（RC）类型。RC类型的QP可以提供可靠的连接和有序的传输。

通过设置 init_attr.cap 结构体的成员，确定了QP的容量和限制。max_send_wr 表示最大发送工作请求数量，max_recv_wr 表示最大接收工作请求数量，max_send_sge 表示每个发送工作请求中的散射/聚集元素数量，max_recv_sge 表示每个接收工作请求中的散射/聚集元素数量。

使用 ibv_create_qp(pd, &init_attr) 创建了一个QP，并将其返回给 qp。pd 是之前通过 ibv_alloc_pd 函数分配的保护域。

接下来，通过 ibv_connect_qp(qp, &addr, NULL) 连接QP到指定的远程地址 addr。这将启动一个连接过程，使得该QP可以与指定的远程QP进行通信。

然后，使用 ibv_post_recv(qp, &r_wr, &bad_wr) 提交一个接收工作请求，将接收缓冲区与QP关联起来。r_wr 是用于接收工作请求的 ibv_recv_wr 结构体，bad_wr 用于记录错误的工作请求。

最后，使用一个循环来等待接收完成，通过检查 r_wr.done 的值来判断接收是否完成。在具体的应用中，可能会采用不同的方式来等待接收完成，这里只是一个简单的示例。

这段代码是创建并配置一个QP，进行连接，并提交一个接收工作请求，以准备后续的通信操作。

5. 在目标主机上写入数据

struct ibv_send_wr s_wr;
struct ibv_sge s_sge;
struct ibv_send_wr *bad_wr = NULL;
int rc;

memset(&s_wr, 0, sizeof(s_wr));
memset(&s_sge, 0, sizeof(s_sge));
void *msg = malloc(size);
memcpy(msg, buf, size);
s_sge.addr = (uintptr_t)msg;
s_sge.length = size;
s_sge.lkey = mr->lkey;
s_wr.wr_id = (uintptr_t)msg;
s_wr.sg_list = &s_sge;
s_wr.num_sge = 1;
s_wr.opcode = IBV_WR_RDMA_WRITE_WITH_IMM;
s_wr.send_flags = IBV_SEND_SIGNALED;
s_wr.wr.rdma.remote_addr = raddr
rc = ibv_post_send(res->qp, &s_wr, &bad_wr);
if (rc) {
    fprintf(stderr, "failed to post SR\n");
}
return rc;

这段代码是配置一个发送工作请求（Send Work Request，WR）并发送一条RDMA写操作的请求。

首先，定义了 ibv_send_wr 结构体 s_wr 和 ibv_sge 结构体 s_sge，用于存储发送工作请求和散射/聚集元素的信息。

通过 memset(&s_wr, 0, sizeof(s_wr)) 和 memset(&s_sge, 0, sizeof(s_sge)) 将这两个结构体的内存清零，以确保所有成员的初始值为0。

然后，使用 malloc(size) 分配了一个大小为 size 的内存块 msg，用于存储要发送的数据。同时使用 memcpy 函数将 buf 中的数据复制到 msg 中。

之后，通过设置 s_sge 结构体的成员，指定了发送数据的地址、长度和内存区域描述符的本地键（lkey）。addr 表示发送数据的地址，length 表示发送数据的长度，lkey 表示内存区域描述符（Memory Region Descriptor，MRD）的本地键。

接下来，将 msg 的地址赋值给 s_wr.wr_id，以便在发送完成后可以用于识别发送工作请求。

通过设置 s_wr.sg_list 为指向 s_sge 的指针，将 s_sge 与 s_wr 关联起来。s_wr.num_sge 设置为 1，表示只有一个散射/聚集元素。

s_wr.opcode 设置为 IBV_WR_RDMA_WRITE_WITH_IMM，表示发送一条带有立即数的RDMA写操作的请求。RDMA写操作是一种通过网络将数据写入远程内存的操作。

s_wr.send_flags 设置为 IBV_SEND_SIGNALED，表示在发送完成后，会发出一个完成事件。

最后，通过 s_wr.wr.rdma.remote_addr 将远程地址 raddr 赋值给 s_wr 的相关字段，以指定写入远程内存的目标地址。

这段代码配置了一个发送工作请求，并指定了要发送的数据、发送数据的地址和长度等信息。通过使用这个发送工作请求，可以发送一条RDMA写操作的请求，将数据写入指定的远程内存地址。

总结

GPU Direct RDMA操作，与普通的RDMA操作没有多大区别，主要是申请的内存是通过cudaMalloc(&dptr, size);申请的GPU内存，还是通过malloc函数申请的主机内存。

参考