llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang 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@-1 {{'a' has type __amdgpu_feature_predicate_t, which is not constructible}} + __amdgpu_feature_predicate_t b = __builtin_amdgcn_processor_is("gfx906"); + // expected-error@-1 {{'b' has type __amdgpu_feature_predicate_t, which is not constructible}} + __amdgpu_feature_predicate_t c = lv; + // expected-error@-1 {{'c' has type __amdgpu_feature_predicate_t, which is not constructible}} + __amdgpu_feature_predicate_t d = rv; + // expected-error@-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@-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@-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@-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@-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@-1 {{no matching function for call to 'predicate'}} + // expected-error@-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@-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@-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@-2 {{valid AMDGCN processor identifiers are: {{.*gfx.*}}}} + if (__builtin_amdgcn_processor_is(str)) return; + // expected-error@-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@-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@-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@-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@-1 {{use of undeclared identifier '__builtin_ia32_pause'}} +} + +bool return_needs_cast() { + return __builtin_amdgcn_processor_is("gfx900"); + // expected-error@-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@-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@-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@-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@-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 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
