[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