1 ; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: --function-signature
2 ; RUN: opt -S -passes=openmp-opt-cgscc -aa-pipeline=basic-aa -openmp-hide-memory-transfer-latency < %s | FileCheck %s
3 target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128"
5 ; CHECK: %struct.__tgt_async_info = type { ptr }
7 %struct.ident_t = type { i32, i32, i32, i32, ptr }
8 %struct.__tgt_offload_entry = type { ptr, ptr, i64, i32, i32 }
10 @.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 35]
11 @.__omp_offloading_heavyComputation1.region_id = weak constant i8 0
12 @.offload_sizes.1 = private unnamed_addr constant [1 x i64] [i64 8]
13 @.offload_maptypes.2 = private unnamed_addr constant [1 x i64] [i64 800]
15 @.__omp_offloading_heavyComputation2.region_id = weak constant i8 0
16 @.offload_maptypes.3 = private unnamed_addr constant [2 x i64] [i64 35, i64 35]
18 @.__omp_offloading_heavyComputation3.region_id = weak constant i8 0
19 @.offload_sizes.2 = private unnamed_addr constant [2 x i64] [i64 4, i64 0]
20 @.offload_maptypes.4 = private unnamed_addr constant [2 x i64] [i64 800, i64 544]
22 @.offload_maptypes.5 = private unnamed_addr constant [1 x i64] [i64 33]
24 @0 = private unnamed_addr global %struct.ident_t { i32 0, i32 34, i32 0, i32 0, ptr @.str0 }, align 8
25 @.str0 = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00", align 1
27 ;double heavyComputation1() {
28 ; double a = rand() % 777;
29 ; double random = rand();
31 ; //#pragma omp target data map(a)
34 ; __tgt_target_data_begin(..., args, ...)
36 ; #pragma omp target teams
37 ; for (int i = 0; i < 1000; ++i) {
43 define dso_local double @heavyComputation1() {
44 ; CHECK-LABEL: define {{[^@]+}}@heavyComputation1() {
46 ; CHECK-NEXT: [[A:%.*]] = alloca double, align 8
47 ; CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8
48 ; CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8
49 ; CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS4:%.*]] = alloca [1 x ptr], align 8
50 ; CHECK-NEXT: [[DOTOFFLOAD_PTRS5:%.*]] = alloca [1 x ptr], align 8
51 ; CHECK-NEXT: [[CALL:%.*]] = tail call i32 (...) @rand()
52 ; CHECK-NEXT: [[REM:%.*]] = srem i32 [[CALL]], 777
53 ; CHECK-NEXT: [[CONV:%.*]] = sitofp i32 [[REM]] to double
54 ; CHECK-NEXT: store double [[CONV]], ptr [[A]], align 8
55 ; CHECK-NEXT: [[CALL1:%.*]] = tail call i32 (...) @rand()
56 ; CHECK-NEXT: store ptr [[A]], ptr [[DOTOFFLOAD_BASEPTRS]], align 8
57 ; CHECK-NEXT: store ptr [[A]], ptr [[DOTOFFLOAD_PTRS]], align 8
58 ; CHECK-NEXT: call void @__tgt_target_data_begin_mapper(ptr @[[GLOB0:[0-9]+]], i64 -1, i32 1, ptr nonnull [[DOTOFFLOAD_BASEPTRS]], ptr nonnull [[DOTOFFLOAD_PTRS]], ptr @.offload_sizes.1, ptr @.offload_maptypes, ptr null, ptr null)
59 ; CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[A]], align 8
60 ; CHECK-NEXT: store i64 [[TMP0]], ptr [[DOTOFFLOAD_BASEPTRS4]], align 8
61 ; CHECK-NEXT: store i64 [[TMP0]], ptr [[DOTOFFLOAD_PTRS5]], align 8
62 ; CHECK-NEXT: [[TMP1:%.*]] = call i32 @__tgt_target_teams_mapper(ptr @[[GLOB0]], i64 -1, ptr nonnull @.__omp_offloading_heavyComputation1.region_id, i32 1, ptr nonnull [[DOTOFFLOAD_BASEPTRS4]], ptr nonnull [[DOTOFFLOAD_PTRS5]], ptr @.offload_sizes.1, ptr @.offload_maptypes.2, ptr null, ptr null, i32 0, i32 0)
63 ; CHECK-NEXT: [[DOTNOT:%.*]] = icmp eq i32 [[TMP1]], 0
64 ; CHECK-NEXT: br i1 [[DOTNOT]], label [[OMP_OFFLOAD_CONT:%.*]], label [[OMP_OFFLOAD_FAILED:%.*]]
65 ; CHECK: omp_offload.failed:
66 ; CHECK-NEXT: call void @heavyComputation1FallBack(i64 [[TMP0]])
67 ; CHECK-NEXT: br label [[OMP_OFFLOAD_CONT]]
68 ; CHECK: omp_offload.cont:
69 ; CHECK-NEXT: [[CONV2:%.*]] = sitofp i32 [[CALL1]] to double
70 ; CHECK-NEXT: call void @__tgt_target_data_end_mapper(ptr @[[GLOB0]], i64 -1, i32 1, ptr nonnull [[DOTOFFLOAD_BASEPTRS]], ptr nonnull [[DOTOFFLOAD_PTRS]], ptr @.offload_sizes.1, ptr @.offload_maptypes, ptr null, ptr null)
71 ; CHECK-NEXT: [[TMP2:%.*]] = load double, ptr [[A]], align 8
72 ; CHECK-NEXT: [[ADD:%.*]] = fadd double [[TMP2]], [[CONV2]]
73 ; CHECK-NEXT: ret double [[ADD]]
82 %a = alloca double, align 8
83 %.offload_baseptrs = alloca [1 x ptr], align 8
84 %.offload_ptrs = alloca [1 x ptr], align 8
85 %.offload_baseptrs4 = alloca [1 x ptr], align 8
86 %.offload_ptrs5 = alloca [1 x ptr], align 8
88 %call = tail call i32 (...) @rand()
89 %rem = srem i32 %call, 777
90 %conv = sitofp i32 %rem to double
91 store double %conv, ptr %a, align 8
93 ; FIXME: call to @__tgt_target_data_begin_mapper_issue(ptr @0, ...) should be moved here.
94 %call1 = tail call i32 (...) @rand()
96 store ptr %a, ptr %.offload_baseptrs, align 8
97 store ptr %a, ptr %.offload_ptrs, align 8
98 call void @__tgt_target_data_begin_mapper(ptr @0, i64 -1, i32 1, ptr nonnull %.offload_baseptrs, ptr nonnull %.offload_ptrs, ptr @.offload_sizes.1, ptr @.offload_maptypes, ptr null, ptr null)
100 %0 = load i64, ptr %a, align 8
101 store i64 %0, ptr %.offload_baseptrs4, align 8
102 store i64 %0, ptr %.offload_ptrs5, align 8
104 ; FIXME: call to @__tgt_target_data_begin_mapper_wait(...) should be moved here.
105 %1 = call i32 @__tgt_target_teams_mapper(ptr @0, i64 -1, ptr nonnull @.__omp_offloading_heavyComputation1.region_id, i32 1, ptr nonnull %.offload_baseptrs4, ptr nonnull %.offload_ptrs5, ptr @.offload_sizes.1, ptr @.offload_maptypes.2, ptr null, ptr null, i32 0, i32 0)
106 %.not = icmp eq i32 %1, 0
107 br i1 %.not, label %omp_offload.cont, label %omp_offload.failed
109 omp_offload.failed: ; preds = %entry
110 call void @heavyComputation1FallBack(i64 %0)
111 br label %omp_offload.cont
113 omp_offload.cont: ; preds = %omp_offload.failed, %entry
114 %conv2 = sitofp i32 %call1 to double
115 call void @__tgt_target_data_end_mapper(ptr @0, i64 -1, i32 1, ptr nonnull %.offload_baseptrs, ptr nonnull %.offload_ptrs, ptr @.offload_sizes.1, ptr @.offload_maptypes, ptr null, ptr null)
116 %2 = load double, ptr %a, align 8
117 %add = fadd double %2, %conv2
121 define internal void @heavyComputation1FallBack(i64 %a) {
122 ; CHECK-LABEL: define {{[^@]+}}@heavyComputation1FallBack
123 ; CHECK-SAME: (i64 [[A:%.*]]) {
125 ; CHECK-NEXT: ret void
128 ; Fallback for offloading function heavyComputation1.
132 ;int heavyComputation2(ptr a, unsigned size) {
133 ; int random = rand() % 7;
135 ; //#pragma omp target data map(a[0:size], size)
139 ; __tgt_target_data_begin(..., args, ...)
141 ; #pragma omp target teams
142 ; for (int i = 0; i < size; ++i) {
143 ; a[i] = ++aptr 3.141624;
148 define dso_local i32 @heavyComputation2(ptr %a, i32 %size) {
149 ; CHECK-LABEL: define {{[^@]+}}@heavyComputation2
150 ; CHECK-SAME: (ptr [[A:%.*]], i32 [[SIZE:%.*]]) {
152 ; CHECK-NEXT: [[SIZE_ADDR:%.*]] = alloca i32, align 4
153 ; CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [2 x ptr], align 8
154 ; CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [2 x ptr], align 8
155 ; CHECK-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [2 x i64], align 8
156 ; CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS2:%.*]] = alloca [2 x ptr], align 8
157 ; CHECK-NEXT: [[DOTOFFLOAD_PTRS3:%.*]] = alloca [2 x ptr], align 8
158 ; CHECK-NEXT: store i32 [[SIZE]], ptr [[SIZE_ADDR]], align 4
159 ; CHECK-NEXT: [[CALL:%.*]] = tail call i32 (...) @rand()
160 ; CHECK-NEXT: [[CONV:%.*]] = zext i32 [[SIZE]] to i64
161 ; CHECK-NEXT: [[TMP0:%.*]] = shl nuw nsw i64 [[CONV]], 3
162 ; CHECK-NEXT: store ptr [[A]], ptr [[DOTOFFLOAD_BASEPTRS]], align 8
163 ; CHECK-NEXT: store ptr [[A]], ptr [[DOTOFFLOAD_PTRS]], align 8
164 ; CHECK-NEXT: store i64 [[TMP0]], ptr [[DOTOFFLOAD_SIZES]], align 8
165 ; CHECK-NEXT: [[TMP1:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i64 0, i64 1
166 ; CHECK-NEXT: store ptr [[SIZE_ADDR]], ptr [[TMP1]], align 8
167 ; CHECK-NEXT: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i64 0, i64 1
168 ; CHECK-NEXT: store ptr [[SIZE_ADDR]], ptr [[TMP2]], align 8
169 ; CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds [2 x i64], ptr [[DOTOFFLOAD_SIZES]], i64 0, i64 1
170 ; CHECK-NEXT: store i64 4, ptr [[TMP3]], align 8
171 ; CHECK-NEXT: call void @__tgt_target_data_begin_mapper(ptr @[[GLOB0]], i64 -1, i32 2, ptr nonnull [[DOTOFFLOAD_BASEPTRS]], ptr nonnull [[DOTOFFLOAD_PTRS]], ptr nonnull [[DOTOFFLOAD_SIZES]], ptr @.offload_maptypes.3, ptr null, ptr null)
172 ; CHECK-NEXT: [[TMP4:%.*]] = load i32, ptr [[SIZE_ADDR]], align 4
173 ; CHECK-NEXT: [[SIZE_CASTED:%.*]] = zext i32 [[TMP4]] to i64
174 ; CHECK-NEXT: store i64 [[SIZE_CASTED]], ptr [[DOTOFFLOAD_BASEPTRS2]], align 8
175 ; CHECK-NEXT: store i64 [[SIZE_CASTED]], ptr [[DOTOFFLOAD_PTRS3]], align 8
176 ; CHECK-NEXT: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i64 0, i64 1
177 ; CHECK-NEXT: store ptr [[A]], ptr [[TMP5]], align 8
178 ; CHECK-NEXT: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i64 0, i64 1
179 ; CHECK-NEXT: store ptr [[A]], ptr [[TMP6]], align 8
180 ; CHECK-NEXT: [[TMP7:%.*]] = call i32 @__tgt_target_teams_mapper(ptr @[[GLOB0]], i64 -1, ptr nonnull @.__omp_offloading_heavyComputation2.region_id, i32 2, ptr nonnull [[DOTOFFLOAD_BASEPTRS2]], ptr nonnull [[DOTOFFLOAD_PTRS3]], ptr @.offload_sizes.2, ptr @.offload_maptypes.4, ptr null, ptr null, i32 0, i32 0)
181 ; CHECK-NEXT: [[DOTNOT:%.*]] = icmp eq i32 [[TMP7]], 0
182 ; CHECK-NEXT: br i1 [[DOTNOT]], label [[OMP_OFFLOAD_CONT:%.*]], label [[OMP_OFFLOAD_FAILED:%.*]]
183 ; CHECK: omp_offload.failed:
184 ; CHECK-NEXT: call void @heavyComputation2FallBack(i64 [[SIZE_CASTED]], ptr [[A]])
185 ; CHECK-NEXT: br label [[OMP_OFFLOAD_CONT]]
186 ; CHECK: omp_offload.cont:
187 ; CHECK-NEXT: [[REM:%.*]] = srem i32 [[CALL]], 7
188 ; CHECK-NEXT: call void @__tgt_target_data_end_mapper(ptr @[[GLOB0]], i64 -1, i32 2, ptr nonnull [[DOTOFFLOAD_BASEPTRS]], ptr nonnull [[DOTOFFLOAD_PTRS]], ptr nonnull [[DOTOFFLOAD_SIZES]], ptr @.offload_maptypes.3, ptr null, ptr null)
189 ; CHECK-NEXT: ret i32 [[REM]]
194 %size.addr = alloca i32, align 4
195 %.offload_baseptrs = alloca [2 x ptr], align 8
196 %.offload_ptrs = alloca [2 x ptr], align 8
197 %.offload_sizes = alloca [2 x i64], align 8
198 %.offload_baseptrs2 = alloca [2 x ptr], align 8
199 %.offload_ptrs3 = alloca [2 x ptr], align 8
201 store i32 %size, ptr %size.addr, align 4
202 %call = tail call i32 (...) @rand()
204 %conv = zext i32 %size to i64
205 %0 = shl nuw nsw i64 %conv, 3
206 store ptr %a, ptr %.offload_baseptrs, align 8
207 store ptr %a, ptr %.offload_ptrs, align 8
208 store i64 %0, ptr %.offload_sizes, align 8
209 %1 = getelementptr inbounds [2 x ptr], ptr %.offload_baseptrs, i64 0, i64 1
210 store ptr %size.addr, ptr %1, align 8
211 %2 = getelementptr inbounds [2 x ptr], ptr %.offload_ptrs, i64 0, i64 1
212 store ptr %size.addr, ptr %2, align 8
213 %3 = getelementptr inbounds [2 x i64], ptr %.offload_sizes, i64 0, i64 1
214 store i64 4, ptr %3, align 8
215 call void @__tgt_target_data_begin_mapper(ptr @0, i64 -1, i32 2, ptr nonnull %.offload_baseptrs, ptr nonnull %.offload_ptrs, ptr nonnull %.offload_sizes, ptr @.offload_maptypes.3, ptr null, ptr null)
217 %4 = load i32, ptr %size.addr, align 4
218 %size.casted = zext i32 %4 to i64
219 store i64 %size.casted, ptr %.offload_baseptrs2, align 8
220 store i64 %size.casted, ptr %.offload_ptrs3, align 8
221 %5 = getelementptr inbounds [2 x ptr], ptr %.offload_baseptrs2, i64 0, i64 1
222 store ptr %a, ptr %5, align 8
223 %6 = getelementptr inbounds [2 x ptr], ptr %.offload_ptrs3, i64 0, i64 1
224 store ptr %a, ptr %6, align 8
226 ; FIXME: call to @__tgt_target_data_begin_mapper_wait(...) should be moved here.
227 %7 = call i32 @__tgt_target_teams_mapper(ptr @0, i64 -1, ptr nonnull @.__omp_offloading_heavyComputation2.region_id, i32 2, ptr nonnull %.offload_baseptrs2, ptr nonnull %.offload_ptrs3, ptr @.offload_sizes.2, ptr @.offload_maptypes.4, ptr null, ptr null, i32 0, i32 0)
228 %.not = icmp eq i32 %7, 0
229 br i1 %.not, label %omp_offload.cont, label %omp_offload.failed
231 omp_offload.failed: ; preds = %entry
232 call void @heavyComputation2FallBack(i64 %size.casted, ptr %a)
233 br label %omp_offload.cont
235 omp_offload.cont: ; preds = %omp_offload.failed, %entry
236 %rem = srem i32 %call, 7
237 call void @__tgt_target_data_end_mapper(ptr @0, i64 -1, i32 2, ptr nonnull %.offload_baseptrs, ptr nonnull %.offload_ptrs, ptr nonnull %.offload_sizes, ptr @.offload_maptypes.3, ptr null, ptr null)
241 define internal void @heavyComputation2FallBack(i64 %size, ptr %a) {
242 ; CHECK-LABEL: define {{[^@]+}}@heavyComputation2FallBack
243 ; CHECK-SAME: (i64 [[SIZE:%.*]], ptr [[A:%.*]]) {
245 ; CHECK-NEXT: ret void
248 ; Fallback for offloading function heavyComputation2.
252 ;int heavyComputation3(ptr restrict a, unsigned size) {
253 ; int random = rand() % 7;
255 ; //#pragma omp target data map(a[0:size], size)
259 ; __tgt_target_data_begin(..., args, ...)
261 ; #pragma omp target teams
262 ; for (int i = 0; i < size; ++i) {
263 ; a[i] = ++aptr 3.141624;
268 define dso_local i32 @heavyComputation3(ptr noalias %a, i32 %size) {
269 ; CHECK-LABEL: define {{[^@]+}}@heavyComputation3
270 ; CHECK-SAME: (ptr noalias [[A:%.*]], i32 [[SIZE:%.*]]) {
272 ; CHECK-NEXT: [[SIZE_ADDR:%.*]] = alloca i32, align 4
273 ; CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [2 x ptr], align 8
274 ; CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [2 x ptr], align 8
275 ; CHECK-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [2 x i64], align 8
276 ; CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS2:%.*]] = alloca [2 x ptr], align 8
277 ; CHECK-NEXT: [[DOTOFFLOAD_PTRS3:%.*]] = alloca [2 x ptr], align 8
278 ; CHECK-NEXT: store i32 [[SIZE]], ptr [[SIZE_ADDR]], align 4
279 ; CHECK-NEXT: [[CALL:%.*]] = tail call i32 (...) @rand()
280 ; CHECK-NEXT: [[CONV:%.*]] = zext i32 [[SIZE]] to i64
281 ; CHECK-NEXT: [[TMP0:%.*]] = shl nuw nsw i64 [[CONV]], 3
282 ; CHECK-NEXT: store ptr [[A]], ptr [[DOTOFFLOAD_BASEPTRS]], align 8
283 ; CHECK-NEXT: store ptr [[A]], ptr [[DOTOFFLOAD_PTRS]], align 8
284 ; CHECK-NEXT: store i64 [[TMP0]], ptr [[DOTOFFLOAD_SIZES]], align 8
285 ; CHECK-NEXT: [[TMP1:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i64 0, i64 1
286 ; CHECK-NEXT: store ptr [[SIZE_ADDR]], ptr [[TMP1]], align 8
287 ; CHECK-NEXT: [[TMP2:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS]], i64 0, i64 1
288 ; CHECK-NEXT: store ptr [[SIZE_ADDR]], ptr [[TMP2]], align 8
289 ; CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds [2 x i64], ptr [[DOTOFFLOAD_SIZES]], i64 0, i64 1
290 ; CHECK-NEXT: store i64 4, ptr [[TMP3]], align 8
291 ; CHECK-NEXT: call void @__tgt_target_data_begin_mapper(ptr @[[GLOB0]], i64 -1, i32 2, ptr nonnull [[DOTOFFLOAD_BASEPTRS]], ptr nonnull [[DOTOFFLOAD_PTRS]], ptr nonnull [[DOTOFFLOAD_SIZES]], ptr @.offload_maptypes.3, ptr null, ptr null)
292 ; CHECK-NEXT: [[TMP4:%.*]] = load i32, ptr [[SIZE_ADDR]], align 4
293 ; CHECK-NEXT: [[SIZE_CASTED:%.*]] = zext i32 [[TMP4]] to i64
294 ; CHECK-NEXT: store i64 [[SIZE_CASTED]], ptr [[DOTOFFLOAD_BASEPTRS2]], align 8
295 ; CHECK-NEXT: store i64 [[SIZE_CASTED]], ptr [[DOTOFFLOAD_PTRS3]], align 8
296 ; CHECK-NEXT: [[TMP5:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_BASEPTRS2]], i64 0, i64 1
297 ; CHECK-NEXT: store ptr [[A]], ptr [[TMP5]], align 8
298 ; CHECK-NEXT: [[TMP6:%.*]] = getelementptr inbounds [2 x ptr], ptr [[DOTOFFLOAD_PTRS3]], i64 0, i64 1
299 ; CHECK-NEXT: store ptr [[A]], ptr [[TMP6]], align 8
300 ; CHECK-NEXT: [[TMP7:%.*]] = call i32 @__tgt_target_teams_mapper(ptr @[[GLOB0]], i64 -1, ptr nonnull @.__omp_offloading_heavyComputation3.region_id, i32 2, ptr nonnull [[DOTOFFLOAD_BASEPTRS2]], ptr nonnull [[DOTOFFLOAD_PTRS3]], ptr @.offload_sizes.2, ptr @.offload_maptypes.4, ptr null, ptr null, i32 0, i32 0)
301 ; CHECK-NEXT: [[DOTNOT:%.*]] = icmp eq i32 [[TMP7]], 0
302 ; CHECK-NEXT: br i1 [[DOTNOT]], label [[OMP_OFFLOAD_CONT:%.*]], label [[OMP_OFFLOAD_FAILED:%.*]]
303 ; CHECK: omp_offload.failed:
304 ; CHECK-NEXT: call void @heavyComputation3FallBack(i64 [[SIZE_CASTED]], ptr [[A]])
305 ; CHECK-NEXT: br label [[OMP_OFFLOAD_CONT]]
306 ; CHECK: omp_offload.cont:
307 ; CHECK-NEXT: [[REM:%.*]] = srem i32 [[CALL]], 7
308 ; CHECK-NEXT: call void @__tgt_target_data_end_mapper(ptr @[[GLOB0]], i64 -1, i32 2, ptr nonnull [[DOTOFFLOAD_BASEPTRS]], ptr nonnull [[DOTOFFLOAD_PTRS]], ptr nonnull [[DOTOFFLOAD_SIZES]], ptr @.offload_maptypes.3, ptr null, ptr null)
309 ; CHECK-NEXT: ret i32 [[REM]]
314 %size.addr = alloca i32, align 4
315 %.offload_baseptrs = alloca [2 x ptr], align 8
316 %.offload_ptrs = alloca [2 x ptr], align 8
317 %.offload_sizes = alloca [2 x i64], align 8
318 %.offload_baseptrs2 = alloca [2 x ptr], align 8
319 %.offload_ptrs3 = alloca [2 x ptr], align 8
320 store i32 %size, ptr %size.addr, align 4
322 ; FIXME: call to @__tgt_target_data_begin_mapper_issue(ptr @0, ...) should be moved here.
323 %call = tail call i32 (...) @rand()
325 %conv = zext i32 %size to i64
326 %0 = shl nuw nsw i64 %conv, 3
327 store ptr %a, ptr %.offload_baseptrs, align 8
328 store ptr %a, ptr %.offload_ptrs, align 8
329 store i64 %0, ptr %.offload_sizes, align 8
330 %1 = getelementptr inbounds [2 x ptr], ptr %.offload_baseptrs, i64 0, i64 1
331 store ptr %size.addr, ptr %1, align 8
332 %2 = getelementptr inbounds [2 x ptr], ptr %.offload_ptrs, i64 0, i64 1
333 store ptr %size.addr, ptr %2, align 8
334 %3 = getelementptr inbounds [2 x i64], ptr %.offload_sizes, i64 0, i64 1
335 store i64 4, ptr %3, align 8
336 call void @__tgt_target_data_begin_mapper(ptr @0, i64 -1, i32 2, ptr nonnull %.offload_baseptrs, ptr nonnull %.offload_ptrs, ptr nonnull %.offload_sizes, ptr @.offload_maptypes.3, ptr null, ptr null)
338 %4 = load i32, ptr %size.addr, align 4
339 %size.casted = zext i32 %4 to i64
340 store i64 %size.casted, ptr %.offload_baseptrs2, align 8
341 store i64 %size.casted, ptr %.offload_ptrs3, align 8
342 %5 = getelementptr inbounds [2 x ptr], ptr %.offload_baseptrs2, i64 0, i64 1
343 store ptr %a, ptr %5, align 8
344 %6 = getelementptr inbounds [2 x ptr], ptr %.offload_ptrs3, i64 0, i64 1
345 store ptr %a, ptr %6, align 8
347 ; FIXME: call to @__tgt_target_data_begin_mapper_wait(...) should be moved here.
348 %7 = call i32 @__tgt_target_teams_mapper(ptr @0, i64 -1, ptr nonnull @.__omp_offloading_heavyComputation3.region_id, i32 2, ptr nonnull %.offload_baseptrs2, ptr nonnull %.offload_ptrs3, ptr @.offload_sizes.2, ptr @.offload_maptypes.4, ptr null, ptr null, i32 0, i32 0)
349 %.not = icmp eq i32 %7, 0
350 br i1 %.not, label %omp_offload.cont, label %omp_offload.failed
352 omp_offload.failed: ; preds = %entry
353 call void @heavyComputation3FallBack(i64 %size.casted, ptr %a)
354 br label %omp_offload.cont
356 omp_offload.cont: ; preds = %omp_offload.failed, %entry
357 %rem = srem i32 %call, 7
358 call void @__tgt_target_data_end_mapper(ptr @0, i64 -1, i32 2, ptr nonnull %.offload_baseptrs, ptr nonnull %.offload_ptrs, ptr nonnull %.offload_sizes, ptr @.offload_maptypes.3, ptr null, ptr null)
362 define internal void @heavyComputation3FallBack(i64 %size, ptr %a) {
363 ; CHECK-LABEL: define {{[^@]+}}@heavyComputation3FallBack
364 ; CHECK-SAME: (i64 [[SIZE:%.*]], ptr [[A:%.*]]) {
366 ; CHECK-NEXT: ret void
369 ; Fallback for offloading function heavyComputation3.
373 ;int dataTransferOnly1(ptr restrict a, unsigned size) {
374 ; // Random computation.
375 ; int random = rand();
377 ; //#pragma omp target data map(to:a[0:size])
380 ; __tgt_target_data_begin(..., args, ...)
382 ; // Random computation.
386 define dso_local i32 @dataTransferOnly1(ptr noalias %a, i32 %size) {
387 ; CHECK-LABEL: define {{[^@]+}}@dataTransferOnly1
388 ; CHECK-SAME: (ptr noalias [[A:%.*]], i32 [[SIZE:%.*]]) {
390 ; CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8
391 ; CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8
392 ; CHECK-NEXT: [[DOTOFFLOAD_SIZES:%.*]] = alloca [1 x i64], align 8
393 ; CHECK-NEXT: [[HANDLE:%.*]] = alloca [[STRUCT___TGT_ASYNC_INFO:%.*]], align 8
394 ; CHECK-NEXT: [[CALL:%.*]] = tail call i32 (...) @rand()
395 ; CHECK-NEXT: [[CONV:%.*]] = zext i32 [[SIZE]] to i64
396 ; CHECK-NEXT: [[TMP0:%.*]] = shl nuw nsw i64 [[CONV]], 3
397 ; CHECK-NEXT: store ptr [[A]], ptr [[DOTOFFLOAD_BASEPTRS]], align 8
398 ; CHECK-NEXT: store ptr [[A]], ptr [[DOTOFFLOAD_PTRS]], align 8
399 ; CHECK-NEXT: store i64 [[TMP0]], ptr [[DOTOFFLOAD_SIZES]], align 8
400 ; CHECK-NEXT: call void @__tgt_target_data_begin_mapper_issue(ptr @[[GLOB0]], i64 -1, i32 1, ptr [[DOTOFFLOAD_BASEPTRS]], ptr [[DOTOFFLOAD_PTRS]], ptr [[DOTOFFLOAD_SIZES]], ptr @.offload_maptypes.5, ptr null, ptr null, ptr [[HANDLE]])
401 ; CHECK-NEXT: [[REM:%.*]] = urem i32 [[CALL]], [[SIZE]]
402 ; CHECK-NEXT: call void @__tgt_target_data_begin_mapper_wait(i64 -1, ptr [[HANDLE]])
403 ; CHECK-NEXT: call void @__tgt_target_data_end_mapper(ptr @[[GLOB0]], i64 -1, i32 1, ptr nonnull [[DOTOFFLOAD_BASEPTRS]], ptr nonnull [[DOTOFFLOAD_PTRS]], ptr nonnull [[DOTOFFLOAD_SIZES]], ptr @.offload_maptypes.5, ptr null, ptr null)
404 ; CHECK-NEXT: ret i32 [[REM]]
413 %.offload_baseptrs = alloca [1 x ptr], align 8
414 %.offload_ptrs = alloca [1 x ptr], align 8
415 %.offload_sizes = alloca [1 x i64], align 8
417 ; FIXME: call to @__tgt_target_data_begin_issue_mapper(...) should be moved here.
418 %call = tail call i32 (...) @rand()
420 %conv = zext i32 %size to i64
421 %0 = shl nuw nsw i64 %conv, 3
422 store ptr %a, ptr %.offload_baseptrs, align 8
423 store ptr %a, ptr %.offload_ptrs, align 8
424 store i64 %0, ptr %.offload_sizes, align 8
425 call void @__tgt_target_data_begin_mapper(ptr @0, i64 -1, i32 1, ptr nonnull %.offload_baseptrs, ptr nonnull %.offload_ptrs, ptr nonnull %.offload_sizes, ptr @.offload_maptypes.5, ptr null, ptr null)
427 %rem = urem i32 %call, %size
429 call void @__tgt_target_data_end_mapper(ptr @0, i64 -1, i32 1, ptr nonnull %.offload_baseptrs, ptr nonnull %.offload_ptrs, ptr nonnull %.offload_sizes, ptr @.offload_maptypes.5, ptr null, ptr null)
433 declare void @__tgt_target_data_begin_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr)
434 declare i32 @__tgt_target_teams_mapper(ptr, i64, ptr, i32, ptr, ptr, ptr, ptr, ptr, ptr, i32, i32)
435 declare void @__tgt_target_data_end_mapper(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr)
437 declare dso_local i32 @rand(...)
439 ; CHECK: declare void @__tgt_target_data_begin_mapper_issue(ptr, i64, i32, ptr, ptr, ptr, ptr, ptr, ptr, ptr)
440 ; CHECK: declare void @__tgt_target_data_begin_mapper_wait(i64, ptr)
442 !llvm.module.flags = !{!0}
444 !0 = !{i32 7, !"openmp", i32 50}