配置P4项目基于Mininet & BMv2

环境说明

我使用配置完成的虚拟机,下载到的是一个2.9G的ova文件(Open Virtual Appliance,就是一种虚拟机文件),打开virtualBox点导入,等它导入,然后选择用户p4,密码是p4,登录。

我也想根P4官方仓库配置,但屡屡失败,我们还是放过自己吧

虚拟机里的终端很难用,如果愿意的话可以使用SSH远程连接,以及设置共享文件夹来用本地ide写代码,不过这并不是本文重点,可自行尝试。

从Makefile了解整个流程

首先要知道一个P4项目分为哪些文件

手动编写的文件(输入):

  • *.p4:数据面逻辑。可以写多个文件,include。
  • topology.json:规定这个模拟的网里有多少主机交换机,如何连线的。
  • s1-commands.txt:静态流表。可不写,平替后文会讲。

编译器生成的文件(中间产物):

  • target.json:BMv2 的处理流程,来自于p4文件的编译。
  • p4info.txt:相当于API的角色,描述了表名、动作名与 ID 的对应关系,供控制器(Python/gRPC)识别。

运行时环境(Mininet所需文件):

  • 需要 topology.json 来创建虚拟网线。
  • 需要 target.json 来初始化每个 simple_switch 进程。

P4的编译

P4编译器是分层的,前端叫做p4c,负责检查语法,后端叫做bm2-ss,负责把语法树转换成特定硬件或软件能懂的格式。bm2-ss 专门针对 BMv2 软件交换机的 Simple Switch 架构。

以上两个合一块就是P4编译器,叫做 p4c-bm2-ss —— p4c-bm2-ss 之于 P4,就像 gcc 之于 C 语言。

在实际开发中,最标准、最常用的写法如下:

p4c-bm2-ss --p4v 16 \
           --p4runtime-files build/target.p4.p4info.txt \
           --p4runtime-format text \
           -o build/target.json \
           ./p4src/source.p4

解释

  • –p4v 16:指定使用的 P4 语言版本。目前主流为 P4-16,若不指定可能按旧版 P4-14 处理。
  • –p4runtime-files :指定生成的 P4Info 文件路径。
  • –p4runtime-format text:规定 P4Info 的排版格式为可读文本。
  • -o :Output。指定生成的 JSON 配置文件路径,这是 BMv2 软件交换机真正运行依赖的逻辑文件。
  • ./p4src/source.p4:主程序入口。编译器会从此文件开始递归处理所有 #include 的头文件。

写到Makefile中,就是

P4C = p4c-bm2-ss
P4_MAIN_SOURCE = ./p4src/source.p4
P4_DEPENDENCIES = (wildcard p4src/*.p4)
BUILD_DIR = build
JSON_FILE =(BUILD_DIR)/target.json
INFO_FILE = (BUILD_DIR)/target.p4.p4info.txt


build:(P4_DEPENDENCIES)
    mkdir -p (BUILD_DIR)(P4C) --p4v 16 --p4runtime-files (INFO_FILE) -o(JSON_FILE) $(P4_MAIN_SOURCE)

P4_MAIN_SOURCE 和 P4_DEPENDENCIES 分开写的原因是要既满足编译器(源文件位置只接收最主要的那个),又要满足 make 的检查逻辑(全部源文件都是依赖,需检查时间戳)

mininet的运行

run_exercise.py 稍后仔细解释,这里只需将其理解为一个黑盒,功能是把之前准备的一些文件整合、运行起来,得到一个动态运行的虚拟网络。其能接受一些命令行参数,包括

  • -t (Topology):指定拓扑 JSON 文件。脚本根据它来创建主机和交换机节点,并像拉网线一样建立连接。
  • -j (JSON):指定默认的 P4 编译 JSON 文件。脚本会将此逻辑加载到拓扑中定义的每一个交换机进程中。
  • -b (Behavioral-model):指定后端交换机执行程序的类型。
    simple_switch: 基础版。
    simple_switch_grpc: 进阶版。支持通过 gRPC 远程控制,稍后解释。

写到Makefile中,就是

UTILS_DIR = /home/p4/tutorials/utils
RUN_SCRIPT = (UTILS_DIR)/run_exercise.py
JSON_FILE =(BUILD_DIR)/myelin.json
TOPO_FILE = ./utils/topology.json

run: build
    sudo python3 (RUN_SCRIPT) -t(TOPO_FILE) -j $(JSON_FILE) -b simple_switch_grpc

从run_exercise.py看topology.json的解析方式

刚才提到,run_exercise.py能把我们编写的静态配置文件转化为一个动态运行的虚拟网络。它通过解析topology.json来自动化构建虚拟链路、分配端口、启动交换机进程并下发初始化流表,一切准备工作就绪后在终端弹出 mininet> 交互式提示符,将控制权移交给开发者,并在实验结束退出时负责清理残留的进程和网络接口。此文件共389行,通常不需要自己写,其已经涵盖了 95% 的实验需求,直接用就行。

run_exercise.py 源码对 topology.json 的解析逻辑非常刻板,因此topology.json一定要按run_exercise.py的解析逻辑去写。

首先,run_exercise.py 期望的 JSON 根对象必须包含以下三个顶级键:

  • hosts: 定义终端节点(即主机)的属性。
  • switches: 定义交换机进程及其实验特定的配置。
  • links: 定义节点间的物理/逻辑连接。

下面讲定义规范,当然是照着run_exercise.py 的解析逻辑来的。

主机配置 (hosts)

  • 每个主机条目支持以下字段:ip、mac、gw、commands
  • ip: 必须带掩码(如 10.0.1.1/24),脚本会据此配置接口。
  • mac: 建议手动指定(如 08:00:00:00:01:01),便于在 P4 代码中匹配。
  • gw: 默认网关。脚本会自动执行 route add default gw 命令。
  • commands (进阶): 一个字符串列表。脚本在启动后会依次在主机命名空间执行这些 Shell 命令(例如:[“arp -s 10.0.1.254 08:00:00:00:01:ff”])。

例子

{
  "hosts": {
    "h1": { "ip": "10.0.1.1/24", "mac": "08:00:00:00:01:01", "gw": "10.0.1.254" },
    "h2": { "ip": "10.0.1.2/24", "mac": "08:00:00:00:01:02", "gw": "10.0.1.254" },
    "h3": { "ip": "10.0.2.3/24", "mac": "08:00:00:00:02:03", "gw": "10.0.2.254" },
    "h4": { "ip": "10.0.2.4/24", "mac": "08:00:00:00:02:04", "gw": "10.0.2.254" }
  },
  "switches": {…},
  "links": […]
}

交换机配置 (switches)

  • 脚本为每个交换机条目启动一个独立的 simple_switch 或 simple_switch_grpc 进程。
  • cli_input: 相当于交换机的“开机初始化脚本”。P4 代码定义了交换机的“功能”(比如它能识别 IPv4 报文),而 cli_input 文件定义了交换机的“数据”(比如具体的路由表项)。
  • runtime_json: 如果使用 P4Runtime (gRPC),指定包含流表定义的 JSON 文件。
  • program 特殊用法 后面讲

例子:

{
  "hosts": {…},
  "switches": {
    "s1": { "cli_input": "utils/s1-commands.txt" },
    "s2": { "cli_input": "utils/s2-commands.txt" }
  },
  "links": […]
}

对应的目录结构

.
├── Makefile
├── p4src/
└── utils/
    ├── topology.json
    ├── s1-commands.txt    <-- 放在这里
    └── s2-commands.txt

s1-commands.txt书写格式就不是由run_exercise.py 规定的,而是由 BMv2 的底层工具 simple_switch_CLI 规定的。可以通过查阅此文档知道有哪些命令可用。

其他下发流表的办法

其实topology.json最简单的写法可以不写s1-commands.txt,即不使用 s1-commands.txt 这种静态文件自动下发流表,那么交换机的转发表(Tables)初始状态是空的。则需要通过Control Plane去下发流表项。两种办法:

  1. 现代SDN,动态下发流表项,通过 gRPC (P4Runtime)

如果在 Makefile 中使用了 -b simple_switch_grpc,交换机会监听 gRPC 端口(s1 通常是 50051,s2 是 50052,以此类推)。自己写一个 Python 脚本,利用 p4runtime_lib,在网络启动后连接 50051 端口,动态地根据网络状况下发流表。此为 P4 进阶功能(如负载均衡、动态路由)的主要配置方式。

  1. 通过 Thrift (CLI) —— 传统调试方式

没用过,不详。

链路配置 (links)

逻辑上很简单,就是连线,语法上有以下规定:

  • 顺序:在定义“主机-交换机”链路时,主机必须放在第一个位置 (node1)。正确:[“h1”, “s1-p1”], 错误:[“s1-p1”, “h1”]
  • 命名与端口规则:脚本使用 node.split(‘-‘) 解析端口。格式必须严格遵守 交换机名-p端口号。脚本会提取 – 后的字符串,并去掉首字母 p,将其转换为整数。例如 s1-p3 对应 BMv2 实例的 port 3。

例子:

{
  "hosts": {…},
  "switches": {…},
  "links": [
    ["h1", "s1-p1"],
    ["h2", "s1-p2"],
    ["h3", "s2-p2"],
    ["h4", "s2-p3"],
    ["s1-p3", "s2-p1"]
  ]
}

把上面三个例子合在一起就是run_exercise.py能解析的topology.json了。

实现异构交换机网络

仅供参考,由于我目前没有用到这个,所以并未亲自验证

现在想实现 s1 和 s2 运行不同的 P4 代码,则需对编译阶段(Makefile) 和 配置阶段(topology.json) 同时进行修改。

修改Makefile

假设已经准备了两个p4源文件,p4src/s1_logic.p4 和 p4src/s1_logic.p4,分别控制 s1 和 s2 的逻辑。那么要让p4c-bm2-ss编译器运行两次,生成两个独立的 JSON 描述文件。

build:
    mkdir -p (BUILD_DIR)(P4C) --p4v 16 -o (S1_JSON)(S1_SRC)
    (P4C) --p4v 16 -o(S2_JSON) $(S2_SRC)

run目标去掉了 -j 参数,其实不去掉也行,因为命令行传入的-j的优先级低于稍后的单独设置。

run: build stop
    mkdir -p logs pcaps
    sudo python3 (RUN_SCRIPT) -t(TOPO_FILE) -b simple_switch_grpc

修改 topology.json

根据 run_exercise.py 的源码逻辑,它会检查 switches 定义中是否有 “program” 字段。如果有,它会覆盖掉全局默认的 JSON。

所以只需

{
  "hosts": {…},
  "switches": {
    "s1": { 
      "program": "build/s1.json",
      …
    },
    "s2": { 
      "program": "build/s2.json",
      …
    }
  },
  "links": […]
}

ODOS-MPI:高性能处理器友好型智能网卡计算/通信内核卸载

DPU=SmartNIC+通用处理,其中的通用处理核给HPC领域带来一些新机遇。但DPU作为新的硬件,支持它的编程框架是没有的。所以这两篇使得OpenMP支持DPU编程。DPU作为能单独运行Linux操作系统的协处理器进行任务卸载。

一些思考

DPU的使用必须涉及并发。而不能像使用GPU那样,把任务卸载后等它干完再返回得数,这是无意义的,因为失去性能。

DPU未来能扮演什么角色(也就是DPU能卸载什么)?已知作为网络处理器,能卸载网络协议,那么能否把计算任务卸载?

两篇论文

  • ODOS:如何让代码在DPU上跑起来
  • ODOS-MPI:如何让代码在DPU上跑起来的同时,还能进行MPI通信
项目 SC23 (ODOS) SC25 (ODOS-MPI)
目标 首次实现OpenMP卸载到DPU 在ODOS基础上支持MPI通信卸载
贡献 实现了基础的OpenMP卸载机制(ODOS) 扩展为MPI+OpenMP混合卸载(ODOS-MPI)
技术支持 LLVM + DOCA SDK 同左 + 扩展Open MPI以支持异构通信
应用场景 纯计算任务卸载(如SPEC ACCEL) 支持通信+计算并发执行(如miniWeather)
性能提升 验证可行性,性能接近MPI基线 在真实应用中实现>18%加速

故事从ODOS讲起

ODOS 论文

ODOS 全称是 OpenMP DOCA Offloading Support。顾名思义,它是一个在LLVM/OpenMP框架内实现的、利用NVIDIA DOCA SDK来实现OpenMP卸载功能的后端。(生态绑定)

它的核心目标是,让程序员能够像写GPU卸载代码一样,使用#pragma omp target等指令,将代码段卸载到DPU上运行,而无需关心底层的复杂通信和异构细节。

在ODOS出现之前,如果想利用DPU的计算能力,主要有两种方式,但都有明显缺点:

  1. 低级别DOCA编程
    • 直接使用DOCA SDK提供的API。
    • 问题:过于底层、复杂,对领域科学家(如物理学家、气象学家)来说编程门槛极高,不利于广泛应用。
  2. MPI多节点编程
    • 将DPU视为一个独立的、能力较弱的计算节点,在上面启动MPI进程。(下图a)
    • 问题:编程模型不自然。程序员需要显式管理主机和DPU之间不同的MPI Rank,处理两者之间巨大的性能差异,并手动进行数据发送/接收,如同在编写分布式程序。

ODOS的价值就在于提供了第三种方式——使用高级、标准化的OpenMP模型,将DPU视为一个本地协处理器,极大地简化了编程。正如上图b,DPU和主机拥有相同Rank号。

ODOS由三个部分组成。

一、交叉编译的修改

负责将包含#pragma omp target的源代码编译成可以在DPU上运行的“胖二进制文件”。

主机通常是x86架构,而DPU是ARM架构。因此需要进行交叉编译,这是我们已经知道的。但是,标准的LLVM交叉编译通常假设目标设备使用与主机相同或类似的基础库。DPU在“独立主机模式”下运行一个完整的、独立的Linux操作系统,拥有自己的一套库和头文件(即独立的sysroot)。所以ODOS修改了Clang和链接器包装器,使其能够同时处理主机(x86_64)和目标设备(aarch64)两个不同的sysroot,从而正确地将DPU所需的库链接到胖二进制文件中。

二、编写OpenMP BlueField插件

是主机端的运行时库,让 OpenMP 运行时能够使用 BlueField DPU 这个硬件设备。

OpenMP运行时有一个抽象的“设备接口”。GPU有它的插件(比如用于NVIDIA GPU的插件),FPGA也有它的插件。但在ODOS之前,DPU没有自己的插件。这个OpenMP BlueField插件是连接标准OpenMP运行时和DOCA的桥梁。其职责包括设备发现、加载二进制文件、数据管理、执行控制、

对上层 OpenMP 运行时隐藏 DPU 的复杂细节(如 DOCA API、PCIe 通信)。特定于 BlueField DPU(及其 DOCA SDK)。

三、DOCA OpenMP服务进程

这是一个运行在DPU本身上的守护进程,接收并执行来自主机插件的请求。

  • 它作为一个服务端,通过DOCA CC监听主机的连接。
  • 它接收主机发送过来的二进制映像,并使用dlopendlsym将其动态加载为共享库。
  • 为了运行目标区域,它使用了libffi库。这是因为目标函数的参数在编译时对服务是未知的,libffi允许在运行时动态地创建和调用函数。

请注意,ODOS的目标硬件(DPU)与传统的卸载目标(GPU)在核心架构上是不同的:

特性 GPU DPU (如 BlueField-3)
核心类型 大量(数千个)轻量级计算核心 少量(16个)功能完整的ARM CPU核心
并行模式 大规模并行:适合在数千个核心上同时执行相同的简单操作 多核并行:适合在十几个核心上执行更复杂、可能有一定分支的任务

所以,最佳并行策略是不同的。GPU偏好的大规模并行语义,而 DPU偏好的传统多核并行语义。所以在具体编程时,GPU:

#pragma omp target teams distribute parallel for map(to: A[0:N], B[0:N]) map(from: C[0:N])
for (int i = 0; i < N; i++) {
    C[i] = A[i] + B[i];
}

DPU:(少了teams distribute,少创建很多很多线程)

#pragma omp target parallel for map(to: A[0:N], B[0:N]) map(from: C[0:N])
for (int i = 0; i < N; i++) {
    C[i] = A[i] + B[i];
}

以上,ODOS是仅仅卸载计算的代表,卸载内容是纯计算内核(无MPI调用),缺点是通信仍需主机CPU处理,通信成为瓶颈。主机在通信时无法计算,DPU在计算时主机可能空闲。

那么能不能将一个既包含计算又包含MPI调用的完整代码块(内核)一起卸载到DPU上执行?就来到今年的《ODOS-MPI》了。

ODOS-MPI 论文

ODOS-MPI 编程模型有如下特点:

  • 将OpenMP的taregt命令所引导的offload region运行在BlueField DPU之上(ODOS已完成)
  • 在offload region之中,使得MPI也能运行(ODOS-MPI的核心)
  • 利用MPI_Info 的提示信息决定,操作的目标在DPU上还是在主机之上(操作比如send/put)

在ODOS基础上,ODOS-MPI进行如下的新增工作(第三列),主要目的是通信和计算的双卸载:

领域 ODOS ODOS-MPI
编译器与工具链 扩展LLVM/Clang,支持DPU交叉编译,生成胖二进制文件。 复用
OpenMP运行时插件 创建了OpenMP BlueField插件,实现设备发现、代码加载、数据传输。 扩展该插件,使其能识别并处理包含MPI调用的目标区域,与新的MPI运行时交互。
DPU端服务 创建了DOCA OpenMP服务,用libffi执行纯计算函数。 增强该服务,使其能与DPU上的MPI运行时交互,支持在目标区域内调用MPI函数。
MPI集成 完全不涉及。卸载区域内不能有MPI调用。 核心新工作
双MPI运行时:在主机和DPU上分别初始化MPI,并让它们共享Rank号。
异构通信:通过扩展UCX,增加主机-设备通信端点。
HETEROGENEOUS提示符:新发明的MPI_Info键,用于明确指定异构通信路径。
编程模型语义 定义了标准的OpenMP卸载语义,仅用于计算。 核心新工作
• 定义了在OpenMP卸载区域内调用MPI的语义(如nowaitMPI_Barrier的交互)。
• 定义了异构通信(主机↔设备)和同构通信(设备↔设备)的规则。
• 定义了集体操作的限制(必须在同一“侧”发起)。
同步模型 使用OpenMP的taskwait进行同步。 扩展同步模型:结合MPI_Barrier和OpenMP taskwait,实现主机与设备间的跨域同步。

ODOS-MPI初始化两个独立的MPI运行时:

  • 主机MPI运行时:运行应用程序主进程
  • DPU MPI运行时:作为后台服务处理卸载请求

两个运行时使用相同的MPI Rank编号,为程序员提供统一的消息传递视图。

异构通信支持:通过扩展Open MPI的UCX组件,ODOS-MPI实现了:

  • 端点扩展:在原有主机-主机端点基础上,增加主机-设备端点
  • 通信模式
    • 同构通信:主机↔主机 或 DPU↔DPU
    • 异构通信:主机↔DPU(通过MPI_Info提示符控制)

小结:设计与实现一个编程模型时需要考虑什么

抽象层次

计算模型

数据模型

并发模型

集成

硬件抽象

工具链

UCC(统一集体通信库)与UCX(统一通信框架)

二者区别

UCX 是底层通信的基础框架(点对点通信),而 UCC 是构建在 UCX 等底层库之上的上层集体通信库(集体通信)。

UCX目标是提供统一接口,让上层(如MPI、UCC)高效利用各种硬件资源。比如,InfiniBand、RoCE、TCP/IP、共享内存等。(UCC、MPI的下层不一定是UCX。 UCX是它们实现时可以选择的其中一种底层通信库,但并不是唯一选择)

UCC专注于集体通信模式,具有后端无关性。可以动态选择 UCX、NCCL、SHARP(InfiniBand的网络内计算)、MPI 甚至 CUDA 等作为执行操作的“引擎”。对于使用者来说,只需要调用如 ucc_collective_initucc_collective_post 等高级API即可。此外,UCC 有一个调度组件,可以根据操作的类型、数据大小、网络拓扑等信息,智能地选择最快、最合适的算法和后端来执行。

UCC是一个比较新的框架,22年4月正式发布1.0版本,在24年HOTI会议上发表论文《Unified Collective Communication (UCC): An Unified Library for CPU, GPU, and DPU Collectives》(统一集体通信UCC:一个面向 CPU、GPU 和 DPU 集体的统一库)。

集体通信

核心特点

  • 涉及一组进程:不止两个进程。
  • 全局操作:操作的结果或过程依赖于组内所有成员。
  • 隐式协调与同步:操作本身通常包含了进程间的协调,有时甚至会引入同步点(即所有进程都到达某个点后才能继续)。

以MPI为例,常见的集体通信操作有

操作名称 功能描述
Broadcast(广播) 一个进程(根进程)将相同的数据发送给组内的所有其他进程。
Scatter(散播) 一个进程将一块数据的不同部分分别发送给组内的所有进程。
Gather(收集) 一个进程从组内所有其他进程那里接收数据,并按照进程号顺序组合成一块数据。
All-Gather(全收集) 每个进程都拥有一个数据块,执行操作后,所有进程都拥有全部数据块的集合。
Reduce(规约) 将组内所有进程的数据通过一个操作(如求和、求最大值、求最小值等)合并为一个结果,存放在一个指定进程(根进程)中。
All-Reduce(全规约) 执行Reduce操作,但结果分发给所有进程。这是机器学习和科学计算中最常用的操作之一。
Barrier(屏障) 组内所有进程在此同步,必须所有进程都到达屏障后,才能继续向下执行。
All-to-All(全交换) 每个进程都向所有其他进程发送不同的消息,同时也从所有其他进程接收消息。

UCC论文(HOTI’24)

Unified Collective Communication (UCC): An Unified Library for CPU, GPU, and DPU Collectives

[推荐阅读](UCC简介 | RedBlog)
[演讲回放]([HOTI2024] 统一集合通信UCC_哔哩哔哩_bilibili)

背景

集体通信用于进程间的数据交换与同步,无论是在HPC领域还是DL领域都非常重要。但现有通信库各自优化专有硬件。UCC希望统一 CPU、GPU、DPU 上的集体通信、支持多种编程模型、整合软件与硬件传输机制。

概述

UCC 的实现精髓在于其模块化分层调度,这使得 UCC 成为一个既能统一接口,又能充分发挥异构硬件性能的通信库。

img

UCC 组件架构如图所示,使用尽可能接近底层的方式实现了集合通信,并向上层应用提供了调用的接口。UCC对集合通信操作的实现分为两层,TL即Transport Layer,主要负责针对不同的底层通信协议实现集合通信操作最基本的功能,如UCP、NCCL、SHARED Memory等。CL层为Collective Layer,在TL层的基础上进一步丰富了其功能,如Hier组件实现了分层的集合通信操作。

设计目标

  • 通用性(既支持MPI等传统HPC负载,又支持PyTorch等AI/DL负载)
  • 兼容性(既能使用UCX传统点对点通信等软件传输,也能支持如[ref14]的硬件传输)
  • 支持并行与重叠(通过推进多个独立的集体操作来实现并行,允许通信与计算重叠来提供重叠能力)
  • 支持异构硬件(CPU、DPU、GPU,集体通信可以由异构硬件中的任何一个发起,或者数据可以驻留在主内存或设备内存中)为支持此设计目标,库的集体调用模型应同时支持紧耦合调用模型(CPU 线程调用并执行集体操作)和松耦合调用模型(CPU 线程调用集体操作,而 GPU 流执行集体操作)。
  • 可扩展性(就是必须模块化)

设计原则

为实现前述目标,UCC 的设计围绕三个核心概念:抽象属性操作,每个抽象有自己的属性与操作

  • 抽象 代表了 UCC 库的基本资源和能力。它们封装了集体操作所需的计算资源、网络资源和网络端点。在规范中,它以 ucc_{抽象名称} 的形式表示。
  • 属性 代表了抽象的语义。它们使用户能够定制抽象以满足各种工作负载的需求。在 UCC 中,它以 ucc_{抽象名称}_params_t 的形式表示。
  • 操作 提供了对这些抽象执行动作的接口。这些操作包括创建和销毁抽象以及更新语义。

UCC 中的主要抽象包括五个:上下文团队端点执行引擎

  1. UCC 库
  • UCC 库是集体通信库所需所有资源的封装。对于一个给定的应用程序,通常只需要一个库实例。
  • 属性:UCC 库的属性包括:
    • 线程模型:UCC 支持多种线程模式(THREAD_SINGLE, THREAD_MULTIPLE, FUNNELED),以适应 MPI 和 OpenSHMEM 编程模型。
    • 同步模型:UCC 支持两种同步模型——同步(SYNC)和异步(NO_SYNC),以满足 MPI、OpenSHMEM 和统一并行 C 编程模型的需求。
    • 支持的集体操作:Barrier、Broadcast、All-gather、Gather、Scatter、Allreduce、AlltoAll 和 Reduce,以及这些集体操作的向量变体。
  • 操作UCC 库 的操作包括初始化、创建和销毁库对象。

  1. UCC 上下文
  • UCC 上下文 封装了本地网络资源,包括软件网络资源(例如,虚拟接口、注入队列、连接端点、共享内存缓冲区)和硬件网络资源(例如,全局资源树)。此抽象有助于管理分配给集体操作的资源。
  • 属性:上下文可以创建为独占类型或共享类型,独占类型提供更高性能,共享类型节省资源。
  • 操作:UCC 上下文操作包括初始化、创建和销毁上下文对象。
  1. UCC Teams
  • UCC Teams 封装了用于集体操作的组资源,类似于 MPI 通信器或 OpenSHMEM Teams 。
  • 属性UCC Teams 的属性包括:
    • (1) 排序模型,它定义了参与集体操作的进程如何调用该集体操作(有序或无序)。
    • (2) 团队的参与者。
    • (3) 为资源分配而设置的未完成集体操作的数量。
    • (4) 集体操作所需的内存资源。
    • (5) 同步模型。
  • 操作UCC 团队 的操作包括提交团队创建请求、测试创建是否完成以及销毁团队对象。
  1. UCC 端点
  • UCC 端点代表了团队中参与者的唯一标识符,允许将端点绑定到并行编程模型中的索引,例如 OpenSHMEM 处理元素(PE)、MPI 排名、操作系统进程 ID 或线程 ID。
  • 属性:端点可以是连续的或非连续的。
  1. 执行引擎(EE)
  • 执行引擎抽象了执行集体操作的计算资源(CPU 或 GPU)。通过与 EE 的事件驱动机制进行交互,可以在调用线程和执行线程之间存在高延迟时实现高效交互。
  • 属性:UCC 支持的 EE 类型包括:CUDA STREAM、ROCM STREAM 和 CPU THREAD。
  • 操作:EE 操作包括创建和销毁,以及事件管理操作,例如 SET、GET 和 ACK。

通过采用这些设计原则,UCC 提供了一个灵活而高效的解决方案,能够满足多样化的工作负载和编程模型的需求。

实现

UCC 的实现像一个高度模块化的工具箱,其核心设计是组件化架构

核心与组件

  • CORE:面向用户的 API 在 ucc.h 头文件中提供。这些接口在 UCC 版本发布之间保持一致性并向后兼容。该 API 被实现为 UCC 的 CORE,其余功能则作为组件实现。
  • 组件:UCC 的大部分功能是作为组件实现的,主要分为CL层、TL层、内存组件、拓扑组件。
    • Collective Layer (CL):CL层组合多个TL层,以提供编程模型所需的、与传输无关的集体实现。它决定集体操作的执行策略。例如,Basic CL 实现直接调用底层传输;Hierachical CL 实现则像先在小团队内完成计算,再在团队代表间汇总,适合大规模系统。
    • Team Layer (TL):这是与具体通信硬件对接的一层。UCC 通过不同的 TL 来调用不同的通信库:
      • UCP TL:使用 UCX 库,适用于 CPU 和多种网络(如 InfiniBand)。
      • NCCL TL:调用 NVIDIA 的 NCCL 库,专门用于 NVIDIA GPU。
      • SHARP TL:使用支持在网络交换机内计算(在网计算)的硬件。
      • CUDA TL:针对单台服务器内 GPU 间的通信。
    • Memory Component:支持 CPU/GPU 内存操作
    • Topology Component:用于分层集体中的端点分组

集体操作的执行流程

概念:一个复杂的【集体操作】(如分层 Allreduce)被拆分成多个小的【任务】(如:节点内规约 -> 节点间全规约 -> 节点内广播)。这些任务及其依赖关系构成一个【调度计划】。UCC 内部通过【进度队列】来管理这些任务,实现并行推进和依赖处理。

UCC 支持阻塞、非阻塞和持久化的调用与执行模型。通过组合以下接口来提供对此的支持:ucc_collective_init, ucc_collective_post, ucc_collective_test, ucc_collective_finalize。

  1. ucc_collective_init:指定计划。程序提供参数(如数据地址、操作类型)。该接口根据这些信息智能地选择(稍后介绍)最合适的CL层、TL层和算法,并规划好执行步骤,生成一个【调度计划】。
  2. ucc_collective_post:开始执行计划。UCC 开始执行调度计划中的【任务】。
  3. ucc_collective_test:程序可以随时询问操作是否完成。UCC会回复完成了或者还在做。
  4. ucc_collective_finalize:操作完成后,释放资源。

关于“智能的选择”:论文的目标是证明这种统一设计的可行性和优势,具体的选择机制官方文档中或许有(或源代码)。

核心任务是为一类特定的集体操作选择一个 “集体层 – 团队层” 的组合。论文中明确提到的选择层次包括:

  • 选择集体层:是使用 基础分层供应商 还是 UROM 实现策略。
  • 选择团队层:根据上面选定的策略,进一步选择使用哪个具体的传输库,例如 UCPNCCLSHARPCUDARCCL 等。

论文中提到的决策因素有:数据大小、Team规模、硬件位置、拓扑、可用资源、工作负载类型。

论文例子:一个典型的分层 Allreduce 集体操作将包括以下调度计划和任务:

  • 任务 1: Intra-node Reduce using TL/UCP
  • 任务 2: Inter-node Allreduce using TL/SHARP
  • 任务 3: Intra-node Broadcast using TL/UCP
  1. 对 GPU 的特别支持

通过执行引擎触发式API优化 GPU 通信

为了让 GPU 也能高效发起和参与集体操作,UCC 引入了执行引擎触发式提交的概念。

  • 执行引擎:抽象了 GPU 的计算上下文(如 CUDA 流)。
  • 触发式提交:程序可以将集体操作”绑定”到一个特定的 GPU 流上,并指定一个事件作为启动信号。当这个事件发生时,集体操作会自动在 GPU 上开始执行。这实现了计算与通信的精细重叠,避免了 GPU 空闲等待。
  1. 对 DPU 的卸载

通过客户端-服务端模型实现向 DPU 的卸载

UCC 可以通过 UROM 组件将集体操作卸载到 DPU 上执行。

  • 工作方式:主机端的 UCC 库作为客户端,将集体操作的命令发送给在 DPU 上运行的 UROM 服务。
  • 流程:DPU 上的服务接收命令,用自己的 UCC 实例在 DPU 之间完成整个集体操作,最后将结果通知主机。
  • 好处:将主机的 CPU/GPU 从通信任务中解放出来,专注于计算,提升整体效率。