[clang] [AMDGPU] Allow w64 ballot to be used on w32 targets (PR #80183)
Yaxun Liu via cfe-commits
cfe-commits at lists.llvm.org
Wed Jan 31 11:39:36 PST 2024
================
@@ -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")
----------------
yxsamliu wrote:
I am not sure whether this will cause the intrinsic being removed by backend if the required target feature is missing. I remember some pass is added for that.
Also removing the check increases chance of miss-use in HIP or OpenCL.
device libs bypass this requirement by adding `__attribute__((target("wavefrontsize64")))` to the callers. can you do the same?
https://github.com/llvm/llvm-project/pull/80183
More information about the cfe-commits
mailing list