1 // expected-no-diagnostics
5 // This file is regex-heavy and takes a long time to execute the test. To speed
6 // testing up, test execution is split over multiple fimes. The RUN commands are
7 // in the corresponding .cpp files now. Do not add them here.
9 // SIMD-ONLY18-NOT: {{__kmpc|__tgt}}
12 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
13 // CK19: [[SIZE00:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
14 // CK19-USE: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i64] [i64 32]
15 // CK19-NOUSE: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer
17 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
18 // CK19: [[SIZE00n:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
19 // CK19-USE: [[MTYPE00n:@.+]] = private {{.*}}constant [1 x i64] [i64 32]
20 // CK19-NOUSE: [[MTYPE00n:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer
22 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
23 // CK19: [[SIZE01:@.+]] = private {{.*}}constant [1 x i64] [i64 400]
24 // CK19-USE: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 33]
25 // CK19-NOUSE: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 1]
27 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
28 // CK19: [[SIZE02:@.+]] = private {{.*}}constant [1 x i64] [i64 240]
29 // CK19-USE: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i64] [i64 34]
30 // CK19-NOUSE: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i64] [i64 2]
32 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
33 // CK19: [[SIZE03:@.+]] = private {{.*}}constant [1 x i64] [i64 240]
34 // CK19-USE: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
35 // CK19-NOUSE: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i64] [i64 3]
37 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
38 // CK19: [[SIZE04:@.+]] = private {{.*}}constant [1 x i64] [i64 400]
39 // CK19-USE: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i64] [i64 32]
40 // CK19-NOUSE: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer
42 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
43 // CK19: [[SIZE05:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
44 // CK19-USE: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i64] [i64 33]
45 // CK19-NOUSE: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i64] [i64 1]
47 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
48 // CK19-USE: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
49 // CK19-NOUSE: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i64] [i64 3]
51 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
52 // CK19-USE: [[MTYPE07:@.+]] = private {{.*}}constant [1 x i64] [i64 32]
53 // CK19-NOUSE: [[MTYPE07:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer
55 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
56 // CK19: [[SIZE08:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
57 // CK19-USE: [[MTYPE08:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
58 // CK19-NOUSE: [[MTYPE08:@.+]] = private {{.*}}constant [1 x i64] [i64 3]
60 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
61 // CK19: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}]
62 // CK19-USE: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 34]
63 // CK19-NOUSE: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 2]
65 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
66 // CK19: [[SIZE10:@.+]] = private {{.*}}constant [1 x i64] [i64 240]
67 // CK19-USE: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
68 // CK19-NOUSE: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i64] [i64 3]
70 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
71 // CK19: [[SIZE11:@.+]] = private {{.*}}constant [1 x i64] [i64 240]
72 // CK19-USE: [[MTYPE11:@.+]] = private {{.*}}constant [1 x i64] [i64 32]
73 // CK19-NOUSE: [[MTYPE11:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer
75 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
76 // CK19: [[SIZE12:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
77 // CK19-USE: [[MTYPE12:@.+]] = private {{.*}}constant [1 x i64] [i64 33]
78 // CK19-NOUSE: [[MTYPE12:@.+]] = private {{.*}}constant [1 x i64] [i64 1]
80 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
81 // CK19-USE: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i64] [i64 32]
82 // CK19-NOUSE: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer
84 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
85 // CK19-USE: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i64] [i64 33]
86 // CK19-NOUSE: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i64] [i64 1]
88 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
89 // CK19: [[SIZE15:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
90 // CK19-USE: [[MTYPE15:@.+]] = private {{.*}}constant [1 x i64] [i64 34]
91 // CK19-NOUSE: [[MTYPE15:@.+]] = private {{.*}}constant [1 x i64] [i64 2]
93 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
94 // CK19-USE: [[SIZE16:@.+]] = private {{.*}}constant [2 x i64] [i64 {{8|4}}, i64 0]
95 // CK19-USE: [[MTYPE16:@.+]] = private {{.*}}constant [2 x i64] [i64 800, i64 33]
96 // CK19-NOUSE: [[MTYPE16:@.+]] = private {{.*}}constant [1 x i64] [i64 1]
98 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
99 // CK19-USE: [[SIZE17:@.+]] = private {{.*}}constant [2 x i64] [i64 {{8|4}}, i64 240]
100 // CK19-USE: [[MTYPE17:@.+]] = private {{.*}}constant [2 x i64] [i64 800, i64 34]
101 // CK19-NOUSE: [[SIZE17:@.+]] = private {{.*}}constant [1 x i64] [i64 240]
102 // CK19-NOUSE: [[MTYPE17:@.+]] = private {{.*}}constant [1 x i64] [i64 2]
104 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
105 // CK19-USE: [[SIZE18:@.+]] = private {{.*}}constant [2 x i64] [i64 {{8|4}}, i64 240]
106 // CK19-USE: [[MTYPE18:@.+]] = private {{.*}}constant [2 x i64] [i64 800, i64 35]
107 // CK19-NOUSE: [[SIZE18:@.+]] = private {{.*}}constant [1 x i64] [i64 240]
108 // CK19-NOUSE: [[MTYPE18:@.+]] = private {{.*}}constant [1 x i64] [i64 3]
110 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
111 // CK19-USE: [[SIZE19:@.+]] = private {{.*}}constant [2 x i64] [i64 {{8|4}}, i64 0]
112 // CK19-USE: [[MTYPE19:@.+]] = private {{.*}}constant [2 x i64] [i64 800, i64 32]
113 // CK19-NOUSE: [[MTYPE19:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer
115 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
116 // CK19-USE: [[SIZE20:@.+]] = private {{.*}}constant [2 x i64] [i64 {{8|4}}, i64 4]
117 // CK19-USE: [[MTYPE20:@.+]] = private {{.*}}constant [2 x i64] [i64 800, i64 33]
118 // CK19-NOUSE: [[SIZE20:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
119 // CK19-NOUSE: [[MTYPE20:@.+]] = private {{.*}}constant [1 x i64] [i64 1]
121 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
122 // CK19-USE: [[SIZE21:@.+]] = private {{.*}}constant [2 x i64] [i64 {{8|4}}, i64 0]
123 // CK19-USE: [[MTYPE21:@.+]] = private {{.*}}constant [2 x i64] [i64 800, i64 35]
124 // CK19-NOUSE: [[MTYPE21:@.+]] = private {{.*}}constant [1 x i64] [i64 3]
126 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
127 // CK19-USE: [[SIZE22:@.+]] = private {{.*}}constant [2 x i64] [i64 {{8|4}}, i64 4]
128 // CK19-USE: [[MTYPE22:@.+]] = private {{.*}}constant [2 x i64] [i64 800, i64 35]
129 // CK19-NOUSE: [[SIZE22:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
130 // CK19-NOUSE: [[MTYPE22:@.+]] = private {{.*}}constant [1 x i64] [i64 3]
132 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
133 // CK19: [[SIZE23:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
134 // CK19-USE: [[MTYPE23:@.+]] = private {{.*}}constant [1 x i64] [i64 39]
135 // CK19-NOUSE: [[MTYPE23:@.+]] = private {{.*}}constant [1 x i64] [i64 7]
137 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
138 // CK19: [[SIZE24:@.+]] = private {{.*}}constant [1 x i64] [i64 480]
139 // CK19-USE: [[MTYPE24:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
140 // CK19-NOUSE: [[MTYPE24:@.+]] = private {{.*}}constant [1 x i64] [i64 3]
142 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
143 // CK19: [[SIZE25:@.+]] = private {{.*}}constant [1 x i64] [i64 16]
144 // CK19-USE: [[MTYPE25:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
145 // CK19-NOUSE: [[MTYPE25:@.+]] = private {{.*}}constant [1 x i64] [i64 3]
147 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
148 // CK19: [[SIZE26:@.+]] = private {{.*}}constant [1 x i64] [i64 24]
149 // CK19-USE: [[MTYPE26:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
150 // CK19-NOUSE: [[MTYPE26:@.+]] = private {{.*}}constant [1 x i64] [i64 3]
152 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
153 // CK19: [[SIZE27:@.+]] = private {{.*}}constant [1 x i64] [i64 4]
154 // CK19-USE: [[MTYPE27:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
155 // CK19-NOUSE: [[MTYPE27:@.+]] = private {{.*}}constant [1 x i64] [i64 3]
157 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
158 // CK19: [[SIZE28:@.+]] = private {{.*}}constant [3 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 16]
159 // CK19-USE: [[MTYPE28:@.+]] = private {{.*}}constant [3 x i64] [i64 35, i64 16, i64 19]
160 // CK19-NOUSE: [[MTYPE28:@.+]] = private {{.*}}constant [3 x i64] [i64 3, i64 16, i64 19]
162 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
163 // CK19: [[SIZE29:@.+]] = private {{.*}}constant [3 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 4]
164 // CK19-USE: [[MTYPE29:@.+]] = private {{.*}}constant [3 x i64] [i64 35, i64 16, i64 19]
165 // CK19-NOUSE: [[MTYPE29:@.+]] = private {{.*}}constant [3 x i64] [i64 3, i64 16, i64 19]
167 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
168 // CK19-USE: [[SIZE30:@.+]] = private {{.*}}constant [4 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 {{8|4}}, i64 0]
169 // CK19-USE: [[MTYPE30:@.+]] = private {{.*}}constant [4 x i64] [i64 800, i64 800, i64 800, i64 35]
170 // CK19-NOUSE: [[MTYPE30:@.+]] = private {{.*}}constant [1 x i64] [i64 3]
172 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
173 // CK19-USE: [[SIZE31:@.+]] = private {{.*}}constant [4 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 {{8|4}}, i64 40]
174 // CK19-USE: [[MTYPE31:@.+]] = private {{.*}}constant [4 x i64] [i64 800, i64 800, i64 800, i64 35]
175 // CK19-NOUSE: [[SIZE31:@.+]] = private {{.*}}constant [1 x i64] [i64 40]
176 // CK19-NOUSE: [[MTYPE31:@.+]] = private {{.*}}constant [1 x i64] [i64 3]
178 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
179 // CK19: [[SIZE32:@.+]] = private {{.*}}constant [1 x i64] [i64 13728]
180 // CK19-USE: [[MTYPE32:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
181 // CK19-NOUSE: [[MTYPE32:@.+]] = private {{.*}}constant [1 x i64] [i64 3]
183 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
184 // CK19: [[SIZE33:@.+]] = private {{.*}}constant [1 x i64] [i64 13728]
185 // CK19-USE: [[MTYPE33:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
186 // CK19-NOUSE: [[MTYPE33:@.+]] = private {{.*}}constant [1 x i64] [i64 3]
188 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
189 // CK19: [[SIZE34:@.+]] = private {{.*}}constant [1 x i64] [i64 13728]
190 // CK19-USE: [[MTYPE34:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
191 // CK19-NOUSE: [[MTYPE34:@.+]] = private {{.*}}constant [1 x i64] [i64 3]
193 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
194 // CK19-USE: [[MTYPE35:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
195 // CK19-NOUSE: [[MTYPE35:@.+]] = private {{.*}}constant [1 x i64] [i64 3]
197 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
198 // CK19: [[SIZE36:@.+]] = private {{.*}}constant [1 x i64] [i64 208]
199 // CK19-USE: [[MTYPE36:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
200 // CK19-NOUSE: [[MTYPE36:@.+]] = private {{.*}}constant [1 x i64] [i64 3]
202 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
203 // CK19-USE: [[SIZE37:@.+]] = private {{.*}}constant [3 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 0]
204 // CK19-USE: [[MTYPE37:@.+]] = private {{.*}}constant [3 x i64] [i64 800, i64 800, i64 35]
205 // CK19-NOUSE: [[MTYPE37:@.+]] = private {{.*}}constant [1 x i64] [i64 3]
207 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
208 // CK19-USE: [[SIZE38:@.+]] = private {{.*}}constant [3 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 0]
209 // CK19-USE: [[MTYPE38:@.+]] = private {{.*}}constant [3 x i64] [i64 800, i64 800, i64 35]
210 // CK19-NOUSE: [[MTYPE38:@.+]] = private {{.*}}constant [1 x i64] [i64 3]
212 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
213 // CK19-USE: [[SIZE39:@.+]] = private {{.*}}constant [3 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 0]
214 // CK19-USE: [[MTYPE39:@.+]] = private {{.*}}constant [3 x i64] [i64 800, i64 800, i64 35]
215 // CK19-NOUSE: [[MTYPE39:@.+]] = private {{.*}}constant [1 x i64] [i64 3]
217 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
218 // CK19-USE: [[SIZE40:@.+]] = private {{.*}}constant [3 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 0]
219 // CK19-USE: [[MTYPE40:@.+]] = private {{.*}}constant [3 x i64] [i64 800, i64 800, i64 35]
220 // CK19-NOUSE: [[MTYPE40:@.+]] = private {{.*}}constant [1 x i64] [i64 3]
222 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
223 // CK19-USE: [[SIZE41:@.+]] = private {{.*}}constant [3 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 208]
224 // CK19-USE: [[MTYPE41:@.+]] = private {{.*}}constant [3 x i64] [i64 800, i64 800, i64 35]
225 // CK19-NOUSE: [[SIZE41:@.+]] = private {{.*}}constant [1 x i64] [i64 208]
226 // CK19-NOUSE: [[MTYPE41:@.+]] = private {{.*}}constant [1 x i64] [i64 3]
228 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
229 // CK19: [[SIZE42:@.+]] = private {{.*}}constant [3 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 104]
230 // CK19-USE: [[MTYPE42:@.+]] = private {{.*}}constant [3 x i64] [i64 35, i64 16, i64 19]
231 // CK19-NOUSE: [[MTYPE42:@.+]] = private {{.*}}constant [3 x i64] [i64 3, i64 16, i64 19]
233 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
234 // CK19-USE: [[MTYPE43:@.+]] = private {{.*}}constant [1 x i64] [i64 35]
235 // CK19-NOUSE: [[MTYPE43:@.+]] = private {{.*}}constant [1 x i64] [i64 3]
237 // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0
238 // CK19: [[SIZE44:@.+]] = private {{.*}}constant [1 x i64] [i64 320]
239 // CK19-USE: [[MTYPE44:@.+]] = private {{.*}}constant [1 x i64] [i64 34]
240 // CK19-NOUSE: [[MTYPE44:@.+]] = private {{.*}}constant [1 x i64] [i64 2]
242 // CK19-LABEL: explicit_maps_single{{.*}}(
243 void explicit_maps_single (int ii){
248 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
249 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
250 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
251 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
252 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
253 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
254 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
256 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
257 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
258 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
259 // CK19-DAG: store ptr [[VAR0]], ptr [[P0]]
261 // CK19-USE: call void [[CALL00:@.+]](ptr {{[^,]+}})
262 // CK19-NOUSE: call void [[CALL00:@.+]]()
263 #pragma omp target map(alloc:a)
270 // Map of a scalar in nested region.
274 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
275 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
276 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
277 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
278 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
279 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
280 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
282 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
283 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
284 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
285 // CK19-DAG: store ptr [[VAR0]], ptr [[P0]]
287 // CK19-USE: call void [[CALL00n:@.+]](ptr {{[^,]+}})
288 // CK19-NOUSE: call void [[CALL00n:@.+]]()
289 #pragma omp target map(alloc:b)
301 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
302 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
303 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
304 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
305 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
306 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
307 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
309 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
310 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
311 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
312 // CK19-DAG: store ptr [[VAR0]], ptr [[P0]]
314 // CK19-USE: call void [[CALL01:@.+]](ptr {{[^,]+}})
315 // CK19-NOUSE: call void [[CALL01:@.+]]()
316 #pragma omp target map(to:arra)
324 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
325 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
326 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
327 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
328 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
329 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
330 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
332 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
333 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
334 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
335 // CK19-DAG: store ptr [[SEC0:%[^,]+]], ptr [[P0]]
336 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0, i{{.+}} 20
338 // CK19-USE: call void [[CALL02:@.+]](ptr {{[^,]+}})
339 // CK19-NOUSE: call void [[CALL02:@.+]]()
340 #pragma omp target map(from:arra[20:60])
348 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
349 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
350 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
351 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
352 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
353 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
354 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
356 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
357 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
358 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
359 // CK19-DAG: store ptr [[SEC0:%[^,]+]], ptr [[P0]]
360 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0, i{{.+}} 0
362 // CK19-USE: call void [[CALL03:@.+]](ptr {{[^,]+}})
363 // CK19-NOUSE: call void [[CALL03:@.+]]()
364 #pragma omp target map(tofrom:arra[:60])
372 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
373 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
374 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
375 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
376 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
377 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
378 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
380 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
381 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
382 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
383 // CK19-DAG: store ptr [[SEC0:%[^,]+]], ptr [[P0]]
384 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0, i{{.+}} 0
386 // CK19-USE: call void [[CALL04:@.+]](ptr {{[^,]+}})
387 // CK19-NOUSE: call void [[CALL04:@.+]]()
388 #pragma omp target map(alloc:arra[:])
396 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
397 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
398 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
399 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
400 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
401 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
402 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
404 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
405 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
406 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
407 // CK19-DAG: store ptr [[SEC0:%[^,]+]], ptr [[P0]]
408 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0, i{{.+}} 15
410 // CK19-USE: call void [[CALL05:@.+]](ptr {{[^,]+}})
411 // CK19-NOUSE: call void [[CALL05:@.+]]()
412 #pragma omp target map(to:arra[15])
420 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
421 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
422 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
423 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
424 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
425 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
426 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]
427 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
428 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
429 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
431 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
432 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
433 // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
434 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
435 // CK19-DAG: store ptr [[SEC0:%[^,]+]], ptr [[P0]]
436 // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], ptr [[S0]]
437 // CK19-DAG: [[CSVAL0]] = {{mul nuw i.+ %.*, 4|sext i32 .+ to i64}}
438 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0, i{{.+}} %{{.*}}
440 // CK19-USE: call void [[CALL06:@.+]](ptr {{[^,]+}})
441 // CK19-NOUSE: call void [[CALL06:@.+]]()
442 #pragma omp target map(tofrom:arra[ii:ii+23])
450 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
451 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
452 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
453 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
454 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
455 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
456 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]
457 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
458 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
459 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
461 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
462 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
463 // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
464 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
465 // CK19-DAG: store ptr [[SEC0:%[^,]+]], ptr [[P0]]
466 // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], ptr [[S0]]
467 // CK19-DAG: [[CSVAL0]] = {{mul nuw i.+ %.*, 4|sext i32 .+ to i64}}
468 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0, i{{.+}} 0
470 // CK19-USE: call void [[CALL07:@.+]](ptr {{[^,]+}})
471 // CK19-NOUSE: call void [[CALL07:@.+]]()
472 #pragma omp target map(alloc:arra[:ii])
480 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
481 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
482 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
483 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
484 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
485 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
486 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
488 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
489 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
490 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
491 // CK19-DAG: store ptr [[SEC0:%[^,]+]], ptr [[P0]]
492 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0, i{{.+}} %{{.*}}
494 // CK19-USE: call void [[CALL08:@.+]](ptr {{[^,]+}})
495 // CK19-NOUSE: call void [[CALL08:@.+]]()
496 #pragma omp target map(tofrom:arra[ii])
507 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
508 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
509 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
510 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
511 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
512 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
513 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
515 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
516 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
517 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
518 // CK19-DAG: store ptr [[VAR0]], ptr [[P0]]
520 // CK19-USE: call void [[CALL09:@.+]](ptr {{[^,]+}})
521 // CK19-NOUSE: call void [[CALL09:@.+]]()
522 #pragma omp target map(from:pa)
530 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
531 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
532 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
533 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
534 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
535 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
536 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
538 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
539 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
540 // CK19-DAG: store ptr [[RVAR0:%.+]], ptr [[BP0]]
541 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
542 // CK19-DAG: [[RVAR0]] = load ptr, ptr [[VAR0:%[^,]+]]
543 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[RVAR00:%.+]], i{{.+}} 20
544 // CK19-DAG: [[RVAR00]] = load ptr, ptr [[VAR0]]
546 // CK19-USE: call void [[CALL10:@.+]](ptr {{[^,]+}})
547 // CK19-NOUSE: call void [[CALL10:@.+]]()
548 #pragma omp target map(tofrom:pa[20:60])
556 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
557 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
558 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
559 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
560 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
561 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
562 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
564 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
565 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
566 // CK19-DAG: store ptr [[RVAR0:%.+]], ptr [[BP0]]
567 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
568 // CK19-DAG: [[RVAR0]] = load ptr, ptr [[VAR0:%[^,]+]]
569 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[RVAR00:%.+]], i{{.+}} 0
570 // CK19-DAG: [[RVAR00]] = load ptr, ptr [[VAR0]]
572 // CK19-USE: call void [[CALL11:@.+]](ptr {{[^,]+}})
573 // CK19-NOUSE: call void [[CALL11:@.+]]()
574 #pragma omp target map(alloc:pa[:60])
582 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
583 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
584 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
585 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
586 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
587 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
588 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
590 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
591 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
592 // CK19-DAG: store ptr [[RVAR0:%.+]], ptr [[BP0]]
593 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
594 // CK19-DAG: [[RVAR0]] = load ptr, ptr [[VAR0:%[^,]+]]
595 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[RVAR00:%.+]], i{{.+}} 15
596 // CK19-DAG: [[RVAR00]] = load ptr, ptr [[VAR0]]
598 // CK19-USE: call void [[CALL12:@.+]](ptr {{[^,]+}})
599 // CK19-NOUSE: call void [[CALL12:@.+]]()
600 #pragma omp target map(to:pa[15])
608 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
609 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
610 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
611 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
612 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
613 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
614 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]
615 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
616 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
617 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
619 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
620 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
621 // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
622 // CK19-DAG: store ptr [[RVAR0:%.+]], ptr [[BP0]]
623 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
624 // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], ptr [[S0]]
625 // CK19-DAG: [[CSVAL0]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}}
626 // CK19-DAG: [[RVAR0]] = load ptr, ptr [[VAR0:%[^,]+]]
627 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[RVAR00:%.+]], i{{.+}} %{{.*}}
628 // CK19-DAG: [[RVAR00]] = load ptr, ptr [[VAR0]]
630 // CK19-USE: call void [[CALL13:@.+]](ptr {{[^,]+}})
631 // CK19-NOUSE: call void [[CALL13:@.+]]()
632 #pragma omp target map(alloc:pa[ii-23:ii])
640 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
641 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
642 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
643 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
644 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
645 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
646 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]
647 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
648 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
649 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
651 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
652 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
653 // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
654 // CK19-DAG: store ptr [[RVAR0:%.+]], ptr [[BP0]]
655 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
656 // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], ptr [[S0]]
657 // CK19-DAG: [[CSVAL0]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}}
658 // CK19-DAG: [[RVAR0]] = load ptr, ptr [[VAR0:%[^,]+]]
659 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[RVAR00:%.+]], i{{.+}} 0
660 // CK19-DAG: [[RVAR00]] = load ptr, ptr [[VAR0]]
662 // CK19-USE: call void [[CALL14:@.+]](ptr {{[^,]+}})
663 // CK19-NOUSE: call void [[CALL14:@.+]]()
664 #pragma omp target map(to:pa[:ii])
672 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
673 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
674 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
675 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
676 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
677 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
678 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
680 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
681 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
682 // CK19-DAG: store ptr [[RVAR0:%.+]], ptr [[BP0]]
683 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
684 // CK19-DAG: [[RVAR0]] = load ptr, ptr [[VAR0:%[^,]+]]
685 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[RVAR00:%.+]], i{{.+}} %{{.*}}
686 // CK19-DAG: [[RVAR00]] = load ptr, ptr [[VAR0]]
688 // CK19-USE: call void [[CALL15:@.+]](ptr {{[^,]+}})
689 // CK19-NOUSE: call void [[CALL15:@.+]]()
690 #pragma omp target map(from:pa[ii+12])
697 // Map of a variable-size array.
701 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
702 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
703 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
704 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
705 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
706 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
707 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]
708 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
709 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
710 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
712 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
713 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
714 // CK19-USE-DAG: store i[[Z:64|32]] {{%.+}}, ptr [[BP0]]
715 // CK19-USE-DAG: store i[[Z]] {{%.+}}, ptr [[P0]]
717 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
718 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
719 // CK19-USE-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
720 // CK19-USE-DAG: store ptr [[VAR1:%.+]], ptr [[BP1]]
721 // CK19-USE-DAG: store ptr [[VAR1]], ptr [[P1]]
722 // CK19-USE-DAG: store i{{.+}} [[CSVAL1:%[^,]+]], ptr [[S1]]
723 // CK19-USE-DAG: [[CSVAL1]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}}
725 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
726 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
727 // CK19-NOUSE-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
728 // CK19-NOUSE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
729 // CK19-NOUSE-DAG: store ptr [[VAR0]], ptr [[P0]]
730 // CK19-NOUSE-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], ptr [[S0]]
731 // CK19-NOUSE-DAG: [[CSVAL0]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}}
733 // CK19-USE: call void [[CALL16:@.+]](i{{.+}} {{[^,]+}}, ptr {{[^,]+}})
734 // CK19-NOUSE: call void [[CALL16:@.+]]()
735 #pragma omp target map(to:va)
743 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
744 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
745 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
746 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
747 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
748 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
749 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
751 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
752 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
753 // CK19-USE-DAG: store i[[Z]] {{%.+}}, ptr [[BP0]]
754 // CK19-USE-DAG: store i[[Z]] {{%.+}}, ptr [[P0]]
756 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
757 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
758 // CK19-USE-DAG: store ptr [[VAR1:%.+]], ptr [[BP1]]
759 // CK19-USE-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]
760 // CK19-USE-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[VAR1]], i{{.+}} 20
762 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
763 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
764 // CK19-NOUSE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
765 // CK19-NOUSE-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
766 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 20
768 // CK19-USE: call void [[CALL17:@.+]](i{{.+}} {{[^,]+}}, ptr {{[^,]+}})
769 // CK19-NOUSE: call void [[CALL17:@.+]]()
770 #pragma omp target map(from:va[20:60])
778 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
779 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
780 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
781 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
782 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
783 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
784 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
786 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
787 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
788 // CK19-USE-DAG: store i[[Z]] {{%.+}}, ptr [[BP0]]
789 // CK19-USE-DAG: store i[[Z]] {{%.+}}, ptr [[P0]]
791 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
792 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
793 // CK19-USE-DAG: store ptr [[VAR1:%.+]], ptr [[BP1]]
794 // CK19-USE-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]
795 // CK19-USE-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[VAR1]], i{{.+}} 0
797 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
798 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
799 // CK19-NOUSE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
800 // CK19-NOUSE-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
801 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0
803 // CK19-USE: call void [[CALL18:@.+]](i{{.+}} {{[^,]+}}, ptr {{[^,]+}})
804 // CK19-NOUSE: call void [[CALL18:@.+]]()
805 #pragma omp target map(tofrom:va[:60])
813 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
814 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
815 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
816 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
817 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
818 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
819 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]
820 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
821 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
822 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
824 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
825 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
826 // CK19-USE-DAG: store i[[Z]] {{%.+}}, ptr [[BP0]]
827 // CK19-USE-DAG: store i[[Z]] {{%.+}}, ptr [[P0]]
829 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
830 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
831 // CK19-USE-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
832 // CK19-USE-DAG: store ptr [[VAR1:%.+]], ptr [[BP1]]
833 // CK19-USE-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]
834 // CK19-USE-DAG: store i{{.+}} [[CSVAL1:%[^,]+]], ptr [[S1]]
835 // CK19-USE-DAG: [[CSVAL1]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}}
836 // CK19-USE-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[VAR1]], i{{.+}} 0
838 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
839 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
840 // CK19-NOUSE-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
841 // CK19-NOUSE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
842 // CK19-NOUSE-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
843 // CK19-NOUSE-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], ptr [[S0]]
844 // CK19-NOUSE-DAG: [[CSVAL0]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}}
845 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0
847 // CK19-USE: call void [[CALL19:@.+]](i{{.+}} {{[^,]+}}, ptr {{[^,]+}})
848 // CK19-NOUSE: call void [[CALL19:@.+]]()
849 #pragma omp target map(alloc:va[:])
857 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
858 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
859 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
860 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
861 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
862 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
863 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
865 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
866 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
867 // CK19-USE-DAG: store i[[Z]] {{%.+}}, ptr [[BP0]]
868 // CK19-USE-DAG: store i[[Z]] {{%.+}}, ptr [[P0]]
870 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
871 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
872 // CK19-USE-DAG: store ptr [[VAR1:%.+]], ptr [[BP1]]
873 // CK19-USE-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]
874 // CK19-USE-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[VAR1]], i{{.+}} 15
876 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
877 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
878 // CK19-NOUSE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
879 // CK19-NOUSE-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
880 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 15
882 // CK19-USE: call void [[CALL20:@.+]](i{{.+}} {{[^,]+}}, ptr {{[^,]+}})
883 // CK19-NOUSE: call void [[CALL20:@.+]]()
884 #pragma omp target map(to:va[15])
892 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
893 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
894 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
895 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
896 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
897 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
898 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]
899 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
900 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
901 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
903 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
904 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
905 // CK19-USE-DAG: store i[[Z]] {{%.+}}, ptr [[BP0]]
906 // CK19-USE-DAG: store i[[Z]] {{%.+}}, ptr [[P0]]
908 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
909 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
910 // CK19-USE-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1
911 // CK19-USE-DAG: store ptr [[VAR1:%.+]], ptr [[BP1]]
912 // CK19-USE-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]
913 // CK19-USE-DAG: store i{{.+}} [[CSVAL1:%[^,]+]], ptr [[S1]]
914 // CK19-USE-DAG: [[CSVAL1]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}}
915 // CK19-USE-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[VAR1]], i{{.+}} %{{.+}}
917 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
918 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
919 // CK19-NOUSE-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
920 // CK19-NOUSE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
921 // CK19-NOUSE-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
922 // CK19-NOUSE-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], ptr [[S0]]
923 // CK19-NOUSE-DAG: [[CSVAL0]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}}
924 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} %{{.+}}
926 // CK19-USE: call void [[CALL21:@.+]](i{{.+}} {{[^,]+}}, ptr {{[^,]+}})
927 // CK19-NOUSE: call void [[CALL21:@.+]]()
928 #pragma omp target map(tofrom:va[ii:ii+23])
936 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
937 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
938 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
939 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
940 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
941 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
942 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
944 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
945 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
946 // CK19-USE-DAG: store i[[Z]] {{%.+}}, ptr [[BP0]]
947 // CK19-USE-DAG: store i[[Z]] {{%.+}}, ptr [[P0]]
949 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
950 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
951 // CK19-USE-DAG: store ptr [[VAR1:%.+]], ptr [[BP1]]
952 // CK19-USE-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]
953 // CK19-USE-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[VAR1]], i{{.+}} %{{.+}}
955 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
956 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
957 // CK19-NOUSE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
958 // CK19-NOUSE-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
959 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} %{{.+}}
961 // CK19-USE: call void [[CALL22:@.+]](i{{.+}} {{[^,]+}}, ptr {{[^,]+}})
962 // CK19-NOUSE: call void [[CALL22:@.+]]()
963 #pragma omp target map(tofrom:va[ii])
972 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
973 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
974 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
975 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
976 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
977 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
978 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
980 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
981 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
982 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
983 // CK19-DAG: store ptr [[VAR0]], ptr [[P0]]
985 // CK19-USE: call void [[CALL23:@.+]](ptr {{[^,]+}})
986 // CK19-NOUSE: call void [[CALL23:@.+]]()
987 #pragma omp target map(always, tofrom: a)
994 // Multidimensional arrays.
999 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
1000 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
1001 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
1002 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
1003 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
1004 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
1005 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
1007 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1008 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1009 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
1010 // CK19-DAG: store ptr [[VAR0]], ptr [[P0]]
1012 // CK19-USE: call void [[CALL24:@.+]](ptr {{[^,]+}})
1013 // CK19-NOUSE: call void [[CALL24:@.+]]()
1014 #pragma omp target map(tofrom: marr)
1022 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
1023 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
1024 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
1025 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
1026 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
1027 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
1028 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
1030 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1031 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1032 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
1033 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
1034 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:[^,]+]], i{{.+}} 0, i{{.+}} 2
1035 // CK19-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[SEC000:[^,]+]], i{{.+}} 0, i{{.+}} 2
1036 // CK19-DAG: [[SEC000]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0, i{{.+}} 1
1038 // CK19-USE: call void [[CALL25:@.+]](ptr {{[^,]+}})
1039 // CK19-NOUSE: call void [[CALL25:@.+]]()
1040 #pragma omp target map(tofrom: marr[1][2][2:4])
1048 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
1049 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
1050 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
1051 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
1052 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
1053 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
1054 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
1056 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1057 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1058 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
1059 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
1060 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:[^,]+]], i{{.+}} 0, i{{.+}} 0
1061 // CK19-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[SEC000:[^,]+]], i{{.+}} 0, i{{.+}} 2
1062 // CK19-DAG: [[SEC000]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0, i{{.+}} 1
1064 // CK19-USE: call void [[CALL26:@.+]](ptr {{[^,]+}})
1065 // CK19-NOUSE: call void [[CALL26:@.+]]()
1066 #pragma omp target map(tofrom: marr[1][2][:])
1074 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
1075 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
1076 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
1077 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
1078 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
1079 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
1080 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
1082 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1083 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1084 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
1085 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
1086 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:[^,]+]], i{{.+}} 0, i{{.+}} 3
1087 // CK19-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[SEC000:[^,]+]], i{{.+}} 0, i{{.+}} 2
1088 // CK19-DAG: [[SEC000]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0, i{{.+}} 1
1090 // CK19-USE: call void [[CALL27:@.+]](ptr {{[^,]+}})
1091 // CK19-NOUSE: call void [[CALL27:@.+]]()
1092 #pragma omp target map(tofrom: marr[1][2][3])
1100 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
1101 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
1102 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
1103 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
1104 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
1105 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
1106 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
1108 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1109 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1110 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
1111 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
1112 // CK19-DAG: [[VAR0]] = load ptr, ptr [[PTR:%[^,]+]],
1113 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:[^,]+]], i{{.+}} 1
1114 // CK19-DAG: [[SEC00]] = load ptr, ptr [[PTR]],
1116 // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
1117 // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
1118 // CK19-DAG: store ptr [[SEC0]], ptr [[BP1]]
1119 // CK19-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]
1120 // CK19-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[SEC11:[^,]+]], i{{.+}} 2
1121 // CK19-DAG: [[SEC11]] = load ptr, ptr [[SEC111:%[^,]+]],
1122 // CK19-DAG: [[SEC111]] = getelementptr {{.*}}ptr [[SEC1111:[^,]+]], i{{.+}} 1
1123 // CK19-DAG: [[SEC1111]] = load ptr, ptr [[PTR]],
1125 // CK19-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
1126 // CK19-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
1127 // CK19-DAG: store ptr [[SEC1]], ptr [[BP2]]
1128 // CK19-DAG: store ptr [[SEC2:%.+]], ptr [[P2]]
1129 // CK19-DAG: [[SEC2]] = getelementptr {{.*}}ptr [[SEC22:[^,]+]], i{{.+}} 2
1130 // CK19-DAG: [[SEC22]] = load ptr, ptr [[SEC222:%[^,]+]],
1131 // CK19-DAG: [[SEC222]] = getelementptr {{.*}}ptr [[SEC2222:[^,]+]], i{{.+}} 2
1132 // CK19-DAG: [[SEC2222]] = load ptr, ptr [[SEC22222:%[^,]+]],
1133 // CK19-DAG: [[SEC22222]] = getelementptr {{.*}}ptr [[SEC222222:[^,]+]], i{{.+}} 1
1134 // CK19-DAG: [[SEC222222]] = load ptr, ptr [[PTR]],
1136 // CK19-USE: call void [[CALL28:@.+]](ptr {{[^,]+}})
1137 // CK19-NOUSE: call void [[CALL28:@.+]]()
1138 #pragma omp target map(tofrom: mptr[1][2][2:4])
1146 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
1147 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
1148 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
1149 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
1150 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
1151 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
1152 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
1154 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1155 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1156 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
1157 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
1158 // CK19-DAG: [[VAR0]] = load ptr, ptr [[PTR:%[^,]+]],
1159 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:[^,]+]], i{{.+}} 1
1160 // CK19-DAG: [[SEC00]] = load ptr, ptr [[PTR]],
1162 // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
1163 // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
1164 // CK19-DAG: store ptr [[SEC0]], ptr [[BP1]]
1165 // CK19-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]
1166 // CK19-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[SEC11:[^,]+]], i{{.+}} 2
1167 // CK19-DAG: [[SEC11]] = load ptr, ptr [[SEC111:%[^,]+]],
1168 // CK19-DAG: [[SEC111]] = getelementptr {{.*}}ptr [[SEC1111:[^,]+]], i{{.+}} 1
1169 // CK19-DAG: [[SEC1111]] = load ptr, ptr [[PTR]],
1171 // CK19-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
1172 // CK19-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
1173 // CK19-DAG: store ptr [[SEC1]], ptr [[BP2]]
1174 // CK19-DAG: store ptr [[SEC2:%.+]], ptr [[P2]]
1175 // CK19-DAG: [[SEC2]] = getelementptr {{.*}}ptr [[SEC22:[^,]+]], i{{.+}} 3
1176 // CK19-DAG: [[SEC22]] = load ptr, ptr [[SEC222:%[^,]+]],
1177 // CK19-DAG: [[SEC222]] = getelementptr {{.*}}ptr [[SEC2222:[^,]+]], i{{.+}} 2
1178 // CK19-DAG: [[SEC2222]] = load ptr, ptr [[SEC22222:%[^,]+]],
1179 // CK19-DAG: [[SEC22222]] = getelementptr {{.*}}ptr [[SEC222222:[^,]+]], i{{.+}} 1
1180 // CK19-DAG: [[SEC222222]] = load ptr, ptr [[PTR]],
1182 // CK19-USE: call void [[CALL29:@.+]](ptr {{[^,]+}})
1183 // CK19-NOUSE: call void [[CALL29:@.+]]()
1184 #pragma omp target map(tofrom: mptr[1][2][3])
1191 // Multidimensional VLA.
1192 double mva[23][ii][ii+5];
1195 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
1196 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
1197 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
1198 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
1199 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
1200 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
1201 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]
1202 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
1203 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
1204 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
1206 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1207 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1208 // CK19-USE-DAG: store i[[Z]] 23, ptr [[BP0]]
1209 // CK19-USE-DAG: store i[[Z]] 23, ptr [[P0]]
1211 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
1212 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
1213 // CK19-USE-DAG: store i[[Z]] [[VAR1:%.+]], ptr [[BP1]]
1214 // CK19-USE-DAG: store i[[Z]] [[VAR11:%.+]], ptr [[P1]]
1215 // CK19-64-USE-DAG: [[VAR1]] = zext i32 %{{[^,]+}} to i64
1216 // CK19-64-USE-DAG: [[VAR11]] = zext i32 %{{[^,]+}} to i64
1218 // CK19-USE-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
1219 // CK19-USE-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
1220 // CK19-USE-DAG: store i[[Z]] [[VAR2:%.+]], ptr [[BP2]]
1221 // CK19-USE-DAG: store i[[Z]] [[VAR22:%.+]], ptr [[P2]]
1222 // CK19-64-USE-DAG: [[VAR2]] = zext i32 %{{[^,]+}} to i64
1223 // CK19-64-USE-DAG: [[VAR22]] = zext i32 %{{[^,]+}} to i64
1225 // CK19-USE-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3
1226 // CK19-USE-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3
1227 // CK19-USE-DAG: [[S3:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 3
1228 // CK19-USE-DAG: store ptr [[VAR3:%.+]], ptr [[BP3]]
1229 // CK19-USE-DAG: store ptr [[VAR3]], ptr [[P3]]
1230 // CK19-USE-DAG: store i64 [[CSVAL3:%[^,]+]], ptr [[S3]]
1231 // CK19-USE-DAG: [[CSVAL3]] = {{mul nuw i64 %[^,]+, 8|sext i32 .+ to i64}}
1233 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1234 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1235 // CK19-NOUSE-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
1236 // CK19-NOUSE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
1237 // CK19-NOUSE-DAG: store ptr [[VAR0]], ptr [[P0]]
1238 // CK19-NOUSE-DAG: store i64 [[CSVAL0:%[^,]+]], ptr [[S0]]
1239 // CK19-NOUSE-DAG: [[CSVAL0]] = {{mul nuw i64 %[^,]+, 8|sext i32 .+ to i64}}
1241 // CK19-USE: call void [[CALL30:@.+]](i[[Z]] 23, i[[Z]] %{{[^,]+}}, i[[Z]] %{{[^,]+}}, ptr %{{[^,]+}})
1242 // CK19-NOUSE: call void [[CALL30:@.+]]()
1243 #pragma omp target map(tofrom: mva)
1251 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
1252 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
1253 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
1254 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
1255 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
1256 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
1257 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
1259 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1260 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1261 // CK19-USE-DAG: store i[[Z]] 23, ptr [[BP0]]
1262 // CK19-USE-DAG: store i[[Z]] 23, ptr [[P0]]
1264 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
1265 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
1266 // CK19-USE-DAG: store i[[Z]] [[VAR1:%.+]], ptr [[BP1]]
1267 // CK19-USE-DAG: store i[[Z]] [[VAR11:%.+]], ptr [[P1]]
1269 // CK19-USE-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
1270 // CK19-USE-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
1271 // CK19-USE-DAG: store i[[Z]] [[VAR2:%.+]], ptr [[BP2]]
1272 // CK19-USE-DAG: store i[[Z]] [[VAR22:%.+]], ptr [[P2]]
1274 // CK19-USE-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3
1275 // CK19-USE-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3
1276 // CK19-USE-DAG: store ptr [[VAR3:%.+]], ptr [[BP3]]
1277 // CK19-USE-DAG: store ptr [[SEC3:%.+]], ptr [[P3]]
1278 // CK19-USE-DAG: [[SEC3]] = getelementptr {{.*}}ptr [[SEC33:%.+]], i[[Z]] 0
1279 // CK19-USE-DAG: [[SEC33]] = getelementptr {{.*}}ptr [[SEC333:%.+]], i[[Z]] [[IDX3:%.+]]
1280 // CK19-USE-DAG: [[IDX3]] = mul nsw i[[Z]] %{{[^,]+}}, %{{[^,]+}}
1281 // CK19-USE-DAG: [[SEC333]] = getelementptr {{.*}}ptr [[VAR3]], i[[Z]] [[IDX33:%.+]]
1282 // CK19-USE-DAG: [[IDX33]] = mul nsw i[[Z]] 1, %{{[^,]+}}
1284 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1285 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1286 // CK19-NOUSE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
1287 // CK19-NOUSE-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
1288 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:%.+]], i[[Z:64|32]] 0
1289 // CK19-NOUSE-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[SEC000:%.+]], i[[Z]] [[IDX0:%.+]]
1290 // CK19-NOUSE-DAG: [[IDX0]] = mul nsw i[[Z]] %{{[^,]+}}, %{{[^,]+}}
1291 // CK19-NOUSE-DAG: [[SEC000]] = getelementptr {{.*}}ptr [[VAR0]], i[[Z]] [[IDX00:%.+]]
1292 // CK19-NOUSE-DAG: [[IDX00]] = mul nsw i[[Z]] 1, %{{[^,]+}}
1294 // CK19-USE: call void [[CALL31:@.+]](i[[Z]] 23, i[[Z]] %{{[^,]+}}, i[[Z]] %{{[^,]+}}, ptr %{{[^,]+}})
1295 // CK19-NOUSE: call void [[CALL31:@.+]]()
1296 #pragma omp target map(tofrom: mva[1][ii-2][:5])
1303 // Multidimensional array sections.
1304 double marras[11][12][13];
1305 double mvlaas[11][ii][13];
1309 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
1310 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
1311 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
1312 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
1313 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
1314 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
1315 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
1317 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1318 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1319 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
1320 // CK19-DAG: store ptr [[VAR0]], ptr [[P0]]
1322 // CK19-USE: call void [[CALL32:@.+]](ptr {{[^,]+}})
1323 // CK19-NOUSE: call void [[CALL32:@.+]]()
1324 #pragma omp target map(marras)
1332 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
1333 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
1334 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
1335 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
1336 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
1337 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
1338 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
1340 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1341 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1342 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
1343 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
1344 // CK19-DAG: [[SEC0]] = getelementptr {{.+}}ptr [[VAR0]], i[[Z]] 0, i[[Z]] 0
1346 // CK19-USE: call void [[CALL33:@.+]](ptr {{[^,]+}})
1347 // CK19-NOUSE: call void [[CALL33:@.+]]()
1348 #pragma omp target map(marras[:])
1356 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
1357 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
1358 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
1359 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
1360 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
1361 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
1362 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
1364 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1365 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1366 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
1367 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
1368 // CK19-DAG: [[SEC0]] = getelementptr {{.+}}ptr [[VAR0]], i[[Z]] 0, i[[Z]] 0
1370 // CK19-USE: call void [[CALL34:@.+]](ptr {{[^,]+}})
1371 // CK19-NOUSE: call void [[CALL34:@.+]]()
1372 #pragma omp target map(marras[:][:][:])
1380 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
1381 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
1382 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
1383 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
1384 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
1385 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
1386 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]
1387 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
1388 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
1389 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
1391 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1392 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1393 // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
1395 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
1396 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
1397 // CK19-DAG: store i64 [[CSVAL0:%[^,]+]], ptr [[S0]]
1398 // CK19-DAG: [[SEC0]] = getelementptr {{.+}}ptr [[SEC00:%[^,]+]], i[[Z]] 0, i[[Z]] 0
1399 // CK19-DAG: [[SEC00]] = getelementptr {{.+}}ptr [[VAR0]], i[[Z]] 0, i[[Z]] 1
1400 // CK19-DAG: [[CSVAL0]] = {{mul nuw i64 %[^,]+, 104|sext i32 .+ to i64}}
1402 // CK19-USE: call void [[CALL35:@.+]](ptr {{[^,]+}})
1403 // CK19-NOUSE: call void [[CALL35:@.+]]()
1404 #pragma omp target map(marras[1][:ii][:])
1412 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
1413 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
1414 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
1415 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
1416 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
1417 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
1418 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
1420 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1421 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1422 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
1423 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
1424 // CK19-DAG: [[SEC0]] = getelementptr {{.+}}ptr [[SEC00:%[^,]+]], i{{.+}} 0
1425 // CK19-DAG: [[SEC00]] = getelementptr {{.+}}ptr [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 0
1426 // CK19-DAG: [[SEC000]] = getelementptr {{.+}}ptr [[VAR0]], i{{.+}} 0, i{{.+}} 0
1428 // CK19-USE: call void [[CALL36:@.+]](ptr {{[^,]+}})
1429 // CK19-NOUSE: call void [[CALL36:@.+]]()
1430 #pragma omp target map(marras[:1][:2][:13])
1438 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
1439 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
1440 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
1441 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
1442 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
1443 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
1444 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]
1445 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
1446 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
1447 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
1449 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1450 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1451 // CK19-USE-DAG: store i[[Z]] 11, ptr [[BP0]]
1452 // CK19-USE-DAG: store i[[Z]] 11, ptr [[P0]]
1454 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
1455 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
1456 // CK19-USE-DAG: store i[[Z]] [[VAR1:%.+]], ptr [[BP1]]
1457 // CK19-USE-DAG: store i[[Z]] [[VAR11:%.+]], ptr [[P1]]
1459 // CK19-USE-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
1460 // CK19-USE-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
1461 // CK19-USE-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
1462 // CK19-USE-DAG: store ptr [[VAR2:%.+]], ptr [[BP2]]
1463 // CK19-USE-DAG: store ptr [[VAR2]], ptr [[P2]]
1464 // CK19-USE-DAG: store i64 [[CSVAL2:%[^,]+]], ptr [[S2]]
1465 // CK19-USE-DAG: [[CSVAL2]] = {{mul nuw i64 %[^,]+, 104|sext i32 .+ to i64}}
1467 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1468 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1469 // CK19-NOUSE-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
1470 // CK19-NOUSE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
1471 // CK19-NOUSE-DAG: store ptr [[VAR0]], ptr [[P0]]
1472 // CK19-NOUSE-DAG: store i64 [[CSVAL0:%[^,]+]], ptr [[S0]]
1473 // CK19-NOUSE-DAG: [[CSVAL0]] = {{mul nuw i64 %[^,]+, 104|sext i32 .+ to i64}}
1475 // CK19-USE: call void [[CALL37:@.+]](i[[Z]] 11, i[[Z]] %{{[^,]+}}, ptr %{{[^,]+}})
1476 // CK19-NOUSE: call void [[CALL37:@.+]]()
1477 #pragma omp target map(mvlaas)
1485 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
1486 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
1487 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
1488 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
1489 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
1490 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
1491 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]
1492 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
1493 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
1494 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
1496 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1497 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1498 // CK19-USE-DAG: store i[[Z]] 11, ptr [[BP0]]
1499 // CK19-USE-DAG: store i[[Z]] 11, ptr [[P0]]
1501 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
1502 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
1503 // CK19-USE-DAG: store i[[Z]] [[VAR1:%.+]], ptr [[BP1]]
1504 // CK19-USE-DAG: store i[[Z]] [[VAR11:%.+]], ptr [[P1]]
1506 // CK19-USE-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
1507 // CK19-USE-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
1508 // CK19-USE-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
1509 // CK19-USE-DAG: store ptr [[VAR2:%.+]], ptr [[BP2]]
1510 // CK19-USE-DAG: store ptr [[SEC2:%.+]], ptr [[P2]]
1511 // CK19-USE-DAG: store i64 [[CSVAL2:%[^,]+]], ptr [[S2]]
1512 // CK19-USE-DAG: [[SEC2]] = getelementptr {{.+}}ptr [[VAR2]], i[[Z]] [[SEC22:%[^,]+]]
1513 // CK19-USE-DAG: [[SEC22]] = mul nsw i[[Z]] 0, %{{[^,]+}}
1514 // CK19-USE-DAG: [[CSVAL2]] = {{mul nuw i64 %[^,]+, 104|sext i32 .+ to i64}}
1516 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1517 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1518 // CK19-NOUSE-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
1519 // CK19-NOUSE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
1520 // CK19-NOUSE-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
1521 // CK19-NOUSE-DAG: store i64 [[CSVAL0:%[^,]+]], ptr [[S0]]
1522 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.+}}ptr [[VAR0]], i[[Z]] [[SEC00:%[^,]+]]
1523 // CK19-NOUSE-DAG: [[SEC00]] = mul nsw i[[Z]] 0, %{{[^,]+}}
1524 // CK19-NOUSE-DAG: [[CSVAL0]] = {{mul nuw i64 %[^,]+, 104|sext i32 .+ to i64}}
1526 // CK19-USE: call void [[CALL38:@.+]](i[[Z]] 11, i[[Z]] %{{[^,]+}}, ptr %{{[^,]+}})
1527 // CK19-NOUSE: call void [[CALL38:@.+]]()
1528 #pragma omp target map(mvlaas[:])
1536 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
1537 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
1538 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
1539 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
1540 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
1541 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
1542 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]
1543 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
1544 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
1545 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
1547 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1548 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1549 // CK19-USE-DAG: store i[[Z]] 11, ptr [[BP0]]
1550 // CK19-USE-DAG: store i[[Z]] 11, ptr [[P0]]
1552 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
1553 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
1554 // CK19-USE-DAG: store i[[Z]] [[VAR1:%.+]], ptr [[BP1]]
1555 // CK19-USE-DAG: store i[[Z]] [[VAR11:%.+]], ptr [[P1]]
1557 // CK19-USE-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
1558 // CK19-USE-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
1559 // CK19-USE-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
1560 // CK19-USE-DAG: store ptr [[VAR2:%.+]], ptr [[BP2]]
1561 // CK19-USE-DAG: store ptr [[SEC2:%.+]], ptr [[P2]]
1562 // CK19-USE-DAG: store i64 [[CSVAL2:%[^,]+]], ptr [[S2]]
1563 // CK19-USE-DAG: [[SEC2]] = getelementptr {{.+}}ptr [[VAR2]], i[[Z]] [[SEC22:%[^,]+]]
1564 // CK19-USE-DAG: [[SEC22]] = mul nsw i[[Z]] 0, %{{[^,]+}}
1565 // CK19-USE-DAG: [[CSVAL2]] = {{mul nuw i64 %[^,]+, 104|sext i32 .+ to i64}}
1567 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1568 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1569 // CK19-NOUSE-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
1570 // CK19-NOUSE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
1571 // CK19-NOUSE-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
1572 // CK19-NOUSE-DAG: store i64 [[CSVAL0:%[^,]+]], ptr [[S0]]
1573 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.+}}ptr [[VAR0]], i[[Z]] [[SEC00:%[^,]+]]
1574 // CK19-NOUSE-DAG: [[SEC00]] = mul nsw i[[Z]] 0, %{{[^,]+}}
1575 // CK19-NOUSE-DAG: [[CSVAL0]] = {{mul nuw i64 %[^,]+, 104|sext i32 .+ to i64}}
1577 // CK19-USE: call void [[CALL39:@.+]](i[[Z]] 11, i[[Z]] %{{[^,]+}}, ptr %{{[^,]+}})
1578 // CK19-NOUSE: call void [[CALL39:@.+]]()
1579 #pragma omp target map(mvlaas[:][:][:])
1587 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
1588 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
1589 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
1590 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
1591 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
1592 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
1593 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]
1594 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
1595 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
1596 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
1598 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1599 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1600 // CK19-USE-DAG: store i[[Z]] 11, ptr [[BP0]]
1601 // CK19-USE-DAG: store i[[Z]] 11, ptr [[P0]]
1603 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
1604 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
1605 // CK19-USE-DAG: store i[[Z]] [[VAR1:%.+]], ptr [[BP1]]
1606 // CK19-USE-DAG: store i[[Z]] [[VAR11:%.+]], ptr [[P1]]
1608 // CK19-USE-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
1609 // CK19-USE-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
1610 // CK19-USE-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2
1611 // CK19-USE-DAG: store ptr [[VAR2:%.+]], ptr [[BP2]]
1612 // CK19-USE-DAG: store ptr [[SEC2:%.+]], ptr [[P2]]
1613 // CK19-USE-DAG: store i64 [[CSVAL2:%[^,]+]], ptr [[S2]]
1614 // CK19-USE-DAG: [[SEC2]] = getelementptr {{.+}}ptr [[SEC22:%[^,]+]], i[[Z]] 0
1615 // CK19-USE-DAG: [[SEC22]] = getelementptr {{.+}}ptr [[VAR2]], i[[Z]] [[SEC222:%[^,]+]]
1616 // CK19-USE-DAG: [[SEC222]] = mul nsw i[[Z]] 1, %{{[^,]+}}
1618 // CK19-NOUSE-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
1619 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1620 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1621 // CK19-NOUSE-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
1622 // CK19-NOUSE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
1623 // CK19-NOUSE-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
1624 // CK19-NOUSE-DAG: store i64 [[CSVAL0:%[^,]+]], ptr [[S0]]
1625 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.+}}ptr [[SEC00:%[^,]+]], i[[Z]] 0
1626 // CK19-NOUSE-DAG: [[SEC00]] = getelementptr {{.+}}ptr [[VAR0]], i[[Z]] [[SEC000:%[^,]+]]
1627 // CK19-NOUSE-DAG: [[SEC000]] = mul nsw i[[Z]] 1, %{{[^,]+}}
1629 // CK19-USE: call void [[CALL40:@.+]](i[[Z]] 11, i[[Z]] %{{[^,]+}}, ptr %{{[^,]+}})
1630 // CK19-NOUSE: call void [[CALL40:@.+]]()
1631 #pragma omp target map(mvlaas[1][:ii][:])
1639 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
1640 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
1641 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
1642 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
1643 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
1644 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
1645 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
1647 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1648 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1649 // CK19-USE-DAG: store i[[Z]] 11, ptr [[BP0]]
1650 // CK19-USE-DAG: store i[[Z]] 11, ptr [[P0]]
1652 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
1653 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
1654 // CK19-USE-DAG: store i[[Z]] [[VAR1:%.+]], ptr [[BP1]]
1655 // CK19-USE-DAG: store i[[Z]] [[VAR11:%.+]], ptr [[P1]]
1657 // CK19-USE-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
1658 // CK19-USE-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
1659 // CK19-USE-DAG: store ptr [[VAR2:%.+]], ptr [[BP2]]
1660 // CK19-USE-DAG: store ptr [[SEC2:%.+]], ptr [[P2]]
1661 // CK19-USE-DAG: [[SEC2]] = getelementptr {{.+}}ptr [[SEC22:%[^,]+]], i[[Z]] 0
1662 // CK19-USE-DAG: [[SEC22]] = getelementptr {{.+}}ptr [[VAR2]], i[[Z]] [[SEC222:%[^,]+]]
1663 // CK19-USE-DAG: [[SEC222]] = mul nsw i[[Z]] 0, %{{[^,]+}}
1664 // CK19-USE-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
1666 // CK19-NO-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1667 // CK19-NO-USE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
1668 // CK19-NO-USE-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
1669 // CK19-NO-USE-DAG: [[SEC0]] = getelementptr {{.+}}ptr [[SEC00:%[^,]+]], i[[Z]] 0
1670 // CK19-NO-USE-DAG: [[SEC00]] = getelementptr {{.+}}ptr [[VAR0]], i[[Z]] [[SEC000:%[^,]+]]
1671 // CK19-NO-USE-DAG: [[SEC000]] = mul nsw i[[Z]] 0, %{{[^,]+}}
1673 // CK19-USE: call void [[CALL41:@.+]](i[[Z]] 11, i[[Z]] %{{[^,]+}}, ptr %{{[^,]+}})
1674 // CK19-NOUSE: call void [[CALL41:@.+]]()
1675 #pragma omp target map(mvlaas[:1][:2][:13])
1683 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
1684 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
1685 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
1686 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
1687 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
1688 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
1689 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
1691 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1692 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1693 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
1694 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
1695 // CK19-DAG: [[VAR0]] = load ptr, ptr [[PTR:%[^,]+]],
1696 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:[^,]+]], i{{.+}} 0
1697 // CK19-DAG: [[SEC00]] = load ptr, ptr [[PTR]],
1699 // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
1700 // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1
1701 // CK19-DAG: store ptr [[SEC0]], ptr [[BP1]]
1702 // CK19-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]
1703 // CK19-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[SEC11:[^,]+]], i{{.+}} 2
1704 // CK19-DAG: [[SEC11]] = load ptr, ptr [[SEC111:%[^,]+]],
1705 // CK19-DAG: [[SEC111]] = getelementptr {{.*}}ptr [[SEC1111:[^,]+]], i{{.+}} 0
1706 // CK19-DAG: [[SEC1111]] = load ptr, ptr [[PTR]],
1708 // CK19-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2
1709 // CK19-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2
1710 // CK19-DAG: store ptr [[SEC1]], ptr [[BP2]]
1711 // CK19-DAG: store ptr [[SEC2:%.+]], ptr [[P2]]
1712 // CK19-DAG: [[SEC2]] = getelementptr {{.*}}ptr [[SEC22:[^,]+]], i{{.+}} 0
1713 // CK19-DAG: [[SEC22]] = load ptr, ptr [[SEC222:%[^,]+]],
1714 // CK19-DAG: [[SEC222]] = getelementptr {{.*}}ptr [[SEC2222:[^,]+]], i{{.+}} 2
1715 // CK19-DAG: [[SEC2222]] = load ptr, ptr [[SEC22222:%[^,]+]],
1716 // CK19-DAG: [[SEC22222]] = getelementptr {{.*}}ptr [[SEC222222:[^,]+]], i{{.+}} 0
1717 // CK19-DAG: [[SEC222222]] = load ptr, ptr [[PTR]],
1719 // CK19-USE: call void [[CALL42:@.+]](ptr {{[^,]+}})
1720 // CK19-NOUSE: call void [[CALL42:@.+]]()
1721 #pragma omp target map(mptras[:1][2][:13])
1728 // Region 43 - the memory is not contiguous for this map - will map the whole last dimension.
1729 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
1730 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
1731 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
1732 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
1733 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
1734 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
1735 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]
1736 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
1737 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
1738 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
1740 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1741 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1742 // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
1744 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
1745 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]
1746 // CK19-DAG: store i64 [[CSVAL0:%[^,]+]], ptr [[S0]]
1747 // CK19-DAG: [[SEC0]] = getelementptr {{.+}}ptr [[SEC00:%[^,]+]], i[[Z]] 0, i[[Z]] 0
1748 // CK19-DAG: [[SEC00]] = getelementptr {{.+}}ptr [[VAR0]], i[[Z]] 0, i[[Z]] 1
1749 // CK19-DAG: [[CSVAL0]] = {{mul nuw i64 %[^,]+, 104|sext i32 .+ to i64}}
1751 // CK19-USE: call void [[CALL43:@.+]](ptr {{[^,]+}})
1752 // CK19-NOUSE: call void [[CALL43:@.+]]()
1753 #pragma omp target map(marras[1][:ii][1:])
1761 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
1762 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
1763 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
1764 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
1765 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
1766 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
1767 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
1769 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
1770 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
1771 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]
1772 // CK19-DAG: store ptr [[SEC0:%[^,]+]], ptr [[P0]]
1773 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0, i{{.+}} 20
1775 // CK19-USE: call void [[CALL44:@.+]](ptr {{[^,]+}})
1776 // CK19-NOUSE: call void [[CALL44:@.+]]()
1777 #pragma omp target map(from:arra[20:])
1786 // CK19: define {{.+}}[[CALL00]]
1787 // CK19: define {{.+}}[[CALL01]]
1788 // CK19: define {{.+}}[[CALL02]]
1789 // CK19: define {{.+}}[[CALL03]]
1790 // CK19: define {{.+}}[[CALL04]]
1791 // CK19: define {{.+}}[[CALL05]]
1792 // CK19: define {{.+}}[[CALL06]]
1793 // CK19: define {{.+}}[[CALL07]]
1794 // CK19: define {{.+}}[[CALL08]]
1795 // CK19: define {{.+}}[[CALL09]]
1796 // CK19: define {{.+}}[[CALL10]]
1797 // CK19: define {{.+}}[[CALL11]]
1798 // CK19: define {{.+}}[[CALL12]]
1799 // CK19: define {{.+}}[[CALL13]]
1800 // CK19: define {{.+}}[[CALL14]]
1801 // CK19: define {{.+}}[[CALL15]]
1802 // CK19: define {{.+}}[[CALL16]]
1803 // CK19: define {{.+}}[[CALL17]]
1804 // CK19: define {{.+}}[[CALL18]]
1805 // CK19: define {{.+}}[[CALL19]]
1806 // CK19: define {{.+}}[[CALL20]]
1807 // CK19: define {{.+}}[[CALL21]]
1808 // CK19: define {{.+}}[[CALL22]]
1809 // CK19: define {{.+}}[[CALL23]]
1810 // CK19: define {{.+}}[[CALL24]]
1811 // CK19: define {{.+}}[[CALL25]]
1812 // CK19: define {{.+}}[[CALL26]]
1813 // CK19: define {{.+}}[[CALL27]]
1814 // CK19: define {{.+}}[[CALL28]]
1815 // CK19: define {{.+}}[[CALL29]]
1816 // CK19: define {{.+}}[[CALL30]]
1817 // CK19: define {{.+}}[[CALL31]]
1818 // CK19: define {{.+}}[[CALL32]]
1819 // CK19: define {{.+}}[[CALL33]]
1820 // CK19: define {{.+}}[[CALL34]]
1821 // CK19: define {{.+}}[[CALL35]]
1822 // CK19: define {{.+}}[[CALL36]]
1823 // CK19: define {{.+}}[[CALL37]]
1824 // CK19: define {{.+}}[[CALL38]]
1825 // CK19: define {{.+}}[[CALL39]]
1826 // CK19: define {{.+}}[[CALL40]]
1827 // CK19: define {{.+}}[[CALL41]]
1828 // CK19: define {{.+}}[[CALL42]]
1829 // CK19: define {{.+}}[[CALL43]]
1830 // CK19: define {{.+}}[[CALL44]]
1833 #endif // HEADER_INC