KhronosGroup / KhronosGroup/SPIRV-LLVM-Translator

SPIR-V instructions mapping is underdefined

Open
#365 0 comments 2 reactions 0 assignees View on GitHub
agenda
Dominant language
LLVM
Stars
625
Forks
279
Avg merge
3d 5h
Merged PRs (30d)
34

Description

The document describing the [SPIR-V representation](https://github.com/KhronosGroup/SPIRV-LLVM-Translator/blob/master/docs/SPIRVRepresentationInLLVM.rst) does not define how builtins should actually be expressed.

The statement: `its name is IA64 mangled` does not give any information as no full prototype is given. Also `SPIR-V builtin functions accepts all argument types accepted by the corresponding SPIR-V instructions` is ambiguous/problematic as the IA64 mangling is affected by the choices made to map those instruction operands.

ATM it seems that things work as the translator does not perform much checks: as long as a function more or less looks like what it is expected, the translator uses it as a builtin.

This causes 2 kind of problems:
- Overload of a builtin (but not a builtin) may be mistaken as the builtin. For instance, for the translator calls to `int get_global_id(int, int)` are equivalent to `int get_global_id(int)`;
- It is not possible to preserve semantic after a round trip. During the SPIR-V -> LLVM IR translation, builtin are mangled in a specific way which may not match the input. So performing a LLVM IR -> SPIR-V -> LLVM IR translation may not preserve semantic as we may end up calling different function (ABI breakage).

Unless I missed something in the document, is that the intention of the translator to roughly accept anything that more or less looks like a builtin ? Can we better define which overload are actual builtins ? IMO, the translator should only accepts as a builtin functions that can be rematerialized (i.e. preserves ABI).

Contributor guide

Open the contributing guide

Research direction

Start with docs/SPIRVRepresentationInLLVM.rst and review the sections describing SPIR-V builtins, IA64 mangling, and operand types. Define the builtin prototypes and overload rules clearly enough to distinguish valid builtins and preserve semantics across LLVM IR to SPIR-V to LLVM IR round trips.

Written by the indexing model from the issue text.

Assessment

Domain
compilers, documentation
Issue type
Documentation
Difficulty
5/5
Estimated time
Over a week
Activity status
Stale
Clarity
Mostly clear
Newbie friendliness
30/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.