xref: /llvm-project/clang/test/OpenMP/target_exit_data_depend_codegen.cpp (revision 94473f4db6a6f5f12d7c4081455b5b596094eac5) 
1 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK1 --check-prefix CK1-64
2 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
3 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s  --check-prefix CK1 --check-prefix CK1-64
4 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s  --check-prefix CK1 --check-prefix CK1-32
5 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
6 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s  --check-prefix CK1 --check-prefix CK1-32
7 
8 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
9 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s
10 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
11 // RUN: %clang_cc1 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
12 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s
13 // RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
14 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
15 
16 // expected-no-diagnostics
17 // CK1: [[ST:%.+]] = type { i32, ptr }
18 // CK1: %struct.kmp_depend_info = type { i[[sz:64|32]],
19 // CK1-SAME: i[[sz]], i8 }
20 #ifndef HEADER
21 #define HEADER
22 
23 template <typename T>
24 struct ST {
25   T a;
26   double *b;
27 };
28 
29 ST<int> gb;
30 double gc[100];
31 
32 // CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 800]
33 // CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 2]
34 
35 // CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i64] [i64 4]
36 // CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 8]
37 
38 // CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 2]
39 
40 // CK1: [[SIZE04:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 24]
41 // CK1: [[MTYPE04:@.+]] = {{.+}}constant [2 x i64] [i64 0, i64 281474976710674]
42 
43 // CK1-LABEL: _Z3fooi
44 void foo(int arg) {
45   int la;
46   float lb[arg];
47 
48   // CK1: alloca [1 x %struct.kmp_depend_info],
49   // CK1: alloca [3 x %struct.kmp_depend_info],
50   // CK1: alloca [4 x %struct.kmp_depend_info],
51   // CK1: alloca [5 x %struct.kmp_depend_info],
52 
53   // Region 00
54   // CK1: [[BP0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP:%.+]], i32 0, i32 0
55   // CK1: store ptr @gc, ptr [[BP0]],
56   // CK1: [[P0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P:%.+]], i32 0, i32 0
57   // CK1: store ptr @gc, ptr [[P0]],
58   // CK1: [[GEPBP0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP]], i32 0, i32 0
59   // CK1: [[GEPP0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P]], i32 0, i32 0
60   // CK1: [[CAP_DEVICE:%.+]] = getelementptr inbounds nuw %struct.anon, ptr [[CAPTURES:%.+]], i32 0, i32 0
61   // CK1: [[DEVICE:%.+]] = load i32, ptr %{{.+}}
62   // CK1: store i32 [[DEVICE]], ptr [[CAP_DEVICE]],
63   // CK1: [[DEV1:%.+]] = load i32, ptr %{{.+}}
64   // CK1: [[DEV2:%.+]] = sext i32 [[DEV1]] to i64
65   // CK1: [[RES:%.+]] = call ptr @__kmpc_omp_target_task_alloc(ptr {{.+}}, i32 {{.+}}, i32 1, i[[sz]] {{64|36}}, i[[sz]] 4, ptr [[TASK_ENTRY0:@.+]], i64 [[DEV2]])
66   // CK1: [[TASK_T:%.+]] = getelementptr inbounds nuw %struct.kmp_task_t_with_privates, ptr [[RES]], i32 0, i32 0
67   // CK1: [[SHAREDS:%.+]] = getelementptr inbounds nuw %struct.kmp_task_t, ptr [[TASK_T]], i32 0, i32 0
68   // CK1: [[SHAREDS_REF:%.+]] = load ptr, ptr [[SHAREDS]],
69   // CK1: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align 4 [[SHAREDS_REF]], ptr align 4 [[CAPTURES]], i[[sz]] 4, i1 false)
70   // CK1: [[PRIVS:%.+]] = getelementptr inbounds nuw %struct.kmp_task_t_with_privates, ptr [[RES]], i32 0, i32 1
71   // CK1-64: [[PRIVS_BASEPTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t, ptr [[PRIVS]], i32 0, i32 0
72   // CK1-64: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_BASEPTRS]], ptr align {{8|4}} [[GEPBP0]], i[[sz]] {{8|4}}, i1 false)
73   // CK1-64: [[PRIVS_PTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t, ptr [[PRIVS]], i32 0, i32 1
74   // CK1-64: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_PTRS]], ptr align {{8|4}} [[GEPP0]], i[[sz]] {{8|4}}, i1 false)
75   // CK1-64: [[PRIVS_SIZES:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t, ptr [[PRIVS]], i32 0, i32 2
76   // CK1-64: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_SIZES]], ptr align {{8|4}} [[SIZE00]], i[[sz]] {{8|4}}, i1 false)
77   // CK1-32: [[PRIVS_SIZES:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t, ptr [[PRIVS]], i32 0, i32 0
78   // CK1-32: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_SIZES]], ptr align {{8|4}} [[SIZE00]], i[[sz]] {{8|4}}, i1 false)
79   // CK1-32: [[PRIVS_BASEPTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t, ptr [[PRIVS]], i32 0, i32 1
80   // CK1-32: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_BASEPTRS]], ptr align {{8|4}} [[GEPBP0]], i[[sz]] {{8|4}}, i1 false)
81   // CK1-32: [[PRIVS_PTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t, ptr [[PRIVS]], i32 0, i32 2
82   // CK1-32: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_PTRS]], ptr align {{8|4}} [[GEPP0]], i[[sz]] {{8|4}}, i1 false)
83   // CK1: [[BC_ADR:%.+]] = ptrtoint ptr %{{.+}} to i[[sz]]
84   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, ptr [[MAIN_DEP:%.+]], i[[sz]] 0
85   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 0
86   // CK1: store i[[sz]] [[BC_ADR]], ptr [[DEP_ADR]],
87   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 1
88   // CK1: store i[[sz]] 4, ptr [[DEP_SIZE]],
89   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 2
90   // CK1: store i8 1, ptr [[DEP_ATTRS]]
91   // CK1: = call i32 @__kmpc_omp_task_with_deps(ptr @{{.+}}, i32 %{{.+}}, ptr [[RES]], i32 1, ptr [[MAIN_DEP]], i32 0, ptr null)
92 
93   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
94   #pragma omp target exit data if(1+3-5) device(arg) map(from:gc) nowait depend(in: arg)
95   {++arg;}
96 
97   // Region 01
98   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
99   #pragma omp target exit data map(release: la) if(1+3-4) depend(in: la) depend(out: arg)
100   {++arg;}
101 
102   // Region 02
103   // CK1: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
104   // CK1: [[IFTHEN]]
105   // CK1: [[BP0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP:%.+]], i32 0, i32 0
106   // CK1: store ptr [[ARG:%.+]], ptr [[BP0]],
107   // CK1: [[P0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P:%.+]], i32 0, i32 0
108   // CK1: store ptr [[ARG]], ptr [[P0]],
109   // CK1: [[GEPBP0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP]], i32 0, i32 0
110   // CK1: [[GEPP0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P]], i32 0, i32 0
111   // CK1: [[IF_DEVICE:%.+]] = getelementptr inbounds nuw %struct.anon{{.+}}, ptr [[CAPTURES:%.+]], i32 0, i32 0
112   // CK1: [[IF:%.+]] = load i8, ptr %{{.+}}
113   // CK1: [[IF_BOOL:%.+]] = trunc i8 [[IF]] to i1
114   // CK1: [[IF:%.+]] = zext i1 [[IF_BOOL]] to i8
115   // CK1: store i8 [[IF]], ptr [[IF_DEVICE]],
116   // CK1: [[RES:%.+]] = call ptr @__kmpc_omp_task_alloc(ptr {{.+}}, i32 {{.+}}, i32 1, i[[sz]] {{64|36}}, i[[sz]] 1, ptr [[TASK_ENTRY2:@.+]])
117   // CK1: [[TASK_T:%.+]] = getelementptr inbounds nuw %struct.kmp_task_t_with_privates{{.+}}, ptr [[RES]], i32 0, i32 0
118   // CK1: [[SHAREDS:%.+]] = getelementptr inbounds nuw %struct.kmp_task_t, ptr [[TASK_T]], i32 0, i32 0
119   // CK1: [[SHAREDS_REF:%.+]] = load ptr, ptr [[SHAREDS]],
120   // CK1: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align 1 [[SHAREDS_REF]], ptr align 1 [[CAPTURES]], i[[sz]] 1, i1 false)
121   // CK1: [[PRIVS:%.+]] = getelementptr inbounds nuw %struct.kmp_task_t_with_privates{{.+}}, ptr [[RES]], i32 0, i32 1
122   // CK1-64: [[PRIVS_BASEPTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 0
123   // CK1-64: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_BASEPTRS]], ptr align {{8|4}} [[GEPBP0]], i[[sz]] {{8|4}}, i1 false)
124   // CK1-64: [[PRIVS_PTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 1
125   // CK1-64: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_PTRS]], ptr align {{8|4}} [[GEPP0]], i[[sz]] {{8|4}}, i1 false)
126   // CK1-64: [[PRIVS_SIZES:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 2
127   // CK1-64: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_SIZES]], ptr align {{8|4}} [[SIZE02]], i[[sz]] {{8|4}}, i1 false)
128   // CK1-32: [[PRIVS_SIZES:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 0
129   // CK1-32: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_SIZES]], ptr align {{8|4}} [[SIZE02]], i[[sz]] {{8|4}}, i1 false)
130   // CK1-32: [[PRIVS_BASEPTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 1
131   // CK1-32: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_BASEPTRS]], ptr align {{8|4}} [[GEPBP0]], i[[sz]] {{8|4}}, i1 false)
132   // CK1-32: [[PRIVS_PTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 2
133   // CK1-32: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_PTRS]], ptr align {{8|4}} [[GEPP0]], i[[sz]] {{8|4}}, i1 false)
134   // CK1: [[BC_ADR:%.+]] = ptrtoint ptr %{{.+}} to i[[sz]]
135   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, ptr [[MAIN_DEP:%.+]], i[[sz]] 0
136   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 0
137   // CK1: store i[[sz]] [[BC_ADR]], ptr [[DEP_ADR]],
138   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 1
139   // CK1: store i[[sz]] 4, ptr [[DEP_SIZE]],
140   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 2
141   // CK1: store i8 3, ptr [[DEP_ATTRS]]
142   // CK1: [[BC_ADR:%.+]] = ptrtoint ptr %{{.+}} to i[[sz]]
143   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, ptr [[MAIN_DEP]], i[[sz]] 1
144   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 0
145   // CK1: store i[[sz]] [[BC_ADR]], ptr [[DEP_ADR]],
146   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 1
147   // CK1: store i[[sz]] 4, ptr [[DEP_SIZE]],
148   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 2
149   // CK1: store i8 3, ptr [[DEP_ATTRS]]
150   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, ptr [[MAIN_DEP]], i[[sz]] 2
151   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 0
152   // CK1: store i[[sz]] ptrtoint (ptr @gc to i[[sz]]), ptr [[DEP_ADR]],
153   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 1
154   // CK1: store i[[sz]] 800, ptr [[DEP_SIZE]],
155   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 2
156   // CK1: store i8 3, ptr [[DEP_ATTRS]]
157   // CK1: call void @__kmpc_omp_taskwait_deps_51(ptr @{{.+}}, i32 %{{.+}}, i32 3, ptr [[MAIN_DEP]], i32 0, ptr null, i32 0)
158   // CK1: call void @__kmpc_omp_task_begin_if0(ptr @{{.+}}, i32 %{{.+}}, ptr [[RES]])
159   // CK1: = call i32 [[TASK_ENTRY2]](i32 %{{.+}}, ptr [[RES]])
160   // CK1: call void @__kmpc_omp_task_complete_if0(ptr @{{.+}}, i32 %{{.+}}, ptr [[RES]])
161 
162   // CK1: br label %[[IFEND:[^,]+]]
163 
164   // CK1: [[IFELSE]]
165   // CK1: br label %[[IFEND]]
166   // CK1: [[IFEND]]
167   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
168   #pragma omp target exit data map(delete: arg) if(arg) device(4) depend(inout: arg, la, gc)
169   {++arg;}
170 
171   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
172   {++arg;}
173 
174   // Region 03
175   // CK1: [[BP0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP:%.+]], i32 0, i32 0
176   // CK1: store ptr [[VLA:%.+]], ptr [[BP0]],
177   // CK1: [[P0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P:%.+]], i32 0, i32 0
178   // CK1: store ptr [[VLA]], ptr [[P0]],
179   // CK1: [[S0:%.+]] = getelementptr inbounds [1 x i64], ptr [[S:%.+]], i32 0, i32 0
180   // CK1: store i64 {{.+}}, ptr [[S0]],
181   // CK1: [[GEPBP0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP]], i32 0, i32 0
182   // CK1: [[GEPP0:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P]], i32 0, i32 0
183   // CK1: [[GEPS0:%.+]] = getelementptr inbounds [1 x i64], ptr [[S]], i32 0, i32 0
184   // CK1: [[RES:%.+]] = call ptr @__kmpc_omp_task_alloc(ptr {{.+}}, i32 {{.+}}, i32 1, i[[sz]] {{64|36}}, i[[sz]] 1, ptr [[TASK_ENTRY3:@.+]])
185   // CK1: [[TASK_T:%.+]] = getelementptr inbounds nuw %struct.kmp_task_t_with_privates{{.+}}, ptr [[RES]], i32 0, i32 0
186   // CK1: [[PRIVS:%.+]] = getelementptr inbounds nuw %struct.kmp_task_t_with_privates{{.+}}, ptr [[RES]], i32 0, i32 1
187   // CK1-64: [[PRIVS_BASEPTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 0
188   // CK1-64: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_BASEPTRS]], ptr align {{8|4}} [[GEPBP0]], i[[sz]] {{8|4}}, i1 false)
189   // CK1-64: [[PRIVS_PTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 1
190   // CK1-64: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_PTRS]], ptr align {{8|4}} [[GEPP0]], i[[sz]] {{8|4}}, i1 false)
191   // CK1-64: [[PRIVS_SIZES:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 2
192   // CK1-64: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_SIZES]], ptr align {{8|4}} [[GEPS0]], i[[sz]] {{8|4}}, i1 false)
193   // CK1-32: [[PRIVS_SIZES:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 0
194   // CK1-32: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_SIZES]], ptr align {{8|4}} [[GEPS0]], i[[sz]] {{8|4}}, i1 false)
195   // CK1-32: [[PRIVS_BASEPTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 1
196   // CK1-32: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_BASEPTRS]], ptr align {{8|4}} [[GEPBP0]], i[[sz]] {{8|4}}, i1 false)
197   // CK1-32: [[PRIVS_PTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 2
198   // CK1-32: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_PTRS]], ptr align {{8|4}} [[GEPP0]], i[[sz]] {{8|4}}, i1 false)
199   // CK1: [[BC_ADR:%.+]] = ptrtoint ptr %{{.+}} to i[[sz]]
200   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, ptr [[MAIN_DEP:%.+]], i[[sz]] 0
201   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 0
202   // CK1: store i[[sz]] [[BC_ADR]], ptr [[DEP_ADR]],
203   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 1
204   // CK1: store i[[sz]] %{{.+}}, ptr [[DEP_SIZE]],
205   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 2
206   // CK1: store i8 3, ptr [[DEP_ATTRS]]
207   // CK1: [[BC_ADR:%.+]] = ptrtoint ptr %{{.+}} to i[[sz]]
208   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, ptr [[MAIN_DEP]], i[[sz]] 1
209   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 0
210   // CK1: store i[[sz]] [[BC_ADR]], ptr [[DEP_ADR]],
211   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 1
212   // CK1: store i[[sz]] 4, ptr [[DEP_SIZE]],
213   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 2
214   // CK1: store i8 3, ptr [[DEP_ATTRS]]
215   // CK1: [[BC_ADR:%.+]] = ptrtoint ptr %{{.+}} to i[[sz]]
216   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, ptr [[MAIN_DEP]], i[[sz]] 2
217   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 0
218   // CK1: store i[[sz]] [[BC_ADR]], ptr [[DEP_ADR]],
219   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 1
220   // CK1: store i[[sz]] 4, ptr [[DEP_SIZE]],
221   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 2
222   // CK1: store i8 3, ptr [[DEP_ATTRS]]
223   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, ptr [[MAIN_DEP]], i[[sz]] 3
224   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 0
225   // CK1: store i[[sz]] ptrtoint (ptr @gc to i[[sz]]), ptr [[DEP_ADR]],
226   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 1
227   // CK1: store i[[sz]] 800, ptr [[DEP_SIZE]],
228   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 2
229   // CK1: store i8 3, ptr [[DEP_ATTRS]]
230   // CK1: call void @__kmpc_omp_taskwait_deps_51(ptr @{{.+}}, i32 %{{.+}}, i32 4, ptr [[MAIN_DEP]], i32 0, ptr null, i32 0)
231   // CK1: call void @__kmpc_omp_task_begin_if0(ptr @{{.+}}, i32 %{{.+}}, ptr [[RES]])
232   // CK1: = call i32 [[TASK_ENTRY3]](i32 %{{.+}}, ptr [[RES]])
233   // CK1: call void @__kmpc_omp_task_complete_if0(ptr @{{.+}}, i32 %{{.+}}, ptr [[RES]])
234   #pragma omp target exit data map(from:lb) depend(out: lb, arg, la, gc)
235   {++arg;}
236 
237   // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
238   {++arg;}
239 
240   // Region 04
241   // CK1: [[DIV:%.+]] = sdiv exact i64 sub (i64 ptrtoint (ptr getelementptr (ptr, ptr getelementptr inbounds nuw (%struct.ST, ptr @gb, i32 0, i32 1), i32 1) to i64), i64 ptrtoint (ptr getelementptr inbounds nuw (%struct.ST, ptr @gb, i32 0, i32 1) to i64)), ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
242   // CK1: [[BP0:%.+]] = getelementptr inbounds [2 x ptr], ptr [[BP:%.+]], i32 0, i32 0
243   // CK1: store ptr @gb, ptr [[BP0]],
244   // CK1: [[P0:%.+]] = getelementptr inbounds [2 x ptr], ptr [[P:%.+]], i32 0, i32 0
245   // CK1: store ptr getelementptr inbounds nuw (%struct.ST, ptr @gb, i32 0, i32 1), ptr [[P0]],
246   // CK1: [[PS0:%.+]] = getelementptr inbounds [2 x i64], ptr [[PS:%.+]], i32 0, i32 0
247   // CK1: store i64 [[DIV]], ptr [[PS0]],
248   // CK1: [[BP1:%.+]] = getelementptr inbounds [2 x ptr], ptr [[BP]], i32 0, i32 1
249   // CK1: store ptr getelementptr inbounds nuw (%struct.ST, ptr @gb, i32 0, i32 1), ptr [[BP1]],
250   // CK1: [[P1:%.+]] = getelementptr inbounds [2 x ptr], ptr [[P]], i32 0, i32 1
251   // CK1: store ptr %{{.+}}, ptr [[P1]],
252   // CK1: [[GEPBP0:%.+]] = getelementptr inbounds [2 x ptr], ptr [[BP]], i32 0, i32 0
253   // CK1: [[GEPP0:%.+]] = getelementptr inbounds [2 x ptr], ptr [[P]], i32 0, i32 0
254   // CK1: [[GEPS0:%.+]] = getelementptr inbounds [2 x i64], ptr [[PS]], i32 0, i32 0
255   // CK1: [[RES:%.+]] = call ptr @__kmpc_omp_task_alloc(ptr {{.+}}, i32 {{.+}}, i32 1, i[[sz]] {{88|52}}, i[[sz]] 1, ptr [[TASK_ENTRY4:@.+]])
256   // CK1: [[TASK_T:%.+]] = getelementptr inbounds nuw %struct.kmp_task_t_with_privates{{.+}}, ptr [[RES]], i32 0, i32 0
257   // CK1: [[PRIVS:%.+]] = getelementptr inbounds nuw %struct.kmp_task_t_with_privates{{.+}}, ptr [[RES]], i32 0, i32 1
258   // CK1-64: [[PRIVS_BASEPTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 0
259   // CK1-64: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_BASEPTRS]], ptr align {{8|4}} [[GEPBP0]], i[[sz]] {{16|8}}, i1 false)
260   // CK1-64: [[PRIVS_PTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 1
261   // CK1-64: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_PTRS]], ptr align {{8|4}} [[GEPP0]], i[[sz]] {{16|8}}, i1 false)
262   // CK1-64: [[PRIVS_SIZES:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 2
263   // CK1-64: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_SIZES]], ptr align {{8|4}} [[GEPS0]], i[[sz]] {{16|8}}, i1 false)
264   // CK1-32: [[PRIVS_SIZES:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 0
265   // CK1-32: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_SIZES]], ptr align {{8|4}} [[GEPS0]], i[[sz]] {{16|8}}, i1 false)
266   // CK1-32: [[PRIVS_BASEPTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 1
267   // CK1-32: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_BASEPTRS]], ptr align {{8|4}} [[GEPBP0]], i[[sz]] {{16|8}}, i1 false)
268   // CK1-32: [[PRIVS_PTRS:%.+]] = getelementptr inbounds nuw %struct..kmp_privates.t{{.+}}, ptr [[PRIVS]], i32 0, i32 2
269   // CK1-32: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{8|4}} [[PRIVS_PTRS]], ptr align {{8|4}} [[GEPP0]], i[[sz]] {{16|8}}, i1 false)
270   // CK1: %{{.+}} = sub nuw
271   // CK1: [[BC_ADR:%.+]] = ptrtoint ptr %{{.+}} to i[[sz]]
272   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, ptr [[MAIN_DEP:%.+]], i[[sz]] 0
273   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 0
274   // CK1: store i[[sz]] [[BC_ADR]], ptr [[DEP_ADR]],
275   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 1
276   // CK1: store i[[sz]] %{{.+}}, ptr [[DEP_SIZE]],
277   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 2
278   // CK1: store i8 1, ptr [[DEP_ATTRS]]
279   // CK1: [[BC_ADR:%.+]] = ptrtoint ptr %{{.+}} to i[[sz]]
280   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, ptr [[MAIN_DEP]], i[[sz]] 1
281   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 0
282   // CK1: store i[[sz]] [[BC_ADR]], ptr [[DEP_ADR]],
283   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 1
284   // CK1: store i[[sz]] 4, ptr [[DEP_SIZE]],
285   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 2
286   // CK1: store i8 1, ptr [[DEP_ATTRS]]
287   // CK1: [[BC_ADR:%.+]] = ptrtoint ptr %{{.+}} to i[[sz]]
288   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, ptr [[MAIN_DEP]], i[[sz]] 2
289   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 0
290   // CK1: store i[[sz]] [[BC_ADR]], ptr [[DEP_ADR]],
291   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 1
292   // CK1: store i[[sz]] %{{.+}}, ptr [[DEP_SIZE]],
293   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 2
294   // CK1: store i8 1, ptr [[DEP_ATTRS]]
295   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, ptr [[MAIN_DEP]], i[[sz]] 3
296   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 0
297   // CK1: store i[[sz]] ptrtoint (ptr @gc to i[[sz]]), ptr [[DEP_ADR]],
298   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 1
299   // CK1: store i[[sz]] 800, ptr [[DEP_SIZE]],
300   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 2
301   // CK1: store i8 1, ptr [[DEP_ATTRS]]
302   // CK1: [[BC_ADR:%.+]] = ptrtoint ptr %{{.+}} to i[[sz]]
303   // CK1: [[DEP:%.+]] = getelementptr %struct.kmp_depend_info, ptr [[MAIN_DEP]], i[[sz]] 4
304   // CK1: [[DEP_ADR:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 0
305   // CK1: store i[[sz]] [[BC_ADR]], ptr [[DEP_ADR]],
306   // CK1: [[DEP_SIZE:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 1
307   // CK1: store i[[sz]] 4, ptr [[DEP_SIZE]],
308   // CK1: [[DEP_ATTRS:%.+]] = getelementptr inbounds nuw %struct.kmp_depend_info, ptr [[DEP]], i32 0, i32 2
309   // CK1: store i8 1, ptr [[DEP_ATTRS]]
310   // CK1: call void @__kmpc_omp_taskwait_deps_51(ptr @{{.+}}, i32 %{{.+}}, i32 5, ptr [[MAIN_DEP]], i32 0, ptr null, i32 0)
311   // CK1: call void @__kmpc_omp_task_begin_if0(ptr @{{.+}}, i32 %{{.+}}, ptr [[RES]])
312   // CK1: = call i32 [[TASK_ENTRY4]](i32 %{{.+}}, ptr [[RES]])
313   // CK1: call void @__kmpc_omp_task_complete_if0(ptr @{{.+}}, i32 %{{.+}}, ptr [[RES]])
314   #pragma omp target exit data map(from:gb.b[:3]) depend(in: gb.b[:3], la, lb, gc, arg)
315   {++arg;}
316 }
317 
318 // CK1: define internal{{.*}} i32 [[TASK_ENTRY0]](i32{{.*}}, ptr noalias noundef %1)
319 // CK1-DAG: call void @__tgt_target_data_end_nowait_mapper(ptr @{{.+}}, i64 [[DEV:%[^,]+]], i32 1, ptr [[BP:%.+]], ptr [[P:%.+]], ptr [[S:%.+]], ptr [[MTYPE00]]{{.+}}, ptr null)
320 // CK1-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
321 // CK1-DAG: [[DEVi32]] = load i32, ptr %{{[^,]+}},
322 // CK1-DAG: [[BP]] = load ptr, ptr [[BP_PRIV:%[^,]+]],
323 // CK1-DAG: [[P]] = load ptr, ptr [[P_PRIV:%[^,]+]],
324 // CK1-DAG: [[S]] = load ptr, ptr [[S_PRIV:%[^,]+]],
325 // CK1-DAG: call void {{%.*}}(ptr %{{[^,]+}}, ptr [[BP_PRIV]], ptr [[P_PRIV]], ptr [[S_PRIV]])
326 // CK1: ret i32 0
327 // CK1: }
328 
329 // CK1: define internal{{.*}} i32 [[TASK_ENTRY2]](i32{{.*}}, ptr noalias noundef %1)
330 // CK1-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 4, i32 1, ptr [[BP:%.+]], ptr [[P:%.+]], ptr [[S:%.+]], ptr [[MTYPE02]]{{.+}}, ptr null)
331 // CK1-DAG: [[BP]] = load ptr, ptr [[BP_PRIV:%[^,]+]],
332 // CK1-DAG: [[P]] = load ptr, ptr [[P_PRIV:%[^,]+]],
333 // CK1-DAG: [[S]] = load ptr, ptr [[S_PRIV:%[^,]+]],
334 // CK1-DAG: call void {{%.*}}(ptr %{{[^,]+}}, ptr [[BP_PRIV]], ptr [[P_PRIV]], ptr [[S_PRIV]])
335 // CK1: ret i32 0
336 // CK1: }
337 
338 // CK1: define internal{{.*}} i32 [[TASK_ENTRY3]](i32{{.*}}, ptr noalias noundef %1)
339 // CK1-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[BP:%.+]], ptr [[P:%.+]], ptr [[S:%.+]], ptr [[MTYPE03]]{{.+}}, ptr null)
340 // CK1-DAG: [[BP]] = load ptr, ptr [[BP_PRIV:%[^,]+]],
341 // CK1-DAG: [[P]] = load ptr, ptr [[P_PRIV:%[^,]+]],
342 // CK1-DAG: [[S]] = load ptr, ptr [[S_PRIV:%[^,]+]],
343 // CK1-DAG: call void {{%.*}}(ptr %{{[^,]+}}, ptr [[BP_PRIV]], ptr [[P_PRIV]], ptr [[S_PRIV]])
344 // CK1-NOT: __tgt_target_data_end_mapper
345 // CK1: ret i32 0
346 // CK1: }
347 
348 // CK1: define internal{{.*}} i32 [[TASK_ENTRY4]](i32{{.*}}, ptr noalias noundef %1)
349 // CK1-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 -1, i32 2, ptr [[BP:%.+]], ptr [[P:%.+]], ptr [[S:%.+]], ptr [[MTYPE04]]{{.+}}, ptr null)
350 // CK1-DAG: [[BP]] = load ptr, ptr [[BP_PRIV:%[^,]+]],
351 // CK1-DAG: [[P]] = load ptr, ptr [[P_PRIV:%[^,]+]],
352 // CK1-DAG: [[S]] = load ptr, ptr [[S_PRIV:%[^,]+]],
353 // CK1-DAG: call void {{%.*}}(ptr %{{[^,]+}}, ptr [[BP_PRIV]], ptr [[P_PRIV]], ptr [[S_PRIV]])
354 // CK1-NOT: __tgt_target_data_end_mapper
355 // CK1: ret i32 0
356 // CK1: }
357 
358 #endif
359