Intel® oneAPI Level Zero Backend Specification

Intel® oneAPI DPC++/C++ Compiler Developer Guide and Reference

Download PDF

ID 767253

Date 11/07/2023

Version

Public

A newer version of this document is available. Customers should click here to go to the newest version.

Visible to Intel only — GUID: GUID-61F09D3E-612A-4DDC-83F6-423BA5A99536

View Details

Intel® oneAPI Level Zero Backend Specification

The Intel® oneAPI Level Zero (Level Zero) extension introduces a Level Zero backend for SYCL. It is built on top of Level Zero runtime enabled with the oneAPI Level Zero Specification. The Level Zero backend aims to provide the best possible performance of SYCL application on a variety of targets supported. The currently supported targets are all Intel GPUs starting with Gen9.

This extension provides a feature-test macro as described in the SYCL spec's section, Feature Test Macros. Any implementation supporting this extension must predefine the macro SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO to one of the values defined in the table below. Applications can test for the existence of this macro to see if the implementation supports this feature, or they can test the macro's value to see the extension APIs the implementation supports:

Value	Description
1	Initial extension version.
2	Added support for the `make_buffer()` API.
3	Added device member to `backend_input_t<backend::ext_oneapi_level_zero, queue>`.
4	Change the definition of `backend_input_t` and `backend_return_t` for the queue object, which changes the API for `make_queue` and `get_native` (when applied to queue).
5	Added support for `make_image()` API.

NOTE:

This extension is following SYCL 2020 backend specification. Prior APIs for interoperability with Level Zero are marked as deprecated and will be removed in the next release.

Prerequisites

The Level Zero loader and drivers must be installed on your system for the SYCL 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 sycl::backend enumeration with:

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

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

Through an Environment Variable

The ONEAPI_DEVICE_SELECTOR environment variable limits the SYCL runtime to use only a subset of the system's devices. By using level_zero for the backend in ONEAPI_DEVICE_SELECTOR, you can select the use of Level Zero as a SYCL 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 ONEAPI_DEVICE_SELECTOR applies filtering to the entire process, this device selector can be used to select the Level Zero backend.

When neither the environment variable nor the filtering device selector is used, the implementation chooses the Level Zero backend for GPU devices supported by the installed Level Zero runtime. The serving backend for a SYCL platform can be queried with the get_backend() member function sycl::platform.

Interoperability with the Level Zero API

The sections below describe the various interoperabilities that are possible between SYCL 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/ext/oneapi/backend/level_zero.hpp"

Mapping of SYCL Objects to Level Zero Handles

These SYCL objects encapsulate the corresponding Level Zero handles:

SYCL Type	backend_return_t <backend::ext_oneapi_level_zero, SyclType>	backend_input_t<backend::ext_oneapi_level_zero, SyclType>
`platform`	`ze_driver_handle_t`	`ze_driver_handle_t`
`device`	`ze_device_handle_t`	`ze_device_handle_t`
`context`	`ze_context_handle_t`	`struct { ze_context_handle_t NativeHandle; std::vector<device> DeviceList; ext::oneapi::level_zero::ownership Ownership{ ext::oneapi::level_zero::ownership::transfer}; }`
`queue`	`ze_command_queue_handle_t`	`struct { ze_command_queue_handle_t NativeHandle; ext::oneapi::level_zero::ownership Ownership{ ext::oneapi::level_zero::ownership::transfer}; }` Deprecated in Version 3 of the Level Zero Backend Specification.
`queue`	`ze_command_queue_handle_t`	`struct { ze_command_queue_handle_t NativeHandle; device Device; ext::oneapi::level_zero::ownership Ownership{ ext::oneapi::level_zero::ownership::transfer}; }` Supported since Version 3 of the Level Zero Backend Specification.
`event`	`ze_event_handle_t`	`struct { ze_event_handle_t NativeHandle; ext::oneapi::level_zero::ownership Ownership{ ext::oneapi::level_zero::ownership::transfer}; }`
`kernel_bundle`	`std::vector<ze_module_handle_t>`	`struct { ze_module_handle_t NativeHandle; ext::oneapi::level_zero::ownership Ownership{ ext::oneapi::level_zero::ownership::transfer}; }`
`kernel`	`ze_kernel_handle_t`	`struct { kernel_bundle<bundle_state::executable> KernelBundle; ze_kernel_handle_t NativeHandle; ext::oneapi::level_zero::ownership Ownership{ ext::oneapi::level_zero::ownership::transfer}; }`
`buffer`	`void *`	`struct { void *NativeHandle; ext::oneapi::level_zero::ownership Ownership{ ext::oneapi::level_zero::ownership::transfer}; }`

Obtaining Native Level Zero Handles from SYCL Objects

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

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

This function is supported for SYCL platform, device, context, queue, event, kernel_bundle, and kernel classes.

The get_native(queue) function returns either ze_command_queue_handle_t or ze_command_list_handle_t depending on the manner in which the input argument queue had been created. Queues created with the SYCL queue constructors have a default setting for whether they use command queues or command lists. The default and how it may be changed is documented in the description for the environment variable SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS. Queues created using make_queue() use either a command list or command queue depending on the input argument to make_queue and are not affected by the default for SYCL queues or the environment variable.

The sycl::get_native<backend::ext_oneapi_level_zero> free-function is not supported for the SYCL buffer class. The native backend object associated with the buffer can be obtained using the interop_hande class as described in the SYCL spec's section, Class interop_handle. The pointer is returned by get_native_mem<backend::ext_oneapi_level_zero> method of the interop_handle class, which is the value returned from a call to zeMemAllocShared(), zeMemAllocDevice(), or zeMemAllocHost() and not directly accessible from the host. You may need to copy your data to the host to access the data. You can get information on the type of the allocation using the type data member of the ze_memory_allocation_properties_t struct that is returned by zeMemGetAllocProperties.

Queue.submit([&](handler &CGH) {
    auto BufferAcc = Buffer.get_access<access::mode::write>(CGH);
    CGH.host_task([=](const interop_handle &IH) {
        void *DevicePtr =
            IH.get_native_mem<backend::ext_oneapi_level_zero>(BufferAcc);
        ze_memory_allocation_properties_t MemAllocProperties{};
        ze_result_t Res = zeMemGetAllocProperties(
            ZeContext, DevicePtr, &MemAllocProperties, nullptr);
        ze_memory_type_t ZeMemType = MemAllocProperties.type;
    });
 }).wait();

Construct a SYCL Object from a Level Zero Handle

The following free functions, defined in the sycl namespace are specialized for the Level Zero backend to allow an application to create a SYCL object that encapsulates a corresponding Level Zero object, see the table below for specific functions.

Level Zero Interoperability Function	Description
`make_platform<backend::ext_oneapi_level_zero>( const backend_input_t< backend::ext_oneapi_level_zero, platform> &)`	Constructs a SYCL platform instance from a Level Zero `ze_driver_handle_t`. The SYCL execution environment contains a fixed number of platforms that are counted with `sycl::platform::get_platforms()`. Calling this function does not create a new platform. Rather it merely creates a `sycl::platform` object that is a copy of one of the platforms from that enumeration.
`make_device<backend::ext_oneapi_level_zero>( const backend_input_t< backend::ext_oneapi_level_zero, device> &)`	Constructs a SYCL device instance from a Level Zero `ze_device_handle_t`. The SYCL execution environment for the Level Zero backend contains a fixed number of devices that are counted with `sycl::device::get_devices()` and a fixed number of sub-devices that are counted with `sycl::device::create_sub_devices(...)`. Calling this function does not create a new device. Rather it merely creates a `sycl::device` object that is a copy of one of the devices from those enumerations.
`make_context<backend::ext_oneapi_level_zero>( const backend_input_t< backend::ext_oneapi_level_zero, context> &)`	Constructs a SYCL context instance from a Level Zero `ze_context_handle_t`. The context is created against the devices passed in a `DeviceList` structure member. There must be at least one device given and all the devices must be from the same SYCL platform and from the same Level Zero driver. The `Ownership` input structure member specifies if the SYCL runtime should take ownership of the passed native handle. The default behavior is to transfer the ownership to the SYCL runtime. See section Level Zero Handle Ownership and Thread-safety for details.
`make_queue<backend::ext_oneapi_level_zero>( const backend_input_t< backend::ext_oneapi_level_zero, queue> &, const context &Context)`	Constructs a SYCL queue instance from a Level Zero `ze_command_queue_handle_t`. The `Context` argument must be a valid SYCL context encapsulating a Level Zero context. The `Device` input structure member specifies the device to create the `queue` against and must be in `Context`. The `Ownership` input structure member specifies if the SYCL runtime should take ownership of the passed native handle. The default behavior is to transfer the ownership to the SYCL runtime. See Level Zero Handle Ownership and Thread-safety for details. If the deprecated variant of `backend_input_t<backend::ext_oneapi_level_zero, queue>` is passed to `make_queue`, the queue is attached to the first device in `Context`. Starting in version 4 of this specification, `make_queue()` can be called by passing either a Level Zero `ze_command_queue_handle_t` or a Level Zero `ze_command_list_handle_t`. Queues created from a Level Zero immediate command list (`ze_command_list_handle_t`) generally perform better than queues created from a standard Level Zero `ze_command_queue_handle_t`. See the Level Zero documentation of these native handles for more details. Also starting in version 4 the `make_queue()` function accepts a Properties member variable. This can contain any of the SYCL properties that are accepted by the SYCL queue constructor, except the `compute_index` property which is built into the command queue or command list.
`make_event<backend::ext_oneapi_level_zero>( const backend_input_t< backend::ext_oneapi_level_zero, event> &, const context &Context)`	Constructs a SYCL event instance from a Level Zero `ze_event_handle_t`. The `Context` argument must be a valid SYCL context encapsulating a Level Zero context. The Level Zero event should be allocated from an event pool created in the same context. The `Ownership` input structure member specifies if the SYCL runtime should take ownership of the passed native handle. The default behavior is to transfer the ownership to the SYCL runtime. See Level Zero Handle Ownership and Thread-safety for details.
`make_kernel_bundle<backend::ext_oneapi_level_zero, bundle_state::executable>( const backend_input_t< backend::ext_oneapi_level_zero, kernel_bundle<bundle_state::executable>> &, const context &Context)`	Constructs a SYCL `kernel_bundle` instance from a Level Zero `ze_module_handle_t`. The `Context` argument must be a valid SYCL context encapsulating a Level Zero context, and the Level Zero module must be created on the same context. The Level Zero module must be fully linked (it cannot require further linking through `zeModuleDynamicLink`). The SYCL `kernel_bundle` is created in the executable state. The `Ownership` input structure member specifies if the SYCL runtime should take ownership of the passed native handle. The default behavior is to transfer the ownership to the SYCL runtime. See Level Zero Handle Ownership and Thread-safety for details. If the behavior is `transfer`, then the runtime is going to destroy the input Level Zero module, and the application must not have any outstanding `ze_kernel_handle_t` handles to the underlying `ze_module_handle_t` by the time this interoperability `kernel_bundle` destructor is called.
`make_kernel<backend::ext_oneapi_level_zero>( const backend_input_t< backend::ext_oneapi_level_zero, kernel> &, const context &Context)`	Constructs a SYCL kernel instance from a Level Zero `ze_kernel_handle_t`. The `KernelBundle` input structure specifies the `kernel_bundle` corresponding to the Level Zero module from which the kernel is created. There must be exactly one Level Zero module in the `KernelBundle`. The `Context` argument must be a valid SYCL context encapsulating a Level Zero context, and the Level Zero module must be created on the same context. The `Ownership` input structure member specifies if the SYCL runtime should take ownership of the passed native handle. The default behavior is to transfer the ownership to the SYCL runtime. See Level Zero Handle Ownership and Thread-safety for details. If the behavior is `transfer`, then the runtime is going to destroy the input Level Zero kernel.
`template <backend Backend, typename T, int Dimensions = 1, typename AllocatorT = buffer_allocator<std::remove_const_t<T>>> buffer<T, Dimensions, AllocatorT> make_buffer( const backend_input_t<Backend, buffer<T, Dimensions, AllocatorT>> &, const context &Context)`	This API is available starting with revision 2 of the Level Zero Backend Specification. Construct a SYCL buffer instance from a pointer to a Level Zero memory allocation. The pointer must be the value returned from a previous call to `zeMemAllocShared()`, `zeMemAllocDevice`(), or `zeMemAllocHost()`. The input SYCL context `Context` must be associated with a single device, matching the device used at the prior allocation. The `Context` argument must be a valid SYCL context encapsulating a Level Zero context, and the Level Zero memory must be allocated on the same context. Created SYCL buffer can be accessed in another contexts, not only in the provided input context. The `Ownership` input structure member specifies if the SYCL runtime should take ownership of the passed native handle. The default behavior is to transfer the ownership to the SYCL runtime. See Level Zero Handle Ownership and Thread-safety for details. If the behavior is `transfer`, then the runtime is going to free the input Level Zero memory allocation. Synchronization rules for a buffer that is created with this API are described in Interoperability Buffer Synchronization Rules.
`template <backend Backend, typename T, int Dimensions = 1, typename AllocatorT = buffer_allocator<std::remove_const_t<T>>> buffer<T, Dimensions, AllocatorT> make_buffer( const backend_input_t<Backend, buffer<T, Dimensions, AllocatorT>> &, const context &Context, event`	This API is available starting with revision 2 of the Level Zero Backend Specification. Construct a SYCL buffer instance from a pointer to a Level Zero memory allocation. Refer to `make_buffer` description above for semantics and restrictions. The additional `AvailableEvent` argument must be a valid SYCL event. The instance of the SYCL buffer class template being constructed must wait for the SYCL event parameter to signal that the memory native handle is ready to be used.
`template<backend Backend, int Dimensions = 1, typename AllocrT = sycl::image_allocator> image<Dimensions, AllocrT> make_image( const backend_input_t<Backend, image<Dimensions, AllocrT>> &backendObject, const context &targetContext);`	This API is available starting with revision 5 of the Level Zero Backend Specification. Construct a SYCL image instance from a `ze_image_handle_t`. Because Level Zero has no way of getting image information from an image, it must be provided. The `backend_input_t` is a `struct` type: `struct type { ze_image_handle_t ZeImageHandle; sycl::image_channel_order ChanOrder; sycl::image_channel_type ChanType; sycl::range<Dimensions> Range; ext::oneapi::level_zero::ownership Ownership{ ext::oneapi::level_zero::ownership::transfer}; };` where the Range should be ordered (width), (width, height), or (width, height, depth) for 1D, 2D and 3D images respectively, with those values matching the dimensions used in the `ze_image_desc` that was used to create the `ze_image_handle_t` initially. Note that the range term ordering (width first, depth last) is true for SYCL 1.2.1 images that are supported here. But future classes like `sampled_image` and `unsampled_image` might have a different ordering. Example: ze_image_handle_t ZeHImage; // ... user provided LevelZero ZeHImage image // handle gotten somehow (possibly zeImageCreate) // the informational data that matches ZeHImage sycl::image_channel_order ChanOrder = sycl::image_channel_order::rgba; sycl::image_channel_type ChanType = sycl::image_channel_type::unsigned_int8; size_t width = 4; size_t height = 2; sycl::range<2> ImgRange_2D(width, height); constexpr sycl::backend BE = sycl::backend::ext_oneapi_level_zero; sycl::backend_input_t<BE, sycl::image<2>> ImageInteropInput{ ZeHImage, ChanOrder, ChanType, ImgRange_2D, sycl::ext::oneapi::level_zero::ownership::transfer }; sycl::image<2> Image_2D = sycl::make_image<BE, 2>(ImageInteropInput, Context); The image can only be used on the single device where it was created. This limitation may be relaxed in the future. The `Context` argument must be a valid SYCL context encapsulating a Level-Zero context, and the Level-Zero image must have been created on the same context. The created SYCL image can only be accessed from kernels that are submitted to a queue using this same context. The `Ownership` input structure member specifies if the SYCL runtime should take ownership of the passed native handle. The default behavior is to transfer the ownership to the SYCL runtime. If the behavior is transfer then the SYCL runtime is going to free the input Level-Zero memory allocation, meaning the memory will be freed when the ~image destructor fires. When using transfer the ~image destructor may not need to block. If the behavior is keep, then the memory will not be freed by the ~image destructor, and the ~image destructor blocks until all work in the queues on the image have been completed. When using keep it is the responsibility of the caller to free the memory appropriately.
`template<backend Backend, int Dimensions = 1, typename AllocrT = sycl::image_allocator> image<Dimensions, AllocrT> make_image( const backend_input_t<Backend, image<Dimensions, AllocrT>> &backendObject, const context &targetContext, event availableEvent);`	This API is available starting with revision 5 of the Level Zero Backend Specification. Construct a SYCL image instance from a pointer to a Level Zero memory allocation. Please refer to `make_image` description above for semantics and restrictions. The additional `AvailableEvent` argument must be a valid SYCL event. The instance of the SYCL image class template being constructed must wait for the SYCL event parameter to signal that the memory native handle is ready to be used.

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. By default, the ownership is transferred to the SYCL runtime, but some interoperability API supports overriding this behavior and keeps the ownership in the application. Use this enumeration for explicit specification of the ownership:

namespace sycl {
namespace ext {
namespace oneapi {
namespace level_zero {

enum class ownership { transfer, keep };

} // namespace level_zero
} // namespace oneapi
} // namespace ext
} // namespace sycl

SYCL Runtime Takes Ownership (default): Whenever the application creates a SYCL object from the corresponding Level Zero handle, with one of the make_* functions, the SYCL runtime takes ownership of the Level Zero handle if no explicit ownership::keep was specified. The application must not use the Level Zero handle after the last host copy of the SYCL object is destroyed. The application must not destroy the Level Zero handle. For more information, see the SYCL Common Reference Semantics section.
Application Keeps Ownership (explicit): If a SYCL object is created with an interoperability API explicitly asking to keep the native handle ownership in the application with ownership::keep, then the SYCL runtime does not take the ownership and will not destroy the Level Zero handle at the destruction of the SYCL object. The application is responsible for destroying the native handle when it no longer needs it, but it must not destroy the handle before the last host copy of the SYCL object is destroyed (as described in the core SYCL specification under SYCL Common Reference Semantics.
Obtaining Native Handle Does Not Change Ownership: The application may call the get_native<backend::ext_oneapi_level_zero> free function on a SYCL object to retrieve the underlying Level Zero handle. Doing so does not change the ownership of the Level Zero handle. The application may not use this handle after the last host copy of the SYCL object is destroyed (as described in the core SYCL specification under SYCL Common Reference Semantics unless the SYCL object was created by the application with ownership::keep.
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 SYCL runtime takes ownership of the Level Zero handles and should not attempt further direct use of those handles.

Interoperability Buffer Synchronization Rules

A SYCL buffer that is constructed with this interop API uses the Level Zero memory allocation for its full lifetime. The contents of the Level Zero memory allocation are unspecified for the lifetime of the SYCL buffer. If the application modifies the contents of that Level Zero memory allocation during the lifetime of the SYCL buffer, the behavior is undefined. The initial contents of the SYCL buffer will be the initial contents of the Level Zero memory allocation at the time of the SYCL buffer's construction.

The behavior of the SYCL buffer destructor depends on the Ownership flag. As with other SYCL buffers, this behavior is triggered only when the last reference count to the buffer is dropped, as described in the SYCL spec's section, Buffer Synchronization Rules.

If the ownership is keep (the application retains ownership of the Level Zero memory allocation), then the SYCL buffer destructor blocks until all work in queues on the buffer have completed. The contents of the buffer is not copied back to the Level Zero memory allocation.
If the ownership is transfer (the SYCL runtime has ownership of the Level Zero memory allocation), then the SYCL buffer destructor does not need to block, even if work on the buffer has not completed. The SYCL runtime frees the Level Zero memory allocation asynchronously when it is no longer in use in queues.

Level Zero Additional Functionality

Device Information Descriptors

The Level Zero backend provides the following device information descriptors that an application can use to query information about a Level Zero device. Applications use these queries with the device::get_backend_info<>() member function as shown in the example below, which illustrates the free_memory query:

sycl::queue Queue;
auto Device = Queue.get_device();

size_t freeMemory =
  Device.get_backend_info<sycl::ext::oneapi::level_zero::info::device::free_memory>();

New descriptors have been added as part of this specification, and are described in the table and example below.

Descriptor	Description
`sycl::ext::oneapi::level_zero::info::device::free_memory`	Returns the number of bytes of free memory for the device.

namespace sycl{
namespace ext {
namespace oneapi {
namespace level_zero {
namespace info {
namespace device {

struct free_memory {
    using return_type = size_t;
};

} // namespace device;
} // namespace info
} // namespace level_zero
} // namespace oneapi
} // namespace ext
} // namespace sycl

Parent topic: Intel® oneAPI Level Zero

Select Your Language

Using Intel.com Search

Quick Links

Recent Searches

Advanced Search

Only search in

Intel® oneAPI DPC++/C++ Compiler Developer Guide and Reference

Intel® oneAPI Level Zero Backend Specification

Prerequisites

User-visible Level Zero Backend Selection and Default Backend

Interoperability with the Level Zero API

Level Zero Additional Functionality