Vulkan driver

This driver uses libvulkan and clspv to run OpenCL code on GPU devices via Vulkan API.

NOTE: THIS DRIVER IS INCOMPLETE, without an active maintainer. Pull Requests welcomed.

Installation

Required:

  • vulkan drivers (on Ubuntu, “mesa-vulkan-drivers” for opensource vulkan drivers)

  • vulkan development files (on Ubuntu, “vulkan-headers” and “libvulkan-dev”)

  • SPIR-V tools (for clspv; on Ubuntu, package “spirv-tools”)

The Vulkan headers, devices and library must support at least Vulkan version 1.1; 1.0 devices may work but are untested. With 1.0 headers, pocl-vulkan won’t compile.

Optional:

  • “vulkan-validationlayers-dev” for vulkan validation layers

  • “vulkan-tools” or “vulkan-utils” package for vulkaninfo

Note that the Vulkan device MUST support the following extensions (clspv requirements):

  • VK_KHR_variable_pointers

  • VK_KHR_storage_buffer_storage_class

  • VK_KHR_shader_non_semantic_info

Optional extensions:

  • VK_EXT_external_memory_host for CL_MEM_USE_HOST_PTR to be useful

  • VK_KHR_16bit_storage, VK_KHR_8bit_storage, VK_KHR_shader_float16_int8 to be able to use 8 bit and 16 bit integers

Easiest to check is with vulkaninfo utility, they must be listed in ‘Device Extensions’ section.

To build the full pocl-vulkan, first you must build the clspv compiler:

git clone https://github.com/google/clspv.git
cd clspv
python utils/fetch_sources.py
mkdir build ; cd build
cmake /path/to/clspv -DCLSPV_BUILD_TESTS=OFF -DCMAKE_BUILD_TYPE=Release
make -jX
make install

… this will take some time and space, because it compiles its own checkout of LLVM.

After the build, copy “clspv” and “clspv-reflection” binaries to some place CLSPV_BIN_DIR

Then build the vulkan driver:

cmake -DENABLE_HOST_CPU_DEVICES=0 -DENABLE_LLVM=0 -DENABLE_VULKAN=1 -DCLSPV_DIR=${CLSPV_BIN_DIR} <path-to-pocl-source-dir>

You may set VULKAN_SDK env variable before running cmake, then it will look for libvulkan in VULKAN_SDK/lib directory.

After build, libpocl can be tested with (run in the build directory):

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

Adding POCL_VULKAN_VALIDATE=1 POCL_DEBUG=vulkan into the environment enables the use of validation layers, this will make output from PoCL much more verbose.

It is possible to build & use pocl-vulkan without clspv, but this limits the usability of the driver to clCreateProgramWithBinaries() with poclbinaries.

What works

  • both integrated and discrete GPUs are supported

  • buffer (cl_mem) kernel arguments

  • POD (plain old data) kernel arguments (int32 and float32; other int/float types are enabled only if indicated by device features; structs with these types)

  • local memory, both as static (in-kernel) and as kernel argument

  • constant memory, both at module-scope and as kernel argument

  • most 1.2 API calls

  • CL_MEM_USE_HOST_PTR with clCreateBuffer(), if the device supports VK_EXT_external_memory_host

  • global offsets to clEnqueueNDRangeKernel

Doesnt work / missing

  • image / sampler support

  • clLinkProgram & clCompileProgram

  • clCreateBuffer(): CL_MEM_USE_HOST_PTR on dGPUs doesn’t work

  • clCreateBuffer(): the CL_MEM_ALLOC_HOST_PTR flag is ignored

Unfinished / non-optimal

  • missing sub-allocator for small allocations

  • statically sized structs that create certain limits

  • descriptor set should be cached (setup once per kernel, then just update)

  • command buffers should be cached

  • kernel library - check what clspv is missing

  • push constants for POD arguments instead of POD UBO

  • stop using deprecated clspv-reflection, instead extract the kernel metadata from the SPIR-V file itself

Known Bugs

Validation layers can print this message:

“After specialization was applied, VkShaderModule 0xXY0000XY[] does not contain valid spirv for stage VK_SHADER_STAGE_COMPUTE_BIT. The Vulkan spec states: module must be a valid VkShaderModule handle (https://www.khronos.org/registry/vulkan/specs/1.1-extensions/html/vkspec.html#VUID-VkPipelineShaderStageCreateInfo-module-parameter)”

This is (AFAIK) caused by Clspv reflection metadata present in SPIR-V, and is harmless.

The pocl vulkan driver will wait indefinitely for a kernel to finish. However GPU drivers have their own “freeze detection” timeouts and could kill the kernel sooner. This would result in PoCL aborting with error -4 (device lost).

Clspv can compile a lot of code, but is still unfinished and has bugs, so pocl-vulkan may fail to compile OpenCL code.

Testing

The tests that should work with Vulkan driver can be run with tools/scripts/run_vulkan_tests.

This driver was tested with these devices:

  • Intel HD 530 integrated GPU

  • AMD Radeon Vega 8 iGPU

  • Nvidia Quadro P600 discrete GPU

  • Raspberry Pi 4 + Ubuntu 22.04