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
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
| // REQUIRES: amdgpu-registered-target
// RUN: %clang_cc1 -triple amdgcn-- -target-cpu tahiti -verify -S -o - %s
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
typedef unsigned long ulong;
typedef unsigned int uint;
// To get all errors for feature checking we need to put them in one function
// since Clang will stop codegen for the next function if it finds error during
// codegen of the previous function.
void test_target_builtin(global int* out, int a)
{
__builtin_amdgcn_s_memrealtime(); // expected-error {{'__builtin_amdgcn_s_memrealtime' needs target feature s-memrealtime}}
*out = __builtin_amdgcn_mov_dpp(a, 0, 0, 0, false); // expected-error {{'__builtin_amdgcn_mov_dpp' needs target feature dpp}}
}
void test_s_sleep(int x)
{
__builtin_amdgcn_s_sleep(x); // expected-error {{argument to '__builtin_amdgcn_s_sleep' must be a constant integer}}
}
void test_s_waitcnt(int x)
{
__builtin_amdgcn_s_waitcnt(x); // expected-error {{argument to '__builtin_amdgcn_s_waitcnt' must be a constant integer}}
}
void test_s_sendmsg(int in)
{
__builtin_amdgcn_s_sendmsg(in, 1); // expected-error {{argument to '__builtin_amdgcn_s_sendmsg' must be a constant integer}}
}
void test_s_sendmsg_var(int in1, int in2)
{
__builtin_amdgcn_s_sendmsg(in1, in2); // expected-error {{argument to '__builtin_amdgcn_s_sendmsg' must be a constant integer}}
}
void test_s_sendmsghalt(int in)
{
__builtin_amdgcn_s_sendmsghalt(in, 1); // expected-error {{argument to '__builtin_amdgcn_s_sendmsghalt' must be a constant integer}}
}
void test_s_sendmsghalt_var(int in1, int in2)
{
__builtin_amdgcn_s_sendmsghalt(in1, in2); // expected-error {{argument to '__builtin_amdgcn_s_sendmsghalt' must be a constant integer}}
}
void test_s_incperflevel(int x)
{
__builtin_amdgcn_s_incperflevel(x); // expected-error {{argument to '__builtin_amdgcn_s_incperflevel' must be a constant integer}}
}
void test_s_decperflevel(int x)
{
__builtin_amdgcn_s_decperflevel(x); // expected-error {{argument to '__builtin_amdgcn_s_decperflevel' must be a constant integer}}
}
void test_sicmp_i32(global ulong* out, int a, int b, uint c)
{
*out = __builtin_amdgcn_sicmp(a, b, c); // expected-error {{argument to '__builtin_amdgcn_sicmp' must be a constant integer}}
}
void test_uicmp_i32(global ulong* out, uint a, uint b, uint c)
{
*out = __builtin_amdgcn_uicmp(a, b, c); // expected-error {{argument to '__builtin_amdgcn_uicmp' must be a constant integer}}
}
void test_sicmp_i64(global ulong* out, long a, long b, uint c)
{
*out = __builtin_amdgcn_sicmpl(a, b, c); // expected-error {{argument to '__builtin_amdgcn_sicmpl' must be a constant integer}}
}
void test_uicmp_i64(global ulong* out, ulong a, ulong b, uint c)
{
*out = __builtin_amdgcn_uicmpl(a, b, c); // expected-error {{argument to '__builtin_amdgcn_uicmpl' must be a constant integer}}
}
void test_fcmp_f32(global ulong* out, float a, float b, uint c)
{
*out = __builtin_amdgcn_fcmpf(a, b, c); // expected-error {{argument to '__builtin_amdgcn_fcmpf' must be a constant integer}}
}
void test_fcmp_f64(global ulong* out, double a, double b, uint c)
{
*out = __builtin_amdgcn_fcmp(a, b, c); // expected-error {{argument to '__builtin_amdgcn_fcmp' must be a constant integer}}
}
void test_ds_swizzle(global int* out, int a, int b)
{
*out = __builtin_amdgcn_ds_swizzle(a, b); // expected-error {{argument to '__builtin_amdgcn_ds_swizzle' must be a constant integer}}
}
void test_s_getreg(global int* out, int a)
{
*out = __builtin_amdgcn_s_getreg(a); // expected-error {{argument to '__builtin_amdgcn_s_getreg' must be a constant integer}}
}
void test_mov_dpp2(global int* out, int a, int b, int c, int d, bool e)
{
*out = __builtin_amdgcn_mov_dpp(a, b, 0, 0, false); // expected-error {{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
*out = __builtin_amdgcn_mov_dpp(a, 0, c, 0, false); // expected-error {{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
*out = __builtin_amdgcn_mov_dpp(a, 0, 0, d, false); // expected-error {{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
*out = __builtin_amdgcn_mov_dpp(a, 0, 0, 0, e); // expected-error {{argument to '__builtin_amdgcn_mov_dpp' must be a constant integer}}
}
void test_update_dpp2(global int* out, int a, int b, int c, int d, int e, bool f)
{
*out = __builtin_amdgcn_update_dpp(a, b, 0, 0, 0, false);
*out = __builtin_amdgcn_update_dpp(a, 0, c, 0, 0, false); // expected-error {{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
*out = __builtin_amdgcn_update_dpp(a, 0, 0, d, 0, false); // expected-error {{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
*out = __builtin_amdgcn_update_dpp(a, 0, 0, 0, e, false); // expected-error {{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
*out = __builtin_amdgcn_update_dpp(a, 0, 0, 0, 0, f); // expected-error {{argument to '__builtin_amdgcn_update_dpp' must be a constant integer}}
}
void test_ds_faddf(local float *out, float src, int a) {
*out = __builtin_amdgcn_ds_faddf(out, src, a, 0, false); // expected-error {{argument to '__builtin_amdgcn_ds_faddf' must be a constant integer}}
*out = __builtin_amdgcn_ds_faddf(out, src, 0, a, false); // expected-error {{argument to '__builtin_amdgcn_ds_faddf' must be a constant integer}}
*out = __builtin_amdgcn_ds_faddf(out, src, 0, 0, a); // expected-error {{argument to '__builtin_amdgcn_ds_faddf' must be a constant integer}}
}
void test_ds_fminf(local float *out, float src, int a) {
*out = __builtin_amdgcn_ds_fminf(out, src, a, 0, false); // expected-error {{argument to '__builtin_amdgcn_ds_fminf' must be a constant integer}}
*out = __builtin_amdgcn_ds_fminf(out, src, 0, a, false); // expected-error {{argument to '__builtin_amdgcn_ds_fminf' must be a constant integer}}
*out = __builtin_amdgcn_ds_fminf(out, src, 0, 0, a); // expected-error {{argument to '__builtin_amdgcn_ds_fminf' must be a constant integer}}
}
void test_ds_fmaxf(local float *out, float src, int a) {
*out = __builtin_amdgcn_ds_fmaxf(out, src, a, 0, false); // expected-error {{argument to '__builtin_amdgcn_ds_fmaxf' must be a constant integer}}
*out = __builtin_amdgcn_ds_fmaxf(out, src, 0, a, false); // expected-error {{argument to '__builtin_amdgcn_ds_fmaxf' must be a constant integer}}
*out = __builtin_amdgcn_ds_fmaxf(out, src, 0, 0, a); // expected-error {{argument to '__builtin_amdgcn_ds_fmaxf' must be a constant integer}}
}
|