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