[clang] [AMDGPU] Allow w64 ballot to be used on w32 targets (PR #80183)

[Joseph Huber via cfe-commits]

================
@@ -151,7 +151,7 @@ BUILTIN(__builtin_amdgcn_mqsad_u32_u8, "V4UiWUiUiV4Ui", "nc")
 //===----------------------------------------------------------------------===//
 
 TARGET_BUILTIN(__builtin_amdgcn_ballot_w32, "ZUib", "nc", "wavefrontsize32")
-TARGET_BUILTIN(__builtin_amdgcn_ballot_w64, "WUib", "nc", "wavefrontsize64")
+BUILTIN(__builtin_amdgcn_ballot_w64, "WUib", "nc")
----------------
jhuber6 wrote:

Just for reference, the main difference is whether or not we read `exec_hi`, here's some code https://godbolt.org/z/z64Yoo7ve. I believe @arsenm said that `exec_hi` is simply zero when read from wave32 mode? I could double check that.

The problem with using the attributes is that it puts the attribute into the function, which prevents it from being called in `wave32` correctly. This would require some compile-time switch like the ROCm libs, or perhaps a runtime check on the wave size.

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