1 // RUN: %clang_cc1 -fopenmp -x c -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
2 // 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
3 // 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
4 // 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
5 // REQUIRES: x86-registered-target
6 // REQUIRES: nvptx-registered-target
7
8 #ifndef DIAGS
9 // expected-no-diagnostics
10 #endif // DIAGS
11
12 #ifdef IMMEDIATE
13 #pragma omp declare target
14 #endif //IMMEDIATE
t1(int r)15 void t1(int r) {
16 #ifdef DIAGS
17 // expected-error@+4 {{invalid input constraint 'mx' in asm}}
18 #endif // DIAGS
19 __asm__("PR3908 %[lf] %[xx] %[li] %[r]"
20 : [ r ] "+r"(r)
21 : [ lf ] "mx"(0), [ li ] "mr"(0), [ xx ] "x"((double)(0)));
22 }
23
t2(signed char input)24 unsigned t2(signed char input) {
25 unsigned output;
26 #ifdef DIAGS
27 // expected-error@+3 {{invalid output constraint '=a' in asm}}
28 #endif // DIAGS
29 __asm__("xyz"
30 : "=a"(output)
31 : "0"(input));
32 return output;
33 }
34
t3(double x)35 double t3(double x) {
36 register long double result;
37 #ifdef DIAGS
38 // expected-error@+3 {{invalid output constraint '=t' in asm}}
39 #endif // DIAGS
40 __asm __volatile("frndint"
41 : "=t"(result)
42 : "0"(x));
43 return result;
44 }
45
t4(unsigned char a,unsigned char b)46 unsigned char t4(unsigned char a, unsigned char b) {
47 unsigned int la = a;
48 unsigned int lb = b;
49 unsigned int bigres;
50 unsigned char res;
51 #ifdef DIAGS
52 // expected-error@+3 {{invalid output constraint '=la' in asm}}
53 #endif // DIAGS
54 __asm__("0:\n1:\n"
55 : [ bigres ] "=la"(bigres)
56 : [ la ] "0"(la), [ lb ] "c"(lb)
57 : "edx", "cc");
58 res = bigres;
59 return res;
60 }
61
t5(void)62 void t5(void) {
63 #ifdef DIAGS
64 // expected-error@+6 {{unknown register name 'st' in asm}}
65 #endif // DIAGS
66 __asm__ __volatile__(
67 "finit"
68 :
69 :
70 : "st", "st(1)", "st(2)", "st(3)",
71 "st(4)", "st(5)", "st(6)", "st(7)",
72 "fpsr", "fpcr");
73 }
74
75 typedef long long __m256i __attribute__((__vector_size__(32)));
t6(__m256i * p)76 void t6(__m256i *p) {
77 #ifdef DIAGS
78 // expected-error@+3 {{unknown register name 'ymm0' in asm}}
79 #endif // DIAGS
80 __asm__ volatile("vmovaps %0, %%ymm0" ::"m"(*(__m256i *)p)
81 : "ymm0");
82 }
83 #ifdef IMMEDIATE
84 #pragma omp end declare target
85 #endif //IMMEDIATE
86
main()87 int main() {
88 #ifdef DELAYED
89 #pragma omp target
90 #endif // DELAYED
91 {
92 #ifdef DELAYED
93 // expected-note@+2 {{called by 'main'}}
94 #endif // DELAYED
95 t1(0);
96 #ifdef DELAYED
97 // expected-note@+2 {{called by 'main'}}
98 #endif // DELAYED
99 t2(0);
100 #ifdef DELAYED
101 // expected-note@+2 {{called by 'main'}}
102 #endif // DELAYED
103 t3(0);
104 #ifdef DELAYED
105 // expected-note@+2 {{called by 'main'}}
106 #endif // DELAYED
107 t4(0, 0);
108 #ifdef DELAYED
109 // expected-note@+2 {{called by 'main'}}
110 #endif // DELAYED
111 t5();
112 #ifdef DELAYED
113 // expected-note@+2 {{called by 'main'}}
114 #endif // DELAYED
115 t6(0);
116 }
117 return 0;
118 }
119