llvm / llvm/llvm-project

[Clang][SPIR-V] Provide builtin function support for SPIR-V instructions required for SYCL

Open
#223,203 2 comments 1 reaction 0 assignees View on GitHub
SPIR-V 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

Open the contributing 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

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.