How to compile OpenCL Kernels to SPIR-V using Clang

Question

I need compile OpenCL kernels in SPIR-V to use with Vulkan, I tried with Google CLSPV https://github.com/google/clspv, but the problem occur with vectorization, functions like vload8 doesn't work. So I need compile OpenCL kernels in SPIR-V using clang.

Friseur

Answer 1

I'm the project lead for Clspv. Jesse is right overall.

The lack of support for vectors of length 8 and 16 is deliberately out of scope for now.
That's because Vulkan itself does not support that.
We haven't added the support to mimic such support, and don't have plans to do so even in the medium term.
There is more info on an old closed issue:
https://github.com/google/clspv/issues/8

Peanut

why did it even allow 16 element vectors in the first place? Did any hardware support that? I'm imagining that was implemented for CPUs with wide SIMD vectors, rather than the GPU focus of Vulkan, was that the reason? I'm suprised that it was a core feature if that was the case, considering IIRC most of these kind of features in OpenCL were extensions. – Ibert

OpenCL and Vulkan have different design philosophies, I guess. OpenCL might emulate operations on 16-element vectors into operations on 4 4-element vectors. That choice yields a simpler programming model. But Vulkan's tries to avoid complexity in the driver, including step that converts the SPIR-V to native GPU code. So that kind of vector emulating code generation isn't included, so you get a simpler, smaller, faster driver. Also OpenCL 1.0 had length-16 vectors since the start. reference – Peanut

Some day Clspv might add the ability to generate the extra code to emulate large vectors with smaller vectors. But that's not a goal at this time. It's an open source project so you could contribute such support. :-) – Peanut

Answer 2

Clspv is the only toolchain I'm aware of that compiles OpenCL C to Vulkan-compatible SPIR-V. You'll need to file an issue against Clspv; attaching a kernel that fails to compile properly would help a lot.

Lumpish

Answer 3

https://github.com/KhronosGroup/SPIR/tree/spirv-1.1

You can follow this Khronos project.

clang -cc1 -emit-spirv -triple=spir-unknown-unknown -cl-std=c++ -I include kernel.cl -o kernel.spv #For OpenCL C++

clang -cc1 -emit-spirv -triple=spir-unknown-unknown -cl-std=CL2.0 -include opencl.h kernel.cl -o kernel.spv #For OpenCL C

Mellins

Clang doesn't have an -emit-spirv option. – Invite

Answer 4

Given

__kernel void add(__global float *a, __global float *result) {
    float8 f = vload8(0, a);
    vstore8(f, 0, result);
}

Run

$ clang -cc1 -emit-llvm-bc -triple spir64-unknown-unknown -cl-std=CL3.0 vadd.cl -finclude-default-header -o vadd.bc
$ llvm-spirv vadd.bc -o vadd.spv
$ spirv-dis vadd.spv

Clang can also generate SPIR-V directly:

$ clang -target spirv64 vadd.cl -Xclang -finclude-default-header -o vadd.spv

Invite