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