1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
|
// REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -std=c++20 -triple amdgcn -target-cpu tahiti -emit-llvm -fcuda-is-device -verify=no-memrealtime -o - %s
// RUN: %clang_cc1 -std=c++20 -triple amdgcn -target-cpu gfx950 -emit-llvm -fcuda-is-device -o - %s
#define __device__ __attribute__((device))
#define __shared__ __attribute__((shared))
struct S {
static constexpr auto memrealtime_lambda = []() {
__builtin_amdgcn_s_memrealtime(); // no-memrealtime-error{{'__builtin_amdgcn_s_memrealtime' needs target feature s-memrealtime}}
};
};
__attribute__((target("s-memrealtime")))
__device__ void test_target_dependant_builtin_attr_fail() {
S::memrealtime_lambda();
}
constexpr auto memrealtime_lambda = []() {
__builtin_amdgcn_s_memrealtime(); // no-memrealtime-error{{'__builtin_amdgcn_s_memrealtime' needs target feature s-memrealtime}}
};
__attribute__((target("s-memrealtime")))
__device__ void global_test_target_dependant_builtin_attr_fail() {
memrealtime_lambda();
}
__attribute__((target("s-memrealtime")))
__device__ void local_test_target_dependant_builtin_attr_fail() {
static constexpr auto f = []() {
__builtin_amdgcn_s_memrealtime(); // no-memrealtime-error{{'__builtin_amdgcn_s_memrealtime' needs target feature s-memrealtime}}
};
f();
}
|
