Sepk

个人博客

0%

NCCL 网络浅析

发表于 更新于 分类于
本文字数： 3.5k 阅读时长 ≈ 3 分钟

NVIDIA 集合通信库（NCCL）是一个实现了多 GPU collective communication 的集合通信库。NCCL 在设计上是以 GPU Driven 为核心的，而非传统的 CPU Driven 代码，以获取更小的传输延迟和更大的传输带宽。在源码中，NCCL 有许多实现上的细节，本文将分享 NCCL 源码中有关网络传输以及 GDR 部分的相关内容。

GPU Driven 模式

在讲解网络以及 RDMA 部分的实现之前，首先需要了解 GPU Driven 这一核心的思路。由于 CPU 与 GPU 之间存在着通信的延迟，如果使用 CPU 进行数据链路的控制，在数据包较小时，CPU - GPU 中间的通信延迟就会显得不可忽略。以一个经典的 CPU 模型来展示：

GPU Driven 就是绕过 CPU 线程的控制，直接使用 GPU 线程去做数据链路的控制。在具体实现上，CPU Driven 其实就是传统的事件\阻塞模型，当 GPU 数据准备完毕后，通知 CPU 去做一些控制工作；而 GPU Driven 则是通过轮询某个标识位，让控制工作去等待数据准备完毕。这样就能够避免线程调度所带来的不可避免的延迟。听起来，GPU Driven 确实延迟会很低，但是不可避免地其 CPU 占用率会较高。但是当数据的传输工作较为密集时，CPU 占用率的提升就显得可以接收，这正是 NCCL 的应用场景。

作为集合通信库，NCCL 完成每一个集合通信操作，例如 allreduce，都是以小块的形式去发送数据，但是当一个集合通信操作被下发到 NCCL 中时，未来所要完成的所有操作都是确定的了。因此，每一个节点所需要接收和发送的数据量都是已经确定好的，这样就能够实现控制事件提前下发并等待数据准备完毕。

网络通信

代理线程

nccl 中所有网络 IO 都是由代理线程完成的，代理线程会以轮询的模式处理已经下发的网络请求。代理线程接收控制消息的方式与接收下发网络请求的方式所有不同。控制消息是通过 socket 接收的，是一个典型的事件通知模式。

网络请求的下发则是通过队列的模式来完成接收的，proxy 会为每一个 rank 创建一个 op 池，网络请求会被下发到对应的 op pool 中。当有可以处理的网络请求时，代理线程将会优先处理这些请求。当已收割的网络请求处理完毕时，代理线程会尝试从 op pool 中收割全部下发的请求，如果 op pool 中也没有请求，则会进入阻塞状态。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
void* ncclProxyProgress(void *proxyState_) {

  ...

  int lastIdle = 0;
  int proxyOpAppendCounter = 0;
  struct ncclProxyArgs profArgs; // Only used for profiling purposes
  while ((state->stop == false || (state->stop == true && state->active)) && *proxyState->abortFlag == 0) {
    int idle = 1;
    // 处理已经完成收割并且转换为 args 的 op
    ncclResult_t ret = progressOps(proxyState, state, state->active, &idle);
      
    // 因为 pool 是一个临界区，所以这里通过调控 pool 的 access 频率来降低临界竞争
    if (idle || (++proxyOpAppendCounter == ncclParamProgressAppendOpFreq())) {
      int added = 0;
      proxyOpAppendCounter = 0;

      // 尝试收割全部下发的 op，并且转换为 args。如果没有 op 下发，则会进入阻塞
      if (state->stop == false)
        ret = ncclProxyGetPostedOps(proxyState, &added);
      
      if (added == 0) {
        sched_yield(); // No request progressed. Let others run.
      }
    }
    lastIdle = idle;
  }
  return NULL;
}

网络收发中的代理线程在总体上是轮询 + 事件通知的模式。当没有下发的任务时，线程会进入阻塞。这里为了避免竞争，NCCL 也做出了优化。

网络事件的下发

NCCL 中的集合操作会经过一个很长的调用链去转换为网络收发操作，这里先忽略其中的转换。在使用以下参数运行all_reduce_perf时，我们观察网络不同的事件下发：

1
./all_reduce_perf -b 4M -e 4M -f 2 -g 1 -t 1 -n 20

这个 allreduce 操作中，每个节点所需要发送的数据量为 4M，数据发送的类型为 float32，因此一共需要发送 1M 个数据点。通过NCCL 中预先埋下的日志分析，我们可以看到，allreduce 最终被分解为两次网络发送和两次网络接收请求。其中，单次网络发送数据的数据量为 2M。allreduce 操作虽然需要经过很多次网络传输，但是 NCCL 只是将其分解为了 reduce-scatter 和 all-gather 这两个操作，并将这两个操作需要的总发送数据量下发到网络发送请求中。下发的网络事件不会立刻进行数据传输，而是由代理线程轮询并等待数据准备完毕。

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
static ncclResult_t sendProxyProgress(struct ncclProxyState* proxyState, struct ncclProxyArgs* args) {
  // --发送完毕的-- done --正在发送的-- transimitted --还没有发送的-- posted ----

    // 检查是否有可以下发的请求
      if (sub->posted < sub->nsteps && sub->posted < sub->done + maxDepth) {
       // Post buffers to the GPU
          ...
      }
    // 如果有下发的网络发送请求，则检查当前是否可以发送数据
    if (sub->transmitted < sub->posted && sub->transmitted < sub->done + NCCL_STEPS) {
		
        // 检查标识位是否允许发送
       if (sizesFifo[buffSlot] != -1 && ((*recvTail > (sub->base+sub->transmitted)) || p == NCCL_PROTO_LL)){ 
       	// STEP1 : 检查数据是否准备完毕
           
        // STEP2 : 发送数据
       
       }
    }
    if (sub->done < sub->transmitted) {
        // 检查是否有数据发送完毕，这一步主要是由于 RDMA 通信的特点
      }
}

这种网络事件的设计模式正是 NCCL 中 GPU Driven 的一个具体体现。当 GPU Kernel 函数完成数据的准备后，就可以通过修改队列尾的方式激活已经提前下发好的网络请求了。

RDMA 通信

与 rsocket 等通用的网络库不同，NCCL 中的 RDMA 通信也针对应用场景采用了不同的设计。使用 RDMA 通信，一个原则性要求是发送方在 post_send 之前，必须要确保接收方已经使用 post_receive_wr 下发了一个接收操作，这样才能够保证消息的正确传递。正常情况下，RDMA 网络库都会在初始化时下发一部分接收操作，并且采用计时的方式进行流量控制。而 NCCL 中则采取了不同的模式——接收方通知发送方进行发送，这是同样也是因为当一个操作被下发之前，所有的网络 IO 次数以及数量都是确定好的。

缓冲区的分配

控制缓冲区

控制缓冲区是用来调控收发请求的，一次操作对应一个槽。这部分是注册在 cudaMallocHost 中的锁页内存中，因为 GPU 和 CPU 都需要去访问这一块缓冲区。

收发缓冲区

这部分缓冲区会根据是否支持 gdr 选择分配在内存或者是显存上。如果分配到显存上面，会使用 gdr ，收发直接到显存上。gdr 优先使用 dmabuf 方式，否则使用 peermem 方式。

RDMA 通信优化

  1. 按组进行接收 comp channel 通知
  2. inline 标识位减少延迟
  3. 减少接收 comp channel 通知次数