聚焦 Level Zero

本文概述了 Level Zero 的应用接口、理念、目的和愿景。我们将探讨Level Zero的基本架构及其在计算单元资源底层访问控制方面的优势。它可以与 OpenMP* 和 SYCL* 等语言扩展一起使用。我们将简要介绍Level Zero与SYCL C++语言扩展抽象层之间的交互如何暴露给应用程序开发者。

Level Zero被设计为一个低级API，用于配置和管理对任意数量卸载设备（offload devices）的访问。在此过程中，它还提供了抽象层，使得C++标准兼容的异构计算无需干扰程序流程即可实现。这使得代码可以在不同的运行时环境中移植。了解Level Zero API后端，可以让你超越SYCL或OpenMP语言扩展的抽象，从而增加你的控制级别。

Level Zero接口 是 oneAPI规范 的一部分。它通过裸金属访问CPU、GPU和加速器，补充了oneAPI的基于API的编程模型和直接编程模型。 Intel的参考实现 针对作为 Intel® oneAPI Base Toolkit 的一部分的Intel® GPU，以及与 Intel® oneAPI DPC++/C++ Compiler 的使用，也都得到了充分的文档记录。阅读本文后，您应该拥有触手可及的资源，可以深入研究并开始使用Level Zero运行时或考虑开发自己的运行时。

释放异构计算的潜力

Intel的Level Zero的首次实现针对Intel GPU。然而，Level Zero的愿景和潜力远不止于此。它有可能为特定设备需求创建定制的抽象。它可以适应以支持更广泛的语言功能，如函数指针、内存和I/O。该API旨在跨多种计算设备工作，包括CPU、GPU、现场可编程门阵列（FPGA）和其他加速器架构。

Level Zero可以与其他低级API（如OpenCL*和Vulkan*）共存。然而，它的目的是独立演进，以允许高级oneAPI和SYCL开发者体验保持硬件无关，并尽可能具有架构独立性和灵活性。它还提供了高级运行时API和利用SYCL的库可能希望获得的显式控制。Level Zero是完全开源的，其 规范 、 源代码库 和针对Intel GPU的 计算运行时实现 都可以在GitHub上轻松访问。

简而言之，Level Zero释放了异构计算的愿景，并提供了灵活的开放后端，以实现真正的卸载计算选择。它通过以下功能提供了显式控制系统级接口的能力：
 

• 设备发现
• 内存分配
• 点对点通信
• 进程间共享
• 内核提交
• 异步执行和调度
• 同步原语
• 指标报告
• 系统管理
   

让我们从Level Zero的关键概念开始。

Level Zero基础

Level Zero位于应用程序层之下。它可以作为C++应用程序和目标设备属性（图 1）之间的抽象接口。这些属性可以包括CPU以及其他计算设备。通过这样做，它使开发人员能够无缝地与共享设备资源交互，并将工作负载调度到特定Level Zero驱动程序支持的设备上。驱动程序将支持的设备添加到可用设备列表中，任何SYCL队列都可以映射到该设备并向其提交工作。如果我们不需要访问特定设备属性，或者多个Level Zero设备之间共享的资源，这些设备就可以像使用不同SYCL后端API的任何其他设备一样运行。

Level Zero的真正强大之处在于其底层控制以及对特定设备内存共享或同步上下文对象（context objects）的支持。这不仅为设备增加了透明度，还为Level Zero API提供了更多针对异构卸载目标设备的配置能力。

Level Zero设备检测和选择的顺序流程如下。

Level Zero加载器

访问卸载设备或加速器始于Level Zero加载器。它会发现系统中的设备所对应的Level Zero驱动程序。加载器项目还包含 Level Zero头文件和库 ，使我们能够构建和与Level Zero实现进行交互。驱动程序初始化和发现将在以下代码示例中说明。

// Initialize the driver
zeInit(0);

// Discover all the driver instances
uint32_t driverCount = 0;
zeDriverGet(&driverCount, nullptr);

ze_driver_handle_t* allDrivers = allocate(driverCount *
                                          sizeof(ze_driver_handle_t));
zeDriverGet(&driverCount, allDrivers);

// Find a driver instance with a GPU device
ze_driver_handle_t hDriver = nullptr;
ze_device_handle_t hDevice = nullptr;
for(i = 0; i < driverCount; ++i)
{
    uint32_t deviceCount = 0;
    zeDeviceGet(allDrivers[i], &deviceCount, nullptr);

    ze_device_handle_t* allDevices = allocate(deviceCount *
                                              sizeof(ze_device_handle_t));
    zeDeviceGet(allDrivers[i], &deviceCount, allDevices);

    for(d = 0; d < deviceCount; ++d)
    {
        ze_device_properties_t device_properties {};
        device_properties.stype = ZE_STRUCTURE_TYPE_DEVICE_PROPERTIES;
        zeDeviceGetProperties(allDevices[d], &device_properties);

        if(ZE_DEVICE_TYPE_GPU == device_properties.type)
        {
            hDriver = allDrivers[i];
            hDevice = allDevices[d];
            break;
        }
    }

    free(allDevices);
    if(nullptr != hDriver)
    {
        break;
    }
}

free(allDrivers);
if(nullptr == hDevice)
    return; // no GPU devices found

接下来是创建一个上下文对象，用于管理内存、命令队列、模块、同步对象等。当管理可能被多个设备共享的系统资源时，上下文的使用尤为重要。以下是一个共享内存场景的简单示例。

// Create context(s)
zeContextCreate(hDriver, &ctxtDesc, &hContextA);
zeContextCreate(hDriver, &ctxtDesc, &hContextB);

zeMemAllocHost(hContextA, &desc, 80, 0, &ptrA);
zeMemAllocHost(hContextB, &desc, 88, 0, &ptrB);

调度模型

任何命令都将按照图 2中的调度图所示，被调度并分发到Level Zero设备。

命令被附加到命令列表中，命令列表代表要在卸载计算单元或加速器上执行的一系列命令。命令列表可以通过重置列表来回收，而无需重新创建。通过多次提交相同的命令序列，可以重复使用命令列表，而无需重新附加命令。

然后将命令列表提交给命令队列以执行。队列是一个逻辑对象，与设备中的物理输入流相关联，它可以配置为同步或异步，并且可以组织在队列组中。这种调度模型转化为图 3所示的源代码流程。

即时命令列表

命令列表处理也可以进行优化，以帮助管理延迟。需要保证响应时间的优先级任务可以通过低延迟的即时命令列表来处理。这是一种特殊的命令列表类型，专用于非常低延迟的提交使用模型。

命令列表及其隐式命令队列是使用命令队列描述符创建的。附加到即时命令列表中的命令将立即在设备上执行。附加到即时命令列表中的命令可能会通过阻塞直到命令完成来进行同步执行。即时命令列表在完成之后不需要关闭或重置。但是，使用将会得到遵循，预期的行为也将得到遵守。以下伪代码演示了即时命令列表创建和使用的基本序列。

// Create an immediate command list
ze_command_queue_desc_t commandQueueDesc = {
    ZE_STRUCTURE_TYPE_COMMAND_QUEUE_DESC,
    nullptr,
    computeQueueGroupOrdinal,
    0, // index
    0, // flags
    ZE_COMMAND_QUEUE_MODE_DEFAULT,
    ZE_COMMAND_QUEUE_PRIORITY_NORMAL
};
ze_command_list_handle_t hCommandList;
zeCommandListCreateImmediate(hContext, hDevice, &commandQueueDesc,
                                                &hCommandList);

// Immediately submit a kernel to the device
zeCommandListAppendLaunchKernel(hCommandList, hKernel, &launchArgs, nullptr, 0, nullptr);
...

设想一下，在确定工业系统错误条件后，可能需要立即采取行动的闭环反馈使用场景的丰富性。即时命令列表的概念将GPU卸载计算引入了需要保证响应时间的使用场景。

应用程序开发者的收益

既然我们已经介绍了Level Zero的关键设计原则，现在让我们看看应用程序开发者如何从他们的SYCL应用程序中与之交互。

设备选择

任何具有Level Zero驱动程序实现的设备都可以由应用程序开发者初始化和使用。Intel的Level Zero实现可在 GitHub 上获得。它可以作为其他希望利用 oneAPI 跨架构异构计算支持的设备的参考。

要选择特定的卸载设备，可以使用 `SYCL_DEVICE_FILTER` 环境变量。使用它会影响所有设备查询函数和设备选择器。要检查正在运行的系统上可用于SYCL的设备的可用性，只需使用 `sycl-ls` 命令，例如：

$ sycl-ls
[opencl:acc:0] Intel(R) FPGA Emulation Platform for OpenCL(TM), Intel(R) FPGA Emulation Device 1.2 [2022.13.3.0.16_160000]
[opencl:cpu:1] Intel(R) OpenCL, 11th Gen Intel(R) Core(TM) i7-1185G7 @ 3.00GHz 3.0 [2022.13.3.0.16_160000]
[opencl:gpu:2] Intel(R) OpenCL HD Graphics, Intel(R) Iris(R) Xe Graphics 3.0 [31.0.101.3358]
[ext_oneapi_level_zero:gpu:0] Intel(R) Level-Zero, Intel(R) Iris(R) Xe Graphics 1.3 [1.3.23828]
[host:host:0] SYCL host platform, SYCL host device 1.2 [1.2]

我们可以这样考虑设备属性。驱动程序对象代表系统中物理设备的集合。可能存在多个驱动程序。例如，一个驱动程序可能支持一个供应商的加速器，而另一个驱动程序可能支持另一个供应商的加速器。上下文对象代表设备或系统资源。它们的主要目的是创建和管理可能被多个设备使用的资源。设备对象代表系统中的物理设备。设备发现API用于枚举系统中的设备。`zeDeviceGet()` 函数用于查询驱动程序支持的Level Zero设备的数量，并获取任何设备对象，这些对象是只读的全局构造。每个设备都有一个分配给它的16字节通用唯一全局标识符（UUID）。设备句柄在创建和管理特定于设备的资源时使用。

如果我们想利用Level Zero的附加功能，我们首先需要启用SYCL支持的C++代码与Level Zero API之间的互操作性。如果我们打算直接从C++应用程序中与Level Zero特定设备上下文对象交互，则需要按照所示顺序在我们的源代码中包含以下头文件：

#include "level_zero/ze_api.h"
#include "sycl/backend/level_zero.hpp"

Level Zero后端已添加到 `sycl::backend` 枚举中，使用：

enum class backend {
  // ...
  ext_oneapi_level_zero,
  // ...
};

这样，您就可以使用SYCL命名空间中的 `sycl:get-native` API来请求SYCL对象底层的Level Zero数据结构。

template <backend BackendName, class SyclObjectT>
auto get_native(const SyclObjectT &Obj)
    -> backend_return_t<BackendName, SyclObjectT>
  ext_oneapi_level_zero,

有关更多详细信息，请参阅 Intel® oneAPI Level Zero后端规范 。给定的SYCL队列将附加到系统中的可用设备。

try {
    vector<device> sub_devices = ...;
    for (auto &d : sub_devices)
    {
      // Each queue is in its own context, no data sharing across them.
      auto q = queue(d);
      q.submit([&](handler& cgh) {...});
    }
}

然后使用该设备，并为其分配一个执行队列，就像任何其他SYCL设备一样。真正有趣的地方在于，当我们想要访问特定于设备的资源，或者主机和卸载执行设备之间共享的资源时。让我们以统一共享内存（USM）为例来访问特定于设备的资源。下面的大纲反映了Intel如何实现USM，但对于其他Level Zero库也可以实现相同的功能。

统一共享内存

内存对上层软件栈可见为统一内存，拥有覆盖CPU和GPU的单个虚拟地址空间。线性的、无格式的内存分配在主机应用程序中表示为指针。主机上的指针大小与设备上的指针大小相同。有三种方法可以使用SYCL命名空间分配内存：

1. sycl::malloc_device
   • 分配只能由指定的设备访问，但不能由主机或其他上下文中的设备访问。
   • 数据始终保留在设备上，并且是内核执行可用的最快数据。
   • 需要显式复制才能将数据传输到主机或其他上下文中的设备。
2. sycl::malloc_host
   • 分配可以由主机和其他上下文中的任何设备访问。
   • 数据始终保留在主机上，并可通过外围组件互连（PCI）从设备访问。
   • 无需显式复制即可将数据与主机或设备同步。
3. sycl::malloc_shared
   • 分配只能由主机和指定设备访问。
   • 数据可以（由Level Zero驱动程序操作）在主机和设备之间迁移，以实现更快的访问。
   • 主机和设备之间的数据同步不需要显式复制，但对于上下文中的其他设备则需要。

通过这种方式调用的等效低级Level Zero调用分别是 `zeMemAllocDevice`、`zeMemAllocHost` 和 `zeMemAlloc_shared`。如果您想更深入地了解其实现，请查看 Level Zero核心API规范 。

USM只是Level Zero提供的先进设备感知能力的一个例子。更多细节可以在 Level Zero核心编程指南 中找到。

总结和后续步骤

Level Zero提供了一套丰富的接口来调度工作并管理计算单元和加速器上的内存。它提供了用于加载和执行程序、分配内存以及管理异构工作负载的所有服务。它通过一个开放的接口来实现这一点，该接口可以根据您的特定硬件配置进行定制，同时保持工作负载在不太专业的设置上运行所需的抽象和灵活性。

Level Zero中定义的 对象，如命令队列和命令列表，允许对底层硬件进行低级控制。通过这些以及可用的优化技术，高级编程语言和应用程序可以以接近金属的延迟执行工作负载，从而获得更高的性能。

结合SYCL，它可以与C++无缝使用和访问。此外，我们目前正在利用我们在Python*、Julia*和Java*方面的经验，为Level Zero在不同语言之间提供更好的语言运行时支持。我们希望邀请开源社区和整个行业为Level Zero做出贡献，使其成为一个更加通用和强大的接口，以实现多架构选择。

额外资源

• Intel® Graphics Compute Runtime for oneAPI Level Zero and OpenCL Driver
• oneAPI DPC++/C++ 编译器开发者指南和参考：Intel oneAPI Level Zero
• oneAPI Level Zero 规范
• oneAPI Level Zero 头文件、加载器和验证层