1 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _ 2 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s 3 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 4 // RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s 5 6 // expected-no-diagnostics 7 #ifndef HEADER 8 #define HEADER 9 10 typedef struct { 11 int a; 12 } C; 13 #pragma omp declare mapper(C s) map(to : s.a) 14 15 typedef struct { 16 int e; 17 C f; 18 int h; 19 } D; 20 21 void foo() { 22 D sa[10]; 23 sa[1].e = 111; 24 sa[1].f.a = 222; 25 26 #pragma omp target map(tofrom : sa [0:2]) 27 { 28 sa[1].e = 333; 29 sa[1].f.a = 444; 30 } 31 } 32 #endif 33 // CHECK-LABEL: define {{[^@]+}}@_Z3foov 34 // CHECK-SAME: () #[[ATTR0:[0-9]+]] { 35 // CHECK-NEXT: entry: 36 // CHECK-NEXT: [[SA:%.*]] = alloca [10 x %struct.D], align 4 37 // CHECK-NEXT: [[DOTOFFLOAD_BASEPTRS:%.*]] = alloca [1 x ptr], align 8 38 // CHECK-NEXT: [[DOTOFFLOAD_PTRS:%.*]] = alloca [1 x ptr], align 8 39 // CHECK-NEXT: [[DOTOFFLOAD_MAPPERS:%.*]] = alloca [1 x ptr], align 8 40 // CHECK-NEXT: [[KERNEL_ARGS:%.*]] = alloca [[STRUCT___TGT_KERNEL_ARGUMENTS:%.*]], align 8 41 // CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x %struct.D], ptr [[SA]], i64 0, i64 1 42 // CHECK-NEXT: [[E:%.*]] = getelementptr inbounds nuw [[STRUCT_D:%.*]], ptr [[ARRAYIDX]], i32 0, i32 0 43 // CHECK-NEXT: store i32 111, ptr [[E]], align 4 44 // CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds [10 x %struct.D], ptr [[SA]], i64 0, i64 1 45 // CHECK-NEXT: [[F:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[ARRAYIDX1]], i32 0, i32 1 46 // CHECK-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_C:%.*]], ptr [[F]], i32 0, i32 0 47 // CHECK-NEXT: store i32 222, ptr [[A]], align 4 48 // CHECK-NEXT: [[ARRAYIDX2:%.*]] = getelementptr inbounds nuw [10 x %struct.D], ptr [[SA]], i64 0, i64 0 49 // CHECK-NEXT: [[TMP0:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 50 // CHECK-NEXT: store ptr [[SA]], ptr [[TMP0]], align 8 51 // CHECK-NEXT: [[TMP1:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 52 // CHECK-NEXT: store ptr [[ARRAYIDX2]], ptr [[TMP1]], align 8 53 // CHECK-NEXT: [[TMP2:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_MAPPERS]], i64 0, i64 0 54 // CHECK-NEXT: store ptr @.omp_mapper._ZTS1D.default, ptr [[TMP2]], align 8 55 // CHECK-NEXT: [[TMP3:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_BASEPTRS]], i32 0, i32 0 56 // CHECK-NEXT: [[TMP4:%.*]] = getelementptr inbounds [1 x ptr], ptr [[DOTOFFLOAD_PTRS]], i32 0, i32 0 57 // CHECK-NEXT: [[TMP5:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 0 58 // CHECK-NEXT: store i32 3, ptr [[TMP5]], align 4 59 // CHECK-NEXT: [[TMP6:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 1 60 // CHECK-NEXT: store i32 1, ptr [[TMP6]], align 4 61 // CHECK-NEXT: [[TMP7:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 2 62 // CHECK-NEXT: store ptr [[TMP3]], ptr [[TMP7]], align 8 63 // CHECK-NEXT: [[TMP8:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 3 64 // CHECK-NEXT: store ptr [[TMP4]], ptr [[TMP8]], align 8 65 // CHECK-NEXT: [[TMP9:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 4 66 // CHECK-NEXT: store ptr @.offload_sizes, ptr [[TMP9]], align 8 67 // CHECK-NEXT: [[TMP10:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 5 68 // CHECK-NEXT: store ptr @.offload_maptypes, ptr [[TMP10]], align 8 69 // CHECK-NEXT: [[TMP11:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 6 70 // CHECK-NEXT: store ptr null, ptr [[TMP11]], align 8 71 // CHECK-NEXT: [[TMP12:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 7 72 // CHECK-NEXT: store ptr [[DOTOFFLOAD_MAPPERS]], ptr [[TMP12]], align 8 73 // CHECK-NEXT: [[TMP13:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 8 74 // CHECK-NEXT: store i64 0, ptr [[TMP13]], align 8 75 // CHECK-NEXT: [[TMP14:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 9 76 // CHECK-NEXT: store i64 0, ptr [[TMP14]], align 8 77 // CHECK-NEXT: [[TMP15:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 10 78 // CHECK-NEXT: store [3 x i32] [i32 -1, i32 0, i32 0], ptr [[TMP15]], align 4 79 // CHECK-NEXT: [[TMP16:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 11 80 // CHECK-NEXT: store [3 x i32] zeroinitializer, ptr [[TMP16]], align 4 81 // CHECK-NEXT: [[TMP17:%.*]] = getelementptr inbounds nuw [[STRUCT___TGT_KERNEL_ARGUMENTS]], ptr [[KERNEL_ARGS]], i32 0, i32 12 82 // CHECK-NEXT: store i32 0, ptr [[TMP17]], align 4 83 // CHECK-NEXT: [[TMP18:%.*]] = call i32 @__tgt_target_kernel(ptr @[[GLOB1:[0-9]+]], i64 -1, i32 -1, i32 0, ptr @.{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l26.region_id, ptr [[KERNEL_ARGS]]) 84 // CHECK-NEXT: [[TMP19:%.*]] = icmp ne i32 [[TMP18]], 0 85 // CHECK-NEXT: br i1 [[TMP19]], label [[OMP_OFFLOAD_FAILED:%.*]], label [[OMP_OFFLOAD_CONT:%.*]] 86 // CHECK: omp_offload.failed: 87 // CHECK-NEXT: call void @{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l26(ptr [[SA]]) #[[ATTR3:[0-9]+]] 88 // CHECK-NEXT: br label [[OMP_OFFLOAD_CONT]] 89 // CHECK: omp_offload.cont: 90 // CHECK-NEXT: ret void 91 // 92 // 93 // CHECK-LABEL: define {{[^@]+}}@{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l26 94 // CHECK-SAME: (ptr noundef nonnull align 4 dereferenceable(120) [[SA:%.*]]) #[[ATTR1:[0-9]+]] { 95 // CHECK-NEXT: entry: 96 // CHECK-NEXT: [[SA_ADDR:%.*]] = alloca ptr, align 8 97 // CHECK-NEXT: store ptr [[SA]], ptr [[SA_ADDR]], align 8 98 // CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[SA_ADDR]], align 8 99 // CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds [10 x %struct.D], ptr [[TMP0]], i64 0, i64 1 100 // CHECK-NEXT: [[E:%.*]] = getelementptr inbounds nuw [[STRUCT_D:%.*]], ptr [[ARRAYIDX]], i32 0, i32 0 101 // CHECK-NEXT: store i32 333, ptr [[E]], align 4 102 // CHECK-NEXT: [[ARRAYIDX1:%.*]] = getelementptr inbounds [10 x %struct.D], ptr [[TMP0]], i64 0, i64 1 103 // CHECK-NEXT: [[F:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[ARRAYIDX1]], i32 0, i32 1 104 // CHECK-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_C:%.*]], ptr [[F]], i32 0, i32 0 105 // CHECK-NEXT: store i32 444, ptr [[A]], align 4 106 // CHECK-NEXT: ret void 107 // 108 // 109 // CHECK-LABEL: define {{[^@]+}}@.omp_mapper._ZTS1D.default 110 // CHECK-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]], ptr noundef [[TMP2:%.*]], i64 noundef [[TMP3:%.*]], i64 noundef [[TMP4:%.*]], ptr noundef [[TMP5:%.*]]) #[[ATTR2:[0-9]+]] { 111 // CHECK-NEXT: entry: 112 // CHECK-NEXT: [[TMP10:%.*]] = udiv exact i64 [[TMP3]], 12 113 // CHECK-NEXT: [[TMP11:%.*]] = getelementptr [[STRUCT_D:%.*]], ptr [[TMP2]], i64 [[TMP10]] 114 // CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY:%.*]] = icmp sgt i64 [[TMP10]], 1 115 // CHECK-NEXT: [[TMP14:%.*]] = and i64 [[TMP4]], 8 116 // CHECK-NEXT: [[TMP15:%.*]] = icmp ne ptr [[TMP1]], [[TMP2]] 117 // CHECK-NEXT: [[TMP16:%.*]] = and i64 [[TMP4]], 16 118 // CHECK-NEXT: [[TMP17:%.*]] = icmp ne i64 [[TMP16]], 0 119 // CHECK-NEXT: [[TMP18:%.*]] = and i1 [[TMP15]], [[TMP17]] 120 // CHECK-NEXT: [[TMP19:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP18]] 121 // CHECK-NEXT: [[DOTOMP_ARRAY__INIT__DELETE:%.*]] = icmp eq i64 [[TMP14]], 0 122 // CHECK-NEXT: [[TMP20:%.*]] = and i1 [[TMP19]], [[DOTOMP_ARRAY__INIT__DELETE]] 123 // CHECK-NEXT: br i1 [[TMP20]], label [[DOTOMP_ARRAY__INIT:%.*]], label [[OMP_ARRAYMAP_HEAD:%.*]] 124 // CHECK: .omp.array..init: 125 // CHECK-NEXT: [[TMP21:%.*]] = mul nuw i64 [[TMP10]], 12 126 // CHECK-NEXT: [[TMP22:%.*]] = and i64 [[TMP4]], -4 127 // CHECK-NEXT: [[TMP23:%.*]] = or i64 [[TMP22]], 512 128 // CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP21]], i64 [[TMP23]], ptr [[TMP5]]) 129 // CHECK-NEXT: br label [[OMP_ARRAYMAP_HEAD]] 130 // CHECK: omp.arraymap.head: 131 // CHECK-NEXT: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP2]], [[TMP11]] 132 // CHECK-NEXT: br i1 [[OMP_ARRAYMAP_ISEMPTY]], label [[OMP_DONE:%.*]], label [[OMP_ARRAYMAP_BODY:%.*]] 133 // CHECK: omp.arraymap.body: 134 // CHECK-NEXT: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END25:%.*]] ] 135 // CHECK-NEXT: [[E:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0 136 // CHECK-NEXT: [[F:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 1 137 // CHECK-NEXT: [[H:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 2 138 // CHECK-NEXT: [[TMP24:%.*]] = getelementptr i32, ptr [[H]], i32 1 139 // CHECK-NEXT: [[TMP25:%.*]] = ptrtoint ptr [[TMP24]] to i64 140 // CHECK-NEXT: [[TMP26:%.*]] = ptrtoint ptr [[E]] to i64 141 // CHECK-NEXT: [[TMP27:%.*]] = sub i64 [[TMP25]], [[TMP26]] 142 // CHECK-NEXT: [[TMP28:%.*]] = sdiv exact i64 [[TMP27]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64) 143 // CHECK-NEXT: [[TMP29:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP0]]) 144 // CHECK-NEXT: [[TMP30:%.*]] = shl i64 [[TMP29]], 48 145 // CHECK-NEXT: [[TMP31:%.*]] = add nuw i64 0, [[TMP30]] 146 // CHECK-NEXT: [[TMP32:%.*]] = and i64 [[TMP4]], 3 147 // CHECK-NEXT: [[TMP33:%.*]] = icmp eq i64 [[TMP32]], 0 148 // CHECK-NEXT: br i1 [[TMP33]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]] 149 // CHECK: omp.type.alloc: 150 // CHECK-NEXT: [[TMP34:%.*]] = and i64 [[TMP31]], -4 151 // CHECK-NEXT: br label [[OMP_TYPE_END:%.*]] 152 // CHECK: omp.type.alloc.else: 153 // CHECK-NEXT: [[TMP35:%.*]] = icmp eq i64 [[TMP32]], 1 154 // CHECK-NEXT: br i1 [[TMP35]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]] 155 // CHECK: omp.type.to: 156 // CHECK-NEXT: [[TMP36:%.*]] = and i64 [[TMP31]], -3 157 // CHECK-NEXT: br label [[OMP_TYPE_END]] 158 // CHECK: omp.type.to.else: 159 // CHECK-NEXT: [[TMP37:%.*]] = icmp eq i64 [[TMP32]], 2 160 // CHECK-NEXT: br i1 [[TMP37]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]] 161 // CHECK: omp.type.from: 162 // CHECK-NEXT: [[TMP38:%.*]] = and i64 [[TMP31]], -2 163 // CHECK-NEXT: br label [[OMP_TYPE_END]] 164 // CHECK: omp.type.end: 165 // CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP34]], [[OMP_TYPE_ALLOC]] ], [ [[TMP36]], [[OMP_TYPE_TO]] ], [ [[TMP38]], [[OMP_TYPE_FROM]] ], [ [[TMP31]], [[OMP_TYPE_TO_ELSE]] ] 166 // CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 [[TMP28]], i64 [[OMP_MAPTYPE]], ptr null) 167 // CHECK-NEXT: [[TMP39:%.*]] = add nuw i64 281474976711171, [[TMP30]] 168 // CHECK-NEXT: [[TMP40:%.*]] = and i64 [[TMP4]], 3 169 // CHECK-NEXT: [[TMP41:%.*]] = icmp eq i64 [[TMP40]], 0 170 // CHECK-NEXT: br i1 [[TMP41]], label [[OMP_TYPE_ALLOC6:%.*]], label [[OMP_TYPE_ALLOC_ELSE7:%.*]] 171 // CHECK: omp.type.alloc1: 172 // CHECK-NEXT: [[TMP42:%.*]] = and i64 [[TMP39]], -4 173 // CHECK-NEXT: br label [[OMP_TYPE_END11:%.*]] 174 // CHECK: omp.type.alloc.else2: 175 // CHECK-NEXT: [[TMP43:%.*]] = icmp eq i64 [[TMP40]], 1 176 // CHECK-NEXT: br i1 [[TMP43]], label [[OMP_TYPE_TO8:%.*]], label [[OMP_TYPE_TO_ELSE9:%.*]] 177 // CHECK: omp.type.to3: 178 // CHECK-NEXT: [[TMP44:%.*]] = and i64 [[TMP39]], -3 179 // CHECK-NEXT: br label [[OMP_TYPE_END11]] 180 // CHECK: omp.type.to.else4: 181 // CHECK-NEXT: [[TMP45:%.*]] = icmp eq i64 [[TMP40]], 2 182 // CHECK-NEXT: br i1 [[TMP45]], label [[OMP_TYPE_FROM10:%.*]], label [[OMP_TYPE_END11]] 183 // CHECK: omp.type.from5: 184 // CHECK-NEXT: [[TMP46:%.*]] = and i64 [[TMP39]], -2 185 // CHECK-NEXT: br label [[OMP_TYPE_END11]] 186 // CHECK: omp.type.end6: 187 // CHECK-NEXT: [[OMP_MAPTYPE12:%.*]] = phi i64 [ [[TMP42]], [[OMP_TYPE_ALLOC6]] ], [ [[TMP44]], [[OMP_TYPE_TO8]] ], [ [[TMP46]], [[OMP_TYPE_FROM10]] ], [ [[TMP39]], [[OMP_TYPE_TO_ELSE9]] ] 188 // CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 4, i64 [[OMP_MAPTYPE12]], ptr null) 189 // CHECK-NEXT: [[TMP47:%.*]] = add nuw i64 281474976711171, [[TMP30]] 190 // CHECK-NEXT: [[TMP48:%.*]] = and i64 [[TMP4]], 3 191 // CHECK-NEXT: [[TMP49:%.*]] = icmp eq i64 [[TMP48]], 0 192 // CHECK-NEXT: br i1 [[TMP49]], label [[OMP_TYPE_ALLOC13:%.*]], label [[OMP_TYPE_ALLOC_ELSE14:%.*]] 193 // CHECK: omp.type.alloc8: 194 // CHECK-NEXT: [[TMP50:%.*]] = and i64 [[TMP47]], -4 195 // CHECK-NEXT: br label [[OMP_TYPE_END18:%.*]] 196 // CHECK: omp.type.alloc.else9: 197 // CHECK-NEXT: [[TMP51:%.*]] = icmp eq i64 [[TMP48]], 1 198 // CHECK-NEXT: br i1 [[TMP51]], label [[OMP_TYPE_TO15:%.*]], label [[OMP_TYPE_TO_ELSE16:%.*]] 199 // CHECK: omp.type.to10: 200 // CHECK-NEXT: [[TMP52:%.*]] = and i64 [[TMP47]], -3 201 // CHECK-NEXT: br label [[OMP_TYPE_END18]] 202 // CHECK: omp.type.to.else11: 203 // CHECK-NEXT: [[TMP53:%.*]] = icmp eq i64 [[TMP48]], 2 204 // CHECK-NEXT: br i1 [[TMP53]], label [[OMP_TYPE_FROM17:%.*]], label [[OMP_TYPE_END18]] 205 // CHECK: omp.type.from12: 206 // CHECK-NEXT: [[TMP54:%.*]] = and i64 [[TMP47]], -2 207 // CHECK-NEXT: br label [[OMP_TYPE_END18]] 208 // CHECK: omp.type.end13: 209 // CHECK-NEXT: [[OMP_MAPTYPE19:%.*]] = phi i64 [ [[TMP50]], [[OMP_TYPE_ALLOC13]] ], [ [[TMP52]], [[OMP_TYPE_TO15]] ], [ [[TMP54]], [[OMP_TYPE_FROM17]] ], [ [[TMP47]], [[OMP_TYPE_TO_ELSE16]] ] 210 // CHECK-NEXT: call void @.omp_mapper._ZTS1C.default(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[F]], i64 4, i64 [[OMP_MAPTYPE19]], ptr null) #[[ATTR3]] 211 // CHECK-NEXT: [[TMP55:%.*]] = add nuw i64 281474976711171, [[TMP30]] 212 // CHECK-NEXT: [[TMP56:%.*]] = and i64 [[TMP4]], 3 213 // CHECK-NEXT: [[TMP57:%.*]] = icmp eq i64 [[TMP56]], 0 214 // CHECK-NEXT: br i1 [[TMP57]], label [[OMP_TYPE_ALLOC20:%.*]], label [[OMP_TYPE_ALLOC_ELSE21:%.*]] 215 // CHECK: omp.type.alloc15: 216 // CHECK-NEXT: [[TMP58:%.*]] = and i64 [[TMP55]], -4 217 // CHECK-NEXT: br label [[OMP_TYPE_END25]] 218 // CHECK: omp.type.alloc.else16: 219 // CHECK-NEXT: [[TMP59:%.*]] = icmp eq i64 [[TMP56]], 1 220 // CHECK-NEXT: br i1 [[TMP59]], label [[OMP_TYPE_TO22:%.*]], label [[OMP_TYPE_TO_ELSE23:%.*]] 221 // CHECK: omp.type.to17: 222 // CHECK-NEXT: [[TMP60:%.*]] = and i64 [[TMP55]], -3 223 // CHECK-NEXT: br label [[OMP_TYPE_END25]] 224 // CHECK: omp.type.to.else18: 225 // CHECK-NEXT: [[TMP61:%.*]] = icmp eq i64 [[TMP56]], 2 226 // CHECK-NEXT: br i1 [[TMP61]], label [[OMP_TYPE_FROM24:%.*]], label [[OMP_TYPE_END25]] 227 // CHECK: omp.type.from19: 228 // CHECK-NEXT: [[TMP62:%.*]] = and i64 [[TMP55]], -2 229 // CHECK-NEXT: br label [[OMP_TYPE_END25]] 230 // CHECK: omp.type.end20: 231 // CHECK-NEXT: [[OMP_MAPTYPE26:%.*]] = phi i64 [ [[TMP58]], [[OMP_TYPE_ALLOC20]] ], [ [[TMP60]], [[OMP_TYPE_TO22]] ], [ [[TMP62]], [[OMP_TYPE_FROM24]] ], [ [[TMP55]], [[OMP_TYPE_TO_ELSE23]] ] 232 // CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[H]], i64 4, i64 [[OMP_MAPTYPE26]], ptr null) 233 // CHECK-NEXT: [[OMP_ARRAYMAP_NEXT]] = getelementptr [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 1 234 // CHECK-NEXT: [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr [[OMP_ARRAYMAP_NEXT]], [[TMP11]] 235 // CHECK-NEXT: br i1 [[OMP_ARRAYMAP_ISDONE]], label [[OMP_ARRAYMAP_EXIT:%.*]], label [[OMP_ARRAYMAP_BODY]] 236 // CHECK: omp.arraymap.exit: 237 // CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY27:%.*]] = icmp sgt i64 [[TMP10]], 1 238 // CHECK-NEXT: [[TMP63:%.*]] = and i64 [[TMP4]], 8 239 // CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP63]], 0 240 // CHECK-NEXT: [[TMP64:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY27]], [[DOTOMP_ARRAY__DEL__DELETE]] 241 // CHECK-NEXT: br i1 [[TMP64]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]] 242 // CHECK: .omp.array..del: 243 // CHECK-NEXT: [[TMP65:%.*]] = mul nuw i64 [[TMP10]], 12 244 // CHECK-NEXT: [[TMP66:%.*]] = and i64 [[TMP4]], -4 245 // CHECK-NEXT: [[TMP67:%.*]] = or i64 [[TMP66]], 512 246 // CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP65]], i64 [[TMP67]], ptr [[TMP5]]) 247 // CHECK-NEXT: br label [[OMP_DONE]] 248 // CHECK: omp.done: 249 // CHECK-NEXT: ret void 250 // 251 // 252 // CHECK-LABEL: define {{[^@]+}}@.omp_mapper._ZTS1C.default 253 // CHECK-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]], ptr noundef [[TMP2:%.*]], i64 noundef [[TMP3:%.*]], i64 noundef [[TMP4:%.*]], ptr noundef [[TMP5:%.*]]) #[[ATTR2]] { 254 // CHECK-NEXT: entry: 255 // CHECK-NEXT: [[TMP10:%.*]] = udiv exact i64 [[TMP3]], 4 256 // CHECK-NEXT: [[TMP11:%.*]] = getelementptr [[STRUCT_C:%.*]], ptr [[TMP2]], i64 [[TMP10]] 257 // CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY:%.*]] = icmp sgt i64 [[TMP10]], 1 258 // CHECK-NEXT: [[TMP14:%.*]] = and i64 [[TMP4]], 8 259 // CHECK-NEXT: [[TMP15:%.*]] = icmp ne ptr [[TMP1]], [[TMP2]] 260 // CHECK-NEXT: [[TMP16:%.*]] = and i64 [[TMP4]], 16 261 // CHECK-NEXT: [[TMP17:%.*]] = icmp ne i64 [[TMP16]], 0 262 // CHECK-NEXT: [[TMP18:%.*]] = and i1 [[TMP15]], [[TMP17]] 263 // CHECK-NEXT: [[TMP19:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP18]] 264 // CHECK-NEXT: [[DOTOMP_ARRAY__INIT__DELETE:%.*]] = icmp eq i64 [[TMP14]], 0 265 // CHECK-NEXT: [[TMP20:%.*]] = and i1 [[TMP19]], [[DOTOMP_ARRAY__INIT__DELETE]] 266 // CHECK-NEXT: br i1 [[TMP20]], label [[DOTOMP_ARRAY__INIT:%.*]], label [[OMP_ARRAYMAP_HEAD:%.*]] 267 // CHECK: .omp.array..init: 268 // CHECK-NEXT: [[TMP21:%.*]] = mul nuw i64 [[TMP10]], 4 269 // CHECK-NEXT: [[TMP22:%.*]] = and i64 [[TMP4]], -4 270 // CHECK-NEXT: [[TMP23:%.*]] = or i64 [[TMP22]], 512 271 // CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP21]], i64 [[TMP23]], ptr [[TMP5]]) 272 // CHECK-NEXT: br label [[OMP_ARRAYMAP_HEAD]] 273 // CHECK: omp.arraymap.head: 274 // CHECK-NEXT: [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP2]], [[TMP11]] 275 // CHECK-NEXT: br i1 [[OMP_ARRAYMAP_ISEMPTY]], label [[OMP_DONE:%.*]], label [[OMP_ARRAYMAP_BODY:%.*]] 276 // CHECK: omp.arraymap.body: 277 // CHECK-NEXT: [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP2]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END:%.*]] ] 278 // CHECK-NEXT: [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_C]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0 279 // CHECK-NEXT: [[TMP24:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP0]]) 280 // CHECK-NEXT: [[TMP25:%.*]] = shl i64 [[TMP24]], 48 281 // CHECK-NEXT: [[TMP26:%.*]] = add nuw i64 1, [[TMP25]] 282 // CHECK-NEXT: [[TMP27:%.*]] = and i64 [[TMP4]], 3 283 // CHECK-NEXT: [[TMP28:%.*]] = icmp eq i64 [[TMP27]], 0 284 // CHECK-NEXT: br i1 [[TMP28]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]] 285 // CHECK: omp.type.alloc: 286 // CHECK-NEXT: [[TMP29:%.*]] = and i64 [[TMP26]], -4 287 // CHECK-NEXT: br label [[OMP_TYPE_END]] 288 // CHECK: omp.type.alloc.else: 289 // CHECK-NEXT: [[TMP30:%.*]] = icmp eq i64 [[TMP27]], 1 290 // CHECK-NEXT: br i1 [[TMP30]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]] 291 // CHECK: omp.type.to: 292 // CHECK-NEXT: [[TMP31:%.*]] = and i64 [[TMP26]], -3 293 // CHECK-NEXT: br label [[OMP_TYPE_END]] 294 // CHECK: omp.type.to.else: 295 // CHECK-NEXT: [[TMP32:%.*]] = icmp eq i64 [[TMP27]], 2 296 // CHECK-NEXT: br i1 [[TMP32]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]] 297 // CHECK: omp.type.from: 298 // CHECK-NEXT: [[TMP33:%.*]] = and i64 [[TMP26]], -2 299 // CHECK-NEXT: br label [[OMP_TYPE_END]] 300 // CHECK: omp.type.end: 301 // CHECK-NEXT: [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP29]], [[OMP_TYPE_ALLOC]] ], [ [[TMP31]], [[OMP_TYPE_TO]] ], [ [[TMP33]], [[OMP_TYPE_FROM]] ], [ [[TMP26]], [[OMP_TYPE_TO_ELSE]] ] 302 // CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[A]], i64 4, i64 [[OMP_MAPTYPE]], ptr null) 303 // CHECK-NEXT: [[OMP_ARRAYMAP_NEXT]] = getelementptr [[STRUCT_C]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 1 304 // CHECK-NEXT: [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr [[OMP_ARRAYMAP_NEXT]], [[TMP11]] 305 // CHECK-NEXT: br i1 [[OMP_ARRAYMAP_ISDONE]], label [[OMP_ARRAYMAP_EXIT:%.*]], label [[OMP_ARRAYMAP_BODY]] 306 // CHECK: omp.arraymap.exit: 307 // CHECK-NEXT: [[OMP_ARRAYINIT_ISARRAY6:%.*]] = icmp sgt i64 [[TMP10]], 1 308 // CHECK-NEXT: [[TMP34:%.*]] = and i64 [[TMP4]], 8 309 // CHECK-NEXT: [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP34]], 0 310 // CHECK-NEXT: [[TMP35:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY6]], [[DOTOMP_ARRAY__DEL__DELETE]] 311 // CHECK-NEXT: br i1 [[TMP35]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]] 312 // CHECK: .omp.array..del: 313 // CHECK-NEXT: [[TMP36:%.*]] = mul nuw i64 [[TMP10]], 4 314 // CHECK-NEXT: [[TMP37:%.*]] = and i64 [[TMP4]], -4 315 // CHECK-NEXT: [[TMP38:%.*]] = or i64 [[TMP37]], 512 316 // CHECK-NEXT: call void @__tgt_push_mapper_component(ptr [[TMP0]], ptr [[TMP1]], ptr [[TMP2]], i64 [[TMP36]], i64 [[TMP38]], ptr [[TMP5]]) 317 // CHECK-NEXT: br label [[OMP_DONE]] 318 // CHECK: omp.done: 319 // CHECK-NEXT: ret void 320 // 321