Baremetal GPU and FPGA programming for Java using Intel's Level Zero API. This project is a Java Native Interface (JNI) binding for Intel's Level Zero. This library is designed to be as closed as possible to the Level Zero API for C++ API, and it supports a subset of the Level Zero 1.4.0 API (Level Zero as in May 2022). This library is used by TornadoVM, and the subset of the Level Zero API 1.4.0 is the minimum required to dispatch SPIR-V kernels on Intel ARC and Intel integrated GPUs under the TornadoVM Runtime. It is, by no means, a complete port of the Level Zero API, but external contributions are very welcome.

git clone https://github.com/oneapi-src/level-zero.git cd level-zero mkdir build cd build cmake .. cmake --build . --config Release cmake --build . --config Release --target package

Tested Configurations:

• Lenovo IdeaPad Gaming 3 15IHU6
• Dell XPS 8950
• Windows 11
• VS Community 2022
  • components: C++, Git, Spectre mitigated libraries
• CMake 3.26.3, Maven 3.9.1, JDK 21+

Run commands in x64 Native Tools Command Prompt for VS 2022.

git clone https://github.com/oneapi-src/level-zero cd level-zero md build cd build cmake .. cmake --build . --config Release rem check .\bin\Release\zello_world.exe

Note: If zello_world.exe fails, search for existing Level Zero API DLLs (file names start with ze_, e.g. ze_tracing_layer.dll) in c:\windows\system32 and move them to another folder.

Set the paths to the directory of Level Zero installation. Here are examples:

scl enable devtoolset-9 bash # << Only for CentOS/RHEL git clone https://github.com/beehive-lab/levelzero-jni export ZE_SHARED_LOADER="<path-to-levelzero>/build/lib/libze_loader.so" export CPLUS_INCLUDE_PATH=<path-to-levelzero>/include:$CPLUS_INCLUDE_PATH export C_INCLUDE_PATH=<path-to-levelzero>/include:$CPLUS_INCLUDE_PATH cd levelzero-jni/levelZeroLib mkdir build cd build cmake .. make

Note: Run commands in x64 Native Tools Command Prompt for VS 2022.

git clone https://github.com/beehive-lab/levelzero-jni set ZE_SHARED_LOADER=%USERPROFILE%\lab\level-zero\build\lib\release\ze_loader.lib set CPLUS_INCLUDE_PATH=%USERPROFILE%\lab\level-zero\include set C_INCLUDE_PATH=%USERPROFILE%\lab\level-zero\include rem add the folder with Intel's Level Zero API DLLs to PATH set PATH=%USERPROFILE%\lab\level-zero\build\bin\release;%PATH% cd levelzero-jni\levelZeroLib md build cd build cmake .. cmake --build . --config Release

In case you want to compile kernels from OpenCL C to SPIR-V and use the Level Zero API, you need to download the llvm-spirv compiler. The implementation we are currently using is Intel LLVM.

mvn clean package ./scripts/run.sh ## < This script compiles an OpenCL C program to SPIR-V using the llvm-spirv compiler (see 2.1)

The OpenCL C kernel provided for this example is as follows:

__kernel void copyData(__global int* input, __global int* output) { uint idx = get_global_id(0); output[idx] = input[idx]; }

To compile from OpenCL C to SPIR-V:

clang -cc1 -triple spir copyData.cl -O0 -finclude-default-header -emit-llvm-bc -o copyData.bc llvm-spirv copyData.bc -o copyData.spv

Note: Java programs that use levelzero-jni are based on DLLs, which are provided by Intel's Level Zero API. For these programs to find these DLLs, the PATH environment variable must contain the folder that contains the DLLs.

mvn clean package rem add the folder with Intel's Level Zero API DLLs to PATH set PATH=%USERPROFILE%\lab\level-zero\build\bin\release;%PATH% rem copyData.spv file expected in CWD .\scripts\run.cmd # more tests .\scripts\copies.cmd .\scripts\events.cmd .\scripts\fences.cmd .\scripts\kernelTimers.cmd .\scripts\transferTimers.cmd .\scripts\largeBuffers.cmd .\scripts\testZelloWorld.cmd

 // Obtain the Level Zero Driver LevelZeroDriver driver = new LevelZeroDriver(); int result = driver.zeInit(ZeInitFlag.ZE_INIT_FLAG_GPU_ONLY); LevelZeroUtils.errorLog("zeInit", result); int[] numDrivers = new int[1]; result = driver.zeDriverGet(numDrivers, null); LevelZeroUtils.errorLog("zeDriverGet", result); ZeDriverHandle driverHandler = new ZeDriverHandle(numDrivers[0]); result = driver.zeDriverGet(numDrivers, driverHandler); LevelZeroUtils.errorLog("zeDriverGet", result);

ZeContextDescriptor contextDescription = new ZeContextDescriptor(); LevelZeroContext context = new LevelZeroContext(driverHandler, contextDescription); result = context.zeContextCreate(driverHandler.getZe_driver_handle_t_ptr()[0]); LevelZeroUtils.errorLog("zeContextCreate", result);

ZeDeviceMemAllocDescriptor deviceMemAllocDesc = new ZeDeviceMemAllocDescriptor(); deviceMemAllocDesc.setFlags(ZeDeviceMemAllocFlags.ZE_DEVICE_MEM_ALLOC_FLAG_BIAS_UNCACHED); ZeHostMemAllocDescriptor hostMemAllocDesc = new ZeHostMemAllocDescriptor(); hostMemAllocDesc.setFlags(ZeHostMemAllocFlags.ZE_HOST_MEM_ALLOC_FLAG_BIAS_UNCACHED); LevelZeroBufferInteger bufferA = new LevelZeroBufferInteger(); result = context.zeMemAllocShared(context.getContextHandle().getContextPtr()[0], deviceMemAllocDesc, hostMemAllocDesc, bufferSize, 1, device.getDeviceHandlerPtr(), bufferA); LevelZeroUtils.errorLog("zeMemAllocShared", result);

See the full list of examples here.

This project is developed at The University of Manchester, and it is fully open source under the MIT license.

The work was partially funded by the EU Horizon 2020 Elegant 957286 project, and Intel Coorporation.