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: [[TTII:%.+]] = type { i32, i32 }
17 // TCHECK-DAG: [[TTIC:%.+]] = type { i8, i8 }
18 // TCHECK-DAG: [[TT:%.+]] = type { i64, i8 }
19 // TCHECK-DAG: [[S1:%.+]] = type { double }
20
foo(int n,double * ptr)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: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
38 // TCHECK-NOT: alloca [[TTII]],
39 // TCHECK: alloca i{{[0-9]+}},
40 // TCHECK: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]],
41 // TCHECK: ret void
42
43 #pragma omp target firstprivate(aa, b, c, d)
44 {
45 aa += 1;
46 b[2] = 1.0;
47 c[1][2] = 1.0;
48 d.X = 1;
49 d.Y = 1;
50 }
51
52 // make sure that firstprivate variables are generated in all cases and that we use those instances for operations inside the
53 // target region
54 // 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:%.+]])
55 // TCHECK: [[A2_ADDR:%.+]] = alloca i{{[0-9]+}},
56 // TCHECK: [[B_ADDR:%.+]] = alloca [10 x float]*,
57 // TCHECK: [[C_ADDR:%.+]] = alloca [5 x [10 x double]]*,
58 // TCHECK: [[D_ADDR:%.+]] = alloca [[TT]]*,
59 // TCHECK-NOT: alloca i{{[0-9]+}},
60 // TCHECK: [[B_PRIV:%.+]] = alloca [10 x float],
61 // TCHECK: [[C_PRIV:%.+]] = alloca [5 x [10 x double]],
62 // TCHECK: [[D_PRIV:%.+]] = alloca [[TT]],
63 // TCHECK: store i{{[0-9]+}} [[A2_IN]], i{{[0-9]+}}* [[A2_ADDR]],
64 // TCHECK: store [10 x float]* [[B_IN]], [10 x float]** [[B_ADDR]],
65 // TCHECK: store [5 x [10 x double]]* [[C_IN]], [5 x [10 x double]]** [[C_ADDR]],
66 // TCHECK: store [[TT]]* [[D_IN]], [[TT]]** [[D_ADDR]],
67 // TCHECK: [[B_ADDR_REF:%.+]] = load [10 x float]*, [10 x float]** [[B_ADDR]],
68 // TCHECK: [[B_ADDR_REF:%.+]] = load [10 x float]*, [10 x float]** %
69 // TCHECK: [[C_ADDR_REF:%.+]] = load [5 x [10 x double]]*, [5 x [10 x double]]** [[C_ADDR]],
70 // TCHECK: [[C_ADDR_REF:%.+]] = load [5 x [10 x double]]*, [5 x [10 x double]]** %
71 // TCHECK: [[D_ADDR_REF:%.+]] = load [[TT]]*, [[TT]]** [[D_ADDR]],
72 // TCHECK: [[D_ADDR_REF:%.+]] = load [[TT]]*, [[TT]]** %
73
74 // firstprivate(aa): a_priv = a_in
75
76 // firstprivate(b): memcpy(b_priv,b_in)
77 // TCHECK: [[B_PRIV_BCAST:%.+]] = bitcast [10 x float]* [[B_PRIV]] to i8*
78 // TCHECK: [[B_ADDR_REF_BCAST:%.+]] = bitcast [10 x float]* [[B_ADDR_REF]] to i8*
79 // TCHECK: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[B_PRIV_BCAST]], i8* align {{[0-9]+}} [[B_ADDR_REF_BCAST]], {{.+}})
80
81 // firstprivate(c)
82 // TCHECK: [[C_PRIV_BCAST:%.+]] = bitcast [5 x [10 x double]]* [[C_PRIV]] to i8*
83 // TCHECK: [[C_IN_BCAST:%.+]] = bitcast [5 x [10 x double]]* [[C_ADDR_REF]] to i8*
84 // TCHECK: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[C_PRIV_BCAST]], i8* align {{[0-9]+}} [[C_IN_BCAST]],{{.+}})
85
86 // firstprivate(d)
87 // TCHECK: [[D_PRIV_BCAST:%.+]] = bitcast [[TT]]* [[D_PRIV]] to i8*
88 // TCHECK: [[D_IN_BCAST:%.+]] = bitcast [[TT]]* [[D_ADDR_REF]] to i8*
89 // TCHECK: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[D_PRIV_BCAST]], i8* align {{[0-9]+}} [[D_IN_BCAST]],{{.+}})
90
91 // TCHECK: load i16, i16* [[A2_ADDR]],
92
93 #pragma omp target firstprivate(ptr)
94 {
95 ptr[0]++;
96 }
97
98 // TCHECK: define weak void @__omp_offloading_{{.+}}(double* [[PTR_IN:%.+]])
99 // TCHECK: [[PTR_ADDR:%.+]] = alloca double*,
100 // TCHECK-NOT: alloca double*,
101 // TCHECK: store double* [[PTR_IN]], double** [[PTR_ADDR]],
102 // TCHECK: [[PTR_IN_REF:%.+]] = load double*, double** [[PTR_ADDR]],
103 // TCHECK-NOT: store double* [[PTR_IN_REF]], double** {{%.+}},
104
105 return a;
106 }
107
108 template <typename tx>
ftemplate(int n)109 tx ftemplate(int n) {
110 tx a = 0;
111 tx b[10];
112
113 #pragma omp target firstprivate(a, b)
114 {
115 a += 1;
116 b[2] += 1;
117 }
118
119 return a;
120 }
121
fstatic(int n)122 static int fstatic(int n) {
123 int a = 0;
124 char aaa = 0;
125 int b[10];
126
127 #pragma omp target firstprivate(a, aaa, b)
128 {
129 a += 1;
130 aaa += 1;
131 b[2] += 1;
132 }
133
134 return a;
135 }
136
137 template <typename tx>
fconst(const tx t)138 void fconst(const tx t) {
139 #pragma omp target firstprivate(t)
140 { }
141 }
142
143 // TCHECK: define {{.*}}void @__omp_offloading_{{.+}}(i{{[0-9]+}}{{.*}} [[A_IN:%.+]], i{{[0-9]+}}{{.*}} [[A3_IN:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]])
144 // TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
145 // TCHECK: [[A3_ADDR:%.+]] = alloca i{{[0-9]+}},
146 // TCHECK: [[B_ADDR:%.+]] = alloca [10 x i{{[0-9]+}}]*,
147 // TCHECK-NOT: alloca i{{[0-9]+}},
148 // TCHECK: [[B_PRIV:%.+]] = alloca [10 x i{{[0-9]+}}],
149 // TCHECK: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]],
150 // TCHECK: store i{{[0-9]+}} [[A3_IN]], i{{[0-9]+}}* [[A3_ADDR]],
151 // TCHECK: store [10 x i{{[0-9]+}}]* [[B_IN]], [10 x i{{[0-9]+}}]** [[B_ADDR]],
152 // TCHECK: [[B_ADDR_REF:%.+]] = load [10 x i{{[0-9]+}}]*, [10 x i{{[0-9]+}}]** [[B_ADDR]],
153 // TCHECK: [[B_ADDR_REF:%.+]] = load [10 x i{{[0-9]+}}]*, [10 x i{{[0-9]+}}]** %
154
155 // firstprivate(a): a_priv = a_in
156
157 // firstprivate(aaa)
158
159 // TCHECK-NOT: store i{{[0-9]+}} %{{.+}}, i{{[0-9]+}}*
160
161 // firstprivate(b)
162 // TCHECK: [[B_PRIV_BCAST:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_PRIV]] to i8*
163 // TCHECK: [[B_IN_BCAST:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_ADDR_REF]] to i8*
164 // TCHECK: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[B_PRIV_BCAST]], i8* align {{[0-9]+}} [[B_IN_BCAST]],{{.+}})
165
166 // TCHECK: ret void
167
168 struct S1 {
169 double a;
170
r1S1171 int r1(int n) {
172 int b = n + 1;
173
174 #pragma omp target firstprivate(b)
175 {
176 this->a = (double)b + 1.5;
177 }
178
179 return (int)b;
180 }
181
182 // TCHECK: define internal void @__omp_offloading_{{.+}}([[S1]]* [[TH:%.+]], i{{[0-9]+}} [[B_IN:%.+]])
183 // TCHECK: [[TH_ADDR:%.+]] = alloca [[S1]]*,
184 // TCHECK: [[B_ADDR:%.+]] = alloca i{{[0-9]+}},
185 // TCHECK-NOT: alloca i{{[0-9]+}},
186
187 // TCHECK: store [[S1]]* [[TH]], [[S1]]** [[TH_ADDR]],
188 // TCHECK: store i{{[0-9]+}} [[B_IN]], i{{[0-9]+}}* [[B_ADDR]],
189 // TCHECK: [[TH_ADDR_REF:%.+]] = load [[S1]]*, [[S1]]** [[TH_ADDR]],
190 // TCHECK-64: [[B_ADDR_CONV:%.+]] = bitcast i{{[0-9]+}}* [[B_ADDR]] to i{{[0-9]+}}*
191
192 // firstprivate(b)
193 // TCHECK-NOT: store i{{[0-9]+}} %{{.+}}, i{{[0-9]+}}*
194
195 // TCHECK: ret void
196 };
197
bar(int n,double * ptr)198 int bar(int n, double *ptr) {
199 int a = 0;
200 a += foo(n, ptr);
201 S1 S;
202 a += S.r1(n);
203 a += fstatic(n);
204 a += ftemplate<int>(n);
205
206 fconst(TT<int, int>{0, 0});
207 fconst(TT<char, char>{0, 0});
208
209 return a;
210 }
211
212 // template
213
214 // TCHECK: define internal void @__omp_offloading_{{.+}}(i{{[0-9]+}} [[A_IN:%.+]], [10 x i{{[0-9]+}}]*{{.+}} [[B_IN:%.+]])
215 // TCHECK: [[A_ADDR:%.+]] = alloca i{{[0-9]+}},
216 // TCHECK: [[B_ADDR:%.+]] = alloca [10 x i{{[0-9]+}}]*,
217 // TCHECK-NOT: alloca i{{[0-9]+}},
218 // TCHECK: [[B_PRIV:%.+]] = alloca [10 x i{{[0-9]+}}],
219 // TCHECK: store i{{[0-9]+}} [[A_IN]], i{{[0-9]+}}* [[A_ADDR]],
220 // TCHECK: store [10 x i{{[0-9]+}}]* [[B_IN]], [10 x i{{[0-9]+}}]** [[B_ADDR]],
221 // TCHECK: [[B_ADDR_REF:%.+]] = load [10 x i{{[0-9]+}}]*, [10 x i{{[0-9]+}}]** [[B_ADDR]],
222 // TCHECK: [[B_ADDR_REF:%.+]] = load [10 x i{{[0-9]+}}]*, [10 x i{{[0-9]+}}]** %
223
224 // firstprivate(a)
225 // TCHECK-NOT: store i{{[0-9]+}} %{{.+}}, i{{[0-9]+}}*
226
227 // firstprivate(b)
228 // TCHECK: [[B_PRIV_BCAST:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_PRIV]] to i8*
229 // TCHECK: [[B_IN_BCAST:%.+]] = bitcast [10 x i{{[0-9]+}}]* [[B_ADDR_REF]] to i8*
230 // TCHECK: call void @llvm.memcpy.{{.+}}(i8* align {{[0-9]+}} [[B_PRIV_BCAST]], i8* align {{[0-9]+}} [[B_IN_BCAST]],{{.+}})
231
232 // TCHECK: ret void
233
234 #endif
235