================
@@ -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
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
