[clang] [llvm] [AMDGPU][clang][CodeGen][opt] Add late-resolved feature identifying predicates (PR #134016)

[Alex Voicu via llvm-commits]

AlexVlx wrote:

> > This change adds two semi-magical builtins for AMDGPU:
> > ```
> > * `__builtin_amdgcn_processor_is`, which is similar in observable behaviour with `__builtin_cpu_is`, except that it is never "evaluated" at run time;
> > 
> > * `__builtin_amdgcn_is_invocable`, which is behaviourally similar with `__has_builtin`, except that it is not a macro (i.e. not evaluated at preprocessing time).
> > ```
> > 
> > 
> >     
> >       
> >     
> > 
> >       
> >     
> > 
> >     
> >   
> > Neither of these are `constexpr`, even though when compiling for concrete (i.e. `gfxXXX` / `gfxXXX-generic`) targets they get evaluated in Clang, so they shouldn't tear the AST too badly / at all for multi-pass compilation cases like HIP. They can only be used in specific contexts (as args to control structures).
> > The motivation for adding these is two-fold:
> > ```
> > * as a nice to have, it provides an AST-visible way to incorporate architecture specific code, rather than having to rely on macros and the preprocessor, which burn in the choice quite early;
> > 
> > * as a must have, it allows featureful AMDGCN flavoured SPIR-V to be produced, where target specific capability is guarded and chosen or discarded when finalising compilation for a concrete target.
> > ```
> > 
> > 
> >     
> >       
> >     
> > 
> >       
> >     
> > 
> >     
> >   
> > I've tried to keep the overall footprint of the change small. The changes to Sema are a bit unpleasant, but there was a strong desire to have Clang validate these, and to constrain their uses, and this was the most compact solution I could come up with (suggestions welcome).
> > In the end, I will note there is nothing that is actually AMDGPU specific here, so it is possible that in the future, assuming interests from other targets / users, we'd just promote them to generic intrinsics.
> 
> First read through this, I find myself wondering WHY these aren't constexpr. They seem exactly the sort of thing that folks would like to use `if constexpr` for.

There are a few reasons, primarily though:

1. at least for builtin checking `__has_builtin` already exists, and that can get lumped into `if constexpr` and pretty much quacks like the same duck, so it felt superfluous (this is not a particularly strong reason though);
2. for an abstract target (in this case AMDGCNSPIRV) it cannob be constexpr/consteval, because when you're doing the initial compilation to SPIR-V you don't actually know which target you'll get eventually finalised for (this is one of the primary motivations for these existing, allowing on to generate "adaptable" AMDGCN SPIR-V that can tightly clamp to target features at finalisation time);
3. if they would be sometimes constexpr i.e. when compiling for a concrete target, and sometimes not constexpr, i.e. when compiling for an abstract one, the user would eventually have to end up doing something like an ifdef guard around their if clauses (to figure out whether or not it is if constexpr), or start playing games defining `MAYBE_CONSTEXPR` macros etc., which would be quite unfortunate.

Concerns around early AST tearing resulting from (3) in multi-pass compilation cases like HIP were also considerable. 

https://github.com/llvm/llvm-project/pull/134016