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