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