https://github.com/jhuber6 created https://github.com/llvm/llvm-project/pull/80183
Summary: Currently we cannot compile `__builtin_amdgcn_ballot_w64` on non-wave64 targets even though it is valid. This is relevant for making library code that can handle both without needing to check the wavefront size. This patch relaxes the semantic check for w64 so it can be used normally. >From 5153ba3b9371cef74ecc6ea4304c28b3faf83f2e Mon Sep 17 00:00:00 2001 From: Joseph Huber <hube...@outlook.com> Date: Wed, 31 Jan 2024 13:18:04 -0600 Subject: [PATCH] [AMDGPU] Allow w64 ballot to be used on w32 targets Summary: Currently we cannot compile `__builtin_amdgcn_ballot_w64` on non-wave64 targets even though it is valid. This is relevant for making library code that can handle both without needing to check the wavefront size. This patch relaxes the semantic check for w64 so it can be used normally. --- clang/include/clang/Basic/BuiltinsAMDGPU.def | 2 +- clang/test/SemaOpenCL/builtins-amdgcn-error-wave64.cl | 7 ++----- 2 files changed, 3 insertions(+), 6 deletions(-) diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index 74dfd1d214e84..a32238381312c 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -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") // Deprecated intrinsics in favor of __builtin_amdgn_ballot_{w32|w64} BUILTIN(__builtin_amdgcn_uicmp, "WUiUiUiIi", "nc") diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-wave64.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error-wave64.cl index 99e93acd9a213..99aa2d634b2a0 100644 --- a/clang/test/SemaOpenCL/builtins-amdgcn-error-wave64.cl +++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-wave64.cl @@ -4,13 +4,10 @@ // RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx1010 -target-feature -wavefrontsize64 -verify -S -o - %s // RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx1010 -verify -S -o - %s +// expected-no-diagnostics + typedef unsigned long ulong; void test_ballot_wave64(global ulong* out, int a, int b) { - *out = __builtin_amdgcn_ballot_w64(a == b); // expected-error {{'__builtin_amdgcn_ballot_w64' needs target feature wavefrontsize64}} -} - -__attribute__((target("wavefrontsize64"))) -void test_ballot_wave64_target_attr(global ulong* out, int a, int b) { *out = __builtin_amdgcn_ballot_w64(a == b); } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits