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

[Alex Voicu via llvm-commits]

================
@@ -9102,6 +9102,15 @@ bool InitializationSequence::Diagnose(Sema &S,
 
   case FK_ConversionFailed: {
     QualType FromType = OnlyArg->getType();
+    // __amdgpu_feature_predicate_t can be explicitly cast to the logical op
+    // type, although this is almost always an error and we advise against it
----------------
AlexVlx wrote:

Not really, although cleanliness, just like beauty, is in the eye of the beholder:) The feature builtins work locally, and naturally controls what might be a single ASM block, or builtin invocation, or mmio access etc. Going with target_clones has a number of implementation problems (it'd require a bunch of other things that are neither there nor impending for AMDGPU / SPIRV to materialise first), it forces outlining / adopting a different software development paradigm, and one that is fairly spammy (do I duplicate the full function (10s / 100s of lines, possibly) just because I might have some target specific behaviour? Do I start sharding out things that are target specific essentially creating a parallel world of `target_clones` based builtins, even though this code might be mature and gnarly etc.). The feature builtins allow a linear, inline transform from:

```cpp
#if defined(__gfx900__)
    __builtin_amdgcn_only_on_gfx900();
#endif
```

to

```cpp
if (__builtin_amdgcn_processor_is("gfx900"))
    __builtin_amdgcn_only_on_gfx900();
```
which is easy to teach / adopt and matches what folks are already doing, and just works as is, both on concrete, compile-time finalised targets, and on abstract run-time finalised ones. Also please consider that we are focusing on the `processor_is` case, however the preferable, more adaptable solution would be to check for the availability of a builtin, rather than an architecture, and that'd be even more spammy (and would make even more sense inline). 

Furthermore, is the current formulation particularly novel? It's simply a builtin that returns an irregular type with an explicit conversion operator for `bool`. Holding it right in this case matches what people already do; none of the exciting indirections we are discussing here would be at all possible today. Finally, we already have e.g. [__builtin_available](https://clang.llvm.org/docs/LanguageExtensions.html#objective-c-available) or `__nvvm_reflect`, which tackle similar problems via similar means, and can be misused in similar ways.

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