[clang] d9e53b2 - clang: Move __builtin_amdgcn_processor_is diagnostic test to sema (#205734)

[via cfe-commits]

Author: Matt Arsenault
Date: 2026-06-25T14:15:28+02:00
New Revision: d9e53b24867e30a9ed513c1253f34f0e8222a8a7

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

LOG: clang: Move __builtin_amdgcn_processor_is diagnostic test to sema (#205734)

This wasn't checking the codegen result, so move it to the right place
and use -verify instead of FileChecking stderr.

Co-authored-by: Claude (Opus 4.8) <noreply at anthropic.com>

Added: 
    clang/test/SemaCXX/amdgpu-feature-builtins-invalid-use.cpp

Modified: 
    

Removed: 
    clang/test/CodeGen/amdgpu-feature-builtins-invalid-use.cpp


################################################################################
diff  --git a/clang/test/CodeGen/amdgpu-feature-builtins-invalid-use.cpp b/clang/test/CodeGen/amdgpu-feature-builtins-invalid-use.cpp
deleted file mode 100644
index 78f18d3a37b46..0000000000000
--- a/clang/test/CodeGen/amdgpu-feature-builtins-invalid-use.cpp
+++ /dev/null
@@ -1,48 +0,0 @@
-// RUN: not %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx900 -emit-llvm %s -o - 2>&1 | FileCheck %s
-// RUN: not %clang_cc1 -triple spirv64-amd-amdhsa -emit-llvm %s -o - 2>&1 | FileCheck %s
-
-bool predicate(bool x);
-void pass_by_value(__amdgpu_feature_predicate_t x);
-
-void invalid_uses(int *p, int x, const __amdgpu_feature_predicate_t &lv,
-                  __amdgpu_feature_predicate_t &&rv) {
-    // CHECK: error: 'a' has type __amdgpu_feature_predicate_t, which is not constructible
-    __amdgpu_feature_predicate_t a;
-    // CHECK: error: 'b' has type __amdgpu_feature_predicate_t, which is not constructible
-    __amdgpu_feature_predicate_t b = __builtin_amdgcn_processor_is("gfx906");
-    // CHECK: error: 'c' has type __amdgpu_feature_predicate_t, which is not constructible
-    __amdgpu_feature_predicate_t c = lv;
-    // CHECK: error: 'd' has type __amdgpu_feature_predicate_t, which is not constructible
-    __amdgpu_feature_predicate_t d = rv;
-    // CHECK: error: '__builtin_amdgcn_processor_is("gfx906")' must be explicitly cast to 'bool'; however, please note that this is almost always an error and that it prevents the effective guarding of target dependent code, and thus should be avoided
-    bool invalid_use_in_init_0 = __builtin_amdgcn_processor_is("gfx906");
-    // CHECK: error: 'x' has type __amdgpu_feature_predicate_t, which is not constructible
-    pass_by_value(__builtin_amdgcn_processor_is("gfx906"));
-    // CHECK: error: '__builtin_amdgcn_is_invocable(__builtin_amdgcn_s_sleep_var)' must be explicitly cast to 'bool'; however, please note that this is almost always an error and that it prevents the effective guarding of target dependent code, and thus should be avoided
-    bool invalid_use_in_init_1 = __builtin_amdgcn_is_invocable(__builtin_amdgcn_s_sleep_var);
-    // CHECK: error: '__builtin_amdgcn_processor_is("gfx906")' must be explicitly cast to 'bool'; however, please note that this is almost always an error and that it prevents the effective guarding of target dependent code, and thus should be avoided
-    if (bool invalid_use_in_init_2 = __builtin_amdgcn_processor_is("gfx906")) return;
-    // CHECK: error: '__builtin_amdgcn_processor_is("gfx1200")' must be explicitly cast to 'bool'; however, please note that this is almost always an error and that it prevents the effective guarding of target dependent code, and thus should be avoided
-    if (predicate(__builtin_amdgcn_processor_is("gfx1200"))) __builtin_amdgcn_s_sleep_var(x);
-}
-
-void invalid_invocations(int x, const char* str) {
-    // CHECK: error: the argument to __builtin_amdgcn_processor_is must be a valid AMDGCN processor identifier; 'not_an_amdgcn_gfx_id' is not valid
-    // CHECK-DAG: note: valid AMDGCN processor identifiers are: {{.*}}gfx{{.*}}
-    if (__builtin_amdgcn_processor_is("not_an_amdgcn_gfx_id")) return;
-    // CHECK: error: the argument to __builtin_amdgcn_processor_is must be a string literal
-    if (__builtin_amdgcn_processor_is(str)) return;
-    // CHECK: error: the argument to __builtin_amdgcn_is_invocable must be either a target agnostic builtin or an AMDGCN target specific builtin; {{.*}}__builtin_amdgcn_s_sleep_var{{.*}} is not valid
-    if (__builtin_amdgcn_is_invocable("__builtin_amdgcn_s_sleep_var")) return;
-    // CHECK: error: the argument to __builtin_amdgcn_is_invocable must be either a target agnostic builtin or an AMDGCN target specific builtin; {{.*}}str{{.*}} is not valid
-    else if (__builtin_amdgcn_is_invocable(str)) return;
-    // CHECK: error: the argument to __builtin_amdgcn_is_invocable must be either a target agnostic builtin or an AMDGCN target specific builtin; {{.*}}x{{.*}} is not valid
-    else if (__builtin_amdgcn_is_invocable(x)) return;
-    // CHECK: error: use of undeclared identifier '__builtin_ia32_pause'
-    else if (__builtin_amdgcn_is_invocable(__builtin_ia32_pause)) return;
-}
-
-bool return_needs_cast() {
-    // CHECK: error: '__builtin_amdgcn_processor_is("gfx900")' must be explicitly cast to 'bool'; however, please note that this is almost always an error and that it prevents the effective guarding of target dependent code, and thus should be avoided
-    return __builtin_amdgcn_processor_is("gfx900");
-}

diff  --git a/clang/test/SemaCXX/amdgpu-feature-builtins-invalid-use.cpp b/clang/test/SemaCXX/amdgpu-feature-builtins-invalid-use.cpp
new file mode 100644
index 0000000000000..e4036ed839cab
--- /dev/null
+++ b/clang/test/SemaCXX/amdgpu-feature-builtins-invalid-use.cpp
@@ -0,0 +1,50 @@
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx900 -fsyntax-only -verify %s
+// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -fsyntax-only -verify=expected,spirv %s
+
+bool predicate(bool x);
+void pass_by_value(__amdgpu_feature_predicate_t x);
+
+void invalid_uses(int *p, int x, const __amdgpu_feature_predicate_t &lv,
+                  __amdgpu_feature_predicate_t &&rv) {
+    __amdgpu_feature_predicate_t a;
+    // expected-error at -1 {{'a' has type __amdgpu_feature_predicate_t, which is not constructible}}
+    __amdgpu_feature_predicate_t b = __builtin_amdgcn_processor_is("gfx906");
+    // expected-error at -1 {{'b' has type __amdgpu_feature_predicate_t, which is not constructible}}
+    __amdgpu_feature_predicate_t c = lv;
+    // expected-error at -1 {{'c' has type __amdgpu_feature_predicate_t, which is not constructible}}
+    __amdgpu_feature_predicate_t d = rv;
+    // expected-error at -1 {{'d' has type __amdgpu_feature_predicate_t, which is not constructible}}
+    bool invalid_use_in_init_0 = __builtin_amdgcn_processor_is("gfx906");
+    // expected-error at -1 {{'__builtin_amdgcn_processor_is("gfx906")' must be explicitly cast to 'bool'; however, please note that this is almost always an error and that it prevents the effective guarding of target dependent code, and thus should be avoided}}
+    pass_by_value(__builtin_amdgcn_processor_is("gfx906"));
+    // expected-error at -1 {{'x' has type __amdgpu_feature_predicate_t, which is not constructible}}
+    bool invalid_use_in_init_1 = __builtin_amdgcn_is_invocable(__builtin_amdgcn_s_sleep_var);
+    // expected-error at -1 {{'__builtin_amdgcn_is_invocable(__builtin_amdgcn_s_sleep_var)' must be explicitly cast to 'bool'; however, please note that this is almost always an error and that it prevents the effective guarding of target dependent code, and thus should be avoided}}
+    if (bool invalid_use_in_init_2 = __builtin_amdgcn_processor_is("gfx906")) return;
+    // expected-error at -1 {{'__builtin_amdgcn_processor_is("gfx906")' must be explicitly cast to 'bool'; however, please note that this is almost always an error and that it prevents the effective guarding of target dependent code, and thus should be avoided}}
+    if (predicate(__builtin_amdgcn_processor_is("gfx1200"))) __builtin_amdgcn_s_sleep_var(x);
+    // expected-error at -1 {{no matching function for call to 'predicate'}}
+    // expected-error at -2 {{'__builtin_amdgcn_processor_is("gfx1200")' must be explicitly cast to 'bool'; however, please note that this is almost always an error and that it prevents the effective guarding of target dependent code, and thus should be avoided}}
+    // spirv-error at -3 {{'__builtin_amdgcn_s_sleep_var' cannot be invoked in the current context, as it requires the 'gfx12-insts' feature(s)}}
+}
+
+void invalid_invocations(int x, const char* str) {
+    if (__builtin_amdgcn_processor_is("not_an_amdgcn_gfx_id")) return;
+    // expected-error at -1 {{the argument to __builtin_amdgcn_processor_is must be a valid AMDGCN processor identifier; 'not_an_amdgcn_gfx_id' is not valid}}
+    // expected-note-re at -2 {{valid AMDGCN processor identifiers are: {{.*gfx.*}}}}
+    if (__builtin_amdgcn_processor_is(str)) return;
+    // expected-error at -1 {{the argument to __builtin_amdgcn_processor_is must be a string literal}}
+    if (__builtin_amdgcn_is_invocable("__builtin_amdgcn_s_sleep_var")) return;
+    // expected-error-re at -1 {{the argument to __builtin_amdgcn_is_invocable must be either a target agnostic builtin or an AMDGCN target specific builtin; {{.*__builtin_amdgcn_s_sleep_var.*}} is not valid}}
+    else if (__builtin_amdgcn_is_invocable(str)) return;
+    // expected-error-re at -1 {{the argument to __builtin_amdgcn_is_invocable must be either a target agnostic builtin or an AMDGCN target specific builtin; {{.*str.*}} is not valid}}
+    else if (__builtin_amdgcn_is_invocable(x)) return;
+    // expected-error-re at -1 {{the argument to __builtin_amdgcn_is_invocable must be either a target agnostic builtin or an AMDGCN target specific builtin; {{.*x.*}} is not valid}}
+    else if (__builtin_amdgcn_is_invocable(__builtin_ia32_pause)) return;
+    // expected-error at -1 {{use of undeclared identifier '__builtin_ia32_pause'}}
+}
+
+bool return_needs_cast() {
+    return __builtin_amdgcn_processor_is("gfx900");
+    // expected-error at -1 {{'__builtin_amdgcn_processor_is("gfx900")' must be explicitly cast to 'bool'; however, please note that this is almost always an error and that it prevents the effective guarding of target dependent code, and thus should be avoided}}
+}