AlexVlx wrote:

> So your users today are building for generic AMDGPU but using builtins that 
> are only available on a specific processor release? Presumably that code is 
> protected _somehow_ and their programs are not simply crashing at runtime. Is 
> that something you'd be able to leverage at all, or is completely ad hoc?

They do crash at run time, except not in the way one would expect - they fail 
when finalising / JIT-ing from SPIRV, which is still a compiler failure, except 
it's a BE / ISEL one. But yes, this is a current problem (which this is 
addressing). Here's an example (there are others):

- client code uses some builtins that are only available on RDNA (GFX10+), `#if 
__has_builtin(something_something)`;
- when targeting AMDGCN-flavoured SPIRV (`amdgcnspirv`), the union of builtins 
is available, since we don't know what the concrete target will end up being, 
and we want to maximally leverage features, so the check is true and the RDNA 
builtin ends up in SPIRV;
- the compiled library / executable gets executed on a CDNA machine;
- depending the nature of the intrinsic a JIT-time error ensues.

What we would like to do is to allow people to handle these cases with a linear 
translation from the above into `if 
(__builtin_amdgcn_is_invocable(something_something)`, which then would lead to 
having code that works everywhere with maximum capability (we don't have to 
reduce things to a common subset), without having to be linear in targets. I 
mention the latter because our device libs (which are not upstream), deal with 
this via a different convoluted mechanism, since there was no amdgcnspirv / 
generic target at the time, which requires generating separate bitcode per 
target, which is not long term viable as we get more and more targets.

https://github.com/llvm/llvm-project/pull/134016
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits

Reply via email to