Sepk

NCCL 的源码中，具有两种方式的 GDR，nvidia-peermem和dmabuf，在注册收发缓冲区时会优先选择后面一种方式。源码注释中的nv_peermem和nvidia-peermem是同一个东西，在不同版本的叫法有所不同。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18

for (int p=0; p<NCCL_NUM_PROTOCOLS; p++) {
    resources->buffers[p] = NCCL_NET_MAP_GET_POINTER(map, cpu, buffs[p]);
    if (resources->buffers[p]) {
#if CUDA_VERSION >= 11070
      /* DMA-BUF support */
      int type = NCCL_NET_MAP_DEV_MEM(map, buffs[p]) ? NCCL_PTR_CUDA : NCCL_PTR_HOST;
      if (type == NCCL_PTR_CUDA && resources->useDmaBuf) {
        int dmabuf_fd;
        CUCHECK(cuMemGetHandleForAddressRange((void *)&dmabuf_fd, (CUdeviceptr)resources->buffers[p], resources->buffSizes[p], CU_MEM_RANGE_HANDLE_TYPE_DMA_BUF_FD, 0));
        NCCLCHECK(proxyState->ncclNet->regMrDmaBuf(resources->netRecvComm, resources->buffers[p], resources->buffSizes[p], type, 0ULL, dmabuf_fd, &resources->mhandles[p]));
        (void)close(dmabuf_fd);
      } else // FALL-THROUGH to nv_peermem GDR path
#endif
      {
        NCCLCHECK(proxyState->ncclNet->regMr(resources->netRecvComm, resources->buffers[p], resources->buffSizes[p], NCCL_NET_MAP_DEV_MEM(map, buffs[p]) ? NCCL_PTR_CUDA : NCCL_PTR_HOST, &resources->mhandles[p]));
      }
    }
  }

中文网站上，有关 dmabuf 和ibv_reg_dmabuf_mr的相关资料比较少。在阅读源码时，笔者不确定这两种方式是否有性能上的差异，因此在 NCCL 的仓库下提出了一个issue，得到的回复是，这两者之间没有性能上的差别：

No, there is no performance difference.

之后，笔者调研了一下相关资料，对比了两者之间的差异，在此作为二手资料进行分享。

虽然 RDMA 这个概念已经很早了，但是 GDR 的提出不过是最近十年左右的事情——早期并没有大量数据直接传输到 GPU 上的需求。因此，GDR 这一方案也并没有一个明确的标准。GDR 与普通 RDMA 的区别在于，计算机访问显存和内存的方式不一样，早期的 libverbs 接口中，并不存在ibv_reg_dmabuf_mr，只存在ibv_reg_mr这一通用的接口。因此cudaMalloc分配出的假显存被注册时，IB core 是不知道这是一片映射得到的内存，在尝试去 pin 这块内存时，内核会报错，具体的原因可以参见pin_user_pages这一函数的设计文档，里面有比较详细的介绍。

nvidia-peermem的这一方案，是 Mallenox 公司推出的。它没有改变 libverbs 的接口，而是修改了自身的驱动，对于不能 pin 住的内存，它会轮询额外加载的peer_memory_client模块，让这些模块尝试翻译并 pin 住这些内存区域。对于 N 卡来说，这一模块就是nvidia-peermem。在忽略 GPU 页表切换时，其实这一模块就是完成了一个地址翻译的工作，数据传输本质上还是依赖于显卡驱动的。

ibv_reg_dmabuf_mr是 OpenFabric 提出的方案，为了反对 Mallenox 公司对 GDR 方案的垄断：nvidia-peermem需要绑定 Mallenox 的网卡。这一方案利用 Kernel 中设计的 dmabuf，将其作为中间值，从而让 RNIC 能够拿到翻译后的物理地址。ibv_reg_dmabuf_mr就是为了这一方案而新引入 libverbs 中的一个接口。这一方案的实现是更加优雅的，并且在注册 mr 时不需要轮询，在注册时的性能“理论上”会好一些。ibv_reg_dmabuf_mr这一接口的具体实现，直至 Linux Kernel 5.12 才被加入内核中，因此对软件栈的要求是比较高的。

如果使用的是 Mallenox 网卡，两种方案并没有本质上的区别，因为数据传输都是由显卡驱动完成的。

libverbs 是 RDMA 通信的通用库，它完成了 RNIC Driver 层的抽象，给用户层提供了一个统一的接口。libverbs 总体上是一个异步通知的队列逻辑。在不考虑 RNIC 网卡性能极限的情况下，一套比较优秀的异步事件交互逻辑能够更加有效地利用网卡的性能，达到更低的延迟和更高的吞吐量。在 NCCL 中源码中，有一些实现上的细节能够提供给我们优化 RDMA 通信逻辑的相关思路。

IBV_SENDLINE 的使用问题

IBV_SEND_LINE是 libverbs 中针对小数据包的优化处理，它能够明显改善 128 字节以下数据包的发送延迟，略微提升吞吐量。这一优化的具体实现是使用 CPU 而非 RNIC 的 DPU 进行发送。但是使用IBV_SEND_INLINE时，必须要使用IBV_SEND_SIGNALED来处理事件的完成通知，否则可能会导致 SQE 用尽，无法发送数据。

在 NCCL 中，会定期使用IBV_SEND_SIGNALED标识位来处理 inline 包的通知问题，即发送一轮 slot 时：

1
2
3
4
5

if (slot == 0) {
    wr.send_flags |= IBV_SEND_SIGNALED;
    wr.wr_id = req - comm->verbs.reqs;
    req->events++;
}

这种思路同样也可以用到其余地方，每隔几次 SR 再带上一次标识位，这样能够降低与 comp channel 之间的通信频率。

多 QP 利用 DPU 性能

在 RC 模式下，为了保证每一个事件的完成是顺序的，单一 SQ 中的事件只能够被 RNIC 中的一个 DPU 处理。当连接较少、数据包较少时，可能无法充分利用 RNIC 的 DPU，导致网络性能下降。 NCCL 中会尝试将一次用户请求的数据分解到多个 QP 当中去，从而更加充分地利用 RNIC 的 DPU 性能：

1
2
3
4
5
6
7
8

for (int q = 0; q < nqps; q++) {

    struct ibv_send_wr *bad_wr;
    // 这里会尝试利用多个 qp 进行数据的发送
    NCCLCHECK(wrap_ibv_post_send(comm->qps[comm->qpIndex], comm->wrs, &bad_wr));
    comm->qpIndex = (comm->qpIndex + 1) % comm->nqps;

}

通信分组通知

RC 模式下，为了检验 SR 是否完成，必须要使用IBV_SEND_SIGNALED与 comp channel 进行交互。但是使用这一标识位意味着应用层需要处理来自于 comp channel 的事件通知，如果每一个请求都使用该标识位，就会导致事件通知过于频繁，应用层需要频繁处理事件完成通知，这会在一定程度上导致 CPU 占用率的上升。由于 NCCL 中分割发送的数据包，在短时间可能就会出现频繁的事件通知，但这种事件通知是不必要的，NCCL 中对此做出了优化：只在分组的最后一个 SR 上使用标识位IBV_SEND_SIGNALED。这样可以利用 RC 的可靠性保证，确保所有请求发送完毕，并且能够有效地减少事件通知的次数，降低 CPU 占用率：

1
2
3
4
5
6
7
8
9
10
11

struct ibv_send_wr *lastWr = comm->wrs + nreqs - 1;
if (nreqs > 1 || (comm->ar && reqs[0]->send.size > ncclParamIbArThreshold())) {
    lastWr++;
    memset(lastWr, 0, sizeof(struct ibv_send_wr));
}
lastWr->wr_id = wr_id;
lastWr->opcode = IBV_WR_RDMA_WRITE_WITH_IMM;
lastWr->imm_data = immData;
lastWr->next = NULL;
// Only use this flag in last WR
lastWr->send_flags = IBV_SEND_SIGNALED;

一些思考

NCCL 所做出的优化，主要还是针对于小数据包的带宽和延迟问题，而针对大数据包的优化比较少。根据我的理解，在数据包较小时，RDMA 的网络通信次数会更加频繁，相应地 comp channel 的通知也比较多。在这种情况下，应用层与 RNIC 驱动之间的交互的耗时占总耗时的比例就比较高，针对这一场景进行优化，效果就比较明显。而数据包较大时，数据的发送耗时大大增加，应用处与 RNIC 驱动交互的耗时就可以忽略不计了。并且当 RNIC 网卡的带宽吃满之后，限制 RDMA 网络通信性能的是网卡本身，而非通信库的处理逻辑。