1 | // Test target codegen - host bc file has to be created first. |
2 | // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc |
3 | // RUN: %clang_cc1 -debug-info-kind=limited -verify -fopenmp -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-64 |
4 | // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc |
5 | // RUN: %clang_cc1 -debug-info-kind=limited -verify -fopenmp -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix TCHECK --check-prefix TCHECK-32 |
6 | // expected-no-diagnostics |
7 | #ifndef HEADER |
8 | #define HEADER |
9 | |
10 | template <typename tx, typename ty> |
11 | struct TT { |
12 | tx X; |
13 | ty Y; |
14 | }; |
15 | |
16 | // TCHECK-DAG: [[TT:%.+]] = type { i64, i8 } |
17 | // TCHECK-DAG: [[TTII:%.+]] = type { i32, i32 } |
18 | // TCHECK-DAG: [[S1:%.+]] = type { double } |
19 | |
20 | // TCHECK: @__omp_offloading_firstprivate__{{.+}}_e_l27 = internal addrspace(4) global [[TTII]] zeroinitializer |
21 | int foo(int n, double *ptr) { |
22 | int a = 0; |
23 | short aa = 0; |
24 | float b[10]; |
25 | double c[5][10]; |
26 | TT<long long, char> d; |
27 | const TT<int, int> e = {n, n}; |
28 | |
29 | #pragma omp target firstprivate(a, e) map(tofrom \ |
30 | : b) |
31 | { |
32 | b[a] = a; |
33 | b[a] += e.X; |
34 | } |
35 | |
36 | // TCHECK: define {{.*}}void @__omp_offloading_{{.+}}([10 x float] addrspace(1)* noalias [[B_IN:%.+]], i{{[0-9]+}} [[A_IN:%.+]], [[TTII]]* noalias [[E_IN:%.+]]) |
37 | // TCHECK-NOT: alloca [[TTII]], |
38 | // TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}}, |
39 | // TCHECK-NOT: alloca [[TTII]], |
40 | // TCHECK-NOT: alloca i{{[0-9]+}}, |
41 | // TCHECK-64: call void @llvm.dbg.declare(metadata [10 x float] addrspace(1)** %{{.+}}, metadata !{{[0-9]+}}, metadata !DIExpression()) |
42 | // TCHECK: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]], |
43 | // TCHECK: ret void |
44 | |
45 | #pragma omp target firstprivate(aa, b, c, d) |
46 | { |
47 | aa += 1; |
48 | b[2] = 1.0; |
49 | c[1][2] = 1.0; |
50 | d.X = 1; |
51 | d.Y = 1; |
52 | } |
53 | |
54 | // make sure that firstprivate variables are generated in all cases and that we use those instances for operations inside the |
55 | // target region |
56 | // TCHECK: define {{.*}}void @__omp_offloading_{{.+}}(i{{[0-9]+}}{{.*}} [[A2_IN:%.+]], [10 x float]*{{.*}} [[B_IN:%.+]], [5 x [10 x double]]*{{.*}} [[C_IN:%.+]], [[TT]]*{{.*}} [[D_IN:%.+]]) |
57 | // TCHECK: [[A2_ADDR:%.+]] = alloca i{{[0-9]+}}, |
58 | // TCHECK: [[B_ADDR:%.+]] = alloca [10 x float]*, |
59 | // TCHECK: [[C_ADDR:%.+]] = alloca [5 x [10 x double]]*, |
60 | // TCHECK: [[D_ADDR:%.+]] = alloca [[TT]]*, |
61 | // TCHECK-NOT: alloca i{{[0-9]+}}, |
62 | // TCHECK: [[B_PRIV:%.+]] = alloca [10 x float], |
63 | // TCHECK: [[C_PRIV:%.+]] = alloca [5 x [10 x double]], |
64 | // TCHECK: [[D_PRIV:%.+]] = alloca [[TT]], |
65 | // TCHECK: store i{{[0-9]+}} [[A2_IN]], i{{[0-9]+}}* [[A2_ADDR]], |
66 | // TCHECK: store [10 x float]* [[B_IN]], [10 x float]** [[B_ADDR]], |
67 | // TCHECK: store [5 x [10 x double]]* [[C_IN]], [5 x [10 x double]]** [[C_ADDR]], |
68 | // TCHECK: store [[TT]]* [[D_IN]], [[TT]]** [[D_ADDR]], |
69 | // TCHECK: [[B_ADDR_REF:%.+]] = load [10 x float]*, [10 x float]** [[B_ADDR]], |
70 | // TCHECK: [[B_ADDR_REF:%.+]] = load [10 x float]*, [10 x float]** % |
71 | // TCHECK: [[C_ADDR_REF:%.+]] = load [5 x [10 x double]]*, [5 x [10 x double]]** [[C_ADDR]], |
72 | // TCHECK: [[C_ADDR_REF:%.+]] = load [5 x [10 x double]]*, [5 x [10 x double]]** % |
73 | // TCHECK: [[D_ADDR_REF:%.+]] = load [[TT]]*, [[TT]]** [[D_ADDR]], |
74 | // TCHECK: [[D_ADDR_REF:%.+]] = load [[TT]]*, [[TT]]** % |
75 | |
76 | // firstprivate(aa): a_priv = a_in |
77 | |
78 | // firstprivate(b): memcpy(b_priv,b_in) |
79 | // TCHECK: [[B_PRIV_BCAST:%.+]] = bitcast [10 x float]* [[B_PRIV]] to i8* |
80 | // TCHECK: [[B_ADDR_REF_BCAST:%.+]] = bitcast [10 x float]* [[B_ADDR_REF]] to i8* |
81 | // TCHECK: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[B_PRIV_BCAST]], i8* align {{[0-9]+}} [[B_ADDR_REF_BCAST]], {{.+}}) |
82 | |
83 | // firstprivate(c) |
84 | // TCHECK: [[C_PRIV_BCAST:%.+]] = bitcast [5 x [10 x double]]* [[C_PRIV]] to i8* |
85 | // TCHECK: [[C_IN_BCAST:%.+]] = bitcast [5 x [10 x double]]* [[C_ADDR_REF]] to i8* |
86 | // TCHECK: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[C_PRIV_BCAST]], i8* align {{[0-9]+}} [[C_IN_BCAST]],{{.+}}) |
87 | |
88 | // firstprivate(d) |
89 | // TCHECK: [[D_PRIV_BCAST:%.+]] = bitcast [[TT]]* [[D_PRIV]] to i8* |
90 | // TCHECK: [[D_IN_BCAST:%.+]] = bitcast [[TT]]* [[D_ADDR_REF]] to i8* |
91 | // TCHECK: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[D_PRIV_BCAST]], i8* align {{[0-9]+}} [[D_IN_BCAST]],{{.+}}) |
92 | |
93 | // TCHECK: load i16, i16* [[A2_ADDR]], |
94 | |
95 | #pragma omp target firstprivate(ptr) |
96 | { |
97 | ptr[0]++; |
98 | } |
99 | |
100 | // TCHECK: define weak void @__omp_offloading_{{.+}}(double* [[PTR_IN:%.+]]) |
101 | // TCHECK: [[PTR_ADDR:%.+]] = alloca double*, |
102 | // TCHECK-NOT: alloca double*, |
103 | // TCHECK: store double* [[PTR_IN]], double** [[PTR_ADDR]], |
104 | // TCHECK: [[PTR_IN_REF:%.+]] = load double*, double** [[PTR_ADDR]], |
105 | // TCHECK-NOT: store double* [[PTR_IN_REF]], double** [[PTR_PRIV]], |
106 | |
107 | return a; |
108 | } |
109 | |
110 | template <typename tx> |
111 | tx ftemplate(int n) { |
112 | tx a = 0; |
113 | tx b[10]; |
114 | |
115 | #pragma omp target firstprivate(a, b) |
116 | { |
117 | a += 1; |
118 | b[2] += 1; |
119 | } |
120 | |
121 | return a; |
122 | } |
123 | |
124 | static int fstatic(int n) { |
125 | int a = 0; |
126 | char aaa = 0; |
127 | int b[10]; |
128 | |
129 | #pragma omp target firstprivate(a, aaa, b) |
130 | { |
131 | a += 1; |
132 | aaa += 1; |
133 | b[2] += 1; |
134 | } |
135 | |
136 | return a; |
137 | } |
138 | |
139 | // TCHECK: define {{.*}}void @__omp_offloading_{{.+}}(i{{[0-9]+}}{{.*}} [[A_IN:%.+]], i{{[0-9]+}}{{.*}} [[A3_IN:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]]) |
140 | // TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}}, |
141 | // TCHECK: [[A3_ADDR:%.+]] = alloca i{{[0-9]+}}, |
142 | // TCHECK: [[B_ADDR:%.+]] = alloca [10 x i{{[0-9]+}}]*, |
143 | // TCHECK-NOT: alloca i{{[0-9]+}}, |
144 | // TCHECK: [[B_PRIV:%.+]] = alloca [10 x i{{[0-9]+}}], |
145 | // TCHECK: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]], |
146 | // TCHECK: store i{{[0-9]+}} [[A3_IN]], i{{[0-9]+}}* [[A3_ADDR]], |
147 | // TCHECK: store [10 x i{{[0-9]+}}]* [[B_IN]], [10 x i{{[0-9]+}}]** [[B_ADDR]], |
148 | // TCHECK: [[B_ADDR_REF:%.+]] = load [10 x i{{[0-9]+}}]*, [10 x i{{[0-9]+}}]** [[B_ADDR]], |
149 | // TCHECK: [[B_ADDR_REF:%.+]] = load [10 x i{{[0-9]+}}]*, [10 x i{{[0-9]+}}]** % |
150 | |
151 | // firstprivate(a): a_priv = a_in |
152 | |
153 | // firstprivate(aaa) |
154 | |
155 | // TCHECK-NOT: store i{{[0-9]+}} %{{.+}}, i{{[0-9]+}}* |
156 | |
157 | // firstprivate(b) |
158 | // TCHECK: [[B_PRIV_BCAST:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_PRIV]] to i8* |
159 | // TCHECK: [[B_IN_BCAST:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_ADDR_REF]] to i8* |
160 | // TCHECK: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[B_PRIV_BCAST]], i8* align {{[0-9]+}} [[B_IN_BCAST]],{{.+}}) |
161 | |
162 | // TCHECK: ret void |
163 | |
164 | struct S1 { |
165 | double a; |
166 | |
167 | int r1(int n) { |
168 | int b = n + 1; |
169 | |
170 | #pragma omp target firstprivate(b) |
171 | { |
172 | this->a = (double)b + 1.5; |
173 | } |
174 | |
175 | return (int)b; |
176 | } |
177 | |
178 | // TCHECK: define internal void @__omp_offloading_{{.+}}([[S1]]* [[TH:%.+]], i{{[0-9]+}} [[B_IN:%.+]]) |
179 | // TCHECK: [[TH_ADDR:%.+]] = alloca [[S1]]*, |
180 | // TCHECK: [[B_ADDR:%.+]] = alloca i{{[0-9]+}}, |
181 | // TCHECK-NOT: alloca i{{[0-9]+}}, |
182 | |
183 | // TCHECK: store [[S1]]* [[TH]], [[S1]]** [[TH_ADDR]], |
184 | // TCHECK: store i{{[0-9]+}} [[B_IN]], i{{[0-9]+}}* [[B_ADDR]], |
185 | // TCHECK: [[TH_ADDR_REF:%.+]] = load [[S1]]*, [[S1]]** [[TH_ADDR]], |
186 | // TCHECK-64: [[B_ADDR_CONV:%.+]] = bitcast i{{[0-9]+}}* [[B_ADDR]] to i{{[0-9]+}}* |
187 | |
188 | // firstprivate(b) |
189 | // TCHECK-NOT: store i{{[0-9]+}} %{{.+}}, i{{[0-9]+}}* |
190 | |
191 | // TCHECK: ret void |
192 | }; |
193 | |
194 | int bar(int n, double *ptr) { |
195 | int a = 0; |
196 | a += foo(n, ptr); |
197 | S1 S; |
198 | a += S.r1(n); |
199 | a += fstatic(n); |
200 | a += ftemplate<int>(n); |
201 | |
202 | return a; |
203 | } |
204 | |
205 | // template |
206 | |
207 | // TCHECK: define internal void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]]) |
208 | // TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}}, |
209 | // TCHECK: [[B_ADDR:%.+]] = alloca [10 x i{{[0-9]+}}]*, |
210 | // TCHECK-NOT: alloca i{{[0-9]+}}, |
211 | // TCHECK: [[B_PRIV:%.+]] = alloca [10 x i{{[0-9]+}}], |
212 | // TCHECK: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]], |
213 | // TCHECK: store [10 x i{{[0-9]+}}]* [[B_IN]], [10 x i{{[0-9]+}}]** [[B_ADDR]], |
214 | // TCHECK: [[B_ADDR_REF:%.+]] = load [10 x i{{[0-9]+}}]*, [10 x i{{[0-9]+}}]** [[B_ADDR]], |
215 | // TCHECK: [[B_ADDR_REF:%.+]] = load [10 x i{{[0-9]+}}]*, [10 x i{{[0-9]+}}]** % |
216 | |
217 | // firstprivate(a) |
218 | // TCHECK-NOT: store i{{[0-9]+}} %{{.+}}, i{{[0-9]+}}* |
219 | |
220 | // firstprivate(b) |
221 | // TCHECK: [[B_PRIV_BCAST:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_PRIV]] to i8* |
222 | // TCHECK: [[B_IN_BCAST:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_ADDR_REF]] to i8* |
223 | // TCHECK: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[B_PRIV_BCAST]], i8* align {{[0-9]+}} [[B_IN_BCAST]],{{.+}}) |
224 | |
225 | // TCHECK: ret void |
226 | |
227 | #endif |
228 | |
229 | // TCHECK-DAG: distinct !DISubprogram(linkageName: "__omp_offloading_{{.+}}_worker", |
230 | // TCHECK-DAG: distinct !DISubprogram(linkageName: "__omp_offloading_{{.+}}_worker", |
231 | // TCHECK-DAG: distinct !DISubprogram(linkageName: "__omp_offloading_{{.+}}_worker", |
232 | // TCHECK-DAG: distinct !DISubprogram(linkageName: "__omp_offloading_{{.+}}_worker", |
233 | // TCHECK-DAG: distinct !DISubprogram(linkageName: "__omp_offloading_{{.+}}_worker", |
234 | |