1 // RUN: %clang_cc1 -no-opaque-pointers -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -debug-info-kind=limited -emit-llvm %s -o - | FileCheck %s --check-prefix DEBUG
2 // RUN: %clang_cc1 -no-opaque-pointers -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK
3 #ifndef HEADER
4 #define HEADER
5 
6 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";d;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
7 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";i;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
8 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";i[1:23];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
9 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";p;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
10 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";p[1:24];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
11 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
12 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.i;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
13 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.s.f;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
14 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.p[:22];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
15 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.ps->s.i;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
16 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.ps->ps;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
17 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.ps->ps->ps;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
18 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.ps->ps->s.f[:22];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
19 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
20 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->i;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
21 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->s.f;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
22 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->p[:22];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
23 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->ps->s.i;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
24 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->ps->ps;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
25 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->ps->ps->ps;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
26 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->ps->ps->s.f[:22];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
27 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.f[:22];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
28 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.p[:33];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
29 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->p[:33];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
30 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.s;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
31 
32 struct S1 {
33     int i;
34     float f[50];
35 };
36 
37 struct S2 {
38     int i;
39     float f[50];
40     S1 s;
41     double *p;
42     struct S2 *ps;
43 };
44 
foo()45 void foo() {
46   double d;
47   int i[100];
48   float *p;
49 
50   S2 s;
51   S2 *ps;
52 
53 #pragma omp target map(d)
54   { }
55 #pragma omp target map(i)
56   { }
57 #pragma omp target map(i[1:23])
58   { }
59 #pragma omp target map(p)
60   { }
61 #pragma omp target map(p[1:24])
62   { }
63 #pragma omp target map(s)
64   { }
65 #pragma omp target map(s.i)
66   { }
67 #pragma omp target map(s.s.f)
68   { }
69 #pragma omp target map(s.p)
70   { }
71 #pragma omp target map(to: s.p[:22])
72   { }
73 #pragma omp target map(s.ps)
74   { }
75 #pragma omp target map(from: s.ps->s.i)
76   { }
77 #pragma omp target map(to: s.ps->ps)
78   { }
79 #pragma omp target map(s.ps->ps->ps)
80   { }
81 #pragma omp target map(to: s.ps->ps->s.f[:22])
82   { }
83 #pragma omp target map(ps)
84   { }
85 #pragma omp target map(ps->i)
86   { }
87 #pragma omp target map(ps->s.f)
88   { }
89 #pragma omp target map(from: ps->p)
90   { }
91 #pragma omp target map(to: ps->p[:22])
92   { }
93 #pragma omp target map(ps->ps)
94   { }
95 #pragma omp target map(from: ps->ps->s.i)
96   { }
97 #pragma omp target map(from: ps->ps->ps)
98   { }
99 #pragma omp target map(ps->ps->ps->ps)
100   { }
101 #pragma omp target map(to: ps->ps->ps->s.f[:22])
102   { }
103 #pragma omp target map(to: s.f[:22]) map(from: s.p[:33])
104   { }
105 #pragma omp target map(from: s.f[:22]) map(to: ps->p[:33])
106   { }
107 #pragma omp target map(from: s.f[:22], s.s) map(to: ps->p[:33])
108   { }
109 }
110 
111 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";B;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
112 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";unknown;unknown;0;0;;\00"
113 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";A;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
114 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";x;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
115 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";fn;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
116 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
117 // DEBUG: @{{.+}} = private constant [7 x i8*] [i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @{{[0-9]+}}, i32 0, i32 0), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @{{[0-9]+}}, i32 0, i32 0), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @{{[0-9]+}}, i32 0, i32 0), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @{{[0-9]+}}, i32 0, i32 0), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @{{[0-9]+}}, i32 0, i32 0), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @{{[0-9]+}}, i32 0, i32 0), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @{{[0-9]+}}, i32 0, i32 0)]
118 
bar(int N)119 void bar(int N) {
120   double B[10];
121   double A[N];
122   double x;
123   S1 s;
124   auto fn = [&x]() { return x; };
125 #pragma omp target
126   {
127     (void)B;
128     (void)A;
129     (void)fn();
130     (void)s.f;
131   }
132 }
133 
134 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";t;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
135 
136 #pragma omp declare target
137 double t;
138 #pragma omp end declare target
139 
baz()140 void baz() {
141 #pragma omp target map(to:t)
142   { }
143 #pragma omp target map(to:t) nowait
144   { }
145 #pragma omp target teams map(to:t)
146   { }
147 #pragma omp target teams map(to:t) nowait
148   { }
149 #pragma omp target data map(to:t)
150   { }
151 #pragma omp target enter data map(to:t)
152 
153 #pragma omp target enter data map(to:t) nowait
154 
155 #pragma omp target exit data map(from:t)
156 
157 #pragma omp target exit data map(from:t) nowait
158 
159 #pragma omp target update from(t)
160 
161 #pragma omp target update to(t)
162 
163 #pragma omp target update from(t) nowait
164 
165 #pragma omp target update to(t) nowait
166 }
167 
168 struct S3 {
S3S3169   S3() {
170 #pragma omp target data map(alloc : Z[0:64])
171     { }
172   }
173   double Z[64];
174 };
175 
176 #pragma omp declare mapper(id: S3 s) map(s.Z[0:64])
177 
qux()178 void qux() {
179   S3 s;
180 #pragma omp target map(mapper(id), to:s)
181   { }
182 }
183 
184 
185 
186 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.Z[0:64];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
187 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";this->Z[0:64];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
188 
189 // Clang used to mistakenly generate the map name "x" for both x and y on this
190 // directive.  Conditions to reproduce the bug: a single map clause has two
191 // variables, and at least the second is used in the associated statement.
192 //
193 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";x;{{.*}}.cpp;[[@LINE+3]];7;;\00"
194 // DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";y;{{.*}}.cpp;[[@LINE+2]];10;;\00"
secondMapNameInClause()195 void secondMapNameInClause() {
196   int x, y;
197   #pragma omp target map(to: x, y)
198   x = y = 1;
199 }
200 
201 // DEBUG: store i8** getelementptr inbounds ([2 x i8*], [2 x i8*]* @[[NAME:.offload_mapnames.[0-9]+]], i32 0, i32 0), i8*** %[[ARG:.+]]
202 // CHECK-NOT: store i8** getelementptr inbounds ([2 x i8*], [2 x i8*]* @[[NAME:.offload_mapnames.[0-9]+]], i32 0, i32 0), i8*** %[[ARG:.+]]
203 
204 // DEBUG: void @.omp_mapper._ZTS2S3.id(i8* {{.*}}, i8* {{.*}}, i8* {{.*}}, i64 {{.*}}, i64 {{.*}}, i8* noundef [[NAME_ARG:%.+]])
205 // DEBUG: store i8* [[NAME_ARG]], i8** [[NAME_STACK:%.+]]
206 // DEBUG: [[MAPPER_NAME:%.+]] = load i8*, i8** [[NAME_STACK]]
207 // DEBUG: call void @__tgt_push_mapper_component(i8* %{{.*}}, i8* %{{.*}}, i8* %{{.*}}, i64 %{{.*}}, i64 %{{.*}}, i8* [[MAPPER_NAME]])
208 
209 #endif
210