使用 oneAPI 进行异构编程

要从当今的硬件中获得最大可实现的性能，需要在最佳利用底层硬件功能与使用可移植、易于维护且能效高的代码之间取得精妙的平衡。这些因素不一定协同工作。它们需要根据用户需求进行优先级排序。用户维护不同架构的独立代码库并非易事。一个通用的、简化的编程模型，可以在标量、矢量、矩阵和空间架构上无缝运行，将通过增加代码重用和减少培训投入来提高开发者的生产力。

oneAPI 是一项行业倡议，旨在实现这些优势。它基于标准和开放规范，包含 Data Parallel C++ (DPC++) 语言以及一套领域库。oneAPI 的目标是让行业内的硬件供应商开发自己的兼容实现，以其 CPU 和加速器为目标。这样，开发者只需用一种语言和一套库 API 即可跨多个架构和多个供应商设备进行编码。

Intel oneAPI 开发者工具的 Beta 版实现，以 Intel® CPU 和加速器为目标，包括 Intel® oneAPI Base Toolkit 以及多个特定领域的工具包—Intel® HPC、IoT、DL Framework Developer 和 Rendering 工具包—以满足不同用户的需求。

图 1 展示了 beta 版 Intel oneAPI 产品以及 Base Toolkit 的不同层，该 Toolkit 包括 Intel oneAPI DPC++ 编译器、Intel® DPC++ 兼容工具、多个优化库以及高级分析和调试工具。跨架构的并行性使用 DPC++ 语言表达，该语言基于 Khronos Group 的 SYCL*。它使用现代 C++ 功能以及 Intel 特定的扩展来高效利用架构。DPC++ 语言功能允许代码在 CPU 上运行，并卸载到可用的加速器—从而可以重用代码。回退属性允许在没有加速器可用时在 CPU 上运行代码。主机和加速器上的执行以及内存依赖性都已明确定义。

用户还可以使用 Intel DPC++ 兼容工具将代码从 CUDA* 移植到 DPC++。它协助开发者进行一次性迁移，通常可以自动迁移 80% 到 90% 的代码。

除了 DPC++ 之外，Intel oneAPI HPC Toolkit 还支持 OpenMP* 5.0 功能，允许将代码卸载到 GPU。用户可以过渡到使用 DPC++，或利用现有 C/C++/Fortran 代码的卸载功能。API 编程通过一套库（例如 Intel® oneAPI Math Kernel Library）提供支持，这些库将针对 Intel GPU 进行优化。

beta 版 Intel oneAPI 产品还在 Intel® VTune™ Profiler1 和 Intel® Advisor2 中提供了新功能，允许用户在代码卸载到加速器时调试代码并查看性能相关指标。

本文介绍 oneAPI 产品 Beta 版，旨在促进异构编程。我们将介绍 oneAPI 软件模型，然后讨论编译模型和二进制生成过程。oneAPI 为所有架构提供单个二进制文件，因此编译和链接步骤与常规二进制生成方法不同。最后，我们将检查一些示例程序。请注意，在本文中，我们互换使用术语“加速器”、“目标”和“设备”。

oneAPI 软件模型

oneAPI 软件模型基于 SYCL 规范，从代码执行和内存使用的角度描述了主机和设备之间的交互。该模型有四个部分：有关编译器优化的更完整信息，请参阅我们的优化声明。订阅未来期号

1. 平台模型，指定主机和设备
2. 执行模型，指定命令队列和将在设备上运行的命令
3. 内存模型，指定主机和设备之间的内存使用
4. 内核模型，将计算内核定位到设备

平台模型

oneAPI 平台模型指定主机和相互通信或与主机通信的多个设备。主机控制设备上的内核执行，并在有多个设备时协调它们。每个设备可以有多个计算单元。每个计算单元可以有多个处理元素。只要平台满足 oneAPI 软件模型的最低要求，oneAPI 规范就可以支持 GPU、FPGA 和 ASIC 等多个设备。（有关平台要求的详细信息，请参阅每个 oneAPI 组件的发行说明。）

执行模型

oneAPI 执行模型指定代码如何在主机和设备上执行。主机执行模型创建命令组，以协调主机和设备之间的内核执行和数据管理。命令组在队列中提交，这些队列可以按顺序或无序策略运行。队列中的命令可以同步，以确保在执行下一个命令之前，设备上的数据更新可供主机使用。

设备执行模型指定计算如何在加速器上完成。执行范围涵盖一组元素，这些元素可以是 一维或多维数据集。此范围被划分为 ND-Range、工作组、子组和工作项的层次结构，如图 2 所示（三维情况）。

请注意，这与 SYCL 模型类似，除了子组（它是 Intel 扩展）。工作项是内核中最小的执行单元。工作组决定了这些工作项之间的数据共享方式。这些分层布局也决定了应使用哪种类型的内存以获得更好的性能。例如，工作项通常在存储在设备内存中的临时数据上操作，而工作组使用全局内存。引入子组分类是为了支持具有矢量单元的硬件资源。这允许对元素进行并行执行。

从图 2 可以清楚地看出，工作组或工作项在 ND-Range 中的位置很重要，因为它决定了计算内核中更新的数据点。每个工作项在其上操作的 ND-Range 的索引是使用 nd_item 类中的内置函数（global_id、work_group_id、sub_group_id 和 local_id）确定的。

内存模型

oneAPI 内存模型定义了主机和设备对内存对象的处理。它帮助用户根据应用程序的需求决定内存分配位置。内存对象分为缓冲区或图像类型。可以使用访问器来指示内存对象的位置和访问模式。访问器为驻留在主机上的对象、设备上的全局内存、设备的本地内存或驻留在主机上的图像提供不同的访问目标。访问类型可以是读、写、原子或读写。

统一共享内存模型允许主机和设备共享内存，而无需使用显式访问器。使用事件的同步管理主机和设备之间的依赖关系。用户可以显式指定一个事件来控制主机或设备更新的数据何时可供重用，或者隐式依赖运行时和设备驱动程序来确定这一点。

内核编程模型

oneAPI 内核编程模型指定在主机和设备上执行的代码。并行性不是自动的。用户需要使用语言结构显式指定它。

DPC++ 语言需要一个能够支持主机端 C++11 及更高版本功能的编译器。然而，设备代码需要一个支持 C++03 功能和某些 C++11 功能（如 lambda 表达式、可变参数模板、右值引用和别名模板）的编译器。它还需要支持 std::string、std::vector 和 std::function。设备代码对某些功能有限制，包括虚函数和虚继承、异常处理、运行时类型信息 (RTTI) 以及使用 new 和 delete 运算符的对象管理。

用户可以选择使用不同的方案来描述主机代码和设备代码之间的分离。lambda 表达式可以将内核代码保留在主机代码行内。函子将主机代码保留在同一个源文件中，但在一个单独的函数中。对于移植 OpenCL 代码的用户，或需要主机代码和设备代码之间显式接口的用户，内核类提供了必要的接口。

用户可以通过三种不同的方式实现并行

• 单个任务，在一个工作项中执行整个内核
• parallel_for 构造，将任务分配给处理元素
• parallel_for_work_group。parallel_for_work_group 构造将任务分配给工作组，并通过使用 barriers 来同步工作组内的所有工作项。

oneAPI 编译模型

oneAPI 编译模型由构建和链接步骤组成。但是，生成的二进制文件需要支持在多个加速器上执行设备代码。这意味着 DPC++ 编译器和链接器必须执行额外的命令来生成二进制文件。这种复杂性通常对用户隐藏，但对于生成特定于目标的二进制文件很有用。

主机代码的编译以标准 x86 架构的默认方式完成。加速器的二进制文件生成更复杂，因为它需要支持单个或多个加速器，以及针对每个加速器的优化。这种称为“fat binary”的加速器二进制文件包含以下组合：

• 中间的标准便携式中间表示 (SPIR-V)，它与设备无关，并在编译过程中生成特定于设备的二进制文件。
• 在编译时生成的特定于目标的二进制文件。由于 oneAPI 旨在支持多个加速器，因此会创建多个代码形式。

多个工具生成这些代码表示，包括 clang 驱动程序、主机和设备 DPC++ 编译器、标准的 Linux* (ld) 或 Windows* (link.exe) 链接器以及用于生成 fat 对象文件的工具。在执行过程中，oneAPI 运行时环境会检查 fat 二进制文件中的特定于设备的映像，如果可用则执行它。否则，将使用 SPIR-V 映像来生成特定于目标的映像。

oneAPI 编程示例

本节将介绍 beta 版 Intel oneAPI DPC++ 编译器、OpenMP 设备卸载和 Intel DPC++ 兼容工具的示例代码。

编写 DPC++ 代码

编写 DPC++ 代码需要用户利用该语言的 API 和语法。清单 1 展示了一些从 C++ (CPU) 代码到 DPC++ (主机和加速器) 代码的示例代码转换。这是在 GitHub4 上发布的 Högbom CLEAN* 算法的实现。该算法迭代地找到图像中的最大值，并减去该点源与观测的 the point spread function 卷积后得到的小增益，直到最大值小于某个阈值。该实现有两个函数：findPeak 和 subtractPSF。这些函数必须按清单 1 和清单 2 所示进行 C++ 到 DPC++ 的移植。

清单 1. subtractPSF 代码的基线和 DPC++ 实现

从 C/C++ 移植到 DPC++ 所需的代码更改包括：

• 引入给定设备的设备队列（使用设备选择器 API）
• 在设备上创建/访问缓冲区（使用 sycl::buffer/get_access API）
• 调用 parallel_for 来启动/执行计算内核
• 等待内核执行完成（并可选地捕获任何异常）
• Intel® DPC++ 编译器和标志：dpcpp -std=c++11 -O2 -lsycl -lOpenCL

清单 2 展示了 findPeak 函数实现的 代码更改。为了更好地利用硬件中的并行性，DPC++ 代码支持 local_work_size、global_id/local_id、workgroup 等许多 API，类似于 OpenCL 和 OpenMP 中使用的构造。

清单 2. findPeak 代码的基线（顶部）和 DPC++（底部）实现。clPeak 是一个包含值和位置数据的结构。通过全局和局部 ID 实现工作组的并发执行，并通过 barriers 实现工作组内多个线程（工作项）的同步。此 parallel_for 执行的结果进一步规约（未显示），以确定跨工作组的最大值和位置。

OpenMP 卸载支持

beta 版 Intel oneAPI HPC Toolkit 提供 OpenMP 卸载支持，使用户能够利用 OpenMP 设备卸载功能。我们来看一个使用 OpenMP pragmas 以 C++ 编写的开源 Jacobi 代码示例3。代码有一个主迭代步骤，该步骤

• 计算 Jacobi 更新
• 计算旧解与新解之间的差异
• 更新旧解
• 计算残差

迭代代码片段在清单 3 中显示。

清单 3. 带有 OpenMP pragmas 的示例 Jacobi 求解器

清单 4 显示了更新后的代码，其中包含 omp target 子句，该子句可用于指定要传输到设备环境的数据，以及一个数据修饰符，可以是 to、from、tofrom 或 alloc。由于数组 b 没有被修改，我们使用 to 子句。由于 x 和 xnew 在卸载指令之前初始化，并在设备环境中更新，因此我们使用 tofrom 子句。规约变量 d 和 r 也在每次迭代中设置和更新，并具有 tofrom 映射子句。

清单 4. 带有 OpenMP 卸载 pragmas 更新的示例 Jacobi 求解器

要使用 oneAPI 编译器编译卸载目标代码，用户需要设置

• 与编译器路径相关的环境变量
• 相关库
• 不同的组件

这些环境变量的路径将取决于用户机器上的 oneAPI 设置。目前，我们将重点介绍编译过程，该过程在不同机器上是相似的，以展示该规范的易用性。要编译代码，请使用基于 LLVM 的 icx 或 icpc -qnextgen 编译器，如下所示：

$ export OMP_TARGET_OFFLOAD=”MANDATORY”
$ export LIBOMPTARGET_DEBUG=1
$./jacobi

-D__STRICT__ANSI 标志确保与 GCC 7.x 及更高版本系统兼容。spir64 标志指的是代码的独立于目标的表示，并在链接阶段或执行阶段移植到特定于目标的代码。要执行代码，请运行这些命令：

OMP_TARGET_OFFLOAD 的 MANDATORY 选项表示卸载必须在 GPU 上运行。默认情况下，它设置为 DEFAULT，表示卸载可以在 CPU 和 GPU 上运行。当设置 LIBOMPTARGET_DEBUG 标志时，它会提供卸载运行时信息，有助于调试。

OpenMP 卸载支持示例是针对 C/C++ 程序的，但 Fortran 卸载也得到支持。这允许拥有 Fortran 代码库的 HPC 用户在 GPU 上运行其代码。

Intel DPC++ 兼容工具

Intel DPC++ 兼容工具是一个基于命令行的代码迁移工具，作为 Intel oneAPI Base Toolkit 的一部分提供。它的主要作用是实现现有 CUDA 源到 DPC++ 的移植。无法自动迁移的源位置将通过适当的错误和警告进行标记。Intel DPC++ 兼容工具还在需要用户干预的源位置插入注释。

图 3 显示了 CUDA 用户可以用来将其源代码移植到 DPC++ 的典型工作流程。Intel DPC++ 兼容工具目前支持 Linux* 和 Windows* 操作系统。本文假设在 Linux 环境下进行。Intel DPC++ 兼容工具目前需要随 CUDA SDK 一起提供的头文件。为了演示迁移过程，我们使用 CUDA SDK 10.1 中的 VectorAdd 示例，通常可以在类似以下位置找到：

$ ls /usr/local/cuda-10.1/samples/0_Simple/vectorAdd
$ icpc -fiopenmp -fopenmp-targets=spir64 -D__STRICT_ANSI__ jacobi.cpp -o
jacobi

VectorAdd 是一个单源示例，代码行数约 150 行。在这种情况下，CUDA 内核设备代码计算数组 A 和 B 与数组 C 的向量加法。

请注意，此处显示的命令、路径和过程在发布时是正确的。最终版本的产品可能会引入一些更改。

要初始化环境以使用 Intel DPC++ 兼容工具，请运行以下命令：

$ source /opt/intel/inteloneapi/setvars.sh

setvars.sh 脚本不仅初始化了 Intel DPC++ 兼容工具的环境，还初始化了 Intel oneAPI Base Toolkit 中提供的所有其他工具。

我们使用一个简化的 CUDA Makefile 版本，如清单 5 所示。

清单 5. 将 CUDA 代码移植到 DPC++ 的 Makefile

下一步拦截 Makefile 执行时发出的命令，并将它们以 JSON 格式存储在编译数据库文件中。Intel DPC++ 兼容工具为此提供了一个名为 intercept-build 的实用程序。以下是一个示例调用：

$ intercept-build make

然后调用实际的转换步骤：

$ dpct -p compile_commands.json --in-root=. --out-root=dpct_output
vectorAdd.cu

--in-root 和 --out-root 标志设置用户程序源文件的位置以及必须写入迁移后的 DPC++ 代码的位置。此步骤生成 ./dpct_output/vectorAdd.dp.cpp。

为了确保向量加法部署到集成 GPU，我们显式指定 GPU 队列，而不是提交到默认队列。通过调用 get_platforms() 和 platform.get_devices() 可以获得支持平台的列表以及每个平台的设备列表。确定目标设备后，为集成 GPU 构建一个队列，并将向量加法内核分派到该队列。可以使用这种方法将多个独立内核定位到连接到同一主机/节点的多个独立设备。

接下来，使用以下命令编译修改后的 DPC++ 代码：

$ dpcpp -std=c++11 -I=/usr/local/cuda-10.1/samples/common/inc
vectorAdd.dp.cpp -lOpenCL

然后调用生成的二进制文件，并确认向量加法正在集成 GPU 上执行，如清单 6 所示。

清单 6. 在集成 GPU 上运行移植后的 DPC++ 代码的输出

有关这些工具的详细信息，请使用以下帮助标志：

$ intercept-build –h
$ dpct –h

为跨多个架构的各种工作负载提供不打折扣的性能

本文介绍了 oneAPI 和 beta 版 Intel oneAPI Toolkits，并概述了 Intel oneAPI Base Toolkit 的组成部分。beta 版本包括帮助用户在 HPC、AI、分析、深度学习、IoT 和视频分析领域过渡到 oneAPI 的工具包。DPC++ 编程指南提供了关于支持优化加速器性能的各种构造的完整详细信息。本文中的 OpenMP 示例是针对 C++ 程序的。但是，GPU 卸载也将支持 C 和 Fortran。oneAPI 提供了您需要在多个加速器上移植和运行代码所需的软件生态系统。

参考文献

1. Intel® VTune™ Profiler
2. Intel® Advisor
3. 使用 OpenMP 的 Jacobi 求解器
4. https://github.com/ATNF/askap-benchmarks/tree/master/tHogbomCleanOMP