[clang] clang: Move __builtin_amdgcn_processor_is diagnostic test to sema (PR #205734)
via cfe-commits
cfe-commits at lists.llvm.org
Thu Jun 25 05:10:50 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-amdgpu
Author: Matt Arsenault (arsenm)
<details>
<summary>Changes</summary>
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@<!-- -->anthropic.com>
---
Full diff: https://github.com/llvm/llvm-project/pull/205734.diff
4 Files Affected:
- (modified) clang/include/clang/Basic/DiagnosticSemaKinds.td (+1-1)
- (removed) clang/test/CodeGen/amdgpu-feature-builtins-invalid-use.cpp (-48)
- (added) clang/test/SemaCXX/amdgpu-feature-builtins-invalid-use.cpp (+50)
- (modified) clang/test/SemaHIP/amdgpu-feature-predicates-guard-use.hip (+2-2)
``````````diff
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index cde99dfb16ec5..e3aa1c7bfeb5b 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -14267,7 +14267,7 @@ def warn_amdgcn_unguarded_asm_stmt
InGroup<UnguardedBuiltinUsageAMDGPU>, DefaultIgnore;
def note_amdgcn_unguarded_asm_silence
: Note<"enclose the '%0' ASM sequence in a scope controlled by a "
- "__builtin_amdgcn_is_processor check to silence this warning">;
+ "__builtin_amdgcn_processor_is check to silence this warning">;
def err_amdgcn_incompatible_builtin
: Error<"%0 cannot be invoked in the current context, as it requires the "
"'%1' feature(s)%select{|, which '%3' does not provide}2">;
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}}
+}
diff --git a/clang/test/SemaHIP/amdgpu-feature-predicates-guard-use.hip b/clang/test/SemaHIP/amdgpu-feature-predicates-guard-use.hip
index 1566bc9aa0be5..345849374e6fd 100644
--- a/clang/test/SemaHIP/amdgpu-feature-predicates-guard-use.hip
+++ b/clang/test/SemaHIP/amdgpu-feature-predicates-guard-use.hip
@@ -10,7 +10,7 @@ __device__ void g();
__device__ void f(int x, bool b) {
long v15_16;
__asm volatile("v_lshlrev_b64 v[15:16], 0, %0" : "={v[15:16]}"(v15_16) : "v"(x)); // expected-warning {{the 'v_lshlrev_b64 v[15:16], 0, $0' ASM sequence might be invalid for some AMDGPU targets}}
- // expected-note at -1 {{enclose the 'v_lshlrev_b64 v[15:16], 0, $0' ASM sequence in a scope controlled by a __builtin_amdgcn_is_processor check to silence this warning}}
+ // expected-note at -1 {{enclose the 'v_lshlrev_b64 v[15:16], 0, $0' ASM sequence in a scope controlled by a __builtin_amdgcn_processor_is check to silence this warning}}
if (__builtin_amdgcn_processor_is("gfx90a")) {
long v15_16;
@@ -20,7 +20,7 @@ __device__ void f(int x, bool b) {
if (!__builtin_amdgcn_processor_is("gfx90a")) {
long v15_16;
__asm volatile("v_lshlrev_b64 v[15:16], 0, %0" : "={v[15:16]}"(v15_16) : "v"(x)); // expected-warning {{the 'v_lshlrev_b64 v[15:16], 0, $0' ASM sequence might be invalid for some AMDGPU targets}}
- // expected-note at -1 {{enclose the 'v_lshlrev_b64 v[15:16], 0, $0' ASM sequence in a scope controlled by a __builtin_amdgcn_is_processor check to silence this warning}}
+ // expected-note at -1 {{enclose the 'v_lshlrev_b64 v[15:16], 0, $0' ASM sequence in a scope controlled by a __builtin_amdgcn_processor_is check to silence this warning}}
}
__builtin_amdgcn_is_invocable(__builtin_amdgcn_s_sleep_var) ? __builtin_amdgcn_s_sleep_var(x) : __builtin_trap();
``````````
</details>
https://github.com/llvm/llvm-project/pull/205734
More information about the cfe-commits
mailing list