18e6958903
This change adds two builtins for AMDGPU: - `__builtin_amdgcn_processor_is`, which is similar in observable behaviour with `__builtin_cpu_is`, except that it is never "evaluated" at run time; - `__builtin_amdgcn_is_invocable`, which is behaviourally similar with `__has_builtin`, except that it is not a macro (i.e. not evaluated at preprocessing time). Neither of these are `constexpr`, even though when compiling for concrete (i.e. `gfxXXX` / `gfxXXX-generic`) targets they get evaluated in Clang, so they shouldn't tear the AST too badly / at all for multi-pass compilation cases like HIP. They can only be used in specific contexts (as args to control structures). The motivation for adding these is two-fold: - as a nice to have, it provides an AST-visible way to incorporate architecture specific code, rather than having to rely on macros and the preprocessor, which burn in the choice quite early; - as a must have, it allows featureful AMDGCN flavoured SPIR-V to be produced, where target specific capability is guarded and chosen or discarded when finalising compilation for a concrete target; this is built atop the Speciali\ation Constant concept which is described in the SPIR-V specification under section [2.12 Specialization](https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#_specialization_2) I've tried to keep the overall footprint of the change small. The changes to Sema are a bit unpleasant, but there was a strong desire to have Clang validate these, and to constrain their uses, and this was the most compact solution I could come up with (suggestions welcome). --------- Co-authored-by: Juan Manuel Martinez Caamaño <jmartinezcaamao@gmail.com> Co-authored-by: Voicu <avoicu@amd.com>
49 lines
4.1 KiB
C++
49 lines
4.1 KiB
C++
// 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");
|
|
}
|