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
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
|
// RUN: %clang_cc1 -fopenmp -x c -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
// RUN: %clang_cc1 -verify -fopenmp -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized
// RUN: %clang_cc1 -verify -DDIAGS -DIMMEDIATE -fopenmp -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized
// RUN: %clang_cc1 -verify -DDIAGS -DDELAYED -fopenmp -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized
// RUN: %clang_cc1 -fopenmp -x c -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -fopenmp-version=50 -emit-llvm-bc %s -o %t-x86-host.bc
// RUN: %clang_cc1 -verify=expected,omp5 -fopenmp -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -fopenmp-version=50 %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized
// RUN: %clang_cc1 -verify=expected,omp5 -DDIAGS -DOMP5 -DIMMEDIATE -fopenmp -fopenmp-version=50 -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized
// RUN: %clang_cc1 -verify=expected,omp5 -DDIAGS -DOMP5 -DDELAYED -fopenmp -fopenmp-version=50 -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized
// REQUIRES: x86-registered-target
// REQUIRES: nvptx-registered-target
#ifndef DIAGS
// expected-no-diagnostics
#endif // DIAGS
#ifdef OMP5
void bar(int r) {
#ifdef IMMEDIATE
// omp5-error@+4 {{invalid input constraint 'mx' in asm}}
#endif // IMMEDIATE
__asm__("PR3908 %[lf] %[xx] %[li] %[r]"
: [ r ] "+r"(r)
: [ lf ] "mx"(0), [ li ] "mr"(0), [ xx ] "x"((double)(0)));
}
#ifdef IMMEDIATE
#pragma omp declare target to(bar) device_type(nohost)
#else
#pragma omp declare target to(bar) device_type(host)
#endif // IMMEDIATE
#endif // OMP5
void foo(int r) {
#ifdef IMMEDIATE
// expected-error@+4 {{invalid input constraint 'mx' in asm}}
#endif // IMMEDIATE
__asm__("PR3908 %[lf] %[xx] %[li] %[r]"
: [ r ] "+r"(r)
: [ lf ] "mx"(0), [ li ] "mr"(0), [ xx ] "x"((double)(0)));
}
#ifdef IMMEDIATE
#pragma omp declare target to(foo)
#endif //IMMEDIATE
#ifdef IMMEDIATE
#pragma omp declare target
#endif //IMMEDIATE
void t1(int r) {
#ifdef DIAGS
// expected-error@+4 {{invalid input constraint 'mx' in asm}}
#endif // DIAGS
__asm__("PR3908 %[lf] %[xx] %[li] %[r]"
: [ r ] "+r"(r)
: [ lf ] "mx"(0), [ li ] "mr"(0), [ xx ] "x"((double)(0)));
}
unsigned t2(signed char input) {
unsigned output;
#ifdef DIAGS
// expected-error@+3 {{invalid output constraint '=a' in asm}}
#endif // DIAGS
__asm__("xyz"
: "=a"(output)
: "0"(input));
return output;
}
double t3(double x) {
register long double result;
#ifdef DIAGS
// expected-error@+3 {{invalid output constraint '=t' in asm}}
#endif // DIAGS
__asm __volatile("frndint"
: "=t"(result)
: "0"(x));
return result;
}
unsigned char t4(unsigned char a, unsigned char b) {
unsigned int la = a;
unsigned int lb = b;
unsigned int bigres;
unsigned char res;
#ifdef DIAGS
// expected-error@+3 {{invalid output constraint '=la' in asm}}
#endif // DIAGS
__asm__("0:\n1:\n"
: [ bigres ] "=la"(bigres)
: [ la ] "0"(la), [ lb ] "c"(lb)
: "edx", "cc");
res = bigres;
return res;
}
void t5(void) {
#ifdef DIAGS
// expected-error@+6 {{unknown register name 'st' in asm}}
#endif // DIAGS
__asm__ __volatile__(
"finit"
:
:
: "st", "st(1)", "st(2)", "st(3)",
"st(4)", "st(5)", "st(6)", "st(7)",
"fpsr", "fpcr");
}
typedef long long __m256i __attribute__((__vector_size__(32)));
void t6(__m256i *p) {
#ifdef DIAGS
// expected-error@+3 {{unknown register name 'ymm0' in asm}}
#endif // DIAGS
__asm__ volatile("vmovaps %0, %%ymm0" ::"m"(*(__m256i *)p)
: "ymm0");
}
#ifdef IMMEDIATE
#pragma omp end declare target
#endif //IMMEDIATE
int main() {
#ifdef DELAYED
#pragma omp target
#endif // DELAYED
{
#ifdef DELAYED
// expected-note@+2 {{called by 'main'}}
#endif // DELAYED
t1(0);
#ifdef DELAYED
// expected-note@+2 {{called by 'main'}}
#endif // DELAYED
t2(0);
#ifdef DELAYED
// expected-note@+2 {{called by 'main'}}
#endif // DELAYED
t3(0);
#ifdef DELAYED
// expected-note@+2 {{called by 'main'}}
#endif // DELAYED
t4(0, 0);
#ifdef DELAYED
// expected-note@+2 {{called by 'main'}}
#endif // DELAYED
t5();
#ifdef DELAYED
// expected-note@+2 {{called by 'main'}}
#endif // DELAYED
t6(0);
}
return 0;
}
|
