1. 版本历史

文档版本

日期

新增功能

V3.1.0

2024年05月21日

支持各种传输通道

V3.2.0

2024年06月20日

支持avg redop

V3.5.1

2025年06月23日

增加enable ll128环境变量

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)

确认是否安装成功， 请执行：

dpkg -l | grep eccl (rpm -qa eccl)

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

ls /usr/include/ | grep eccl.h

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

ls /usr/lib/ | grep libeccl.so

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 卡

下列代码将展示如何使用 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);

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

代码展示如下：

#include <cstdio>
#include <thread>
#include <cassert>
#include "eccl.h"

#define ECCLCHECK(call) do { \
  ecclResult_t res = call; \
  if (res != ecclSuccess) { \
    return res; \
  } \
} while (0);

struct ecclTestArgs {
  int rank;
  int nranks;
  ecclUniqueId id;
};

#define N_TEST_RANKS 4
thread_local pthread_t ecclTestThreads[N_TEST_RANKS];
thread_local struct ecclTestArgs ecclThreadArgs[N_TEST_RANKS];

void* ecclTestThread(void* args_) {
  struct ecclTestArgs* args = (struct ecclTestArgs*)args_;
  int myRank = args->rank;
  int nRanks = args->nranks;
  ecclUniqueId id;
  ecclComm_t comm;

  memcpy(&id, &args->id, sizeof(id));

  printf("ecclTestThread myRank %d, nRanks%d\n", myRank, nRanks);
  assert(topsSetDevice(myRank) == topsSuccess);
  assert(ecclCommInitRank(&comm, nRanks, id, myRank) == ecclSuccess);

  using T = int;
  const size_t TEST_SIZE = 4 * 1024 * 1024;
  size_t count = TEST_SIZE / sizeof(T);

  T *srcH, *dstH;
  // Alloc host buffer
  srcH = (T *)malloc(TEST_SIZE);
  dstH = (T *)malloc(TEST_SIZE);

  // Init host buffer
  for (int i = 0; i < count; ++i) {
    srcH[i] = i;
    dstH[i] = -1;
  }

  // Alloc device buffer
  void *sendBuff, *recvBuff;
  assert(topsMalloc(&sendBuff, TEST_SIZE) == topsSuccess);
  assert(topsMalloc(&recvBuff, TEST_SIZE) == topsSuccess);

  assert(
        topsMemcpy(sendBuff, srcH, TEST_SIZE, topsMemcpyHostToDevice) == topsSuccess);
  assert(
        topsMemcpy(recvBuff, dstH, TEST_SIZE, topsMemcpyHostToDevice) == topsSuccess);

  topsStream_t topsStream;
  assert(topsStreamCreate(&topsStream) == topsSuccess);
  ecclDataType_t dataType = ecclInt32;
  ecclRedOp_t redOp = ecclSum;

  assert(ecclAllReduce(sendBuff, recvBuff, count, dataType, redOp,
                            comm, topsStream) == ecclSuccess);
  assert(topsStreamSynchronize(topsStream) == topsSuccess);

  assert(
        topsMemcpy(dstH, recvBuff, TEST_SIZE, topsMemcpyDeviceToHost) == topsSuccess);
  for (int j = 0; j < count; ++j) {
    if (abs(dstH[j] - srcH[j] * N_TEST_RANKS) > 0.00001) {
      printf("Check failed: at srcH[%d]=%d, dstH[%d]=%d\n",
              j, srcH[j], j, dstH[j]);
      assert(false);
    }
  }

  // Release resources
  assert(topsStreamSynchronize(topsStream) == topsSuccess);
  assert(topsFree(sendBuff) == topsSuccess);
  assert(topsFree(recvBuff) == topsSuccess);
  free(srcH);
  free(dstH);
  assert(topsStreamDestroy(topsStream) == topsSuccess);
  assert(ecclCommDestroy(comm) == ecclSuccess);

  printf("Test single process multiple thread allreduce passed!\n");
  return args;
}

int main(int argc, char **argv) {

  ecclUniqueId id;
  ECCLCHECK(ecclGetUniqueId(&id));

  for(int i= 0 ; i < N_TEST_RANKS; i++){
    ecclThreadArgs[i].rank = i;
    ecclThreadArgs[i].nranks = N_TEST_RANKS;
    memcpy(&ecclThreadArgs[i].id, &id, sizeof(id));
    pthread_create(ecclTestThreads+i, NULL, ecclTestThread, &ecclThreadArgs[i]);
  }

  for(int i= 0 ; i < N_TEST_RANKS; i++){
    int err = pthread_join(ecclTestThreads[i], NULL);
    if (err != 0) {
      printf("Error waiting for %d pthread_join : %s\n", i, strerror(errno));
      return -1;
    }
  }

  return 0;
}

  • 单进程单线程启动，单个线程对应多张 GCU 卡

代码展示如下：

#include <cstdio>
#include <thread>
#include "eccl.h"

#define TOPSCHECK(cmd)                                      \
{                                                            \
  topsError_t error  = cmd;                                  \
  if (error != topsSuccess) {                                \
    fprintf(stderr, "error: '%s'(%d) at %s:%d\n",            \
      topsGetErrorString(error), error, __FILE__, __LINE__); \
    exit(EXIT_FAILURE);                                      \
    exit(EXIT_FAILURE);                                      \
  }                                                          \
}

#define ECCLCHECK(call) do { \
  ecclResult_t res = call; \
  if (res != ecclSuccess) { \
    return res; \
  } \
} while (0);

int main(int argc, char **argv) {

  const int N_TEST_RANKS = 4;

  // Init eccl library
  ecclComm_t comm[N_TEST_RANKS];
  ecclUniqueId id;
  ECCLCHECK(ecclGetUniqueId(&id));

  ECCLCHECK(ecclGroupStart());
  for (int i = 0; i < N_TEST_RANKS; ++i) {
    TOPSCHECK(topsSetDevice(i));
    ECCLCHECK(ecclCommInitRank(&comm[i], N_TEST_RANKS, id, i));
  }
  ECCLCHECK(ecclGroupEnd());

  using T = int;
  const size_t TEST_SIZE = 4 * 1024 * 1024;
  size_t count = TEST_SIZE / sizeof(T);

  T *srcH, *dstH;
  // Alloc host buffer
  srcH = (T *)malloc(TEST_SIZE);
  dstH = (T *)malloc(TEST_SIZE);

  // Init host buffer
  for (int i = 0; i < count; ++i) {
    srcH[i] = i;
    dstH[i] = -1;
  }

  // Alloc device buffer
  void *sendBuffs[N_TEST_RANKS], *recvBuffs[N_TEST_RANKS];
  for (int i = 0; i < N_TEST_RANKS; ++i) {
    TOPSCHECK(topsSetDevice(i));
    TOPSCHECK(topsMalloc(&sendBuffs[i], TEST_SIZE));
    TOPSCHECK(topsMalloc(&recvBuffs[i], TEST_SIZE));
  }

  // Init device buffer, copy data from host to device
  for (int i = 0; i < N_TEST_RANKS; ++i) {
    TOPSCHECK(topsSetDevice(i));
    TOPSCHECK(
        topsMemcpy(sendBuffs[i], srcH, TEST_SIZE, topsMemcpyHostToDevice));
    TOPSCHECK(
        topsMemcpy(recvBuffs[i], dstH, TEST_SIZE, topsMemcpyHostToDevice));
  }

  topsStream_t topsStream[N_TEST_RANKS];
  for(int i = 0; i < N_TEST_RANKS; ++i) {
    TOPSCHECK(topsSetDevice(i));
    TOPSCHECK(topsStreamCreate(&topsStream[i]));
  }

  ecclDataType_t dataType = ecclInt32;
  ecclRedOp_t redOp = ecclSum;

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

  for (int i = 0; i < N_TEST_RANKS; ++i) {
    TOPSCHECK(topsSetDevice(i));
    TOPSCHECK(topsStreamSynchronize(topsStream[i]));
  }

  // Check result
  for (int i = 0; i < N_TEST_RANKS; ++i) {
    TOPSCHECK(topsSetDevice(i));
    TOPSCHECK(
        topsMemcpy(dstH, recvBuffs[i], TEST_SIZE, topsMemcpyDeviceToHost));
    for (int j = 0; j < count; ++j) {
      if (abs(dstH[j] - srcH[j] * N_TEST_RANKS) > 0.00001) {
        printf("Check failed: at recvBuffs[%d] srcH[%d]=%d, dstH[%d]=%d\n", i,
               j, srcH[j], j, dstH[j]);
        TOPSCHECK(topsErrorUnknown);
      }
    }
  }

  // Release resources
  for (int i = 0; i < N_TEST_RANKS; ++i) {
    TOPSCHECK(topsFree(sendBuffs[i]));
    TOPSCHECK(topsFree(recvBuffs[i]));
    TOPSCHECK(topsStreamDestroy(topsStream[i]));
    ECCLCHECK(ecclCommDestroy(comm[i]));
  }
  free(srcH);
  free(dstH);

  printf("Test single process single thread allreduce passed!\n");

  return 0;
}

组操作

组操作（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_PROTO

用于设置 ECCL 使用的协议，目前支持 LL128，Simple。

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

注意： 由于当前LL128使用会存在一些限制，且在某些CPU系统下可能有潜在的mismatch问题, 所以LL128 协议是默认disable的。如果用户经过严格的验证明确在某些CPU下不会出现问题。那么推荐用户设置 export ECCL_PROTO="LL128" 使能LL128协议。对于不支持的场景，在ECCL内部会基于当前的环境变量设置自动tune到一个合适的算法协议。

LL128只支持allreduce原语，只支持单机场景，不支持SHM/NET链路，不支持较大数据量的传输。对于不支持的场景，会添加如下log的打印:

WARN : no algorithm/protocol available for function ...
Redirect to  ...

...表示这里省略了其他内容。

允许传入的值

  • LL128

  • Simple

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 to 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_GID_INDEX

该变量用于设置RoCE模式中使用的Global ID index（可通过show_gids命令查看），每个端口具有一个GUID（Globally Unique Identifier），GUID是全局唯一的，类似于以太网MAC地址。

允许传入的值

默认值为-1，可设置为相应的Global ID index。

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。

ECCL_SOCKET_NTHREADS

该变量用于设置使用socket连接时CPU创建的线程个数，有可能会提高socket传输性能，以cpu利用率为代价。该环境变量会影响socket通信的性能。

允许传入的值

与ECCL_NSOCKS_PERTHREAD配合使用，取值范围是1至16，默认值为4。

使用示例

使用socket连接时CPU创建的线程个数2：

export ECCL_SOCKET_NTHREADS=2

ECCL_NSOCKS_PERTHREAD

该变量用于设置使用socket连接时单个线程创建的socket个数。在单个socket传输速率限制的情况下，设置>1可能会提高性能。该环境变量会影响socket通信的性能。

允许传入的值

与ECCL_SOCKET_NTHREADS配合使用，默认值为1。总的socket个数不超过64(即ECCL_NSOCKS_PERTHREAD*ECCL_SOCKET_NTHREADS<=64)。

使用示例

使用socket连接时CPU的单个线程创建的socket个数2：

export ECCL_NSOCKS_PERTHREAD=2

ECCL_SOCKET_IFNAME

该变量用于设置socket通信使用的ip interface。

允许传入的值

用于筛选ip interface的字符串前缀。"^"用于排除，"="用于精确选择。默认值为空。

使用示例

选择所有以eth开头的ip interface，例如eth0、eth1等：

export ECCL_SOCKET_IFNAME="eth"

只选择eth0：

export ECCL_SOCKET_IFNAME="=eth0"

不选择所有以eth开头的ip interface：

export ECCL_SOCKET_IFNAME="^eth"

不选择eth0：

export ECCL_SOCKET_IFNAME="^=eth0"

5. eccl-tests

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

5.1. eccl-tests安装

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

eccl-tests安装

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

确认是否安装成功， 请执行：

dpkg -l | grep eccl-tests (rpm -qa eccl-tests)

检查eccl_all_reduce_perf、eccl_all_gather_perf、eccl_broadcast_perf、eccl_reduce_perf、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，在一个进程上有一个线程。eccl-tests目前支持多种通信算子的性能测试，包括AllReduce、Broadcast、ReduceScatter、AllGather、Reduce、Gather、Scatter、SendRecv、AlltoAll，当前AllReduce、Broadcast、ReduceScatter、AllGather、Reduce、Gather、Scatter支持单机和多机，SendRecv、AlltoAll仅支持单机。

5.3. 命令参数

Number of GCUs

表 5.3.5 GCU及线程数量

参数名称

使用方法

默认值

参数说明

-t

--nthreads <num threads>

1（仅支持1）

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

-g

--ngcus <GCUs per thread>

1（仅支持1）

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

Attention

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

Sizes to scan

表 5.3.6 需要扫描的尺寸

参数名称

使用方法

默认值

参数说明

-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.3.7 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.3.8 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.3.9 Test operation

参数名称

使用方法

默认值

参数说明

-c

--check <0/1>

1

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

-z

--blocking <0/1>

0

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

5.4. 单机测试示例

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.5. 多机测试准备以及示例

eccl-tests同样支持多机测试，推荐在单机环境搭建测试成功后，再进行多机的相关测试。多机测试依赖测试环境的配置，以及设置一些环境变量，可按照如下步骤进行检查、测试（以下均以两机十六卡测试为例）：

检查eccl、eccl-tests安装情况

多机测试需要待测的所有机器都已经安装TopsPlatform、eccl以及eccl-tests，安装以及检查可参考本文档ECCL安装以及eccl-tests安装部分，若要使能peer mem测试，则需要安装enflame_peer_mem，可以参考《topsplatform 安装说明手册》进行安装。

检查网络配置情况

eccl以及eccl-tests的测试依赖环境的网络配置，用户需进行一些相关检查以保证网络正常。目前燧原公司机器默认安装Mellanox网卡，测试前需确认网卡个数、驱动安装、以及网卡状态等情况，其他厂商网卡命令会有所不同；推荐不同机器使用相同个数的网卡，若不同，需要配合环境变量ECCL_IB_HCA使用，以下是一些基本检查步骤:

网卡的物理情况

以Mellanox网卡为例，检查当前设备Mellanox网卡情况，命令：

lspci -nn | grep Eth | grep Mellanox
#部分参考输出，以下结果表示设备具备两张Mellanox系列网卡
#型号是ConnectX-5；若无显示，则需要插上相关网卡
Ethernet controller : Mellanox Technologies MT28800 Family [ConnectX-5 Ex]
Ethernet controller : Mellanox Technologies MT28800 Family [ConnectX-5 Ex]

也可通过：

lspci -t -v

来查看网卡设备在PCI总线上的位置、个数、是否有多品牌网卡，以及包括在具体哪个numa节点，连接CPU、PCIe switch的情况，网卡的不同位置对实际性能会有较大影响。

网卡驱动

以Mellanox网卡为例，确认是否已经正确安装OFED驱动，命令：

ofed_info -s
#参考输出，以下结果表示版本为5.4-3.1.0.0；若无显示，则需要安装OFED驱动
MLNX_OFED_LINUX-5.4-3.1.0.0:

网卡接口

可通过以下命令查看所有网络接口，包含普通以太网口，以及rdma相关网口；

ifconfig

网卡状态

#查看网口与网卡映射情况
ibdev2netdev
#参考输出，以下结果表示mlx5_0和mlx5_2对应端口为ib0和ib2，均为up状态
mlx5_0 port 1 ==> ib0 (Up)
mlx5_2 port 1 ==> ib2 (Up)

#查看网卡工作情况，命令：
ibstatus
#参考输出，以下结果表示mlx5_0和mlx5_2均为ACTIVE状态，速率为100Gb/s
#工作在Ethernet模式（也可工作在InfiniBand模式，相关网卡需要保持一致）
Infiniband device 'mlx5_0' port 1 status:
        default gid:     fe80:0000:0000:0000:0e42:a1ff:fe60:0a88
        base lid:        0x0
        sm lid:          0x0
        state:           4: ACTIVE
        phys state:      5: LinkUp
        rate:            100 Gb/sec (4X EDR)
        link_layer:      Ethernet

Infiniband device 'mlx5_2' port 1 status:
        default gid:     fe80:0000:0000:0000:0e42:a1ff:fe60:0408
        base lid:        0x0
        sm lid:          0x0
        state:           4: ACTIVE
        phys state:      5: LinkUp
        rate:            100 Gb/sec (4X EDR)
        link_layer:      Ethernet

若网口以及网卡工作状态不符合预期，首先检查相关物理链路是否连接正常，以及检查是否为端口设置ip地址，ifconfig可查看，若无ip地址，则需要配置ip地址。

网卡实际速率测试

可通过ib_send_bw、ib_write_bw等命令测试网卡的实际收发速率：

#server端命令，mlx5_1为server端待测网卡
ib_write_bw --ib-dev=mlx5_1 -a  --report_gbits

#client端命令，mlx5_2为client端待测网卡
#a.b.c.d为server端的待测网卡对应网口ip地址
ib_write_bw --ib-dev=mlx5_2 -a  --report_gbits a.b.c.d -F

实际的收发速率若与网卡标称速率不符，或远低于标称速率则需要进行相关检查或者更换网卡；若测试失败，则需要检查物理连接以及网络配置。

mpirun相关配置

本手册使用mpirun启动进行测试，涉及到mpirun启动命令，以下是部分必须的设置，测试当中更多的mpirun启动问题可自行搜索mpirun相关文档。

免密登录配置

查看待测试的机器是否可以互相ssh免密登录，若无法免密登录，则：

#生成密钥
ssh-keygen
#拷贝至对应服务器，username为用户名，a.b.c.d为对应ip地址
ssh-copy-id username@a.b.c.d

hostfile配置

多机测试需要使用hostfile文件，用于指定slot的数量，在每台测试机的相同路径创建该文件，文件名为hostfile（文件名可自定义），并输入如下内容：

a.b.c.d slots=8
a.b.c.e slots=8

以上hostfile文件是以两机十六卡为例，a.b.c.d和a.b.c.e分别为两台机器的ip地址，slots=8指定每台机器使用八张GCU，也可根据实际情况设置GCU个数。

指定网络接口

mpirun通常不需要直接指定使用哪个物理网卡，会自动选择最优的通信路径。但推荐通过环境变量来设置，以下环境变量会限制TCP通信仅使用特定网络接口：

export OMPI_MCA_btl_tcp_if_include=${name}

这里的${name}是假设的网卡接口名，实际使用中需替换为系统中对应的接口名。

推荐配置的环境变量

多机测试默认需要配置一些网络相关环境变量，否则会导致测试失败，如下列举了推荐使用的环境变量，环境变量详细使用可参考环境变量章节。

ECCL_SOCKET_IFNAME

该变量用于设置socket通信使用的ip interface，使用mpirun进行eccl初始化时候会搜索可使用的ip interface，并且选择搜索到的第一个端口，不同机器搜索到的端口可能会不一致，从而导致初始化失败。可以ifconfig查看所有机器上都有的ip interface，并设置进行测试；例如所有机器上都存在端口eth0，则设置：

export ECCL_SOCKET_IFNAME="eth0"

ECCL_IB_HCA

该变量用于设置RDMA传输使用的网口，不同机器上的网卡数目、网卡设备名可能不一致，推荐设置从而精确选择网卡；例如所有机器上都存在网卡mlx5_0，则设置：

export ECCL_IB_HCA="mlx5_0"

ECCL_IB_GID_INDEX

该变量用于设置RoCE模式中使用的Global ID index，可通过show_gids命令查看，推荐使用RoCEv2对应的index（依赖物理环境的组网，需要先确认网络可通）；例如RoCEv2对应的index为3，则设置：

export ECCL_IB_GID_INDEX=3

执行测试

在其中一台待测试机器执行命令，相关测试的具体命令如下，其中： "-np 16" 中的16表示测试一共使用16张卡，其中每台机器各八张卡； "--mca btl_tcp_if_include ${name}" 表示使用名为${name}的ip interface作为mpi tcp的通信端口，需保证${name}为两台机器的可用端口，可与ECCL_SOCKET_IFNAME环境变量设置一致，若不同机器的端口不一致，则需要按照mpirun指定网络接口章节设置； "-hostfile hostfile" 中第二个hostfile为设置所用hostfile文件的文件名，用户可自定义。

AllReduce

mpirun -np 16 -hostfile hostfile --allow-run-as-root -bind-to none -map-by slot --mca plm_rsh_args "-p 2223" --mca btl_tcp_if_include ${name} eccl_all_reduce_perf -b 1k -e 128M -f 2

Broadcast

mpirun -np 16 -hostfile hostfile --allow-run-as-root -bind-to none -map-by slot --mca plm_rsh_args "-p 2223" --mca btl_tcp_if_include ${name} eccl_broadcast_perf -b 1k -e 128M -f 2

ReduceScatter

mpirun -np 16 -hostfile hostfile --allow-run-as-root -bind-to none -map-by slot --mca plm_rsh_args "-p 2223" --mca btl_tcp_if_include ${name} eccl_reduce_scatter_perf -b 1k -e 128M -f 2

AllGather

mpirun -np 16 -hostfile hostfile --allow-run-as-root -bind-to none -map-by slot --mca plm_rsh_args "-p 2223" --mca btl_tcp_if_include ${name} eccl_all_gather_perf -b 1k -e 128M -f 2

Reduce

mpirun -np 16 -hostfile hostfile --allow-run-as-root -bind-to none -map-by slot --mca plm_rsh_args "-p 2223" --mca btl_tcp_if_include ${name} eccl_reduce_perf -b 1k -e 128M -f 2

Gather

mpirun -np 16 -hostfile hostfile --allow-run-as-root -bind-to none -map-by slot --mca plm_rsh_args "-p 2223" --mca btl_tcp_if_include ${name} eccl_gather_perf -b 1k -e 128M -f 2

Scatter

mpirun -np 16 -hostfile hostfile --allow-run-as-root -bind-to none -map-by slot --mca plm_rsh_args "-p 2223" --mca btl_tcp_if_include ${name} eccl_scatter_perf -b 1k -e 128M -f 2

5.6. 结果

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

参数名称

参数说明

size

数据的大小(S)

count

数据的计数

type

数据的类型

redop

操作的类型

time

操作的执行时间(t)

algbw

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

busbw

总线带宽，根据数据链路上实际传输的数据大小以及操作执行时间计算得到。实际是根据操作类型，将算法带宽乘以一个系数，详见下表
表 5.6.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.7. 注意事项

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