卸载设计与内部机制

简介

本文档描述了 Clang 驱动程序和代码生成步骤，用于创建卸载应用程序。Clang 支持使用 CUDA、HIP 和 OpenMP 等编程模型将应用程序卸载到各种架构。本文档的目的是说明使用 Clang 创建卸载应用程序所需的步骤。

OpenMP 卸载

Clang 支持将 OpenMP 目标卸载到多个不同的架构，例如 NVPTX、AMDGPU、X86_64、Arm 和 PowerPC。卸载代码由 Clang 生成，然后使用 libomptarget 运行时和与目标架构相关的插件（例如 libomptarget.rtl.cuda）执行。本节描述创建可由 OpenMP 运行时加载的功能性设备映像所需的步骤。有关 OpenMP 运行时的更多信息，请参阅 OpenMP 文档页面。

卸载概述

卸载编译的目标是创建一个可以在目标设备上运行的可执行设备映像。OpenMP 卸载通过分别为主机和目标设备编译输入文件来创建可执行映像。然后需要将设备阶段的输出嵌入到主机中以创建胖对象。然后，一个特殊的工具需要从胖对象中提取设备代码，运行设备链接步骤，并将最终映像嵌入到主机运行时库可以用来注册库并访问设备上符号的符号中。

编译过程

编译器执行以下高级操作来生成 OpenMP 卸载代码

• 编译主机输入文件以生成一个位代码文件。将 #pragma omp target 声明降低为 卸载条目，并创建元数据以指示哪些条目位于设备上。

• 使用主机创建的 卸载条目 元数据，为目标 设备 编译输入文件。

• 链接 OpenMP 设备运行时库并运行后端以创建设备目标文件。

• 在主机位代码文件上运行后端，并使用设备目标文件创建 胖目标文件。

• 将胖目标文件传递给 链接器包装器工具 并提取设备对象。对提取的对象运行设备链接操作。

• 将 设备映像 和 卸载条目 封装 到主机可以访问的符号中。

• 将 封装后的二进制文件 添加到链接器输入并运行主机链接操作。与 libomptarget 链接以注册和执行映像。

生成卸载条目

编译的第一步是为主机生成卸载条目。这些信息用于识别将由设备提供的函数内核或全局值。包含在 #pragma omp target 中的块或 #pragma omp declare target 指令内的符号将生成卸载条目。下表显示了 卸载条目结构。

__tgt_offload_entry 结构

类型	标识符	描述
void*	addr	设备映像中全局符号的地址（函数或全局变量）
char*	name	符号的名称
size_t	size	条目信息的长度（如果为函数，则为 0）
int32_t	flags	与条目关联的标志（请参阅 目标区域条目标志）
int32_t	reserved	保留，供运行时库使用。

一旦设备映像加载，运行时将把全局符号的地址设置为设备指针值。标志设置为指示卸载条目所需的处理。如果卸载条目是目标区域的条目，则它可以具有以下 条目标志 之一。

目标区域条目标志

名称	值	描述
OMPTargetRegionEntryTargetRegion	0x00	将条目标记为通用目标区域
OMPTargetRegionEntryCtor	0x02	将条目标记为全局构造函数
OMPTargetRegionEntryDtor	0x04	将条目标记为全局析构函数

如果卸载条目是全局变量（由非零大小指示），则它将具有以下 全局 标志之一。

目标区域全局

名称	值	描述
OMPTargetGlobalVarEntryTo	0x00	将条目标记为“to”属性（相对于 to 子句）
OMPTargetGlobalVarEntryLink	0x01	将条目标记为“link”属性（相对于 link 子句）

运行时使用目标卸载条目访问设备内核和将由最终设备映像提供的全局变量。每个卸载条目都设置为使用 omp_offloading_entries 部分。当创建最终应用程序时，链接器将提供 __start_omp_offloading_entries 和 __stop_omp_offloading_entries 符号，这些符号用于创建 最终映像。

设备编译阶段使用这些信息来确定需要从设备导出哪些符号。我们使用 omp_offload.info 元数据节点将这些信息传递给设备编译阶段。

在设备上访问条目

通过使用 卸载条目 中的 address 字段来访问设备中的条目。运行时将在运行时初始化期间将地址设置为与设备映像关联的指针。这用于在进入 #pragma omp target 区域时调用相应的内核函数。对于变量，运行时维护一个将主机指针映射到设备指针的表。在 #pragma omp target declare 指令中的全局变量首先初始化为主机的地址。一旦设备地址初始化，我们将其插入到表中以将主机地址映射到设备地址。

调试信息

我们生成结构来保存传递给 libomptarget 的调试信息。这允许前端生成运行时库用于更具信息性的错误消息的信息。这是使用 libomp 和 libomptarget 中使用的标准 标识符结构 完成的。这用于将信息和源位置传递给运行时。

ident_t 结构

类型	标识符	描述
int32_t	reserved	保留，供运行时库使用。
int32_t	flags	用于指示某些功能的标志，大多数未使用。
int32_t	reserved	保留，供运行时库使用。
int32_t	reserved	保留，供运行时库使用。
char*	psource	程序源信息，存储为“;filename;function;line;column;;\0”

如果启用调试信息，我们还将创建字符串以指示在目标区域中映射的变量的名称和声明。它们具有与 标识符结构 中的源位置相同的格式，但函数名称将被替换为变量名称。

卸载设备编译

输入文件为每个活动设备工具链编译。设备编译阶段与主机阶段的执行方式不同。也就是说，我们不会生成任何卸载条目。这是通过将 -fopenmp-is-target-device 标志传递给前端来设置的。我们使用主机位代码来确定要从设备导出哪些符号。位代码文件使用 -fopenmp-host-ir-file-path 标志从上一阶段传递进来。否则，编译将与对任何其他目标三元组的编译方式相同。

在为 OpenMP 设备编译时，我们默认将所有设备符号的可见性设置为 protected。这提高了性能，并防止了一类错误，在这种错误中，目标设备中的符号可能会抢占主机库。

OpenMP 运行时库在编译期间链接进来，以提供标准 OpenMP 功能的实现。对于 GPU 目标，这是通过在编译期间链接一个特殊的位码库来完成的（例如，libomptarget-nvptx64-sm_70.bc），使用 -mlink-builtin-bitcode 标志。其他设备库，如 CUDA 的 libdevice，也以这种方式链接。如果目标是具有现有 libomp 实现的标准体系结构，则将链接该库。最后，设备工具用于创建可重定位的设备对象文件，该文件可以嵌入到主机中。

创建胖对象

胖二进制文件是包含用于另一个设备的信息的二进制文件。我们通过将设备编译阶段的输出嵌入到主机中作为命名部分来创建胖对象。来自设备编译的输出通过 -fembed-offload-object 标志传递给主机后端。这将设备映像嵌入到 .llvm.offloading 部分中，使用一种特殊的二进制格式，该格式的行为类似于字符串映射。这种二进制格式用于捆绑有关映像的元数据，以便链接器可以将正确的设备链接操作与映像关联起来。每个设备映像都将以魔数 0x10FF10AD 开头。

@llvm.embedded.object = private constant [1 x i8] c"\00", section ".llvm.offloading"

然后，在后端在主机上运行时，设备代码将被放置在相应的部分中，从而创建胖对象。使用胖对象使我们能够将卸载对象视为标准主机对象。最终的对象文件应包含以下 卸载部分。我们将在 链接目标设备代码 时使用此信息。

卸载部分

部分	描述
omp_offloading_entries	卸载条目信息（参见 __tgt_offload_entry 结构）
.llvm.offloading	嵌入的目标设备和体系结构的设备对象文件

链接目标设备代码

包含 卸载部分 的对象需要特殊处理才能创建可执行的设备映像。这是使用 Clang 工具完成的，有关更多信息，请参见 Clang 链接器包装器。此工具充当主机链接作业的包装器。它扫描输入对象文件以查找卸载部分 .llvm.offloading。然后提取存储在此部分中的设备文件并传递给相应的链接作业。然后将链接的设备映像 包装，以创建用于加载设备映像并将其与主机链接的符号。

链接器包装器工具支持通过链接时优化 (LTO) 链接位码文件。当嵌入到主机中的对象文件包含 LLVM 位码时，就会使用它。对于不支持可重定位对象格式的体系结构（如 AMDGPU 或 SPIR-V），或者如果用户使用 -foffload-lto 标志请求了它，将嵌入位码。

设备二进制文件包装

各种结构和函数用于创建在设备上卸载代码所需的信息。我们使用 链接的设备可执行文件 和相应的卸载条目来创建加载和执行设备映像所需的符号。

结构类型

使用几种不同的结构来存储卸载信息。设备映像结构 存储单个链接的设备映像及其关联的卸载条目。卸载条目使用链接器使用 __tgt_offload_entry 结构 生成的 __start_omp_offloading_entries 和 __stop_omp_offloading_entries 符号存储。

__tgt_device_image 结构

类型	标识符	描述
void*	ImageStart	指向目标代码开头的指针
void*	ImageEnd	指向目标代码结尾的指针
__tgt_offload_entry*	EntriesBegin	包含所有目标条目的表的开头
__tgt_offload_entry*	EntriesEnd	表的结尾（不包含）

目标 目标二进制描述符 用于在一个数组中存储所有二进制映像和卸载条目。

__tgt_bin_desc 结构

类型	标识符	描述
int32_t	NumDeviceImages	支持的设备类型的数量
__tgt_device_image*	DeviceImages	设备映像数组（每个设备类型 1 个）
__tgt_offload_entry*	HostEntriesBegin	包含所有主机条目的表的开头
__tgt_offload_entry*	HostEntriesEnd	表的结尾（不包含）

全局变量

全局变量 列出了各种全局变量，以及它们的类型和它们的显式 ELF 部分，这些部分用于存储设备映像和相关符号。

全局变量

变量	类型	ELF 部分	描述
__start_omp_offloading_entries	__tgt_offload_entry	.omp_offloading_entries	卸载条目表的开头符号。
__stop_omp_offloading_entries	__tgt_offload_entry	.omp_offloading_entries	卸载条目表的结尾符号。
__dummy.omp_offloading.entry	__tgt_offload_entry	.omp_offloading_entries	卸载条目部分中的虚拟零大小对象，用于强制链接器定义上面定义的开头/结尾符号。
.omp_offloading.device_image	__tgt_device_image	.omp_offloading_entries	第一个映像的 ELF 设备代码对象。
.omp_offloading.device_image.N	__tgt_device_image	.omp_offloading_entries	第 (N+1) 个映像的 ELF 设备代码对象。
.omp_offloading.device_images	__tgt_device_image	.omp_offloading_entries	映像数组。
.omp_offloading.descriptor	__tgt_bin_desc	.omp_offloading_entries	二进制描述符对象（参见 设备映像的二进制描述符）

设备映像的二进制描述符

此对象在程序启动时传递给卸载运行时，它描述了可执行文件或共享库中可用的所有设备映像。它的定义如下：

__attribute__((visibility("hidden")))
extern __tgt_offload_entry *__start_omp_offloading_entries;
__attribute__((visibility("hidden")))
extern __tgt_offload_entry *__stop_omp_offloading_entries;
static const char Image0[] = { <Bufs.front() contents> };
...
static const char ImageN[] = { <Bufs.back() contents> };
static const __tgt_device_image Images[] = {
  {
    Image0,                            /*ImageStart*/
    Image0 + sizeof(Image0),           /*ImageEnd*/
    __start_omp_offloading_entries,    /*EntriesBegin*/
    __stop_omp_offloading_entries      /*EntriesEnd*/
  },
  ...
  {
    ImageN,                            /*ImageStart*/
    ImageN + sizeof(ImageN),           /*ImageEnd*/
    __start_omp_offloading_entries,    /*EntriesBegin*/
    __stop_omp_offloading_entries      /*EntriesEnd*/
  }
};
static const __tgt_bin_desc BinDesc = {
  sizeof(Images) / sizeof(Images[0]),  /*NumDeviceImages*/
  Images,                              /*DeviceImages*/
  __start_omp_offloading_entries,      /*HostEntriesBegin*/
  __stop_omp_offloading_entries        /*HostEntriesEnd*/
};

全局构造函数和析构函数

全局构造函数 (.omp_offloading.descriptor_reg()) 通过调用 __tgt_register_lib() 运行时函数向运行时注册设备映像。构造函数在 .text.startup 部分中显式定义，并在程序启动时运行一次。类似地，全局析构函数 (.omp_offloading.descriptor_unreg()) 为析构函数调用 __tgt_unregister_lib()，并且也定义在 .text.startup 部分中，并在程序退出时运行。

卸载示例

本节包含使用 OpenMP 卸载生成卸载代码的简单示例。我们将使用一个简单的 ZAXPY BLAS 例程。

#include <complex>

using complex = std::complex<double>;

void zaxpy(complex *X, complex *Y, complex D, std::size_t N) {
#pragma omp target teams distribute parallel for
  for (std::size_t i = 0; i < N; ++i)
    Y[i] = D * X[i] + Y[i];
}

int main() {
  const std::size_t N = 1024;
  complex X[N], Y[N], D;
#pragma omp target data map(to:X[0 : N]) map(tofrom:Y[0 : N])
  zaxpy(X, Y, D, N);
}

此代码使用以下 Clang 标志编译。

$ clang++ -fopenmp -fopenmp-targets=nvptx64 -O3 zaxpy.cpp -c

可以使用 readelf 实用程序查看对象文件中的输出部分。 .llvm.offloading 部分具有 SHF_EXCLUDE 标志，因此它将由链接器从最终的可执行文件或共享库中删除。

$ llvm-readelf -WS zaxpy.o
Section Headers:
[Nr] Name                   Type     Address          Off    Size   ES Flg Lk Inf Al
[11] omp_offloading_entries PROGBITS 0000000000000000 0001f0 000040 00   A  0   0  1
[12] .llvm.offloading       PROGBITS 0000000000000000 000260 030950 00   E  0   0  8

再次编译此文件将调用 clang-linker-wrapper 实用程序来提取和链接存储在名为 .llvm.offloading 的部分中的设备代码，然后使用存储在名为 omp_offloading_entries 的部分中的条目来创建 libomptarget 注册设备映像并调用入口函数所需的符号。

$ clang++ -fopenmp -fopenmp-targets=nvptx64 zaxpy.o -o zaxpy
$ ./zaxpy

我们可以使用 Clang 中的 -ccc-print-phases 选项查看 clang 生成卸载代码创建的步骤。这与 卸载概述 中的描述相匹配。

$ clang++ -fopenmp -fopenmp-targets=nvptx64 -ccc-print-phases zaxpy.cpp
# "x86_64-unknown-linux-gnu" - "clang", inputs: ["zaxpy.cpp"], output: "/tmp/zaxpy-host.bc"
# "nvptx64-nvidia-cuda" - "clang", inputs: ["zaxpy.cpp", "/tmp/zaxpy-e6a41b.bc"], output: "/tmp/zaxpy-07f434.s"
# "nvptx64-nvidia-cuda" - "NVPTX::Assembler", inputs: ["/tmp/zaxpy-07f434.s"], output: "/tmp/zaxpy-0af7b7.o"
# "x86_64-unknown-linux-gnu" - "clang", inputs: ["/tmp/zaxpy-e6a41b.bc", "/tmp/zaxpy-0af7b7.o"], output: "/tmp/zaxpy-416cad.o"
# "x86_64-unknown-linux-gnu" - "Offload::Linker", inputs: ["/tmp/zaxpy-416cad.o"], output: "a.out"

可重定位链接

卸载编译管道通常会将最终的设备链接和运行时注册推迟到运行 clang-linker-wrapper 创建可执行文件时。这是在 -fgpu-rdc 模式下为 OpenMP 卸载或 CUDA 和 HIP 编译时的标准行为。但是，在某些情况下，用户可能希望过早地执行此设备处理。这在 链接器包装器文档 中进行了描述。

实际上，这允许用户在分发对象或静态库时提前处理特定于卸载的链接。可以将其视为在对象文件子集上执行标准 -fno-gpu-rdc 编译。这对于减少链接时间、防止用户与库的设备代码交互或用于将库分发到不兼容的编译器很有用。

通常，如果使用 clang -r 完成可重定位链接，它只会合并 .llvm.offloading 部分，这些部分将在稍后创建可执行文件时链接。但是，如果 -r 标志与卸载工具链一起使用，它将执行设备链接和注册阶段，然后将注册代码合并到最终的可重定位对象文件中。

以下示例展示了如何使用可重定位链接与卸载管道一起创建包含卸载代码的静态库，该库可以重新分发，而无需任何额外的处理。

$ clang++ -fopenmp -fopenmp-targets=nvptx64 foo.cpp -c
$ clang++ -lomptarget.devicertl --offload-link -r foo.o -o merged.o
$ llvm-ar rcs libfoo.a merged.o
# g++ app.cpp -L. -lfoo