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

[Alex Voicu via llvm-commits]

AlexVlx wrote:

> So in short: what you're trying to prevent is "this was stored in a variable, then checked later when we are no longer on the device, thus the answer is different". 

Not quite, although that is definitely an interesting consideration. What I am trying to address here is the not invalid concern that if you allow these to compose in arbitrary ways, with arbitrary values, stash and retrieve them to the point of being unable to trace them back, people will do it. And then they will build a rather complex contraption that makes perfect sense for them, is extremely valid and useful, but also ends up either failing to fold (not very helpful error message, they'll probably be unhappy with the compiler) or, worse, folds into something that completely subverted the value of the predicate (what was meant to be false is now true), and an ASM sequence that melts their GPU or launches nuclear missiles gets through (definitely unhappy). I am slightly more ambivalent than some of my colleagues towards this, but I cannot outright discard the concern - hence the awkward design.
 
>Your solution doesn't actually DO that, and acts in a way that is inconsistent with the language. Your attempts here are >defeated by common patterns, including once where variables are declared or altered inside of a condition statement. So >any attempts here are pretty fraught with errors. Consider:
> 
> ```
> if (auto x = <builtin-call>) {
> // value is here
> }
> 

Indeed, with the minor nit that with the current PR that'd actually not work, but rather fail, since the innermost context for the BI's use would be the initialisation of X, and for initialisations it is just a void returning function; we only expand/promote in a boolean condition.

> // Or even:
> bool b = false;
> if (<builtin-call>) b = true;
> ```
> 

Completely agreed, with the slight objection that this is subversive. The goal is not to outsmart really clever users, that's not tractable, but merely to prevent enthusiasm driven misuse that could lead to extremely sub-optimal outcomes.

> So any attempt to do so is, honestly, partial at best, and confusing for no good reason. Values change in a program, and these are no different, so that sounds like a common programming mistake. Also, for some reason, this disallows the conditional-operator as well? Which is another confusing thing for users, as they consider that to be a shortcut for if/else assignments
>

The conditional-operator was considered "too confusing / why would you ever do that / what use could it be" material when I suggested it should work, hence it ended up as disallowed. Adding it back wouldn't be a problem.

> IMO, we are better served by a warning diagnostic if we detect these are misused. `ParseCXXCondition` (or the C equivalent, but since you are returning bool it seems you're not concerned about C?) might be a good place to set a variable to enable the warning.

This is a very good suggestion, thank you very much for it - it might well be where we end up. My worry is that ignoring warning and diagnostics is rather common. If I may be so bold as to inquire: would you and @AaronBallman be slightly less horrified if the return type variance would be replaced with returning an odd type that only knows how to `bool`ify itself in conditions? More explicitly, if instead of `void __builtin_amdgcn_processor_is(const char*)` what we see is `__amdgpu_predicate_t __builtin_amdgcn_processor_is(const char*)`, would that be somewhat less bad? There is precedent for special-ish builtins returning special-ish target types (please consider `__amdgpu_buffer_rsrc_t` for `__builtin_amdgcn_make_buffer_rsrc` or `svcount_t`)



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