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