Level Zero driver

This driver uses libze and LLVM/Clang to run OpenCL code on GPU devices via Level Zero API.

The implementation is now (as of Mar 2025) able to pass OpenCL-CTS tests, however improvements to performance and functions are still ongoing.

Installation

Required:

  • Clang+LLVM: 18, 19 and 20 should work, older may work but are untested. It’s highly recommended to use a LLVM built with static component libraries (this is the default)

  • Level Zero ICD + development files (level-zero and level-zero-devel)

  • Level Zero drivers (on Ubuntu, intel-level-zero-gpu)

  • SPIRV-LLVM-Translator from Khronos (https://github.com/KhronosGroup/SPIRV-LLVM-Translator) Must be built for the corresponding Clang/LLVM branch. PoCL will utilize libLLVMSPIRV library if it’s available, otherwise it will fallback to llvm-spirv binary

Note the driver is known to fail more tests with LLVM 16 because of unsolved bugs related to opaque-pointer changes in LLVM and SPIRV-LLVM-Translator.

The libze_loader + headers should support at least Level Zero specification version 1.11; older may work but are untested.

To build the Level Zero driver driver (with recommended options):

cmake -DENABLE_LEVEL0=1 -DENABLE_LLVM=1 -DSTATIC_LLVM=ON -DWITH_LLVM_CONFIG=/path/to/bin/llvm-config -DLLVM_SPIRV=/path/to/llvm-spirv <path-to-pocl-source-dir>

Using STATIC_LLVM is recommended to avoid problems with symbol conflicts (the LevelZero driver uses its own version of LLVM), even though some LLVM versions might work when linked as shared libraries.

For additional CMake variables see CMake options.

After build, it can be tested without installation (in the build directory):

OCL_ICD_VENDORS=$PWD/ocl-vendors/pocl-tests.icd POCL_BUILDING=1 POCL_DEVICES=level0 ./examples/example1/example1

or using the test runner script:

./tools/scripts/run_level0_tests

This assumes that libOpenCL.so is the opensource ocl-icd loader; for other ICD loaders you will need to somehow point them to the built libpocl.so. For the meaning of environment variables, see Tuning pocl behavior with ENV variables.

Known issues / broken features

These features are currently disabled when PoCL is built with ENABLE_CONFORMANCE=ON.

  • Images: host synchronization when CL_MEM_USE_HOST_PTR is used works with buffers, but doesn’t work with Images.

  • Subgroups: require device queries which aren’t yet available through L0 API

  • Program-scope variables: fails a particular use case from OpenCL-cts

  • FP64: this is emulated on most consumer hardware, and fails math tests

  • FP16: this is available on consumer hardware, but fails some CTS math tests

  • 64bit atomics: can cause a GPU hang

The last two features may work on some hardware but fail on other; the other features require fixes in software.

Extra features / tunables

POCL_LEVEL0_JIT env variable can be used to enable JIT compilation (kernels are compiled lazily only when launched via clEnqueueNDRangeKernel, instead of eagerly at clBuildProgram-time). Useful with programs that have thousands of kernels (e.g. from heavily templated code). See Tuning pocl behavior with ENV variables for accepted values.

Testing

The tests that should work with Level Zero driver can be run with tools/scripts/run_level0_tests.

This driver was tested with these de0vices:

  • Intel Tiger Lake integrated GPU

  • Intel Alder Lake integrated GPU

  • Intel Arc A750