[clang] [llvm] [AMDGPU] Add `wave_id` and `wave_shuffle` Clang builtins. (PR #179492)

[Alex Voicu via cfe-commits]

================
@@ -229,3 +229,10 @@ void test_atomic_dec64() {
   __INT64_TYPE__ signedVal = 15;
   signedVal = __builtin_amdgcn_atomic_dec64(&signedVal, signedVal, __ATOMIC_ACQUIRE, ""); // expected-warning {{passing '__private long *' to parameter of type 'volatile __private unsigned long *' converts between pointers to integer types with different sign}}
 }
+
+void test_wave_shuffle(double d, int i, long long lli) {
+  struct S { int x; } s;
+  int x = __builtin_amdgcn_wave_shuffle(lli, i); // expected-error {{'lli' argument must be a scalar 'int' or 16 or 32 bit floating-point type (was '__private long long')}}
+  int y = __builtin_amdgcn_wave_shuffle(i, lli); // expected-error {{'lli' argument must be a scalar 'int' type (was '__private long long')}}
+  float z = __builtin_amdgcn_wave_shuffle(s, i); // expected-error {{'s' argument must be a scalar 'int' or 16 or 32 bit floating-point type (was '__private struct S')}}
+}
----------------
AlexVlx wrote:

I'm not entirely certain we should reject narrower types, as the upcast is information preserving, so it should be safe - it appears potentially more helpful to work in handling for them (when emitting the builtin, just do the implicit casting). Vectors (<4 x i8>, <2 x i16> wouldn't work as is without an explicit cast, but we could handle that here. I believe @b-sumner expressed a desire to add wider type support - to me it seems as if we should / could match all the types that wave_reduce supports (but this'd be separate work on the LLVM side, and not part of this change).

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