Intel® oneAPI Level Zero Backend Specification

Introduction

This extension introduces a Level Zero backend for Data Parallel C++ (DPC++), which is built on top of Level Zero runtime enabled with the oneAPI Level Zero Specification. The supported targets are Intel GPUs, starting with Gen9.

Note

This specification is a draft. It is not complete or exhaustive in its descriptions. More information, including explanations on mapping the Data Parallel C++ (DPC++) programming model to a Level Zero API, is forthcoming. In the future, it will conform to the SYCL* 2020 spec.

Prerequisites

The Level Zero loader and drivers must be installed on your system for the DPC++ runtime to recognize and enable the Level Zero backend. Visit Intel® oneAPI DPC++/C++ Compiler System Requirements for specific instructions.

User-visible Level Zero Backend Selection and Default Backend

The Level Zero backend is added to the cl::sycl::backend enumeration with:

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

The sections below explain the different ways the Level Zero backend can be selected.

Through an Environment Variable

The SYCL_DEVICE_FILTER environment variable limits the DPC++ runtime to use only a subset of the system's devices. By using level_zero for the backend in SYCL_DEVICE_FILTER, you can select the use of Level Zero as a DPC++ backend. For more information, see the Environment Variables.

Through a Programming API

The Filter Selector extension is described in SYCL* Proposals: Filter Selector. Similar to how the SYCL_DEVICE_FILTER applies filtering to the entire process, this device selector can be used to programmatically select the Level Zero backend.

If the environment variable or filtering device selector is NOT used, the implementation chooses the Level Zero backend for GPU devices that are supported by the installed Level Zero runtime. The serving backend for a DPC++ platform can be queried with the get_backend() member function of the cl::sycl::platform command.

Interoperability with the Level Zero API

The sections below describe the various interoperabilities that are possible between DPC++ and Level Zero. The application must include the following headers to use any of the inter-operation APIs described in this section. These headers must be included in the order shown:

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

Mapping of DPC++ Objects to Level Zero Handles

These DPC++ objects encapsulate the corresponding Level Zero handles:

DPC++ Object	Level Zero Handle
Platform	`ze_driver_handle_t`
Device	`ze_device_handle_t`
Context	`ze_context_handle_t`
Queue	`ze_command_queue_handle_t`
Program	`ze_module_handle_t`

Obtaining Built-in Level Zero Handles from DPC++ Objects

The get_native<cl::sycl::backend::level_zero>() member function is how you can use a raw native Level Zero handle to obtain a specific DPC++ object. The function is supported for the DPC++platform, device, context, queue, event and program classes. You can use a free-function defined in the cl::sycl namespace instead of the member function with:

template <backend BackendName, class SyclObjectT>
auto get_native(const SyclObjectT &Obj) ->
	    typename interop<BackendName, SyclObjectT>::type;

Construct a DPC++ Object from a Level Zero Handle

The following free functions, defined in the cl::sycl::level_zero namespace, allow an application to create a DPC++ object that encapsulates a corresponding Level Zero object:

Level Zero Interoperability Function	Description
`make<platform>(ze_driver_handle_t);`	Constructs a DPC++ platform instance with `ze_driver_handle_t`.
`make<device>(const platform &, ze_device_handle_t);`	Constructs a DPC++ device instance with `ze_device_handle_t`. The platform argument gives a DPC++ platform, which encapsulates a Level Zero driver that supports the passed Level Zero device.
`make<context>(const vector_class<device> &, ze_context_handle_t);`	Constructs a DPC++ context instance with `ze_context_handle_t`. The context is created against the devices that are passed in. You must give at least one device and all the devices must be from the same DPC++ platform (from the same Level Zero driver).
`make<queue>(const context &, ze_command_queue_handle_t);`	Constructs a DPC++ queue instance with `ze_command_queue_handle_t`. The context argument must be a valid DPC++ context that encapsulates a Level Zero context. The queue is attached to the first device in the passed DPC++ context.
`make<program>(const context &, ze_module_handle_t);`	Constructs a DPC++ program instance with `ze_module_handle_t`. The context argument must be a valid DPC++ that encapsulates a Level Zero context. The Level Zero module must be fully linked (example: it does not require further linking through `zeModuleDynamicLink`) and then the DPC++ program is created in the linked state.

Level Zero Handle Ownership and Thread-safety

The Level Zero runtime does not do reference-counting of its objects, so it is crucial to adhere to these practices of how Level Zero handles are managed:

DPC++ Runtime Takes Ownership: Whenever the application creates a DPC++ object from the corresponding Level Zero handle, via one of the make<T>() functions, the DPC++ runtime takes ownership of the Level Zero handle. The application must not use the Level Zero handle after the last host copy of the DPC++ object is destroyed. The application must not destroy the Level Zero handle. For more information see the SYCL Common Reference Semantics section: https://www.khronos.org/registry/SYCL/specs/sycl-2020/pdf/sycl-2020.pdf
DPC++ Runtime Assumes Ownership: The application may call the get_native<T>() member function of a DPC++ object to retrieve the underlying Level Zero handle, however, the DPC++ runtime continues to retain ownership of this handle. The application must not use this handle after the last host copy of the DPC++ object is destroyed. The application must not destroy the Level Zero handle. For more information see the SYCL Common Reference Semantics section: https://www.khronos.org/registry/SYCL/specs/sycl-2020/pdf/sycl-2020.pdf
Considerations for Multi-threaded Environment: The Level Zero API is not thread-safe, refer to Multithreading and Concurrency for more information. Applications must make sure that the Level Zero handles are not used simultaneously from different threads. The DPC++ runtime takes ownership of the Level Zero handles and should not attempt further direct use of those handles.