clang/test/OpenMP/target_map_nest_defalut_mapper_codegen.cpp

   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:    [[DOTADDR:%.*]] = alloca ptr, align 8
 113 // CHECK-NEXT:    [[DOTADDR1:%.*]] = alloca ptr, align 8
 114 // CHECK-NEXT:    [[DOTADDR2:%.*]] = alloca ptr, align 8
 115 // CHECK-NEXT:    [[DOTADDR3:%.*]] = alloca i64, align 8
 116 // CHECK-NEXT:    [[DOTADDR4:%.*]] = alloca i64, align 8
 117 // CHECK-NEXT:    [[DOTADDR5:%.*]] = alloca ptr, align 8
 118 // CHECK-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 8
 119 // CHECK-NEXT:    store ptr [[TMP1]], ptr [[DOTADDR1]], align 8
 120 // CHECK-NEXT:    store ptr [[TMP2]], ptr [[DOTADDR2]], align 8
 121 // CHECK-NEXT:    store i64 [[TMP3]], ptr [[DOTADDR3]], align 8
 122 // CHECK-NEXT:    store i64 [[TMP4]], ptr [[DOTADDR4]], align 8
 123 // CHECK-NEXT:    store ptr [[TMP5]], ptr [[DOTADDR5]], align 8
 124 // CHECK-NEXT:    [[TMP6:%.*]] = load i64, ptr [[DOTADDR3]], align 8
 125 // CHECK-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[DOTADDR]], align 8
 126 // CHECK-NEXT:    [[TMP8:%.*]] = load ptr, ptr [[DOTADDR1]], align 8
 127 // CHECK-NEXT:    [[TMP9:%.*]] = load ptr, ptr [[DOTADDR2]], align 8
 128 // CHECK-NEXT:    [[TMP10:%.*]] = udiv exact i64 [[TMP6]], 12
 129 // CHECK-NEXT:    [[TMP11:%.*]] = getelementptr [[STRUCT_D:%.*]], ptr [[TMP9]], i64 [[TMP10]]
 130 // CHECK-NEXT:    [[TMP12:%.*]] = load i64, ptr [[DOTADDR4]], align 8
 131 // CHECK-NEXT:    [[TMP13:%.*]] = load ptr, ptr [[DOTADDR5]], align 8
 132 // CHECK-NEXT:    [[OMP_ARRAYINIT_ISARRAY:%.*]] = icmp sgt i64 [[TMP10]], 1
 133 // CHECK-NEXT:    [[TMP14:%.*]] = and i64 [[TMP12]], 8
 134 // CHECK-NEXT:    [[TMP15:%.*]] = icmp ne ptr [[TMP8]], [[TMP9]]
 135 // CHECK-NEXT:    [[TMP16:%.*]] = and i64 [[TMP12]], 16
 136 // CHECK-NEXT:    [[TMP17:%.*]] = icmp ne i64 [[TMP16]], 0
 137 // CHECK-NEXT:    [[TMP18:%.*]] = and i1 [[TMP15]], [[TMP17]]
 138 // CHECK-NEXT:    [[TMP19:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP18]]
 139 // CHECK-NEXT:    [[DOTOMP_ARRAY__INIT__DELETE:%.*]] = icmp eq i64 [[TMP14]], 0
 140 // CHECK-NEXT:    [[TMP20:%.*]] = and i1 [[TMP19]], [[DOTOMP_ARRAY__INIT__DELETE]]
 141 // CHECK-NEXT:    br i1 [[TMP20]], label [[DOTOMP_ARRAY__INIT:%.*]], label [[OMP_ARRAYMAP_HEAD:%.*]]
 142 // CHECK:       .omp.array..init:
 143 // CHECK-NEXT:    [[TMP21:%.*]] = mul nuw i64 [[TMP10]], 12
 144 // CHECK-NEXT:    [[TMP22:%.*]] = and i64 [[TMP12]], -4
 145 // CHECK-NEXT:    [[TMP23:%.*]] = or i64 [[TMP22]], 512
 146 // CHECK-NEXT:    call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[TMP8]], ptr [[TMP9]], i64 [[TMP21]], i64 [[TMP23]], ptr [[TMP13]])
 147 // CHECK-NEXT:    br label [[OMP_ARRAYMAP_HEAD]]
 148 // CHECK:       omp.arraymap.head:
 149 // CHECK-NEXT:    [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP9]], [[TMP11]]
 150 // CHECK-NEXT:    br i1 [[OMP_ARRAYMAP_ISEMPTY]], label [[OMP_DONE:%.*]], label [[OMP_ARRAYMAP_BODY:%.*]]
 151 // CHECK:       omp.arraymap.body:
 152 // CHECK-NEXT:    [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP9]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END25:%.*]] ]
 153 // CHECK-NEXT:    [[E:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
 154 // CHECK-NEXT:    [[F:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 1
 155 // CHECK-NEXT:    [[H:%.*]] = getelementptr inbounds nuw [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 2
 156 // CHECK-NEXT:    [[TMP24:%.*]] = getelementptr i32, ptr [[H]], i32 1
 157 // CHECK-NEXT:    [[TMP25:%.*]] = ptrtoint ptr [[TMP24]] to i64
 158 // CHECK-NEXT:    [[TMP26:%.*]] = ptrtoint ptr [[E]] to i64
 159 // CHECK-NEXT:    [[TMP27:%.*]] = sub i64 [[TMP25]], [[TMP26]]
 160 // CHECK-NEXT:    [[TMP28:%.*]] = sdiv exact i64 [[TMP27]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
 161 // CHECK-NEXT:    [[TMP29:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP7]])
 162 // CHECK-NEXT:    [[TMP30:%.*]] = shl i64 [[TMP29]], 48
 163 // CHECK-NEXT:    [[TMP31:%.*]] = add nuw i64 0, [[TMP30]]
 164 // CHECK-NEXT:    [[TMP32:%.*]] = and i64 [[TMP12]], 3
 165 // CHECK-NEXT:    [[TMP33:%.*]] = icmp eq i64 [[TMP32]], 0
 166 // CHECK-NEXT:    br i1 [[TMP33]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]]
 167 // CHECK:       omp.type.alloc:
 168 // CHECK-NEXT:    [[TMP34:%.*]] = and i64 [[TMP31]], -4
 169 // CHECK-NEXT:    br label [[OMP_TYPE_END:%.*]]
 170 // CHECK:       omp.type.alloc.else:
 171 // CHECK-NEXT:    [[TMP35:%.*]] = icmp eq i64 [[TMP32]], 1
 172 // CHECK-NEXT:    br i1 [[TMP35]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]]
 173 // CHECK:       omp.type.to:
 174 // CHECK-NEXT:    [[TMP36:%.*]] = and i64 [[TMP31]], -3
 175 // CHECK-NEXT:    br label [[OMP_TYPE_END]]
 176 // CHECK:       omp.type.to.else:
 177 // CHECK-NEXT:    [[TMP37:%.*]] = icmp eq i64 [[TMP32]], 2
 178 // CHECK-NEXT:    br i1 [[TMP37]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]]
 179 // CHECK:       omp.type.from:
 180 // CHECK-NEXT:    [[TMP38:%.*]] = and i64 [[TMP31]], -2
 181 // CHECK-NEXT:    br label [[OMP_TYPE_END]]
 182 // CHECK:       omp.type.end:
 183 // CHECK-NEXT:    [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP34]], [[OMP_TYPE_ALLOC]] ], [ [[TMP36]], [[OMP_TYPE_TO]] ], [ [[TMP38]], [[OMP_TYPE_FROM]] ], [ [[TMP31]], [[OMP_TYPE_TO_ELSE]] ]
 184 // CHECK-NEXT:    call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 [[TMP28]], i64 [[OMP_MAPTYPE]], ptr null)
 185 // CHECK-NEXT:    [[TMP39:%.*]] = add nuw i64 281474976711171, [[TMP30]]
 186 // CHECK-NEXT:    [[TMP40:%.*]] = and i64 [[TMP12]], 3
 187 // CHECK-NEXT:    [[TMP41:%.*]] = icmp eq i64 [[TMP40]], 0
 188 // CHECK-NEXT:    br i1 [[TMP41]], label [[OMP_TYPE_ALLOC6:%.*]], label [[OMP_TYPE_ALLOC_ELSE7:%.*]]
 189 // CHECK:       omp.type.alloc6:
 190 // CHECK-NEXT:    [[TMP42:%.*]] = and i64 [[TMP39]], -4
 191 // CHECK-NEXT:    br label [[OMP_TYPE_END11:%.*]]
 192 // CHECK:       omp.type.alloc.else7:
 193 // CHECK-NEXT:    [[TMP43:%.*]] = icmp eq i64 [[TMP40]], 1
 194 // CHECK-NEXT:    br i1 [[TMP43]], label [[OMP_TYPE_TO8:%.*]], label [[OMP_TYPE_TO_ELSE9:%.*]]
 195 // CHECK:       omp.type.to8:
 196 // CHECK-NEXT:    [[TMP44:%.*]] = and i64 [[TMP39]], -3
 197 // CHECK-NEXT:    br label [[OMP_TYPE_END11]]
 198 // CHECK:       omp.type.to.else9:
 199 // CHECK-NEXT:    [[TMP45:%.*]] = icmp eq i64 [[TMP40]], 2
 200 // CHECK-NEXT:    br i1 [[TMP45]], label [[OMP_TYPE_FROM10:%.*]], label [[OMP_TYPE_END11]]
 201 // CHECK:       omp.type.from10:
 202 // CHECK-NEXT:    [[TMP46:%.*]] = and i64 [[TMP39]], -2
 203 // CHECK-NEXT:    br label [[OMP_TYPE_END11]]
 204 // CHECK:       omp.type.end11:
 205 // CHECK-NEXT:    [[OMP_MAPTYPE12:%.*]] = phi i64 [ [[TMP42]], [[OMP_TYPE_ALLOC6]] ], [ [[TMP44]], [[OMP_TYPE_TO8]] ], [ [[TMP46]], [[OMP_TYPE_FROM10]] ], [ [[TMP39]], [[OMP_TYPE_TO_ELSE9]] ]
 206 // CHECK-NEXT:    call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[E]], i64 4, i64 [[OMP_MAPTYPE12]], ptr null)
 207 // CHECK-NEXT:    [[TMP47:%.*]] = add nuw i64 281474976711171, [[TMP30]]
 208 // CHECK-NEXT:    [[TMP48:%.*]] = and i64 [[TMP12]], 3
 209 // CHECK-NEXT:    [[TMP49:%.*]] = icmp eq i64 [[TMP48]], 0
 210 // CHECK-NEXT:    br i1 [[TMP49]], label [[OMP_TYPE_ALLOC13:%.*]], label [[OMP_TYPE_ALLOC_ELSE14:%.*]]
 211 // CHECK:       omp.type.alloc13:
 212 // CHECK-NEXT:    [[TMP50:%.*]] = and i64 [[TMP47]], -4
 213 // CHECK-NEXT:    br label [[OMP_TYPE_END18:%.*]]
 214 // CHECK:       omp.type.alloc.else14:
 215 // CHECK-NEXT:    [[TMP51:%.*]] = icmp eq i64 [[TMP48]], 1
 216 // CHECK-NEXT:    br i1 [[TMP51]], label [[OMP_TYPE_TO15:%.*]], label [[OMP_TYPE_TO_ELSE16:%.*]]
 217 // CHECK:       omp.type.to15:
 218 // CHECK-NEXT:    [[TMP52:%.*]] = and i64 [[TMP47]], -3
 219 // CHECK-NEXT:    br label [[OMP_TYPE_END18]]
 220 // CHECK:       omp.type.to.else16:
 221 // CHECK-NEXT:    [[TMP53:%.*]] = icmp eq i64 [[TMP48]], 2
 222 // CHECK-NEXT:    br i1 [[TMP53]], label [[OMP_TYPE_FROM17:%.*]], label [[OMP_TYPE_END18]]
 223 // CHECK:       omp.type.from17:
 224 // CHECK-NEXT:    [[TMP54:%.*]] = and i64 [[TMP47]], -2
 225 // CHECK-NEXT:    br label [[OMP_TYPE_END18]]
 226 // CHECK:       omp.type.end18:
 227 // CHECK-NEXT:    [[OMP_MAPTYPE19:%.*]] = phi i64 [ [[TMP50]], [[OMP_TYPE_ALLOC13]] ], [ [[TMP52]], [[OMP_TYPE_TO15]] ], [ [[TMP54]], [[OMP_TYPE_FROM17]] ], [ [[TMP47]], [[OMP_TYPE_TO_ELSE16]] ]
 228 // CHECK-NEXT:    call void @.omp_mapper._ZTS1C.default(ptr [[TMP7]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[F]], i64 4, i64 [[OMP_MAPTYPE19]], ptr null) #[[ATTR3]]
 229 // CHECK-NEXT:    [[TMP55:%.*]] = add nuw i64 281474976711171, [[TMP30]]
 230 // CHECK-NEXT:    [[TMP56:%.*]] = and i64 [[TMP12]], 3
 231 // CHECK-NEXT:    [[TMP57:%.*]] = icmp eq i64 [[TMP56]], 0
 232 // CHECK-NEXT:    br i1 [[TMP57]], label [[OMP_TYPE_ALLOC20:%.*]], label [[OMP_TYPE_ALLOC_ELSE21:%.*]]
 233 // CHECK:       omp.type.alloc20:
 234 // CHECK-NEXT:    [[TMP58:%.*]] = and i64 [[TMP55]], -4
 235 // CHECK-NEXT:    br label [[OMP_TYPE_END25]]
 236 // CHECK:       omp.type.alloc.else21:
 237 // CHECK-NEXT:    [[TMP59:%.*]] = icmp eq i64 [[TMP56]], 1
 238 // CHECK-NEXT:    br i1 [[TMP59]], label [[OMP_TYPE_TO22:%.*]], label [[OMP_TYPE_TO_ELSE23:%.*]]
 239 // CHECK:       omp.type.to22:
 240 // CHECK-NEXT:    [[TMP60:%.*]] = and i64 [[TMP55]], -3
 241 // CHECK-NEXT:    br label [[OMP_TYPE_END25]]
 242 // CHECK:       omp.type.to.else23:
 243 // CHECK-NEXT:    [[TMP61:%.*]] = icmp eq i64 [[TMP56]], 2
 244 // CHECK-NEXT:    br i1 [[TMP61]], label [[OMP_TYPE_FROM24:%.*]], label [[OMP_TYPE_END25]]
 245 // CHECK:       omp.type.from24:
 246 // CHECK-NEXT:    [[TMP62:%.*]] = and i64 [[TMP55]], -2
 247 // CHECK-NEXT:    br label [[OMP_TYPE_END25]]
 248 // CHECK:       omp.type.end25:
 249 // CHECK-NEXT:    [[OMP_MAPTYPE26:%.*]] = phi i64 [ [[TMP58]], [[OMP_TYPE_ALLOC20]] ], [ [[TMP60]], [[OMP_TYPE_TO22]] ], [ [[TMP62]], [[OMP_TYPE_FROM24]] ], [ [[TMP55]], [[OMP_TYPE_TO_ELSE23]] ]
 250 // CHECK-NEXT:    call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[H]], i64 4, i64 [[OMP_MAPTYPE26]], ptr null)
 251 // CHECK-NEXT:    [[OMP_ARRAYMAP_NEXT]] = getelementptr [[STRUCT_D]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 1
 252 // CHECK-NEXT:    [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr [[OMP_ARRAYMAP_NEXT]], [[TMP11]]
 253 // CHECK-NEXT:    br i1 [[OMP_ARRAYMAP_ISDONE]], label [[OMP_ARRAYMAP_EXIT:%.*]], label [[OMP_ARRAYMAP_BODY]]
 254 // CHECK:       omp.arraymap.exit:
 255 // CHECK-NEXT:    [[OMP_ARRAYINIT_ISARRAY27:%.*]] = icmp sgt i64 [[TMP10]], 1
 256 // CHECK-NEXT:    [[TMP63:%.*]] = and i64 [[TMP12]], 8
 257 // CHECK-NEXT:    [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP63]], 0
 258 // CHECK-NEXT:    [[TMP64:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY27]], [[DOTOMP_ARRAY__DEL__DELETE]]
 259 // CHECK-NEXT:    br i1 [[TMP64]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]]
 260 // CHECK:       .omp.array..del:
 261 // CHECK-NEXT:    [[TMP65:%.*]] = mul nuw i64 [[TMP10]], 12
 262 // CHECK-NEXT:    [[TMP66:%.*]] = and i64 [[TMP12]], -4
 263 // CHECK-NEXT:    [[TMP67:%.*]] = or i64 [[TMP66]], 512
 264 // CHECK-NEXT:    call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[TMP8]], ptr [[TMP9]], i64 [[TMP65]], i64 [[TMP67]], ptr [[TMP13]])
 265 // CHECK-NEXT:    br label [[OMP_DONE]]
 266 // CHECK:       omp.done:
 267 // CHECK-NEXT:    ret void
 268 //
 269 //
 270 // CHECK-LABEL: define {{[^@]+}}@.omp_mapper._ZTS1C.default
 271 // CHECK-SAME: (ptr noundef [[TMP0:%.*]], ptr noundef [[TMP1:%.*]], ptr noundef [[TMP2:%.*]], i64 noundef [[TMP3:%.*]], i64 noundef [[TMP4:%.*]], ptr noundef [[TMP5:%.*]]) #[[ATTR2]] {
 272 // CHECK-NEXT:  entry:
 273 // CHECK-NEXT:    [[DOTADDR:%.*]] = alloca ptr, align 8
 274 // CHECK-NEXT:    [[DOTADDR1:%.*]] = alloca ptr, align 8
 275 // CHECK-NEXT:    [[DOTADDR2:%.*]] = alloca ptr, align 8
 276 // CHECK-NEXT:    [[DOTADDR3:%.*]] = alloca i64, align 8
 277 // CHECK-NEXT:    [[DOTADDR4:%.*]] = alloca i64, align 8
 278 // CHECK-NEXT:    [[DOTADDR5:%.*]] = alloca ptr, align 8
 279 // CHECK-NEXT:    store ptr [[TMP0]], ptr [[DOTADDR]], align 8
 280 // CHECK-NEXT:    store ptr [[TMP1]], ptr [[DOTADDR1]], align 8
 281 // CHECK-NEXT:    store ptr [[TMP2]], ptr [[DOTADDR2]], align 8
 282 // CHECK-NEXT:    store i64 [[TMP3]], ptr [[DOTADDR3]], align 8
 283 // CHECK-NEXT:    store i64 [[TMP4]], ptr [[DOTADDR4]], align 8
 284 // CHECK-NEXT:    store ptr [[TMP5]], ptr [[DOTADDR5]], align 8
 285 // CHECK-NEXT:    [[TMP6:%.*]] = load i64, ptr [[DOTADDR3]], align 8
 286 // CHECK-NEXT:    [[TMP7:%.*]] = load ptr, ptr [[DOTADDR]], align 8
 287 // CHECK-NEXT:    [[TMP8:%.*]] = load ptr, ptr [[DOTADDR1]], align 8
 288 // CHECK-NEXT:    [[TMP9:%.*]] = load ptr, ptr [[DOTADDR2]], align 8
 289 // CHECK-NEXT:    [[TMP10:%.*]] = udiv exact i64 [[TMP6]], 4
 290 // CHECK-NEXT:    [[TMP11:%.*]] = getelementptr [[STRUCT_C:%.*]], ptr [[TMP9]], i64 [[TMP10]]
 291 // CHECK-NEXT:    [[TMP12:%.*]] = load i64, ptr [[DOTADDR4]], align 8
 292 // CHECK-NEXT:    [[TMP13:%.*]] = load ptr, ptr [[DOTADDR5]], align 8
 293 // CHECK-NEXT:    [[OMP_ARRAYINIT_ISARRAY:%.*]] = icmp sgt i64 [[TMP10]], 1
 294 // CHECK-NEXT:    [[TMP14:%.*]] = and i64 [[TMP12]], 8
 295 // CHECK-NEXT:    [[TMP15:%.*]] = icmp ne ptr [[TMP8]], [[TMP9]]
 296 // CHECK-NEXT:    [[TMP16:%.*]] = and i64 [[TMP12]], 16
 297 // CHECK-NEXT:    [[TMP17:%.*]] = icmp ne i64 [[TMP16]], 0
 298 // CHECK-NEXT:    [[TMP18:%.*]] = and i1 [[TMP15]], [[TMP17]]
 299 // CHECK-NEXT:    [[TMP19:%.*]] = or i1 [[OMP_ARRAYINIT_ISARRAY]], [[TMP18]]
 300 // CHECK-NEXT:    [[DOTOMP_ARRAY__INIT__DELETE:%.*]] = icmp eq i64 [[TMP14]], 0
 301 // CHECK-NEXT:    [[TMP20:%.*]] = and i1 [[TMP19]], [[DOTOMP_ARRAY__INIT__DELETE]]
 302 // CHECK-NEXT:    br i1 [[TMP20]], label [[DOTOMP_ARRAY__INIT:%.*]], label [[OMP_ARRAYMAP_HEAD:%.*]]
 303 // CHECK:       .omp.array..init:
 304 // CHECK-NEXT:    [[TMP21:%.*]] = mul nuw i64 [[TMP10]], 4
 305 // CHECK-NEXT:    [[TMP22:%.*]] = and i64 [[TMP12]], -4
 306 // CHECK-NEXT:    [[TMP23:%.*]] = or i64 [[TMP22]], 512
 307 // CHECK-NEXT:    call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[TMP8]], ptr [[TMP9]], i64 [[TMP21]], i64 [[TMP23]], ptr [[TMP13]])
 308 // CHECK-NEXT:    br label [[OMP_ARRAYMAP_HEAD]]
 309 // CHECK:       omp.arraymap.head:
 310 // CHECK-NEXT:    [[OMP_ARRAYMAP_ISEMPTY:%.*]] = icmp eq ptr [[TMP9]], [[TMP11]]
 311 // CHECK-NEXT:    br i1 [[OMP_ARRAYMAP_ISEMPTY]], label [[OMP_DONE:%.*]], label [[OMP_ARRAYMAP_BODY:%.*]]
 312 // CHECK:       omp.arraymap.body:
 313 // CHECK-NEXT:    [[OMP_ARRAYMAP_PTRCURRENT:%.*]] = phi ptr [ [[TMP9]], [[OMP_ARRAYMAP_HEAD]] ], [ [[OMP_ARRAYMAP_NEXT:%.*]], [[OMP_TYPE_END:%.*]] ]
 314 // CHECK-NEXT:    [[A:%.*]] = getelementptr inbounds nuw [[STRUCT_C]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 0, i32 0
 315 // CHECK-NEXT:    [[TMP24:%.*]] = call i64 @__tgt_mapper_num_components(ptr [[TMP7]])
 316 // CHECK-NEXT:    [[TMP25:%.*]] = shl i64 [[TMP24]], 48
 317 // CHECK-NEXT:    [[TMP26:%.*]] = add nuw i64 1, [[TMP25]]
 318 // CHECK-NEXT:    [[TMP27:%.*]] = and i64 [[TMP12]], 3
 319 // CHECK-NEXT:    [[TMP28:%.*]] = icmp eq i64 [[TMP27]], 0
 320 // CHECK-NEXT:    br i1 [[TMP28]], label [[OMP_TYPE_ALLOC:%.*]], label [[OMP_TYPE_ALLOC_ELSE:%.*]]
 321 // CHECK:       omp.type.alloc:
 322 // CHECK-NEXT:    [[TMP29:%.*]] = and i64 [[TMP26]], -4
 323 // CHECK-NEXT:    br label [[OMP_TYPE_END]]
 324 // CHECK:       omp.type.alloc.else:
 325 // CHECK-NEXT:    [[TMP30:%.*]] = icmp eq i64 [[TMP27]], 1
 326 // CHECK-NEXT:    br i1 [[TMP30]], label [[OMP_TYPE_TO:%.*]], label [[OMP_TYPE_TO_ELSE:%.*]]
 327 // CHECK:       omp.type.to:
 328 // CHECK-NEXT:    [[TMP31:%.*]] = and i64 [[TMP26]], -3
 329 // CHECK-NEXT:    br label [[OMP_TYPE_END]]
 330 // CHECK:       omp.type.to.else:
 331 // CHECK-NEXT:    [[TMP32:%.*]] = icmp eq i64 [[TMP27]], 2
 332 // CHECK-NEXT:    br i1 [[TMP32]], label [[OMP_TYPE_FROM:%.*]], label [[OMP_TYPE_END]]
 333 // CHECK:       omp.type.from:
 334 // CHECK-NEXT:    [[TMP33:%.*]] = and i64 [[TMP26]], -2
 335 // CHECK-NEXT:    br label [[OMP_TYPE_END]]
 336 // CHECK:       omp.type.end:
 337 // CHECK-NEXT:    [[OMP_MAPTYPE:%.*]] = phi i64 [ [[TMP29]], [[OMP_TYPE_ALLOC]] ], [ [[TMP31]], [[OMP_TYPE_TO]] ], [ [[TMP33]], [[OMP_TYPE_FROM]] ], [ [[TMP26]], [[OMP_TYPE_TO_ELSE]] ]
 338 // CHECK-NEXT:    call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], ptr [[A]], i64 4, i64 [[OMP_MAPTYPE]], ptr null)
 339 // CHECK-NEXT:    [[OMP_ARRAYMAP_NEXT]] = getelementptr [[STRUCT_C]], ptr [[OMP_ARRAYMAP_PTRCURRENT]], i32 1
 340 // CHECK-NEXT:    [[OMP_ARRAYMAP_ISDONE:%.*]] = icmp eq ptr [[OMP_ARRAYMAP_NEXT]], [[TMP11]]
 341 // CHECK-NEXT:    br i1 [[OMP_ARRAYMAP_ISDONE]], label [[OMP_ARRAYMAP_EXIT:%.*]], label [[OMP_ARRAYMAP_BODY]]
 342 // CHECK:       omp.arraymap.exit:
 343 // CHECK-NEXT:    [[OMP_ARRAYINIT_ISARRAY6:%.*]] = icmp sgt i64 [[TMP10]], 1
 344 // CHECK-NEXT:    [[TMP34:%.*]] = and i64 [[TMP12]], 8
 345 // CHECK-NEXT:    [[DOTOMP_ARRAY__DEL__DELETE:%.*]] = icmp ne i64 [[TMP34]], 0
 346 // CHECK-NEXT:    [[TMP35:%.*]] = and i1 [[OMP_ARRAYINIT_ISARRAY6]], [[DOTOMP_ARRAY__DEL__DELETE]]
 347 // CHECK-NEXT:    br i1 [[TMP35]], label [[DOTOMP_ARRAY__DEL:%.*]], label [[OMP_DONE]]
 348 // CHECK:       .omp.array..del:
 349 // CHECK-NEXT:    [[TMP36:%.*]] = mul nuw i64 [[TMP10]], 4
 350 // CHECK-NEXT:    [[TMP37:%.*]] = and i64 [[TMP12]], -4
 351 // CHECK-NEXT:    [[TMP38:%.*]] = or i64 [[TMP37]], 512
 352 // CHECK-NEXT:    call void @__tgt_push_mapper_component(ptr [[TMP7]], ptr [[TMP8]], ptr [[TMP9]], i64 [[TMP36]], i64 [[TMP38]], ptr [[TMP13]])
 353 // CHECK-NEXT:    br label [[OMP_DONE]]
 354 // CHECK:       omp.done:
 355 // CHECK-NEXT:    ret void
 356 //