[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.

Author: jhuber6

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")

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);
 }