03-RDMA与InfiniBand原理
概述
RDMA (Remote Direct Memory Access) 是高性能计算和大规模AI训练的核心网络技术。本章深入解析RDMA原理、InfiniBand架构、以及GPUDirect RDMA技术，帮助理解分布式训练的网络底层实现。
RDMA 基础原理
传统网络 vs RDMA
┌─────────────────────────────────────────────────────────────────────────┐
│ 传统 TCP/IP 网络栈 │
├─────────────────────────────────────────────────────────────────────────┤
│ │
│ 应用程序 │
│ │ │
│ ▼ ←── 系统调用 (用户态 → 内核态切换) │
│ ┌──────────────────────────────────────────┐ │
│ │ Socket 层 │ │
│ │ ├─ send()/recv() │ │
│ │ └─ 数据拷贝: 用户空间 → 内核缓冲区 │ ←── 第1次拷贝 │
│ └──────────────────────────────────────────┘ │
│ │ │
│ ▼ │
│ ┌──────────────────────────────────────────┐ │
│ │ TCP/UDP 层 │ │
│ │ ├─ 分段/重组 │ │
│ │ ├─ 流控/拥塞控制 │ │
│ │ └─ 可靠性保证 (ACK/重传) │ │
│ └──────────────────────────────────────────┘ │
│ │ │
│ ▼ │
│ ┌──────────────────────────────────────────┐ │
│ │ IP 层 │ │
│ │ ├─ 路由选择 │ │
│ │ └─ 分片/重组 │ │
│ └──────────────────────────────────────────┘ │
│ │ │
│ ▼ │
│ ┌──────────────────────────────────────────┐ │
│ │ 网卡驱动 │ │
│ │ └─ 数据拷贝: 内核缓冲区 → DMA 缓冲区 │ ←── 第2次拷贝 │
│ └──────────────────────────────────────────┘ │
│ │ │
│ ▼ │
│ ┌──────────────────────────────────────────┐ │
│ │ 网卡 (NIC) │ │
│ │ └─ DMA 传输到网络 │ │
│ └──────────────────────────────────────────┘ │
│ │
│ 问题: │
│ • 多次内存拷贝 (至少2次) │
│ • 多次上下文切换 (用户态 ↔ 内核态) │
│ • CPU 参与每个数据包处理 │
│ • 高延迟 (~10-100μs) │
└─────────────────────────────────────────────────────────────────────────┘
┌─────────────────────────────────────────────────────────────────────────┐
│ RDMA 网络 │
├─────────────────────────────────────────────────────────────────────────┤
│ │
│ 应用程序 │
│ │ │
│ ▼ ←── 直接操作 (无系统调用) │
│ ┌──────────────────────────────────────────┐ │
│ │ RDMA Verbs API │ │
│ │ ├─ ibv_post_send() │ 用户态直接操作 │
│ │ ├─ ibv_post_recv() │ 无需内核参与 │
│ │ └─ ibv_poll_cq() │ │
│ └──────────────────────────────────────────┘ │
│ │ │
│ ▼ ←── 零拷贝 (应用内存直接注册到网卡) │
│ ┌──────────────────────────────────────────┐ │
│ │ RDMA 网卡 (RNIC/HCA) │ │
│ │ ├─ 硬件实现协议栈 │ │
│ │ ├─ 内存注册 (Memory Registration) │ │
│ │ ├─ DMA 直接读写应用内存 │ │
│ │ └─ 硬件可靠性保证 │ │
│ └──────────────────────────────────────────┘ │
│ │
│ 优势: │
│ • 零拷贝 (Zero-Copy) │
│ • 内核旁路 (Kernel Bypass) │
│ • CPU 卸载 (协议栈在硬件实现) │
│ • 超低延迟 (~1-2μs) │
└─────────────────────────────────────────────────────────────────────────┘
RDMA 三大特性
// RDMA 核心特性实现原理
// 1. 零拷贝 (Zero-Copy)
// =====================
// 传统方式:
// 用户缓冲区 → 内核缓冲区 → DMA缓冲区 → 网卡
//
// RDMA 方式:
// 用户缓冲区 ────────────────────────→ 网卡 (直接DMA)
struct ibv_mr *mr = ibv_reg_mr(
pd, // Protection Domain
buffer, // 用户空间缓冲区
size, // 大小
IBV_ACCESS_LOCAL_WRITE | // 本地写权限
IBV_ACCESS_REMOTE_WRITE | // 远程写权限
IBV_ACCESS_REMOTE_READ // 远程读权限
);
// 注册后,网卡可以直接DMA访问这块内存
// 2. 内核旁路 (Kernel Bypass)
// ==========================
// 传统方式: 每次I/O都需要系统调用
// 用户态 → 系统调用 → 内核态 → 系统调用返回 → 用户态
//
// RDMA 方式: 用户态直接操作
// 用户态 → 直接操作QP → 完成
// 发送数据 - 完全在用户态完成
struct ibv_send_wr wr = {
.wr_id = id,
.sg_list = &sge, // 指向已注册内存
.num_sge = 1,
.opcode = IBV_WR_SEND,
.send_flags = IBV_SEND_SIGNALED,
};
// 直接写入QP的发送队列,无系统调用
ibv_post_send(qp, &wr, &bad_wr);
// 3. CPU 卸载 (Protocol Offload)
// ==============================
// 传统方式: CPU处理所有协议栈
// TCP/IP 分段、校验和、重传、流控全部CPU处理
//
// RDMA 方式: 硬件处理协议
// HCA/RNIC 硬件实现:
// - 数据分段
// - CRC校验
// - 可靠传输 (ACK/NAK/重传)
// - 流量控制
// - 拥塞控制
// CPU只需要发起操作,其他由硬件完成
RDMA 操作类型
┌─────────────────────────────────────────────────────────────────────────┐
│ RDMA 操作类型 │
├─────────────────────────────────────────────────────────────────────────┤
│ │
│ 1. Send/Receive (双边操作) │
│ ═══════════════════════════ │
│ │
│ 发送方 接收方 │
│ ┌────────┐ ┌────────┐ │
│ │ Buffer │ │ Buffer │ │
│ └───┬────┘ └───▲────┘ │
│ │ ibv_post_send() │ ibv_post_recv() │
│ ▼ │ │
│ ┌────────┐ 数据传输 ┌────────┐ │
│ │ QP │ ─────────────────────→ │ QP │ │
│ └────────┘ └────────┘ │
│ │
│ • 需要双方参与 (发送方post_send, 接收方post_recv) │
│ • 接收方需要预先post receive buffer │
│ • 类似传统socket语义 │
│ │
│ 2. RDMA Write (单边操作) │
│ ═══════════════════════ │
│ │
│ 本地 远程 │
│ ┌────────┐ ┌────────┐ │
│ │ Buffer │ │ Buffer │ │
│ │ (src) │ │ (dst) │ │
│ └───┬────┘ └───▲────┘ │
│ │ │ │
│ │ ibv_post_send(RDMA_WRITE) │ │
│ └───────────────────────────────────┘ │
│ 直接写入远程内存,远程CPU不感知 │
│ │
│ • 只需本地发起,远程CPU完全不参与 │
│ • 需要知道远程地址和rkey │
│ • 远程无任何通知 (除非使用RDMA Write with Immediate) │
│ │
│ 3. RDMA Read (单边操作) │
│ ══════════════════════ │
│ │
│ 本地 远程 │
│ ┌────────┐ ┌────────┐ │
│ │ Buffer │ │ Buffer │ │
│ │ (dst) │ │ (src) │ │
│ └───▲────┘ └───┬────┘ │
│ │ │ │
│ │ ibv_post_send(RDMA_READ) │ │
│ └───────────────────────────────────┘ │
│ 直接读取远程内存,远程CPU不感知 │
│ │
│ 4. Atomic 操作 │
│ ════════════ │
│ │
│ • Compare and Swap (CAS) │
│ • Fetch and Add (FAA) │
│ • 原子性由硬件保证 │
│ • 用于分布式锁、计数器等 │
│ │
└─────────────────────────────────────────────────────────────────────────┘
InfiniBand 架构
InfiniBand 网络拓扑
┌─────────────────────────────────────────────────────────────────────────┐
│ InfiniBand 网络架构 │
├─────────────────────────────────────────────────────────────────────────┤
│ │
│ ┌───────────────┐ │
│ │ Subnet │ │
│ │ Manager │ │
│ │ (SM) │ │
│ └───────┬───────┘ │
│ │ 管理所有子网 │
│ │ │
│ ┌────────────────────────────┼────────────────────────────┐ │
│ │ │ │ │
│ ▼ ▼ ▼ │
│ ┌──────┐ ┌──────┐ ┌──────┐ │
│ │ Leaf │ │ Leaf │ │ Leaf │ │
│ │Switch│ │Switch│ │Switch│ │
│ └──┬───┘ └──┬───┘ └──┬───┘ │
│ │ │ │ │
│ ┌──┴───────────┐ ┌────┴──────────┐ ┌────┴────────┐ │
│ │ │ │ │ │ │ │ │ │ │ │ │ │ │ │ │ │ │ │
│ ▼ ▼ ▼ ▼ ▼ ▼ ▼ ▼ ▼ ▼ ▼ ▼ ▼ ▼ ▼ ▼ ▼ ▼ │
│ ┌─┐┌─┐┌─┐┌─┐┌─┐┌─┐ ┌─┐┌─┐┌─┐┌─┐┌─┐┌─┐ ┌─┐┌─┐┌─┐┌─┐┌─┐┌─┐ │
│ │H││H││H││H││H││H│ │H││H││H││H││H││H│ │H││H││H││H││H││H│ │
│ │C││C││C││C││C││C│ │C││C││C││C││C││C│ │C││C││C││C││C││C│ │
│ │A││A││A││A││A││A│ │A││A││A││A││A││A│ │A││A││A││A││A││A│ │
│ └─┘└─┘└─┘└─┘└─┘└─┘ └─┘└─┘└─┘└─┘└─┘└─┘ └─┘└─┘└─┘└─┘└─┘└─┘ │
│ │ │ │ │ │ │ │ │ │ │ │ │ │ │ │ │ │ │ │
│ ▼ ▼ ▼ ▼ ▼ ▼ ▼ ▼ ▼ ▼ ▼ ▼ ▼ ▼ ▼ ▼ ▼ ▼ │
│ 计算节点/GPU服务器 计算节点/GPU服务器 计算节点/GPU服务器 │
│ │
├─────────────────────────────────────────────────────────────────────────┤
│ │
│ HCA (Host Channel Adapter) - InfiniBand 网卡 │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ │ │
│ │ ┌──────────┐ ┌──────────┐ ┌──────────┐ ┌──────────┐ │ │
│ │ │ Port │ │ Port │ │ Port │ │ Port │ │ │
│ │ │ 200Gbps │ │ 200Gbps │ │ 200Gbps │ │ 200Gbps │ │ │
│ │ └────┬─────┘ └────┬─────┘ └────┬─────┘ └────┬─────┘ │ │
│ │ │ │ │ │ │ │
│ │ └─────────────┴──────┬──────┴─────────────┘ │ │
│ │ │ │ │
│ │ ┌────────▼────────┐ │ │
│ │ │ IB Transport │ │ │
│ │ │ Engine │ │ │
│ │ │ ├─ RC/UC/UD │ │ │
│ │ │ ├─ 重传逻辑 │ │ │
│ │ │ └─ 流控 │ │ │
│ │ └────────┬────────┘ │ │
│ │ │ │ │
│ │ ┌────────▼────────┐ │ │
│ │ │ DMA Engine │ │ │
│ │ │ ├─ 零拷贝 │ │ │
│ │ │ └─ 内存注册 │ │ │
│ │ └────────┬────────┘ │ │
│ │ │ │ │
│ │ ┌────────▼────────┐ │ │
│ │ │ PCIe Gen4/5 │ │ │
│ │ │ x16 Interface │ │ │
│ │ └─────────────────┘ │ │
│ │ │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ │
└─────────────────────────────────────────────────────────────────────────┘
InfiniBand 传输类型
// InfiniBand 传输服务类型
// 1. RC (Reliable Connection) - 可靠连接
// ======================================
// 特点:
// - 面向连接,一对一通信
// - 硬件保证消息顺序和可靠性
// - 支持所有RDMA操作 (Send/Recv, Read, Write, Atomic)
// - AI训练最常用
struct ibv_qp_init_attr init_attr = {
.send_cq = cq,
.recv_cq = cq,
.cap = {
.max_send_wr = 128,
.max_recv_wr = 128,
.max_send_sge = 1,
.max_recv_sge = 1,
},
.qp_type = IBV_QPT_RC, // Reliable Connection
};
// RC 优点:
// - 硬件自动处理重传、去重、排序
// - 编程简单,语义清晰
// - 延迟最低 (~1μs)
// RC 缺点:
// - 每对节点需要一个QP (N^2 扩展性问题)
// - 资源消耗大 (QP、CQ等)
// 2. UC (Unreliable Connection) - 不可靠连接
// ==========================================
// 特点:
// - 面向连接,一对一
// - 不保证可靠性 (可能丢包)
// - 支持 Send/Recv, RDMA Write
// - 不支持 RDMA Read, Atomic
struct ibv_qp_init_attr uc_attr = {
.qp_type = IBV_QPT_UC, // Unreliable Connection
// ... 其他配置
};
// UC 适用场景:
// - 可以容忍丢包的流式数据
// - 上层有自己的可靠性机制
// 3. UD (Unreliable Datagram) - 不可靠数据报
// ==========================================
// 特点:
// - 无连接,类似UDP
// - 一个QP可以与任意节点通信
// - 只支持 Send/Recv
// - 消息大小限制 (MTU)
struct ibv_qp_init_attr ud_attr = {
.qp_type = IBV_QPT_UD, // Unreliable Datagram
// ... 其他配置
};
// UD 优点:
// - 扩展性好 (一个QP对所有节点)
// - 资源消耗小
// UD 缺点:
// - 不支持RDMA Read/Write
// - 消息大小限制
// 4. XRC (Extended Reliable Connection) - 扩展可靠连接
// ===================================================
// 特点:
// - 解决RC的扩展性问题
// - 多个QP共享SRQ (Shared Receive Queue)
// - 减少资源消耗
// 比较表:
// ┌─────────┬────────┬────────┬────────┬─────────┐
// │ 类型 │ 可靠性 │ 连接性 │ RDMA │ 扩展性 │
// ├─────────┼────────┼────────┼────────┼─────────┤
// │ RC │ ✓ │ 连接 │ 全支持 │ 差 │
// │ UC │ ✗ │ 连接 │ 部分 │ 差 │
// │ UD │ ✗ │ 无连接 │ 不支持 │ 好 │
// │ XRC │ ✓ │ 连接 │ 全支持 │ 较好 │
// └─────────┴────────┴────────┴────────┴─────────┘
IB Verbs API 详解
// InfiniBand Verbs API 完整流程
#include <infiniband/verbs.h>
// ============================================================
// 步骤1: 获取设备列表
// ============================================================
struct ibv_device **dev_list = ibv_get_device_list(&num_devices);
if (!dev_list) {
perror("Failed to get IB devices list");
return -1;
}
// 打印所有设备
for (int i = 0; i < num_devices; i++) {
printf("Device %d: %s\n", i, ibv_get_device_name(dev_list[i]));
}
// ============================================================
// 步骤2: 打开设备
// ============================================================
struct ibv_context *ctx = ibv_open_device(dev_list[0]);
if (!ctx) {
perror("Failed to open device");
return -1;
}
// 查询设备属性
struct ibv_device_attr dev_attr;
ibv_query_device(ctx, &dev_attr);
printf("Max QP: %d\n", dev_attr.max_qp);
printf("Max CQ: %d\n", dev_attr.max_cq);
printf("Max MR: %d\n", dev_attr.max_mr);
// ============================================================
// 步骤3: 分配 Protection Domain
// ============================================================
struct ibv_pd *pd = ibv_alloc_pd(ctx);
if (!pd) {
perror("Failed to allocate PD");
return -1;
}
// PD用于:
// - 隔离不同应用的资源
// - 关联 QP、MR、AH 等资源
// ============================================================
// 步骤4: 注册内存 (Memory Registration)
// ============================================================
void *buf = malloc(BUF_SIZE);
struct ibv_mr *mr = ibv_reg_mr(
pd,
buf,
BUF_SIZE,
IBV_ACCESS_LOCAL_WRITE | // 本地写
IBV_ACCESS_REMOTE_WRITE | // 远程写 (RDMA Write)
IBV_ACCESS_REMOTE_READ // 远程读 (RDMA Read)
);
if (!mr) {
perror("Failed to register MR");
return -1;
}
// mr->lkey: 本地访问密钥
// mr->rkey: 远程访问密钥 (需要发送给对端)
printf("lkey: 0x%x, rkey: 0x%x\n", mr->lkey, mr->rkey);
// ============================================================
// 步骤5: 创建 Completion Queue
// ============================================================
struct ibv_cq *cq = ibv_create_cq(
ctx,
128, // CQ 大小
NULL, // CQ context
NULL, // Completion channel (用于事件通知)
0 // Completion vector
);
if (!cq) {
perror("Failed to create CQ");
return -1;
}
// ============================================================
// 步骤6: 创建 Queue Pair
// ============================================================
struct ibv_qp_init_attr qp_init_attr = {
.send_cq = cq,
.recv_cq = cq,
.cap = {
.max_send_wr = 128, // 发送队列深度
.max_recv_wr = 128, // 接收队列深度
.max_send_sge = 1, // 每个WR的SGE数
.max_recv_sge = 1,
.max_inline_data = 64, // 内联数据大小
},
.qp_type = IBV_QPT_RC, // 可靠连接
};
struct ibv_qp *qp = ibv_create_qp(pd, &qp_init_attr);
if (!qp) {
perror("Failed to create QP");
return -1;
}
printf("QP number: %u\n", qp->qp_num);
// ============================================================
// 步骤7: QP 状态转换 (RESET → INIT → RTR → RTS)
// ============================================================
// 7.1 RESET → INIT
struct ibv_qp_attr init_attr = {
.qp_state = IBV_QPS_INIT,
.pkey_index = 0,
.port_num = 1,
.qp_access_flags = IBV_ACCESS_REMOTE_WRITE |
IBV_ACCESS_REMOTE_READ |
IBV_ACCESS_LOCAL_WRITE,
};
ibv_modify_qp(qp, &init_attr,
IBV_QP_STATE |
IBV_QP_PKEY_INDEX |
IBV_QP_PORT |
IBV_QP_ACCESS_FLAGS);
// 7.2 INIT → RTR (Ready to Receive)
struct ibv_qp_attr rtr_attr = {
.qp_state = IBV_QPS_RTR,
.path_mtu = IBV_MTU_4096,
.dest_qp_num = remote_qpn, // 对端QP号
.rq_psn = 0, // 接收PSN
.max_dest_rd_atomic = 1,
.min_rnr_timer = 12,
.ah_attr = {
.is_global = 0,
.dlid = remote_lid, // 对端LID
.sl = 0,
.src_path_bits = 0,
.port_num = 1,
},
};
ibv_modify_qp(qp, &rtr_attr,
IBV_QP_STATE |
IBV_QP_AV |
IBV_QP_PATH_MTU |
IBV_QP_DEST_QPN |
IBV_QP_RQ_PSN |
IBV_QP_MAX_DEST_RD_ATOMIC |
IBV_QP_MIN_RNR_TIMER);
// 7.3 RTR → RTS (Ready to Send)
struct ibv_qp_attr rts_attr = {
.qp_state = IBV_QPS_RTS,
.timeout = 14,
.retry_cnt = 7,
.rnr_retry = 7,
.sq_psn = 0, // 发送PSN
.max_rd_atomic = 1,
};
ibv_modify_qp(qp, &rts_attr,
IBV_QP_STATE |
IBV_QP_TIMEOUT |
IBV_QP_RETRY_CNT |
IBV_QP_RNR_RETRY |
IBV_QP_SQ_PSN |
IBV_QP_MAX_QP_RD_ATOMIC);
// ============================================================
// 步骤8: 发送数据 (RDMA Write)
// ============================================================
struct ibv_sge sge = {
.addr = (uintptr_t)buf,
.length = data_size,
.lkey = mr->lkey,
};
struct ibv_send_wr wr = {
.wr_id = 0,
.sg_list = &sge,
.num_sge = 1,
.opcode = IBV_WR_RDMA_WRITE,
.send_flags = IBV_SEND_SIGNALED,
.wr = {
.rdma = {
.remote_addr = remote_addr, // 远程地址
.rkey = remote_rkey, // 远程key
},
},
};
struct ibv_send_wr *bad_wr;
int ret = ibv_post_send(qp, &wr, &bad_wr);
if (ret) {
perror("Failed to post send");
return -1;
}
// ============================================================
// 步骤9: 轮询完成队列
// ============================================================
struct ibv_wc wc;
int ne;
do {
ne = ibv_poll_cq(cq, 1, &wc);
} while (ne == 0);
if (ne < 0) {
perror("Failed to poll CQ");
return -1;
}
if (wc.status != IBV_WC_SUCCESS) {
printf("Work completion error: %s\n",
ibv_wc_status_str(wc.status));
return -1;
}
printf("RDMA Write completed successfully\n");
// ============================================================
// 步骤10: 清理资源
// ============================================================
ibv_destroy_qp(qp);
ibv_destroy_cq(cq);
ibv_dereg_mr(mr);
ibv_dealloc_pd(pd);
ibv_close_device(ctx);
ibv_free_device_list(dev_list);
free(buf);
GPUDirect 技术
GPUDirect 演进
┌─────────────────────────────────────────────────────────────────────────┐
│ GPUDirect 技术演进 │
├─────────────────────────────────────────────────────────────────────────┤
│ │
│ 1. GPUDirect Shared Access (2010, CUDA 3.1) │
│ ════════════════════════════════════════════ │
│ │
│ 传统方式: │
│ CPU Memory ←──→ GPU Memory ←──→ CPU Memory ←──→ NIC │
│ 3次拷贝 │
│ │
│ GPUDirect: │
│ CPU Memory (pinned) ←────────────────────────→ NIC │
│ ↑ │
│ └── GPU 可以直接访问 pinned memory │
│ │
│ 2. GPUDirect P2P (2011, CUDA 4.0) │
│ ═════════════════════════════════ │
│ │
│ 同一PCIe树内的GPU直接通信: │
│ │
│ ┌─────────────────────┐ │
│ │ PCIe Switch │ │
│ └──────┬────┬─────────┘ │
│ │ │ │
│ ┌──────┴─┐ ┌┴───────┐ │
│ │ GPU0 │ │ GPU1 │ │
│ │ Memory │ │ Memory │ │
│ └────────┘ └────────┘ │
│ ↑ ↑ │
│ └────┬────┘ │
│ │ │
│ cudaMemcpyPeer (不经过CPU) │
│ │
│ 3. GPUDirect RDMA (2013, CUDA 5.0) │
│ ══════════════════════════════════ │
│ │
│ GPU Memory 直接与网络通信: │
│ │
│ ┌──────────┐ ┌──────────┐ │
│ │ GPU 0 │ │ GPU 1 │ │
│ │ Memory │ │ Memory │ │
│ └────┬─────┘ └────┬─────┘ │
│ │ │ │
│ │ DMA │ DMA │
│ ▼ ▼ │
│ ┌─────────┐ 网络 ┌─────────┐ │
│ │ NIC │ ◄────────────────────► │ NIC │ │
│ └─────────┘ └─────────┘ │
│ Node A Node B │
│ │
│ • 完全绕过CPU和系统内存 │
│ • 延迟降低 ~50% │
│ • 带宽提升 ~30% │
│ │
│ 4. NVLink 与 GPUDirect (2016+) │
│ ════════════════════════════════ │
│ │
│ ┌─────────┐ NVLink (600GB/s) ┌─────────┐ │
│ │ GPU 0 │ ◄────────────────► │ GPU 1 │ │
│ └─────────┘ └─────────┘ │
│ ▲ ▲ │
│ │ NVLink │ NVLink │
│ ▼ ▼ │
│ ┌─────────┐ NVLink ┌─────────┐ │
│ │ GPU 2 │ ◄────────────────► │ GPU 3 │ │
│ └─────────┘ └─────────┘ │
│ │
│ NVSwitch: 所有GPU全互联, 900GB/s 双向带宽 │
│ │
└─────────────────────────────────────────────────────────────────────────┘
GPUDirect RDMA 实现
// GPUDirect RDMA 编程示例
#include <cuda_runtime.h>
#include <infiniband/verbs.h>
// ============================================================
// 步骤1: 分配 GPU 内存
// ============================================================
void *gpu_buf;
size_t size = 1024 * 1024; // 1MB
cudaError_t err = cudaMalloc(&gpu_buf, size);
if (err != cudaSuccess) {
fprintf(stderr, "cudaMalloc failed: %s\n", cudaGetErrorString(err));
return -1;
}
// 获取GPU指针属性,确认是设备内存
cudaPointerAttributes attr;
cudaPointerGetAttributes(&attr, gpu_buf);
printf("Memory type: %d (device=%d)\n", attr.type, cudaMemoryTypeDevice);
// ============================================================
// 步骤2: 检查 GPUDirect RDMA 支持
// ============================================================
// 检查 nvidia_peermem 模块是否加载
// $ lsmod | grep nvidia_peermem
// nvidia_peermem 16384 0
// 检查 IB 设备是否支持
struct ibv_device_attr_ex attr_ex;
ibv_query_device_ex(ctx, NULL, &attr_ex);
// 检查 IBV_DEVICE_MEM_WINDOW_TYPE_2B 等标志
// ============================================================
// 步骤3: 注册 GPU 内存到 RDMA
// ============================================================
// 方法1: 使用标准 ibv_reg_mr (需要 nvidia_peermem 支持)
struct ibv_mr *gpu_mr = ibv_reg_mr(
pd,
gpu_buf,
size,
IBV_ACCESS_LOCAL_WRITE |
IBV_ACCESS_REMOTE_WRITE |
IBV_ACCESS_REMOTE_READ
);
if (!gpu_mr) {
// 可能的错误:
// - nvidia_peermem 未加载
// - GPU驱动版本不兼容
// - CUDA版本问题
perror("Failed to register GPU memory");
// 检查错误原因
printf("Check: lsmod | grep nvidia_peermem\n");
printf("Check: nvidia-smi -q | grep 'Persistence Mode'\n");
return -1;
}
printf("GPU memory registered: lkey=0x%x, rkey=0x%x\n",
gpu_mr->lkey, gpu_mr->rkey);
// ============================================================
// 步骤4: 使用 GPU 内存进行 RDMA 操作
// ============================================================
// 发送 GPU 数据
struct ibv_sge sge = {
.addr = (uintptr_t)gpu_buf,
.length = size,
.lkey = gpu_mr->lkey,
};
struct ibv_send_wr wr = {
.wr_id = 0,
.sg_list = &sge,
.num_sge = 1,
.opcode = IBV_WR_RDMA_WRITE,
.send_flags = IBV_SEND_SIGNALED,
.wr.rdma = {
.remote_addr = remote_gpu_addr,
.rkey = remote_rkey,
},
};
struct ibv_send_wr *bad_wr;
ibv_post_send(qp, &wr, &bad_wr);
// 数据路径:
// Local GPU Memory → PCIe → NIC → Network → NIC → PCIe → Remote GPU Memory
// 完全不经过CPU和系统内存!
// ============================================================
// 步骤5: 同步 GPU 和 RDMA 操作
// ============================================================
// 问题: GPU计算和RDMA传输如何同步?
// 方案1: 显式同步
cudaDeviceSynchronize(); // 等待GPU计算完成
ibv_post_send(...); // 然后发送
// 等待RDMA完成
struct ibv_wc wc;
while (ibv_poll_cq(cq, 1, &wc) == 0);
cudaKernel<<<...>>>(); // 然后启动下一个计算
// 方案2: 使用CUDA Stream和RDMA协调
cudaStream_t stream;
cudaStreamCreate(&stream);
// 在stream中启动kernel
myKernel<<<grid, block, 0, stream>>>(gpu_buf, ...);
// 记录event
cudaEvent_t event;
cudaEventCreate(&event);
cudaEventRecord(event, stream);
// 等待event完成
cudaEventSynchronize(event);
// 发送数据
ibv_post_send(qp, &wr, &bad_wr);
// 方案3: 使用GDRCopy优化小数据
// GDRCopy允许CPU直接访问GPU内存映射
// 适合小数据的低延迟传输
// ============================================================
// 性能对比
// ============================================================
/*
传输场景: 节点间 GPU-to-GPU 1GB 数据
传统方式 (GPU→CPU→NIC→NIC→CPU→GPU):
- GPU→CPU: 12GB/s (PCIe Gen3)
- CPU→NIC: 6GB/s
- NIC→NIC: 12GB/s (100Gbps IB)
- NIC→CPU: 6GB/s
- CPU→GPU: 12GB/s
- 总耗时: ~500ms
- CPU利用率: 高
GPUDirect RDMA (GPU→NIC→NIC→GPU):
- GPU→NIC: 12GB/s (PCIe Gen3)
- NIC→NIC: 12GB/s (100Gbps IB)
- NIC→GPU: 12GB/s
- 总耗时: ~250ms (2x加速)
- CPU利用率: 几乎为0
使用NVLink + IB HDR:
- 带宽可达 200Gbps
- 延迟 < 2μs
*/
GDRCopy 详解
// GDRCopy: 让CPU低延迟访问GPU内存
// https://github.com/NVIDIA/gdrcopy
#include <gdrapi.h>
// ============================================================
// GDRCopy 原理
// ============================================================
/*
传统CPU访问GPU内存:
CPU → cudaMemcpy() → Driver → PCIe DMA → GPU Memory
延迟: 数十微秒
GDRCopy方式:
CPU → 直接内存映射 → GPU Memory (通过BAR1)
延迟: ~1-2微秒
原理:
GPU的BAR (Base Address Register) 可以映射到CPU地址空间
GDRCopy利用这个机制创建CPU可访问的映射
*/
// ============================================================
// GDRCopy 使用示例
// ============================================================
// 1. 初始化
gdr_t g = gdr_open();
if (!g) {
fprintf(stderr, "gdr_open failed\n");
return -1;
}
// 2. 分配GPU内存 (必须是page-aligned)
void *gpu_ptr;
size_t size = 1024 * 1024; // 1MB
cudaMalloc(&gpu_ptr, size);
// 3. Pin GPU内存
gdr_mh_t mh;
int ret = gdr_pin_buffer(g, (unsigned long)gpu_ptr, size, 0, 0, &mh);
if (ret) {
fprintf(stderr, "gdr_pin_buffer failed: %d\n", ret);
return -1;
}
// 4. 映射到CPU地址空间
void *cpu_ptr;
ret = gdr_map(g, mh, &cpu_ptr, size);
if (ret) {
fprintf(stderr, "gdr_map failed: %d\n", ret);
return -1;
}
// 5. CPU直接读写GPU内存 (低延迟!)
// 写入GPU
gdr_copy_to_mapping(mh, cpu_ptr, src_data, size);
// 读取GPU
gdr_copy_from_mapping(mh, dst_data, cpu_ptr, size);
// 甚至可以直接使用memcpy (但gdr_copy更优化)
// memcpy(cpu_ptr, src_data, size);
// 6. 清理
gdr_unmap(g, mh, cpu_ptr, size);
gdr_unpin_buffer(g, mh);
gdr_close(g);
cudaFree(gpu_ptr);
// ============================================================
// GDRCopy 在 NCCL 中的应用
// ============================================================
/*
NCCL使用GDRCopy:
1. 小消息优化:
- 对于小于阈值(通常几KB)的消息
- 使用GDRCopy直接CPU访问,而不是GPU kernel
- 减少kernel launch开销
2. 控制路径:
- 同步信号、标志位等小数据
- 使用GDRCopy低延迟访问
3. 与GPUDirect RDMA配合:
- GDRCopy处理小数据
- GPUDirect RDMA处理大数据
*/
// NCCL中的相关代码 (简化)
if (size < GDRCOPY_THRESHOLD) {
// 小消息用GDRCopy
gdr_copy_to_mapping(mh, cpu_ptr, src, size);
// 然后通过socket/IB发送cpu_ptr指向的数据
} else {
// 大消息用GPUDirect RDMA
ibv_post_send(qp, &rdma_wr, &bad_wr);
}
RoCE (RDMA over Converged Ethernet)
RoCE vs InfiniBand
┌─────────────────────────────────────────────────────────────────────────┐
│ RoCE vs InfiniBand 对比 │
├─────────────────────────────────────────────────────────────────────────┤
│ │
│ RoCE (RDMA over Converged Ethernet) │
│ ═══════════════════════════════════ │
│ │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ 应用层 (Verbs API) │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ IB Transport (RoCE v2) │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ UDP (RoCE v2) / GRH (RoCE v1) │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ IP (RoCE v2) / Ethernet (RoCE v1) │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ 以太网 (标准交换机) │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ │
│ 优点: 缺点: │
│ • 使用标准以太网交换机 • 性能略低于原生IB │
│ • 成本较低 • 需要无损网络配置 │
│ • 可跨子网路由 (v2) • PFC/ECN配置复杂 │
│ • 运维团队熟悉 • 大规模部署挑战 │
│ │
├─────────────────────────────────────────────────────────────────────────┤
│ │
│ InfiniBand │
│ ═══════════ │
│ │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ 应用层 (Verbs API) │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ IB Transport │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ IB Network Layer │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ IB Link Layer │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ ┌─────────────────────────────────────────────────────────────┐ │
│ │ IB 专用交换机 │ │
│ └─────────────────────────────────────────────────────────────┘ │
│ │
│ 优点: 缺点: │
│ • 最高性能和最低延迟 • 专用硬件成本高 │
│ • 原生无损网络 • 需要IB专业知识 │
│ • 成熟的拥塞控制 • 交换机vendor lock-in │
│ • SM统一管理 • 与以太网隔离 │
│ │
├─────────────────────────────────────────────────────────────────────────┤
│ │
│ 性能对比 (典型值) │
│ ┌─────────────┬────────────────┬────────────────┐ │
│ │ 指标 │ InfiniBand │ RoCE v2 │ │
│ ├─────────────┼────────────────┼────────────────┤ │
│ │ 带宽 │ 400Gbps (NDR) │ 400Gbps │ │
│ │ 延迟 │ 0.6μs │ 1-2μs │ │
│ │ 成本 │ 高 │ 中 │ │
│ │ 扩展性 │ 中 (子网限制) │ 高 (IP路由) │ │
│ │ 配置复杂度 │ 中 │ 高 (PFC/ECN) │ │
│ └─────────────┴────────────────┴────────────────┘ │
│ │
└─────────────────────────────────────────────────────────────────────────┘
RoCE 无损网络配置
# RoCE v2 需要无损以太网,通过 PFC 和 ECN 实现
# ============================================================
# PFC (Priority Flow Control) 配置
# ============================================================
# 检查PFC状态
mlnx_qos -i eth0
# 启用PFC (优先级3用于RoCE)
mlnx_qos -i eth0 --pfc 0,0,0,1,0,0,0,0
# 配置信任模式
mlnx_qos -i eth0 --trust dscp
# ============================================================
# ECN (Explicit Congestion Notification) 配置
# ============================================================
# 启用ECN
echo 1 > /sys/class/net/eth0/ecn/roce_np/enable/3
echo 1 > /sys/class/net/eth0/ecn/roce_rp/enable/3
# 配置ECN阈值
# 当队列深度超过阈值时标记ECN
echo 150000 > /sys/class/net/eth0/ecn/roce_np/min_time_between_cnps
# ============================================================
# DSCP 映射配置
# ============================================================
# RoCE使用DSCP标记流量优先级
# DSCP 26 (AF31) 是常见的RoCE配置
# 配置DSCP到TC映射
mlnx_qos -i eth0 --dscp2prio set,26,3
# 配置TC到队列映射
tc qdisc add dev eth0 root mqprio \
num_tc 8 \
map 0 1 2 3 4 5 6 7 \
queues 1@0 1@1 1@2 1@3 1@4 1@5 1@6 1@7 \
hw 1
# ============================================================
# 验证配置
# ============================================================
# 检查RoCE模式
cat /sys/class/infiniband/mlx5_0/ports/1/gid_attrs/types/0
# 输出: RoCE v2
# 检查GID表
ibv_devinfo -v | grep GID
# 运行诊断
ibdiagnet -r
# 性能测试
ib_write_bw -d mlx5_0 -a -F # 服务端
ib_write_bw -d mlx5_0 -a -F <server_ip> # 客户端
NCCL 中的 RDMA 实现
NCCL IB Transport
// NCCL InfiniBand 传输层实现分析
// 文件位置: nccl/src/transport/net_ib.cc
// ============================================================
// IB Transport 初始化
// ============================================================
ncclResult_t ncclIbInit(ncclDebugLogger_t logFunction) {
// 1. 获取IB设备列表
struct ibv_device** devices;
devices = ibv_get_device_list(&nIbDevs);
// 2. 遍历并初始化每个设备
for (int d = 0; d < nIbDevs; d++) {
struct ncclIbDev* dev = ncclIbDevs + ncclNIbDevs;
// 打开设备
dev->context = ibv_open_device(devices[d]);
// 查询设备属性
ibv_query_device(dev->context, &dev->deviceAttr);
// 查询端口属性
ibv_query_port(dev->context, port, &dev->portAttr);
// 检查端口状态
if (dev->portAttr.state != IBV_PORT_ACTIVE) {
continue; // 跳过非活动端口
}
// 分配PD
dev->pd = ibv_alloc_pd(dev->context);
ncclNIbDevs++;
}
return ncclSuccess;
}
// ============================================================
// 连接建立
// ============================================================
ncclResult_t ncclIbConnect(int dev, void* opaqueHandle,
void** sendComm) {
struct ncclIbSendComm* comm;
// 1. 创建QP
struct ibv_qp_init_attr qpInitAttr = {
.send_cq = ibCq,
.recv_cq = ibCq,
.qp_type = IBV_QPT_RC,
.cap = {
.max_send_wr = MAX_REQUESTS,
.max_recv_wr = MAX_REQUESTS,
.max_send_sge = 1,
.max_recv_sge = 1,
.max_inline_data = ncclParamIbUseInline() ?
sizeof(struct ncclIbSendFifo) : 0,
},
};
comm->qp = ibv_create_qp(ncclIbDevs[dev].pd, &qpInitAttr);
// 2. 交换QP信息 (通过带外通道)
struct ncclIbQpInfo localQpInfo = {
.lid = portAttr.lid,
.qpn = comm->qp->qp_num,
.psn = 0,
// GID for RoCE
.gid = gid,
};
// 发送本地信息,接收远程信息
exchangeQpInfo(&localQpInfo, &remoteQpInfo);
// 3. 转换QP状态: RESET → INIT → RTR → RTS
ibv_modify_qp(comm->qp, &initAttr, ...); // → INIT
ibv_modify_qp(comm->qp, &rtrAttr, ...); // → RTR
ibv_modify_qp(comm->qp, &rtsAttr, ...); // → RTS
*sendComm = comm;
return ncclSuccess;
}
// ============================================================
// 内存注册 (支持GPUDirect RDMA)
// ============================================================
ncclResult_t ncclIbRegMr(void* comm, void* data, int size,
int type, void** mhandle) {
struct ncclIbMrCache* cache = &ncclIbDevs[dev].mrCache;
struct ibv_mr* mr;
// 检查MR缓存
mr = ncclIbMrCacheGet(cache, data, size);
if (mr) {
*mhandle = mr;
return ncclSuccess;
}
// 注册新的MR
int access = IBV_ACCESS_LOCAL_WRITE |
IBV_ACCESS_REMOTE_WRITE |
IBV_ACCESS_REMOTE_READ;
// 对于GPU内存,需要nvidia_peermem支持
mr = ibv_reg_mr(pd, data, size, access);
if (!mr) {
// 注册失败,可能是GPU内存但没有GPUDirect支持
WARN("IB MR registration failed for %p, size %d", data, size);
return ncclSystemError;
}
// 添加到缓存
ncclIbMrCacheAdd(cache, mr);
*mhandle = mr;
return ncclSuccess;
}
// ============================================================
// 数据发送 (RDMA Write)
// ============================================================
ncclResult_t ncclIbIsend(void* sendComm, void* data, int size,
int tag, void* mhandle, void** request) {
struct ncclIbSendComm* comm = sendComm;
struct ncclIbRequest* req;
// 获取空闲请求槽
req = ncclIbGetRequest(comm);
// 准备发送
struct ibv_sge sge = {
.addr = (uintptr_t)data,
.length = size,
.lkey = ((struct ibv_mr*)mhandle)->lkey,
};
struct ibv_send_wr wr = {
.wr_id = (uintptr_t)req,
.sg_list = &sge,
.num_sge = 1,
.opcode = IBV_WR_RDMA_WRITE_WITH_IMM, // 带立即数的RDMA Write
.send_flags = IBV_SEND_SIGNALED,
.imm_data = tag, // 使用立即数传递tag
.wr.rdma = {
.remote_addr = comm->remoteAddr,
.rkey = comm->remoteRkey,
},
};
// 发送
struct ibv_send_wr* bad_wr;
ibv_post_send(comm->qp, &wr, &bad_wr);
*request = req;
return ncclSuccess;
}
// ============================================================
// 完成检查
// ============================================================
ncclResult_t ncclIbTest(void* request, int* done, int* size) {
struct ncclIbRequest* req = request;
struct ncclIbSendComm* comm = req->comm;
// 轮询CQ
struct ibv_wc wc;
int ne = ibv_poll_cq(comm->cq, 1, &wc);
if (ne == 0) {
*done = 0;
return ncclSuccess;
}
if (wc.status != IBV_WC_SUCCESS) {
WARN("IB completion error: %s", ibv_wc_status_str(wc.status));
return ncclSystemError;
}
*done = 1;
if (size) *size = wc.byte_len;
// 回收请求
ncclIbFreeRequest(req);
return ncclSuccess;
}
性能调优指南
RDMA 性能优化
# ============================================================
# 系统级优化
# ============================================================
# 1. 禁用CPU频率调节
cpupower frequency-set -g performance
# 2. 禁用NUMA自动平衡
echo 0 > /proc/sys/kernel/numa_balancing
# 3. 设置大页内存
echo 4096 > /proc/sys/vm/nr_hugepages
# 4. 调整网络参数
sysctl -w net.core.rmem_max=67108864
sysctl -w net.core.wmem_max=67108864
sysctl -w net.core.rmem_default=67108864
sysctl -w net.core.wmem_default=67108864
# ============================================================
# IB 参数调优
# ============================================================
# 1. 增加QP深度
export NCCL_IB_QPS_PER_CONNECTION=8
# 2. 调整MTU
ibv_devinfo -d mlx5_0 | grep active_mtu
# 设置为4K: mlxconfig -d mlx5_0 set LINK_TYPE_P1=1
# 3. 启用自适应路由 (仅限某些交换机)
export NCCL_IB_AR_THRESHOLD=0
# 4. 多轨道配置
export NCCL_IB_HCA=mlx5_0,mlx5_1,mlx5_2,mlx5_3
# ============================================================
# GPUDirect RDMA 优化
# ============================================================
# 1. 确保nvidia_peermem加载
modprobe nvidia_peermem
# 2. 检查GPU和NIC的PCIe拓扑
nvidia-smi topo -m
# 最佳拓扑: GPU和NIC在同一PCIe switch下
# GPU0 <---> NIC0 (同一PCIe switch)
# GPU1 <---> NIC1 (同一PCIe switch)
# 3. 绑定进程到正确的NUMA节点
numactl --cpunodebind=0 --membind=0 ./my_app
# 4. 设置GPU和NIC亲和性
export CUDA_VISIBLE_DEVICES=0,1
export NCCL_IB_HCA=mlx5_0,mlx5_1
# ============================================================
# 性能诊断
# ============================================================
# 1. 检查IB统计
perfquery -x
# 2. 检查错误计数
cat /sys/class/infiniband/mlx5_0/ports/1/counters/*error*
# 3. 带宽测试
ib_write_bw -d mlx5_0 -s 1048576 -n 1000 <server>
# 4. 延迟测试
ib_write_lat -d mlx5_0 -s 8 -n 10000 <server>
# 5. NCCL测试
nccl-tests/all_reduce_perf -b 1M -e 1G -f 2 -g 8
常见问题排查
# ============================================================
# 问题1: ibv_reg_mr 失败
# ============================================================
# 症状: GPU内存注册失败
# 检查1: nvidia_peermem是否加载
lsmod | grep nvidia_peermem
# 如果没有: modprobe nvidia_peermem
# 检查2: CUDA和驱动版本
nvidia-smi
nvcc --version
# 检查3: GPU persistence mode
nvidia-smi -pm 1
# ============================================================
# 问题2: 连接超时
# ============================================================
# 症状: QP状态转换失败或连接超时
# 检查1: 端口状态
ibstatus
# 检查2: 子网管理器
sminfo
# 检查3: 路由表
ibroute
# ============================================================
# 问题3: 性能低于预期
# ============================================================
# 检查1: PCIe带宽
nvidia-smi -q | grep -A 10 "PCI"
# 检查2: 是否有PCIe降速
lspci -vv | grep -E "LnkSta:|Speed"
# 检查3: 拥塞控制
# RoCE: 检查PFC和ECN
mlnx_qos -i eth0
# 检查4: 重传统计
perfquery -x | grep Retry
# ============================================================
# 问题4: GPU Direct不工作
# ============================================================
# 检查ACS (Access Control Services)
# ACS会阻止GPU Direct
setpci -s <pcie_switch_bdf> ECAP_ACS+6.w
# 禁用ACS
setpci -s <pcie_switch_bdf> ECAP_ACS+6.w=0
总结
RDMA 核心概念
┌─────────────────────────────────────────────────────────────────────────┐
│ RDMA 核心概念总结 │
├─────────────────────────────────────────────────────────────────────────┤
│ │
│ 三大特性 │
│ ──────── │
│ • 零拷贝: 数据直接从用户空间DMA到网络 │
│ • 内核旁路: 数据路径不经过内核 │
│ • CPU卸载: 协议栈由硬件实现 │
│ │
│ 操作类型 │
│ ──────── │
│ • Send/Receive: 双边操作,类似socket │
│ • RDMA Write: 单边写,远程CPU不参与 │
│ • RDMA Read: 单边读,远程CPU不参与 │
│ • Atomic: 原子操作 (CAS, FAA) │
│ │
│ 传输类型 │
│ ──────── │
│ • RC: 可靠连接,AI训练首选 │
│ • UC: 不可靠连接,流媒体场景 │
│ • UD: 数据报,高扩展性场景 │
│ │
│ 网络技术 │
│ ──────── │
│ • InfiniBand: 最高性能,专用硬件 │
│ • RoCE v2: 以太网上的RDMA,需要无损配置 │
│ │
│ GPUDirect │
│ ───────── │
│ • GPUDirect P2P: GPU间直接通信 │
│ • GPUDirect RDMA: GPU内存直接网络传输 │
│ • GDRCopy: CPU低延迟访问GPU内存 │
│ │
└─────────────────────────────────────────────────────────────────────────┘
AI Infra 架构师必知
- RDMA 是大规模训练的基础: 理解零拷贝、内核旁路原理
- GPUDirect RDMA 实现: 掌握GPU内存直接网络传输
- NCCL 网络层: 理解IB Transport的QP管理和数据传输
- 性能调优: PCIe拓扑、NUMA亲和性、IB参数
- 故障排查: 连接问题、性能问题、GPUDirect问题
面试常见问题
- RDMA相比TCP的优势是什么?
- RC和UD传输类型的区别和适用场景?
- GPUDirect RDMA的工作原理?
- RoCE需要哪些网络配置?
- 如何排查RDMA性能问题?