1. 版本历史

文档版本

版本

作者

V3.2.0

ECCL 用户使用手册

Enflame Tech

2. 词汇表

名词

描述

ECCL

Enflame Collective Communications Library,燧原科技集合通信库

GCU

General Compute Unit, 燧原科技通用计算单元

CPU

Central Processing Unit,中央处理器

NUMA

Non Uniform Memory Access,非统一内存访问

SMP

Symmetric Multi-Processing,对称多处理器

QPI

Quick Path Interconnect,快速通道互联

UPI

Ultra Path Interconnect,超路径互联

SHM

Shared Memory,共享内存

NIC

Network Interface Controller,网络适配器

RDMA

Remote Direct Memory Access,远程直接数据存取

IB

InfiniBand,无限带宽技术

RoCE

RDMA over Converged Ethernet,基于以太网的远程直接数据存取

3. ECCL 简介

燧原公司集合通信库(Enflame Collective Communications Library, 简称为 ECCL)在基于公司自研的芯片 GCU 上提供了卡间通信的功能。

其主要提供了以下多个通信原语:

  • ecclAllReduce:将多卡上的数据进行规约操作,并将得到的结果保存在每张卡上

  • ecclBroadcast:将一张卡上的数据广播给多张卡

  • ecclReduceScatter:将多卡上的数据进行规约,并将得到的结果均匀分片后分发到每张卡上

  • ecclAllGather:将多卡上的数据聚合一起,每张卡上的原始数据大小需一致,并将得到的结果保存在每张卡上

ECCL 同样支持点对点通信功能:

  • ecclSend:将卡上的数据发送给另一张指定的卡

  • ecclRecv:从指定的卡上接收数据到本卡

4. ECCL 用户使用手册

4.1. ECCL 使用前置准备

  • TopsPlatform,使用燧原公司提供《topsplatform 安装说明手册》进行安装

4.2. ECCL 安装

安装 ECCL 提供的 deb 包(rpm 包):

dpkg -i eccl_*_amd64.deb (rpm -ivh eccl-*.x86_64.rpm)

确认是否安装成功, 请执行:

ls /usr/include/ | grep eccl.h

检查 “eccl.h” 头文件是否存在于目录 /usr/include/ 下

ls /usr/lib/ | grep libeccl.so

检查 “libeccl.so” 是否存在于目录 /usr/lib/ 下

4.3. Docker使用注意事项

通常情况下,docker 容器会默认限制共享内存的申请上限。当在容器中使用 ECCL 的时候,请按需调整容器内共享内存的大小。

下面展示如何在 Docker 启动命令行中添加参数来实现增加共享内存大小,推荐设置为4G的容量:

--shm-size=4g --ulimit memlock=-1

4.4. ECCL 使用方式

以下内容描述了如何使用 ECCL 进行集合通信操作。包括了 ECCL 接口描述、使用概述以及使用示例介绍

ECCL 对外接口

描述了 ECCL 集合通信库的对外接口以及相关的参数信息。包括了通信域生成与管理类接口、通信算子类接口、组操作接口以及点对点通信接口。详情可见《ECCL API 参考》文档。

ECCL 使用

ECCL 支持多种启动方式:

  • 多进程启动,每个进程对应单张 GCU 卡

  • 单进程多线程启动,每个线程对应单张 GCU 卡,详细使用方法可以参考 test_single_process_multi_thread sample 用例

  • 单进程单线程启动,单个线程对应多张 GCU 卡,详细使用方法可以参考 test_single_process_single_thread sample 或者 test_commInitAll 用例(test_commInitAll 用例展示了单线程初始化多卡接口 ecclCommInitAll 的使用方法)

下列代码将会展示如何使用 MPI 启动多进程,每个进程对应单张 GCU 卡,创建一个通信域并在通信域上运行集合通信算子。

  • 检索有关线程的 MPI 信息:

int myRank, nRanks;
MPI_Comm_rank(MPI_COMM_WORLD, &myRank);
MPI_Comm_size(MPI_COMM_WORLD, &nRanks);
  • 选择任一个rank生成一个通信域的 UniqueId,并把 Id 广播给域内其余的卡:

ecclUniqueId id;
if (myRank == 0) ecclGetUniqueId(&id);
MPI_Bcast((void *)&id, sizeof(id), MPI_BYTE, 0, MPI_COMM_WORLD);
  • 根据以上信息可以生成一个通信域:

ecclComm_t comm;
ecclCommInitRank(&comm, nRanks, id, myRank);
  • 在生成的通信域上调用相应的通信算子进行集合通信(以 ecclAllReduce 为例):

ecclAllReduce(..., comm);
  • 集合通信结束时,释放相应的资源并摧毁通信域:

ecclCommDestroy(comm);

组操作

组操作(ecclGroupStart/ecclGroupEnd)可以被用来将多次调用合并为单次调用。当前 ECCL 仅支持在同一个通信域中聚合相同的通信操作以提高性能。

组操作使用示例

聚合通信操作意味着在单次任务下发中,将执行多次通信操作(当前仅支持聚合多个相同通信操作)。组操作对于减少下发开销效果明显,因为其在执行多个通信操作中仅会下发一次。

对通信操作进行聚合可以通过在 ecclGroupStart 和 ecclGroupEnd 区域内多次调用 ECCL 来完成,下列代码将会展示如何在单次下发中聚合多次 AllReduce 操作:

ecclGroupStart();
for (int i = 0; i < nlayers; ++i) {
  ecclAllReduce(sendBuffs[i], recvBuffs[i], count, dataType, redOp,
                comm, topsStream);
}
ecclGroupEnd();

点对点通信

点对点通信可以指定通信域内任意两张卡之间进行数据接发。该通信方式由两个举动组成:发送端调用 ecclSend(),相应的接收端调用 ecclRecv()。此时两端传入的数据类型和位数参数需要一致。

将针对不同节点的多次点对点通信调用跟组操作结合在一起,能够完成更加复杂的通信操作,下列代码将会展示如何在单次下发中聚合多次点对点通信:

ecclGroupStart();
for (int r=0; r<nlayers; r++) {
  ecclSend(sendbuff[r], sendcount, sendtype, peer, comm, topsStream);
  ecclRecv(recvbuff[r], recvcount, recvtype, peer, comm, topsStream);
}
ecclGroupEnd();

Attention

点对点通信中发送端和接收端需要严格对应,即发送端发送的次数要等于接收端接收的次数

运行 ECCL sample

ECCL 软件包安装完成后,对应的 sample 会被一并安装至 /usr/src/eccl_samples 目录下。

此次提供了下列多个 demo sample:

  • test_sendrecv

该 demo 提供了单次点对点通信的测试方法,编译命令如下:

cd test_sendrecv && make

运行单机4卡测试命令如下,由参数np指定参与的卡数:

mpirun -np 4 ./test_sendrecv
  • test_group_sendrecv

该 demo 提供了组操作下点对点通信的测试方法,编译命令如下:

cd test_group_sendrecv && make

运行单机4卡测试命令如下,由参数np指定参与的卡数:

mpirun -np 4 ./test_group_sendrecv
  • test_single_process_single_thread

该 demo 提供了单进程单线程模式下集合通信的测试方法,编译命令如下:

cd test_single_process_single_thread && make

运行测试命令如下(测试内部指定运行的卡数,此测试运行卡数为4):

./test_single_process_single_thread
  • test_single_process_multi_thread

该 demo 提供了单进程多线程模式下集合通信的测试方法,编译命令如下:

cd test_single_process_multi_thread && make

运行测试命令如下(测试内部指定运行的卡数,此测试运行卡数为4):

./test_single_process_multi_thread
  • test_commInitAll

该 demo 提供了单线程模式下使用 ecclCommInitAll 接口进行初始化并进行集合通信的测试方法,编译命令如下:

cd test_commInitAll && make

运行测试命令如下(测试内部指定运行的卡数,此测试运行卡数为4):

./test_commInitAll

ECCL 使用示例

以下代码描述了一个完整的工作示例,其中包含多个 MPI 进程,每个进程管理一个 GCU 设备。

示例代码可以使用”gcc”编译器进行编译。需要在编译时链接下列的库:

  • Openmpi,推荐版本4.0.5。

  • ECCL

  • topsruntime

#include <stdio.h>
#include "tops/tops_runtime.h"
#include "eccl.h"
#include "mpi.h"
#include <unistd.h>
#include <stdint.h>
#include <stdlib.h>

#define MPICHECK(cmd) do {                          \
  int e = cmd;                                      \
  if( e != MPI_SUCCESS ) {                          \
    printf("Failed: MPI error %s:%d '%d'\n",        \
        __FILE__,__LINE__, e);   \
    exit(EXIT_FAILURE);                             \
  }                                                 \
} while(0)

#define TOPSCHECK(cmd) do {                         \
  topsError_t e = cmd;                              \
  if( e != topsSuccess ) {                          \
    printf("Failed: Tops error %s:%d '%s'\n",       \
        __FILE__,__LINE__,topsGetErrorString(e));   \
    exit(EXIT_FAILURE);                             \
  }                                                 \
} while(0)


#define ECCLCHECK(cmd) do {                         \
  ecclResult_t r = cmd;                             \
  if (r!= ecclSuccess) {                            \
    printf("Failed, ECCL error %s:%d '%d'\n",       \
        __FILE__,__LINE__,r);   \
    exit(EXIT_FAILURE);                             \
  }                                                 \
} while(0)


static uint64_t getHostHash(const char* string) {
  // Based on DJB2a, result = result * 33 ^ char
  uint64_t result = 5381;
  for (int c = 0; string[c] != '\0'; c++){
    result = ((result << 5) + result) ^ string[c];
  }
  return result;
}

static void getHostName(char* hostname, int maxlen) {
  gethostname(hostname, maxlen);
  for (int i=0; i< maxlen; i++) {
    if (hostname[i] == '.') {
        hostname[i] = '\0';
        return;
    }
  }
}

int main(int argc, char* argv[])
{
  int size = 32*1024*1024;

  int myRank, nRanks, localRank = 0;

  // 初始化MPI
  MPICHECK(MPI_Init(&argc, &argv));
  MPICHECK(MPI_Comm_rank(MPI_COMM_WORLD, &myRank));
  MPICHECK(MPI_Comm_size(MPI_COMM_WORLD, &nRanks));

  // 基于hostname计算localRank, localRank用于指定使用哪个GCU
  uint64_t hostHashs[nRanks];
  char hostname[1024];
  getHostName(hostname, 1024);
  hostHashs[myRank] = getHostHash(hostname);
  MPICHECK(MPI_Allgather(MPI_IN_PLACE, 0, MPI_DATATYPE_NULL, hostHashs, sizeof(uint64_t), MPI_BYTE, MPI_COMM_WORLD));
  for (int p=0; p<nRanks; p++) {
     if (p == myRank) break;
     if (hostHashs[p] == hostHashs[myRank]) localRank++;
  }

  ecclUniqueId id;
  ecclComm_t comm;
  float *sendbuff, *recvbuff;
  topsStream_t s;

  // 获取rank 0的ECCL unique ID 然后广播到其他rank
  if (myRank == 0) ECCLCHECK(ecclGetUniqueId(&id));
  MPICHECK(MPI_Bcast((void *)&id, sizeof(id), MPI_BYTE, 0, MPI_COMM_WORLD));

  // 依据localRank选择一个GCU, 分配资源
  TOPSCHECK(topsSetDevice(localRank));
  TOPSCHECK(topsMalloc(&sendbuff, size * sizeof(float)));
  TOPSCHECK(topsMalloc(&recvbuff, size * sizeof(float)));
  TOPSCHECK(topsStreamCreate(&s));

  // 初始化ECCL
  ECCLCHECK(ecclCommInitRank(&comm, nRanks, id, myRank));

  // 使用ECCL进行集合通信
  ECCLCHECK(ecclAllReduce((const void*)sendbuff, (void*)recvbuff, size, ecclFloat, ecclSum,
        comm, s));

  // stream 同步ECCL集合操作
  TOPSCHECK(topsStreamSynchronize(s));

  // 释放空间
  TOPSCHECK(topsFree(sendbuff));
  TOPSCHECK(topsFree(recvbuff));

  // 销毁 ECCL
  ecclCommDestroy(comm);

  // 销毁MPI
  MPICHECK(MPI_Finalize());

  printf("[MPI Rank %d] Success \n", myRank);
  return 0;
}

4.5. ECCL 环境变量

ECCL_ALGO

用于设置 ECCL 使用的算法,目前支持 Ring, Mesh 和 MeshDirect 三种算法。

该环境变量会影响分布式的性能,建议用户不额外设置此环境变量,ECCL会自适应选择合适的算法类型。

允许传入的值

  • Ring

  • Mesh

  • MeshDirect

ECCL_DEBUG

该变量用于设置 ECCL 打印的 log 级别,WARN、INFO 两种级别。WARN 级别在业务逻辑出现异常时打印,INFO 级别用于用户行为留痕、调用链跟踪。其中 INFO 需与定义子模块的环境变量 ECCL_DEBUG_SUBSYS 配合使用,WARN 会打印所有子模块。

允许传入的值

  • WARN:打开 WARN 级别 log

  • INFO:打开 INFO 级别 log

使用示例

打印 INFO 级别的初始化以及环境变量相关 log (即默认值):

export ECCL_DEBUG=INFO
export ECCL_DEBUG_SUBSYS=INIT,ENV

打印 WARN 级别的所有 log:

export ECCL_DEBUG=WARN
export ECCL_DEBUG_SUBSYS=ALL

ECCL_DEBUG_SUBSYS

用于设置需要打印的子模块,主要包括初始化、环境变量、调用等等子模块,可以设置多个子模块,多个子模块间用”,”隔开,默认值为”INIT,ENV”。关闭子模块,在子模块前加”^”。ECCL_DEBUG 设置为 INFO 时需与 ECCL_DEBUG_SUBSYS 配合使用。

允许传入的值

  • INIT:初始化阶段信息,默认值

  • ENV:环境变量信息,默认值

  • CALL:ECCL 对外接口调用信息

  • ALLOC:ECCL 内部资源分配信息

  • TRANSPORT:ECCL 内部 transport 信息

  • NET:网络相关信息,调用 socket、ibverbs 接口信息

  • TOPO:ECCL 内部拓扑相关信息

  • P2P:ECCL 点对点通信相关信息

  • ALL:ECCL 所有模块信息

使用示例

打印 INFO 级别的初始化以及环境变量相关 log:

export ECCL_DEBUG=INFO
export ECCL_DEBUG_SUBSYS=INIT,ENV

打印 WARN 级别的所有 log:

export ECCL_DEBUG=WARN
export ECCL_DEBUG_SUBSYS=ALL

关闭初始化模块打印:

export ECCL_DEBUG_SUBSYS=^INIT

ECCL_TOPO_DUMP_FILE

给定一个存储XML文件的路径,用于转存探测的物理拓扑链接。

允许传入的值

一个指向某文件的路径,这个文件可能会被创建,或者会被覆盖原有内容。

ECCL_GRAPH_DUMP_FILE

给定一个存储XML文件的路径,用于转存搜索graph的信息。

允许传入的值

一个指向某文件的路径,这个文件可能会被创建,或者会被覆盖原有内容。

ECCL_P2P_DISABLE

该变量用于设置关闭 peer ot peer (P2P) transport。P2P 通信将使用 PCI/ESL 等方式直接访问来实现通信的功能。若关闭 P2P,ECCL 将转而使用 Shared Memory (SHM) 进行通信。

允许传入的值

默认值为0,设置为1则表示关闭。

ECCL_SHM_DISABLE

该变量用于设置关闭 Shared Memory (SHM) transport。SHM 是在设备之间无法进行 P2P 通信时使用的(此时 ECCL_P2P_DISABLE 被设置为 1),其将使用主机内存进行通信。若关闭 SHM,ECCL 将转而使用 IB/RoCE 进行通信。

允许传入的值

默认值为0,设置为1则表示关闭。

ECCL_IB_DISABLE

该变量用于设置关闭IB/RoCE transport,若关闭,ECCL将转而使用IP sockets。

允许传入的值

默认值为0,设置为1则表示关闭。

Attention

transport的优先级为 P2P > SHM > IB/RoCE > IP sockets。

ECCL_IB_HCA

该变量用于设置RDMA传输使用的网口。

允许传入的值

用于筛选RDMA网口的字符串,存在多个筛选条件时用”,”隔开。默认值为空。 接受以下三种形式的字符串(以网卡mlx5_1为例):

  • 网卡名(例:”mlx5_1”)。

  • 网卡名的部分前缀(例:”mlx5”)。

  • 网卡名加物理口序号,网卡与物理口之间用”:”隔开(例:”mlx5_1:1”表示网卡mlx5_1的第一个口)。

字符串中可增加”^”用于排除,”=”用于精确选择。在字符串为网卡名时,”=”可忽略。

使用示例

选择所有以mlx5开头的口,例如mlx5_0、mlx5_1等:

export ECCL_IB_HCA="mlx5"

选择网卡mlx5_1上的所有的口,例如mlx5_1:1、mlx5_1:2等:

export ECCL_IB_HCA="mlx5_1"ECCL_IB_HCA="=mlx5_1"

使用网卡mlx5_0、mlx5_1的第一个口:

export ECCL_IB_HCA="=mlx5_0:1,mlx5_1:1"

不选择mlx5_1上的口:

export ECCL_IB_HCA="^mlx5_1"ECCL_IB_HCA="^=mlx5_1"

不选择mlx5_1上的第一个口:

export ECCL_IB_HCA="^=mlx5_1:1"

ECCL_NET_GDR_LEVEL

该变量定义NIC和GCU之间的最大距离,用于精细控制NIC和GCU之间何时使用GCU Direct RDMA。如用户不指定,ECCL将根据架构以及环境选择最优的值。

允许传入的值

允许传入字符串形式的值:

  • LOC : 不使用GCU Direct RDMA。

  • PIX : GCU和NIC连接同一个PCI交换机,使用GCU Direct RDMA。

  • PXB : GCU和NIC通过多个PCI交换机相连(多跳),使用GCU Direct RDMA。

  • PHB : GCU和NIC位于同一个NUMA节点下,使用GCU Direct RDMA,此时流量将通过CPU。

  • SYS : GCU和NIC跨NUMA节点通过SMP互联(如QPI、UPI),也使用GCU Direct RDMA。

也允许传入数字形式的值,与字符串的对应关系如下:

  • 0 : LOC

  • 1 : PIX

  • 2 : PXB

  • 3 : PHB

  • 4 : SYS

Attention

设置的值大于4时,同样按照”SYS”处理。

ECCL_NET_GDR_READ

该变量用于GCU和NIC在ECCL_NET_GDR_LEVEL指定的距离内时,发送数据开启GCU Direct RDMA。

Attention

已知在某些平台(如:PCI-E)上,发送数据直接从GCU内存读取比从CPU内存读取稍慢。

允许传入的值

0或者1。设置为1时使用GCU Direct RDMA将数据直接发送到NIC(即绕过CPU),设置为0时关闭使用GCU Direct RDMA。

5. eccl-tests

用来测试检查ECCL操作的性能和正确性。

5.1. 环境构建

eccl-tests安装包依赖eccl安装包,在安装eccl-tests之前需要先安装eccl包,见本文档 ECCL安装 部分。

eccl-tests安装

dpkg -i eccl-tests_*_amd64.deb (rpm -ivh eccl-tests-*.x86_64.rpm)

确认是否安装成功, 请执行:

ls /usr/local/bin/ | grep eccl

检查eccl_all_reduce_perf、eccl_all_gather_perf、eccl_broadcast_perf、eccl_reduce_per、eccl_reduce_scatter_perf、eccl_sendrecv_perf、eccl_gather_perf、eccl_scatter_perf、eccl_alltoall_perf 可执行文件是否存在于目录 /usr/local/bin 下。

5.2. 测试使用

eccl-tests运行在多张GCU卡上,当前支持通过mpirun命令启动程序并指定启动的进程的个数,每个进程对应着一个唯一rank。目前支持一个进程管理一块GCU,在一个进程上有一个线程。

5.3. 使用方法

eccl-tests可以用mpirun中的参数np来控制参与的卡数,并支持多种通信算子的性能测试,包括AllReduce、Broadcast、ReduceScatter、AllGather、Reduce、SendRecv、Gather、Scatter、AlltoAll。当前集合通信只支持单机,不支持多机。

下面提供了测试各种通信算子的命令格式。

AllReduce

单机8卡:

mpirun -np 8 eccl_all_reduce_perf -b 1k -e 128M -f 2

Broadcast

单机8卡:

mpirun -np 8 eccl_broadcast_perf -b 1k -e 128M -f 2

ReduceScatter

单机8卡:

mpirun -np 8 eccl_reduce_scatter_perf -b 1k -e 128M -f 2

AllGather

单机8卡:

mpirun -np 8 eccl_all_gather_perf -b 1k -e 128M -f 2

Reduce

单机8卡:

mpirun -np 8 eccl_reduce_perf -b 1k -e 128M -f 2

SendRecv

单机8卡:

mpirun -np 8 eccl_sendrecv_perf -b 1k -e 128M -f 2

Gather

单机8卡:

mpirun -np 8 eccl_gather_perf -b 1k -e 128M -f 2

Scatter

单机8卡:

mpirun -np 8 eccl_scatter_perf -b 1k -e 128M -f 2

AlltoAll

单机8卡:

mpirun -np 8 eccl_alltoall_perf -b 1k -e 128M -f 2

5.4. 命令参数

Number of GCUs

表 5.4.2 GCU及线程数量

参数名称

使用方法

默认值

参数说明

-t

--nthreads <num threads>

1(仅支持1)

运行在每一个进程上的线程数

-g

--ngcus <GCUs per thread>

1(仅支持1)

运行在每一个线程上的GCU个数

Attention

目前仅支持一个进程对应一个GCU启动

Sizes to scan

表 5.4.3 需要扫描的尺寸

参数名称

使用方法

默认值

参数说明

-b

--minbytes <min size in bytes>

32M

启动时的最小尺寸

-e

--maxbytes <max size in bytes>

32M

结束时的最大的尺寸

-i

--stepbytes <increment size>

1M

在一定的尺度等量选择增量

-f

--stepfactor <increment factor>

不使用

乘法的使用

Attention

-i和-f只使用一个,默认使用-i

ECCL operations arguments

表 5.4.4 ECCL operations argument

参数名称

使用方法

默认值

参数说明

-o

--op <sum/prod/min/max/all>

sum

指定要执行的规约操作类型。仅与AllReduce、Reduce或ReduceScatter等源操作相关

-d

--datatype <eccltype>

float32

指定数据类型

-r

--root <root/all>

0

指定 root GCU,应用于Broadcast或者Reduce相关操作

Attention

-d 可指定的类型有 <int8/uint8/int32/uint32/int64/uint64/float16/float32/float64/bfloat16>。AllReduce,ReduceScatter和Reduce当前不支持 64位的数据类型: int64,uint64和float64

Performance

表 5.4.5 Performance

参数名称

使用方法

默认值

参数说明

-n

--iters <iteration count>

20

循环迭代次数

-w

--warmup_iters <warmup iteration count>

5

预热迭代次数(不是时间)

-m

--agg_iters <aggregation count>

1

每次迭代中聚合在一起的操作数

-a

--average <0/1/2/3>

1

报告所有rank的平均性能(仅MPI=1)<0=Rank0,1=Avg,2=Min,3=Max>

Test operation

表 5.4.6 Test operation

参数名称

使用方法

默认值

参数说明

-c

--check <0/1>

1

检查结果的正确性。在大量GCU上可能较慢,默认检查

-z

--blocking <0/1>

0

进行ECCL集合通信阻塞,即让CPU在每次集体阻塞后等待和同步

5.5. 结果

表 5.5.1 eccl-tests性能相关名词解释

参数名称

参数说明

size

数据的大小(S)

count

数据的计数

type

数据的类型

redop

操作的类型

time

操作的执行时间(t)

algbw

算法带宽,根据数据大小以及操作执行时间计算得到,algbw=S/t

busbw

总线带宽,根据数据链路上实际传输的数据大小以及操作执行时间计算得到。实际是根据操作类型,将算法带宽乘以一个系数,详见下表

表 5.5.2 不同操作类型总线带宽与算法带宽的关系

操作的类型

计算关系

AllReduce

busbw=algbw*2*(n-1)/n

ReduceScatter

busbw=algbw*(n-1)/n

AllGather

busbw=algbw*(n-1)/n

Broadcast

busbw=algbw

Reduce

busbw=algbw

SendRecv

busbw=algbw

Gather

busbw=algbw*(n-1)/n

Scatter

busbw=algbw*(n-1)/n

AlltoAll

busbw=algbw*(n-1)/n

5.6. 注意事项

打开调试开关(ECCL_DEBUG,ECCL_DEBUG_SUBSYS)可能导致eccl-tests测得的性能低于实际性能,建议在进行eccl-tests性能测试时不要设置调试相关环境变量。