[Clang][SPIR-V] Provide builtin function support for SPIR-V instructions required for SYCL
- Dominant language
- LLVM
- Stars
- 40.5k
- Forks
- 18.7k
- PR merge metrics
- PR metrics pending
Description
Many SPIR-V instructions that will be needed by the SYCL runtime library are not currently exposed as Clang builtin functions. This issue tracks identifying which ones are needed, providing LLVM intrinsic support where it is not currently available (likely via delegation to other issues), and exposing Clang builtin functions when LLVM intrinsics are (or become) available.
LLVM defines intrinsics for some SPIR-V instructions in [`llvm/include/llvm/IR/IntrinsicsSPIRV.td`](https://github.com/llvm/llvm-project/blob/387c965bb013d14dbcee52098e714c27d6f98249/llvm/include/llvm/IR/IntrinsicsSPIRV.td). For example, that file [defines `int_spv_workgroup_size`](https://github.com/llvm/llvm-project/blob/387c965bb013d14dbcee52098e714c27d6f98249/llvm/include/llvm/IR/IntrinsicsSPIRV.td#L85) which corresponds to the `WorkgroupSize` builtin from [section 3.2.20, "BuiltIn" of the SPIR-V specification](https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#BuiltIn). That intrinsic gets associated with the `llvm.spv.workgroup.size` name through `llvm-tblgen` magic which makes the intrinsic available for use in LLVM IR `call` instructions. Lowering to SPIR-V is accomplished by [mapping that intrinsic, in `SPIRVInstructionSelector::selectIntrinsic()`](https://github.com/llvm/llvm-project/blob/387c965bb013d14dbcee52098e714c27d6f98249/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp#L5523-L5525), to a [builtin operand, `WorkgroupSize`, defined in `llvm/lib/Target/SPIRV/SPIRVSymbolicOperands.td`](https://github.com/llvm/llvm-project/blob/387c965bb013d14dbcee52098e714c27d6f98249/llvm/lib/Target/SPIRV/SPIRVSymbolicOperands.td#L1491) for use in `OpDecorate %N BuiltIn WorkgroupSize` instructions and subsequent `OpVariable` instructions.
A few of the LLVM intrinsics are exposed as Clang builtins via [`clang/include/clang/Basic/BuiltinsSPIRVCommon.td`](https://github.com/llvm/llvm-project/blob/387c965bb013d14dbcee52098e714c27d6f98249/clang/include/clang/Basic/BuiltinsSPIRVCommon.td) and [`clang/include/clang/Basic/BuiltinsSPIRVCL.td`](https://github.com/llvm/llvm-project/blob/387c965bb013d14dbcee52098e714c27d6f98249/clang/include/clang/Basic/BuiltinsSPIRVCL.td). Those builtins are mapped to the corresponding LLVM intrinsics in [`CodeGenFunction::EmitSPIRVBuiltinExpr()`](https://github.com/llvm/llvm-project/blob/387c965bb013d14dbcee52098e714c27d6f98249/clang/lib/CodeGen/TargetBuiltins/SPIR.cpp#L22-L177). For example, [`workgroup_size`](https://github.com/llvm/llvm-project/blob/387c965bb013d14dbcee52098e714c27d6f98249/clang/include/clang/Basic/BuiltinsSPIRVCommon.td#L12), which obtains a `__builtin_spirv_` prefix via [`SPIRVBuiltin`](https://github.com/llvm/llvm-project/blob/387c965bb013d14dbcee52098e714c27d6f98249/clang/include/clang/Basic/BuiltinsSPIRVBase.td#L11-L15), is mapped to the `llvm.spv.workgroup.size` intrinsic [here](https://github.com/llvm/llvm-project/blob/387c965bb013d14dbcee52098e714c27d6f98249/clang/lib/CodeGen/TargetBuiltins/SPIR.cpp#L126-L131).
At present, [libsycl](https://github.com/llvm/llvm-project/tree/387c965bb013d14dbcee52098e714c27d6f98249/libsycl) is using a manual method to execute SPIR-V instructions by declaring corresponding functions in [`libsycl/include/sycl/__spirv/spirv_vars.hpp`](https://github.com/llvm/llvm-project/blob/387c965bb013d14dbcee52098e714c27d6f98249/libsycl/include/sycl/__spirv/spirv_vars.hpp). These declarations rely on the demangled name of the declared functions (including those of function template specializations) matching names with `__spirv_` prefixes known to the LLVM SPIR-V translator and native SPIR-V backend. See documentation [here](https://llvm.org/docs/SPIRVUsage.html#builtin-functions). For example, the `__spirv_BuiltInWorkgroupSize()` function can be declared and called to retrieve the workgroup size.
See https://godbolt.org/z/4Y54xhqP7 for an example of the `__builtin_spirv_workgroup_size()` and `__spirv_BuiltInWorkgroupSize()` builtin functions working as intended.
Newly added LLVM intrinsics should be evaluated for potential documentation at https://llvm.org/docs/SPIRVUsage.html#target-intrinsics.
The builtin form that uses an LLVM intrinsic is preferred for several reasons.
1. The programmer is not required to provide explicit (and possibly incorrect) declarations of functions corresponding to required SPIR-V instructions.
2. Compilation time is improved by not having to spend time parsing such declarations.
3. SPIR-V instruction requirements, such as parameters that are required to have constant argument values, can be enforced and specialized diagnostics provided for incorrect use.
4. Builtin functions can be suppressed, or diagnosed as unavailable, when the target is known to lack support for the corresponding SPIR-V instructions.
Contributor guide
Research direction
Start by comparing llvm/include/llvm/IR/IntrinsicsSPIRV.td with clang/include/clang/Basic/BuiltinsSPIRVCommon.td and BuiltinsSPIRVCL.td, then trace mappings in clang/lib/CodeGen/TargetBuiltins/SPIR.cpp and SPIRVInstructionSelector::selectIntrinsic(). Review libsycl/include/sycl/__spirv/spirv_vars.hpp and SPIRVUsage.html for gaps. Done means required SYCL instructions are identified and supported or explicitly delegated, with relevant documentation updated.
Written by the indexing model from the issue text.
Assessment
- Tech stack
- cpp
- Domain
- compilers
- Issue type
- Feature
- Difficulty
- 5/5
- Estimated time
- Over a week
- Activity status
- Active
- Clarity
- Needs clarification
- Newbie friendliness
- 35/100