Clang Project

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