Vulkan

Experimental Vulkan compute backend. This backend is under active development and not all SYCL features are supported, or all possible Vulkan backend devices tested.

Vulkan consumes SPIR-V kernels which are required to match the Vulkan SPIR-V environment requirements, rather than the OpenCL SPIR-V environment requirements like for the Level-Zero and OpenCL backends. In order to turn SYCL kernels into Vulkan consumable SPIR-V the clspv tool is used as part of the generic SSCP compilation flow.

Requirements

AdaptiveCpp should be built against LLVM version 19 or higher.
For a Vulkan backend device to be reported through SYCL a Vulkan 1.2 or later physical device is required. Vulkan 1.2 is the minimum version as it is the release where timeline semaphore and buffer device address features are made core (optional) features rather than through VK_KHR_timeline_semaphore and VK_KHR_buffer_device_address respectively.
The LunarG Vulkan SDK versions 1.4 or later for Vulkan loader, layers, headers, and other tools.
A clspv executable from commit 8286ad49be1fa8708f07aa6852e1b84ed4dd0861.
Linux, macOS, and Windows operating systems are tested in CI. Ubuntu 22.04 and later are the tested distributions for Linux. MacOS 15.7.4 is tested CI with MoltenVK. Windows server 2022 is used with llvmpipe in CI. However on Windows there are sporadic failures in the group_functions suite of SYCL tests, during ~buffer_impl there is a crash in the loop calling event::wait() on pending events, so these are filtered out of the test executable invocation. Adding support for Android is future work.
A Vulkan driver, use vulkaninfo from the Vulkan SDK to see the available devices on your system.

Capabilities

The following physical device features are required to support for a Vulkan device being reported as a SYCL backend:

bufferDeviceAddress - Enables device USM implementation by allowing the runtime to get a device pointer to return to the user.
timelineSemaphore - Used to implement synchronization between command submissions.
shaderSubgroupExtendedTypes - Enables support for subgroup operations in kernels using i8, i16, i64, and f16 types. There is no SPIR-V capability to check for this lazily.

Additionally, these shader capabilities are lazily checked against device support when a SPIR-V kernel is loaded, and an error printed if the device doesn't support them:

shaderInt8 - Enables 8-bit types in SPIR-V, allows SYCL kernels to use char and uchar types.
shaderFloat16 - Enables 16-bit floating point types in SPIR-V, allows SYCL kernels to use half types.
shaderInt16 - Enables 16-bit integers in SPIR-V, allows SYCL kernels to use short and ushort types.
shaderInt64 - Enables 64-bit integers in SPIR-V, allows SYCL kernels to use long and ulong types.
storagePushConstant8 - Enables 8-bit integers to be passed in push constants, allowing kernels to pass 8-bit arguments by value.
storagePushConstant16 - Enables 16-bit integers to be passed in push constants, allowing kernels to pass 16-bit arguments by value.
variablePointers - Allows a SPIR-V pointer to not know what statically what object is comes from. Enables lowering of LLVM OpPhi instructions.
variablePointersStorageBuffer - Allows SPIR-V pointer to cross blocks, generated by clspv but may not be require as all arguments are passed in a single buffer due to single kernel per TU compilation model.
physicalStorageBufferAddresses - Use PhysicalStorageBuffer64 memory model where physical pointers are allowed.
VK_SUBGROUP_FEATURE_BASIC_BIT - Allows SPIR-V CapabilityGroupNonUniform instructions to be used in the implementation of SYCL sub-group builtins.
VK_SUBGROUP_FEATURE_VOTE_BIT - Allows SPIR-V CapabilityGroupNonUniformVote instructions to be used in the implementation of SYCL sub-group builtins.
VK_SUBGROUP_FEATURE_SHUFFLE_BIT - Allows SPIR-V CapabilityGroupNonUniformShuffle instructions to be used in the implementation of SYCL sub-group builtins.

Building

The LunarG Vulkan SDK is a dependency for building the Vulkan backend, as it provides the Vulkan layers & loader along with SPIR-V Tools, and should be available in your system path after using the setup-env.sh script it ships. A clspv binary that is detectable by find_program is also a dependency.

Combined with the -DWITH_VULKAN_BACKEND=ON option for enabling the Vulkan backend in the build, the relevant parts of CMake invocation are:

$ cmake -DWITH_VULKAN_BACKEND=ON -DCMAKE_PROGRAM_PATH=<path/to/clspv/build/bin/dir>

The Vulkan validation layer is enabled in debug builds.

Running

acpp-info can be used to list all of the Vulkan devices available on your system, alongside the devices for other backends. For example,

$ acpp-info -l
=================Backend information===================
Loaded backend 0: OpenMP
  Found device: AdaptiveCpp OpenMP host device
Loaded backend 1: Vulkan
  Found device: AMD Radeon Graphics (RADV PHOENIX)
  Found device: llvmpipe (LLVM 20.1.2, 256 bits)

The SYCL device selectors will tend to prefer GPU devices where available, but to narrow down the selection of visible devices the ACPP_VISIBILITY_MASK=vk environment variable can be used. Optionally setting a individual device, for example in the output above ACPP_VISIBILITY_MASK=vk:0 would select the "AMD Radeon Graphics" Vulkan device ACPP_VISIBILITY_MASK=vk:1 would select the llvmpipe Vulkan device.

Limitations

Due to the challenge of supporting a generic address space in SYCL/PCUDA on Vulkan SPIR-V PhysicalStorageBuffer64 memory model, kernels cannot currently chase pointers in a linked list fashion or store to global memory the addresses of variables.

All the runtime functionality and synchronization of SYCL is expected to work, although some optional features are not supported. Most notably only device USM is available, and not host or shared USM. However, there is no support of backend interop at this time.

The holes in support arise from deficiencies in the compilation of kernel code. Features which are not supported for use in kernels targeting Vulkan SPIR-V are:

Floating point atomics - Not supported, will generate incorrect results rather than error.
marray - Will fail to compile. SPIR-V does not support non-standard integer types without the SPV_INTEL_arbitrary_precision_integers extension which clspv does not support. However, using marray can generate such types in LLVM IR, for example an i48 for marray<short, 3>. See the chipStar HipPromoteInts LLVM pass as a possible way to resolve this.
group function variants - Unstable support, will fail to compile or generate incorrect results.

Device Support

In additional to the restrictions around supported features individual devices will have their own bugs that need addressed (which may turn out to be not device specific once investigated).

Devices known to be well supported by the backend include:

Driver Name	Driver Version	Device Type	Status
llvmpipe	Mesa 25.0.7	CPU	Well supported
RADV Pheonix	Mesa 25.0.7	AMD Integrated GPU	Well supported
RADV MI100	Mesa 23.2.1	AMD Discrete GPU	Well Supported
RADV MI210	Mesa 23.2.1	AMD Discrete GPU	Many Issues
Arc MTL	Mesa 25.0.7	Intel Integrated GPU	Issues with sub-groups
RTX 500	NVIDIA 580.95.5.0	NVIDIA Discrete GPU	Tests pass in isolation, but device stops being detected when running full sycl suite
MoltenVK	Khronos 1.4.1	Apple Integrated GPU	CI testing with Macos 15.7.4, some tests disabled
Swiftshader	Google 5.0.0	CPU	Poor support, no `Int64` or `VariablePointer` capabilities

Other devices are untested and support status is unknown.

Benchmarks

The pcuda branch of illuhad/benchmark-scipts contains benchmarks that work with SYCL and PCUDA. No thorough benchmarking has been done but an initial assessment of the status of the benchmarks is as follows on llvmpipe 25.0.7.

Benchmark	Status
HeCbench dslash	Completes successfully on SYCL and PCUDA
HeCbench fdtd3d	Can't compile kernel for SYCL or PCUDA
HeCbench FFT	SYCL & PCUDA compiles kernel, but fails verification
HeCbench ising	Completes successfully on SYCL and PCUDA
HeCbench mandlebrot	Completes successfully on SYCL and PCUDA
HeCbench nbody	Completes successfully on SYCL and PCUDA
HeCbench rsbench	Can't compile kernel for SYCL or PCUDA
HeCbench sph	Completes successfully on SYCL and PCUDA
Bude	SYCL verification fails, PCUDA device not compatible
Cloverleaf	SYCL verification fails and PCUDA raises an error

Compilation fail investigations:

HecBench fdtd3d - LLVM-IR contains llvm.memmove intrinsic which is not implemented in clspv.
Hecbench rsbench - Error processing GEP into alloca array of type { double, double } struct.

Benchmark Results

Initial benchmarking results are below. Comparing the 5 correctly running HeCbench benchmarks across different AdaptiveCpp SYCL backends, using an absolute time metric so "lower is better". We can see that when comparing against OpenCL or CUDA on the same hardware (2 v 3, 4 v 5, 6 v 7) Vulkan is always less performant. In some cases, such as dslash on NVIDIA Ada, it is an order of magnitude slower. These cases need to be analysed and improved as future work.

benchmark	OMP[1]	OCL CPU[2]	VK CPU[3]	OCL GPU[4]	VK MTL GPU[5]	CUDA GPU[6]	VK Ada GPU[7]
nbody	48.6	9.8	35.4	1.6	3.7	1.3	4.2
mandlebrot	22.3 / 21	4.6 / 3.3	36 / 27	2.5 / 0.7	42.1 / 2.5	2.6 / 0.2	63 / 5.5
ising	3.2	2.1	3.7	0.3	0.8	0.1	3.4
dslash	2.3	2	3.9	2.2	2.3	1.4	160
sph	28.2	53	36	22.7	38	35.7	64.8

SYCL Backends: 1. OpenMP host device 2. OpenCL CPU driver for Intel Meteorlake 3. Vulkan on llvmpipe backend (mesa) 4. OpenCL GPU driver for Intel Meteorlake (Arc GPU) 5. Vulkan on Intel Meteorlake GPU (mesa) 6. Cuda backend on NVIDIA RTX 500 Ada 7. Vulkan backend on NVIDIA RTX 500 Ada (NVIDIA proprietary driver)

Configs: * ndbody - ./sycl-generic 128000 10, time is total time in seconds. * mandlebrot - ./sycl-generic 1000, time is average parallel time / average kernel execution time in ms. * ising - ./sycl-generic -x 5120 -y 5120 -w 10 -n 100, time is elapsed time in seconds. * dslash - ./sycl-generic 256, time is total execution time in seconds. * sph - ./sycl-generic sph, time is average execution time of sph kernels in ms.

PCUDA

Supporting PCUDA is not a primary goal of this backend as it can only support pcudaMalloc, which is equivalent to device USM, and not pcudaHostManaged or pcudaMallocHost which are equivalent to shared/host USM. However in CUDA this functionality is not optional and many tests & applications make use of the managed/host allocation APIs.

Design

Runtime

USM

The use of the VK_KHR_buffer_device_address extension is a key part of the implementation as it allows us to implement device USM. By using vkGetBufferDeviceAddress we can get a device addressable to give to the user that they can use in a kernel. However, host/shared USM isn't supportable as this pointer isn't valid on host, the underlying Vulkan buffer needs memory mapped to a make that possible. This mapping is an implementation detail that the SYCL runtime performs when it needs to access the data at a device USM pointer on host, for example when implementing sycl::buffer.

In order to implement memcpy using the asynchronous command-buffer API, we need to use the VK buffer objects tied to those device addressable USM pointers. However, if the source or destination operand is a host pointer then there won't be an associated VK buffer. We therefore need to create temporary VK buffers for these operands which are tied to the lifetime of the async command. If a source operand is a host pointer, then we do a memcopy from the host pointer into the temporary VK buffer before submitting the async operation. When the destination operand is a host pointer we submit an async host thread that executes after the command-buffer has completed which copies the temporary VK buffer back to the host pointer.

In future VK_KHR_device_address_commands could be used to implement memcpy.

Compute Pipeline

A compute pipeline is created for each kernel JIT compiled for a specific workgroup size. When we JIT compiling kernels with adaptivity level 1 and higher, each translation unit (TU) that the clspv compiler sees only contains a single kernel in it. Whereas on adaptivity level 0 the TU contains mulitple kernels. The clspv tool will create a single descriptor set for all the kernels in the TU, with each kernel only caring about arguments at it's specific offsets in the buffer bound to that descriptor set.

When a kernel is re-submitted in the SYCL application the existing vk_executable_object is found in the cache with the JIT compiled SPIR-V kernel code. This will already have a vk_kernel_object attached representing the kernel function in that binary. If the workgroup size of the kernel enqueue is the same as previously then the old cached vk_kernel_pipeline instance will be used with the Vulkan compute pipeline created for that specific workgroup size. Otherwise, if the workgroup size is different, a new vk_kernel_pipeline object will be created to instantiate a Vulkan compute pipeline for that workgroup size. At adaptivity level 1 (the default) and above the AdaptiveCpp SSCP compiler creates a new executable for each work-group size with the information embedded in the required work-group size metadata. In this case a new vk_executable_object is created for each work-group size variant of a kernel. At adaptivity level 0, a specialization constant is used to set the work-group size on the same vk_executable_object object for each work-group size variant.

For small numbers of kernel arguments, no buffer descriptors may be required and all arguments can be set pushing push constants before the invocation of the Vulkan command-buffer containing the kernel. For larger numbers of kernel arguments or struct arguments that are decomposed, a single uniform buffer is used with binding 0, and each argument is an offset into that uniform buffer.

Queue

Vulkan devices are conceptually split into physical devices and logical devices, where a logical devices is created by the user based on the available physical devices and their capabilities. During the creation of the logical device the SYCL runtime finds the first Vulkan queue family which supports compute, and creates a single single queue to tie to the logical device. All sycl::queue objects created map to that single queue, but each have control over their own command-buffers to use with that queue.

Each queue submission is in it's own command-buffer, recorded with a single-submit bit. And incrementing a timeline semaphore owned by the queue such that the submission depends on the previous submission and the signal value can be used to identify the command. This combination of a signal value and queue uniquely identifying a command leads them to the implementation basis of the event class.

Host Threads

Not all SYCL commands are implementation in command-buffer commands, e.g. memset needs to work with less than 4-bytes. In which case asynchronous host threads are used to perform the work, waiting for and signalling the queue semaphore from host.

Host threads are also used to perform work that needs to be done after the device command has completed, but before we tell other SYCL commands the DAG node has completed. For example, freeing temporary allocations created to implement memcopy.

Compiler

clspv

Kernels are compiled by using the LLVM-IR input to clspv to create Vulkan consumable SPIR-V version 1.5 (supported by Vulkan 1.2 drivers onward). This involves the SSCP compiler flavouring the LLVM-IR to resemble OpenCL-C generated IR, which is then lowered appropriately by clspv. Transformations are defined in the llvm-to-clspv tool to achieve this. For more information on the transformations see the llvm-to-clspv doc

The SPIR-V reflection non-semantic instructions are then used to provide the Vulkan runtime with the appropriate information on how to set the kernel arguments. For example as push constants, or via a uniform buffer.

See clspv doc on OpenCL-C restrictions for inherited restrictions, such as the assumption of no aliasing.

Physical Addressing

A key part of the SPIR-V generated is that we can request clspv uses the Physical64 memory model, as opposed to Logical addressing, which allows for the use of device USM allocated pointers.

Specialization constants

A specialization constant will be created by clspv if the required workgroup size attribute was not set in the kernel IR, which is part of the reflection information that the runtime can detect and set.

libKernel

To implement the kernel functions declared by libkernel many of them are mapped to OpenCL-C builtin functions, which clspv then knows how to implement correctly for SPIR-V. However, clspv only supports OpenCL-C version 1.2 and some builtins (e.g. work-group collectives) require builtins from OpenCL 3.0. In such cases a more generic implementation is fallen back to.

The existing spirv libkernel library is not used because it is implemented in terms of the SPIRV-LLVM Translator IR expectations which are not recognized by other tools like clspv. clspv however can represent spirv builtins in IR by using a function call to a mangled name prefixed with convention _Z8spirv.op.<SPIRV OpCode>.<Mangled parameters>.