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

[via cfe-commits]

Author: Joseph Huber
Date: 2024-02-05T08:42:28-06:00
New Revision: 5249379d742148728f654665e113084c6b93cdf2

URL: https://github.com/llvm/llvm-project/commit/5249379d742148728f654665e113084c6b93cdf2
DIFF: https://github.com/llvm/llvm-project/commit/5249379d742148728f654665e113084c6b93cdf2.diff

LOG: [AMDGPU] Allow w64 ballot to be used on w32 targets (#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.

Added: 
    

Modified: 
    clang/include/clang/Basic/BuiltinsAMDGPU.def
    clang/test/SemaOpenCL/builtins-amdgcn-error-wave64.cl

Removed: 
    


################################################################################
diff  --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def
index e9dd8dcd0b60e..5f8001e61a028 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..1bbd84d09bf94 100644
--- a/clang/test/SemaOpenCL/builtins-amdgcn-error-wave64.cl
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-wave64.cl
@@ -1,13 +1,13 @@
 // RUN: %clang_cc1 -triple amdgcn-- -verify -S -o - %s
-// RUN: %clang_cc1 -triple amdgcn-- -target-feature +wavefrontsize32 -verify -S -o - %s
-// RUN: %clang_cc1 -triple amdgcn-- -target-cpu gfx1010 -target-feature +wavefrontsize32 -verify -S -o - %s
 // 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}}
+  *out = __builtin_amdgcn_ballot_w64(a == b);
 }
 
 __attribute__((target("wavefrontsize64")))