1 // expected-no-diagnostics
2 #ifndef HEADER
3 #define HEADER
4
5 ///==========================================================================///
6 // RUN: %clang_cc1 -DCK24 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-64
7 // RUN: %clang_cc1 -DCK24 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
8 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-64
9 // RUN: %clang_cc1 -DCK24 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-32
10 // RUN: %clang_cc1 -DCK24 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
11 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-32
12
13 // RUN: %clang_cc1 -DCK24 -verify -fopenmp -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-64
14 // RUN: %clang_cc1 -DCK24 -fopenmp -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
15 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-64
16 // RUN: %clang_cc1 -DCK24 -verify -fopenmp -fopenmp-version=45 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-32
17 // RUN: %clang_cc1 -DCK24 -fopenmp -fopenmp-version=45 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
18 // RUN: %clang_cc1 -fopenmp -fopenmp-version=45 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-32
19
20 // RUN: %clang_cc1 -DCK24 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-64
21 // RUN: %clang_cc1 -DCK24 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
22 // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-64
23 // RUN: %clang_cc1 -DCK24 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-32
24 // RUN: %clang_cc1 -DCK24 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
25 // RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-32
26
27 // RUN: %clang_cc1 -DCK24 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY23 %s
28 // RUN: %clang_cc1 -DCK24 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
29 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY23 %s
30 // RUN: %clang_cc1 -DCK24 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY23 %s
31 // RUN: %clang_cc1 -DCK24 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
32 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY23 %s
33 // SIMD-ONLY23-NOT: {{__kmpc|__tgt}}
34 #ifdef CK24
35
36 // CK24-DAG: [[SC:%.+]] = type { i32, [[SB:%.+]], ptr, [10 x i32] }
37 // CK24-DAG: [[SB]] = type { i32, [[SA:%.+]], [10 x [[SA:%.+]]], [10 x ptr], ptr }
38 // CK24-DAG: [[SA]] = type { i32, ptr, [10 x i32] }
39
40 struct SA{
41 int a;
42 struct SA *p;
43 int b[10];
44 };
45 struct SB{
46 int a;
47 struct SA s;
48 struct SA sa[10];
49 struct SA *sp[10];
50 struct SA *p;
51 };
52 struct SC{
53 int a;
54 struct SB s;
55 struct SB *p;
56 int b[10];
57 };
58
59 // CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
60 // CK24: [[SIZE01:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
61 // CK24: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
62
63 // CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
64 // CK24: [[SIZE13:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
65 // CK24: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
66
67 // CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
68 // CK24: [[SIZE14:@.+]] = private {{.*}}constant [1 x i64] [i64 {{48|56}}]
69 // CK24: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
70
71 // CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
72 // CK24: [[SIZE15:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
73 // CK24: [[MTYPE15:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
74
75 // CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
76 // CK24: [[SIZE16:@.+]] = private {{.*}}constant [2 x i64] [i64 0, i64 20]
77 // CK24: [[MTYPE16:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710659]
78
79 // CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
80 // CK24: [[SIZE17:@.+]] = private {{.*}}constant [2 x i64] [i64 0, i64 {{3560|2880}}]
81 // CK24: [[MTYPE17:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710675]
82
83 // CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
84 // CK24: [[SIZE18:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
85 // CK24: [[MTYPE18:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
86
87 // CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
88 // CK24: [[SIZE19:@.+]] = private {{.*}}constant [3 x i64] [i64 0, i64 {{8|4}}, i64 4]
89 // CK24: [[MTYPE19:@.+]] = private {{.*}}constant [3 x i64] [i64 32, i64 281474976710659, i64 281474976710675]
90
91 // CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
92 // CK24: [[SIZE20:@.+]] = private {{.*}}constant [2 x i64] [i64 0, i64 4]
93 // CK24: [[MTYPE20:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710675]
94
95 // CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
96 // CK24: [[SIZE21:@.+]] = private {{.*}}constant [3 x i64] [i64 0, i64 {{8|4}}, i64 4]
97 // CK24: [[MTYPE21:@.+]] = private {{.*}}constant [3 x i64] [i64 32, i64 281474976710659, i64 281474976710675]
98
99 // CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
100 // CK24: [[SIZE22:@.+]] = private {{.*}}constant [2 x i64] [i64 0, i64 8]
101 // CK24: [[MTYPE22:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710659]
102
103 // CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
104 // CK24: [[SIZE23:@.+]] = private {{.*}}constant [3 x i64] [i64 0, i64 {{8|4}}, i64 8]
105 // CK24: [[MTYPE23:@.+]] = private {{.*}}constant [3 x i64] [i64 32, i64 281474976710659, i64 281474976710675]
106
107 // CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
108 // CK24: [[SIZE24:@.+]] = private {{.*}}constant [4 x i64] [i64 0, i64 {{8|4}}, i64 {{8|4}}, i64 4]
109 // CK24: [[MTYPE24:@.+]] = private {{.*}}constant [4 x i64] [i64 32, i64 281474976710672, i64 16, i64 19]
110
111 // CK24-LABEL: explicit_maps_struct_fields
explicit_maps_struct_fields(int a)112 int explicit_maps_struct_fields(int a){
113 SC s;
114 SC *p;
115
116 // Region 01
117 // CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
118 // CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
119 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
120 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
121 // CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
122 // CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
123 // CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
124
125 // CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
126 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
127 // CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
128 // CK24-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
129 // CK24-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0, i{{.+}} 0
130
131 // CK24: call void [[CALL01:@.+]](ptr {{[^,]+}})
132 #pragma omp target map(s.a)
133 { s.a++; }
134
135 //
136 // Same thing but starting from a pointer.
137 //
138 // Region 13
139 // CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
140 // CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
141 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
142 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
143 // CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
144 // CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
145 // CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
146
147 // CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
148 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
149 // CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
150 // CK24-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
151 // CK24-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 0
152
153 // CK24-DAG: [[VAR0]] = load ptr, ptr %{{.+}}
154 // CK24-DAG: [[VAR00]] = load ptr, ptr %{{.+}}
155
156 // CK24: call void [[CALL13:@.+]](ptr {{[^,]+}})
157 #pragma omp target map(p->a)
158 { p->a++; }
159
160 // Region 14
161 // CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
162 // CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
163 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
164 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
165 // CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
166 // CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
167 // CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
168
169 // CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
170 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
171 // CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
172 // CK24-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
173 // CK24-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 1
174 // CK24-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1
175
176 // CK24-DAG: [[VAR0]] = load ptr, ptr %{{.+}}
177 // CK24-DAG: [[VAR00]] = load ptr, ptr %{{.+}}
178
179 // CK24: call void [[CALL14:@.+]](ptr {{[^,]+}})
180 #pragma omp target map(p->s.s)
181 { p->a++; }
182
183 // Region 15
184 // CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
185 // CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
186 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
187 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
188 // CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
189 // CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
190 // CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
191
192 // CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
193 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
194 // CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
195 // CK24-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
196 // CK24-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0
197 // CK24-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 1
198 // CK24-DAG: [[SEC000]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1
199
200 // CK24-DAG: [[VAR0]] = load ptr, ptr %{{.+}}
201 // CK24-DAG: [[VAR00]] = load ptr, ptr %{{.+}}
202
203 // CK24: call void [[CALL15:@.+]](ptr {{[^,]+}})
204 #pragma omp target map(p->s.s.a)
205 { p->a++; }
206
207 // Region 16
208 // CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
209 // CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
210 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
211 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
212 // CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
213 // CK24-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
214 // CK24-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]
215 // CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
216 // CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
217 // CK24-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
218
219 // CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
220 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
221 // CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
222 // CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
223 // CK24-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
224 // CK24-DAG: store i64 {{%.+}}, ptr [[S0]]
225 // CK24-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0
226 // CK24-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 3
227
228 // CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
229 // CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
230 // CK24-DAG: store ptr [[VAR0]], ptr [[BP1]]
231 // CK24-DAG: store ptr [[SEC0]], ptr [[P1]]
232
233 // CK24-DAG: [[VAR0]] = load ptr, ptr %{{.+}}
234 // CK24-DAG: [[VAR00]] = load ptr, ptr %{{.+}}
235
236 // CK24: call void [[CALL16:@.+]](ptr {{[^,]+}})
237 #pragma omp target map(p->b[:5])
238 { p->a++; }
239
240 // Region 17
241 // CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
242 // CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
243 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
244 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
245 // CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
246 // CK24-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
247 // CK24-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]
248 // CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
249 // CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
250 // CK24-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
251
252 // CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
253 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
254 // CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
255 // CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
256 // CK24-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
257 // CK24-DAG: store i64 {{%.+}}, ptr [[S0]]
258 // CK24-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 2
259
260 // CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
261 // CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
262 // CK24-DAG: store ptr [[SEC0]], ptr [[BP1]]
263 // CK24-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]
264 // CK24-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[SEC11:%[^,]+]], i{{.+}} 0
265 // CK24-DAG: [[SEC11]] = load ptr, ptr [[SEC111:%[^,]+]],
266 // CK24-DAG: [[SEC111]] = getelementptr {{.*}}ptr [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 2
267
268 // CK24-DAG: [[VAR0]] = load ptr, ptr %{{.+}}
269 // CK24-DAG: [[VAR00]] = load ptr, ptr %{{.+}}
270 // CK24-DAG: [[VAR000]] = load ptr, ptr %{{.+}}
271
272 // CK24: call void [[CALL17:@.+]](ptr {{[^,]+}})
273 #pragma omp target map(p->p[:5])
274 { p->a++; }
275
276 // Region 18
277 // CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
278 // CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
279 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
280 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
281 // CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
282 // CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
283 // CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
284
285 // CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
286 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
287 // CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
288 // CK24-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
289 // CK24-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0
290 // CK24-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 3
291 // CK24-DAG: [[SEC000]] = getelementptr {{.*}}ptr [[SEC0000:%[^,]+]], i{{.+}} 0, i{{.+}} 2
292 // CK24-DAG: [[SEC0000]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1
293
294 // CK24-DAG: [[VAR0]] = load ptr, ptr %{{.+}}
295 // CK24-DAG: [[VAR00]] = load ptr, ptr %{{.+}}
296
297 // CK24: call void [[CALL18:@.+]](ptr {{[^,]+}})
298 #pragma omp target map(p->s.sa[3].a)
299 { p->a++; }
300
301 // Region 19
302 // CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
303 // CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
304 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
305 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
306 // CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
307 // CK24-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
308 // CK24-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]
309 // CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
310 // CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
311 // CK24-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
312
313 // CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
314 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
315 // CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
316 // CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
317 // CK24-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
318 // CK24-DAG: store i64 {{%.+}}, ptr [[S0]]
319 // CK24-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 3
320 // CK24-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 3
321 // CK24-DAG: [[SEC000]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1
322
323 // CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
324 // CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
325 // CK24-DAG: store ptr [[VAR0]], ptr [[BP1]]
326 // CK24-DAG: store ptr [[SEC0]], ptr [[P1]]
327
328 // CK24-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
329 // CK24-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
330 // CK24-DAG: store ptr [[SEC0]], ptr [[BP2]]
331 // CK24-DAG: store ptr [[SEC1:%.+]], ptr [[P2]]
332 // CK24-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[SEC11:%[^,]+]], i{{.+}} 0, i{{.+}} 0
333 // CK24-DAG: [[SEC11]] = load ptr, ptr [[SEC111:%[^,]+]],
334 // CK24-DAG: [[SEC111]] = getelementptr {{.*}}ptr [[SEC1111:%[^,]+]], i{{.+}} 0, i{{.+}} 3
335 // CK24-DAG: [[SEC1111]] = getelementptr {{.*}}ptr [[SEC11111:%[^,]+]], i{{.+}} 0, i{{.+}} 3
336 // CK24-DAG: [[SEC11111]] = getelementptr {{.*}}ptr [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 1
337
338 // CK24-DAG: [[VAR0]] = load ptr, ptr %{{.+}}
339 // CK24-DAG: [[VAR00]] = load ptr, ptr %{{.+}}
340 // CK24-DAG: [[VAR000]] = load ptr, ptr %{{.+}}
341
342 // CK24: call void [[CALL19:@.+]](ptr {{[^,]+}})
343 #pragma omp target map(p->s.sp[3]->a)
344 { p->a++; }
345
346 // Region 20
347 // CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
348 // CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
349 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
350 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
351 // CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
352 // CK24-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
353 // CK24-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]
354 // CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
355 // CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
356 // CK24-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
357
358 // CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
359 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
360 // CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
361 // CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
362 // CK24-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
363 // CK24-DAG: store i64 {{%.+}}, ptr [[S0]]
364 // CK24-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 2
365
366 // CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
367 // CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
368 // CK24-DAG: store ptr [[SEC0]], ptr [[BP1]]
369 // CK24-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]
370 // CK24-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[SEC11:%[^,]+]], i{{.+}} 0
371 // CK24-DAG: [[SEC11]] = load ptr, ptr [[SEC111:%[^,]+]],
372 // CK24-DAG: [[SEC111]] = getelementptr {{.*}}ptr [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 2
373
374 // CK24-DAG: [[VAR0]] = load ptr, ptr %{{.+}}
375 // CK24-DAG: [[VAR00]] = load ptr, ptr %{{.+}}
376 // CK24-DAG: [[VAR000]] = load ptr, ptr %{{.+}}
377
378 // CK24: call void [[CALL20:@.+]](ptr {{[^,]+}})
379 #pragma omp target map(p->p->a)
380 { p->a++; }
381
382 // Region 21
383 // CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
384 // CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
385 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
386 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
387 // CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
388 // CK24-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
389 // CK24-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]
390 // CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
391 // CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
392 // CK24-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
393
394 // CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
395 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
396 // CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
397 // CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
398 // CK24-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
399 // CK24-DAG: store i64 {{%.+}}, ptr [[S0]]
400 // CK24-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:[^,]+]], i{{.+}} 0, i{{.+}} 4
401 // CK24-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1
402
403 // CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
404 // CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
405 // CK24-DAG: store ptr [[VAR0]], ptr [[BP1]]
406 // CK24-DAG: store ptr [[SEC0]], ptr [[P1]]
407
408 // CK24-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
409 // CK24-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
410 // CK24-DAG: store ptr [[SEC0]], ptr [[BP2]]
411 // CK24-DAG: store ptr [[SEC1:%.+]], ptr [[P2]]
412 // CK24-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[SEC11:%[^,]+]], i{{.+}} 0
413 // CK24-DAG: [[SEC11]] = load ptr, ptr [[SEC111:%[^,]+]],
414 // CK24-DAG: [[SEC111]] = getelementptr {{.*}}ptr [[SEC1111:[^,]+]], i{{.+}} 0, i{{.+}} 4
415 // CK24-DAG: [[SEC1111]] = getelementptr {{.*}}ptr [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 1
416
417 // CK24-DAG: [[VAR0]] = load ptr, ptr %{{.+}}
418 // CK24-DAG: [[VAR00]] = load ptr, ptr %{{.+}}
419 // CK24-DAG: [[VAR000]] = load ptr, ptr %{{.+}}
420
421 // CK24: call void [[CALL21:@.+]](ptr {{[^,]+}})
422 #pragma omp target map(p->s.p->a)
423 { p->a++; }
424
425 // Region 22
426 // CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
427 // CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
428 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
429 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
430 // CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
431 // CK24-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
432 // CK24-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]
433 // CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
434 // CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
435 // CK24-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
436
437 // CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
438 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
439 // CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
440 // CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
441 // CK24-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
442 // CK24-DAG: store i64 {{%.+}}, ptr [[S0]]
443 // CK24-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0
444 // CK24-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 2
445 // CK24-DAG: [[SEC000]] = getelementptr {{.*}}ptr [[SEC0000:%[^,]+]], i{{.+}} 0, i{{.+}} 1
446 // CK24-DAG: [[SEC0000]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1
447
448 // CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
449 // CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
450 // CK24-DAG: store ptr [[VAR0]], ptr [[BP1]]
451 // CK24-DAG: store ptr [[SEC0]], ptr [[P1]]
452
453 // CK24-DAG: [[VAR0]] = load ptr, ptr %{{.+}}
454 // CK24-DAG: [[VAR00]] = load ptr, ptr %{{.+}}
455
456 // CK24: call void [[CALL22:@.+]](ptr {{[^,]+}})
457 #pragma omp target map(p->s.s.b[:2])
458 { p->a++; }
459
460 // Region 23
461 // CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
462 // CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
463 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
464 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
465 // CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
466 // CK24-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
467 // CK24-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]
468 // CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
469 // CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
470 // CK24-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
471
472 // CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
473 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
474 // CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
475 // CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
476 // CK24-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
477 // CK24-DAG: store i64 {{%.+}}, ptr [[S0]]
478 // CK24-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 4
479 // CK24-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1
480
481 // CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
482 // CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
483 // CK24-DAG: store ptr [[VAR0]], ptr [[BP1]]
484 // CK24-DAG: store ptr [[SEC0]], ptr [[P1]]
485
486 // CK24-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
487 // CK24-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
488 // CK24-DAG: store ptr [[SEC0]], ptr [[BP2]]
489 // CK24-DAG: store ptr [[SEC1:%.+]], ptr [[P2]]
490 // CK24-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[SEC11:%[^,]+]], i{{.+}} 0, i{{.+}} 0
491 // CK24-DAG: [[SEC11]] = getelementptr {{.*}}ptr [[SEC111:%[^,]+]], i{{.+}} 0, i{{.+}} 2
492 // CK24-DAG: [[SEC111]] = load ptr, ptr [[SEC1111:%[^,]+]],
493 // CK24-DAG: [[SEC1111]] = getelementptr {{.*}}ptr [[SEC11111:%[^,]+]], i{{.+}} 0, i{{.+}} 4
494 // CK24-DAG: [[SEC11111]] = getelementptr {{.*}}ptr [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 1
495
496 // CK24-DAG: [[VAR0]] = load ptr, ptr %{{.+}}
497 // CK24-DAG: [[VAR00]] = load ptr, ptr %{{.+}}
498 // CK24-DAG: [[VAR000]] = load ptr, ptr %{{.+}}
499
500 // CK24: call void [[CALL23:@.+]](ptr {{[^,]+}})
501 #pragma omp target map(p->s.p->b[:2])
502 { p->a++; }
503
504 // Region 24
505 // CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
506 // CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
507 // CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
508 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
509 // CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
510 // CK24-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
511 // CK24-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]
512 // CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
513 // CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
514 // CK24-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
515
516 // CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
517 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
518 // CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
519 // CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
520 // CK24-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
521 // CK24-DAG: store i64 {{%.+}}, ptr [[S0]]
522 // CK24-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 2
523
524 // CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
525 // CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
526 // CK24-DAG: store ptr [[SEC0]], ptr [[BP1]]
527 // CK24-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]
528 // CK24-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[SEC11:%[^,]+]], i{{.+}} 0, i{{.+}} 4
529 // CK24-DAG: [[SEC11]] = load ptr, ptr [[SEC111:%[^,]+]],
530 // CK24-DAG: [[SEC111]] = getelementptr {{.*}}ptr [[VAR000:%.+]], i{{.+}} 0, i{{.+}} 2
531
532 // CK24-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
533 // CK24-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
534 // CK24-DAG: store ptr [[SEC1]], ptr [[BP2]]
535 // CK24-DAG: store ptr [[SEC2:%.+]], ptr [[P2]]
536 // CK24-DAG: [[SEC2]] = getelementptr {{.*}}ptr [[SEC22:%[^,]+]], i{{.+}} 0, i{{.+}} 1
537 // CK24-DAG: [[SEC22]] = load ptr, ptr [[SEC222:%[^,]+]],
538 // CK24-DAG: [[SEC222]] = getelementptr {{.*}}ptr [[SEC2222:%[^,]+]], i{{.+}} 0, i{{.+}} 4
539 // CK24-DAG: [[SEC2222]] = load ptr, ptr [[SEC22222:%[^,]+]],
540 // CK24-DAG: [[SEC22222]] = getelementptr {{.*}}ptr [[VAR0000:%.+]], i{{.+}} 0, i{{.+}} 2
541
542 // CK24-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3
543 // CK24-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3
544 // CK24-DAG: store ptr [[SEC2]], ptr [[BP3]]
545 // CK24-DAG: store ptr [[SEC3:%.+]], ptr [[P3]]
546 // CK24-DAG: [[SEC3]] = getelementptr {{.*}}ptr [[SEC33:%[^,]+]], i{{.+}} 0, i{{.+}} 0
547 // CK24-DAG: [[SEC33]] = load ptr, ptr [[SEC333:%[^,]+]],
548 // CK24-DAG: [[SEC333]] = getelementptr {{.*}}ptr [[SEC3333:%[^,]+]], i{{.+}} 0, i{{.+}} 1
549 // CK24-DAG: [[SEC3333]] = load ptr, ptr [[SEC33333:%[^,]+]],
550 // CK24-DAG: [[SEC33333]] = getelementptr {{.*}}ptr [[SEC333333:%[^,]+]], i{{.+}} 0, i{{.+}} 4
551 // CK24-DAG: [[SEC333333]] = load ptr, ptr [[SEC3333333:%[^,]+]],
552 // CK24-DAG: [[SEC3333333]] = getelementptr {{.*}}ptr [[VAR00000:%.+]], i{{.+}} 0, i{{.+}} 2
553
554 // CK24-DAG: [[VAR0]] = load ptr, ptr %{{.+}}
555 // CK24-DAG: [[VAR00]] = load ptr, ptr %{{.+}}
556 // CK24-DAG: [[VAR000]] = load ptr, ptr %{{.+}}
557 // CK24-DAG: [[VAR0000]] = load ptr, ptr %{{.+}}
558 // CK24-DAG: [[VAR00000]] = load ptr, ptr %{{.+}}
559
560 // CK24: call void [[CALL24:@.+]](ptr {{[^,]+}})
561 #pragma omp target map(p->p->p->p->a)
562 { p->a++; }
563
564 return s.a;
565 }
566
567 // CK24: define {{.+}}[[CALL01]]
568 // CK24: define {{.+}}[[CALL13]]
569 // CK24: define {{.+}}[[CALL14]]
570 // CK24: define {{.+}}[[CALL15]]
571 // CK24: define {{.+}}[[CALL16]]
572 // CK24: define {{.+}}[[CALL17]]
573 // CK24: define {{.+}}[[CALL18]]
574 // CK24: define {{.+}}[[CALL19]]
575 // CK24: define {{.+}}[[CALL20]]
576 // CK24: define {{.+}}[[CALL21]]
577 // CK24: define {{.+}}[[CALL22]]
578 // CK24: define {{.+}}[[CALL23]]
579 // CK24: define {{.+}}[[CALL24]]
580 #endif // CK24
581 #endif
582