clang/test/OpenMP/declare_mapper_codegen.cpp

   1 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
   2
   3 // expected-no-diagnostics
   4 #ifndef HEADER
   5 #define HEADER
   6
   7 ///==========================================================================///
   8 // RUN: %clang_cc1 -DCK0 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK0 --check-prefix CK0-64 %s
   9 // RUN: %clang_cc1 -DCK0 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
  10 // RUN: %clang_cc1 -DCK0 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK0 --check-prefix CK0-64 %s
  11 // RUN: %clang_cc1 -DCK0 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK0 --check-prefix CK0-32 %s
  12 // RUN: %clang_cc1 -DCK0 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
  13 // RUN: %clang_cc1 -DCK0 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK0 --check-prefix CK0-32 %s
  14
  15 // RUN: %clang_cc1 -DCK0 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
  16 // RUN: %clang_cc1 -DCK0 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
  17 // RUN: %clang_cc1 -DCK0 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
  18 // RUN: %clang_cc1 -DCK0 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
  19 // RUN: %clang_cc1 -DCK0 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
  20 // RUN: %clang_cc1 -DCK0 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
  21
  22 #ifdef CK0
  23 // Mapper function code generation and runtime interface.
  24
  25 // CK0: [[IDENT_T:%.+]] = type { i32, i32, i32, i32, ptr }
  26 // CK0: [[ENTRY:%.+]] = type { ptr, ptr, i[[SZ:32|64]], i32, i32 }
  27 // CK0: [[ANON_T:%.+]] = type { ptr }
  28 // CK0: [[ANON_T_0:%.+]] = type { ptr }
  29 // CK0: [[KMP_TASK_T_WITH_PRIVATES:%.+]] = type { [[KMP_TASK_T:%[^,]+]], [[KMP_PRIVATES_T:%.+]] }
  30 // CK0: [[KMP_TASK_T]] = type { ptr, ptr, i32, %{{[^,]+}}, %{{[^,]+}} }
  31 // CK0-32: [[KMP_PRIVATES_T]] = type { [1 x i64], [1 x ptr], [1 x ptr], [1 x ptr] }
  32 // CK0-64: [[KMP_PRIVATES_T]] = type { [1 x ptr], [1 x ptr], [1 x i64], [1 x ptr] }
  33 // CK0: [[KMP_TASK_T_WITH_PRIVATES_1:%.+]] = type { [[KMP_TASK_T]], [[KMP_PRIVATES_T_2:%.+]] }
  34 // CK0-32: [[KMP_PRIVATES_T_2]] = type { [1 x i64], [1 x ptr], [1 x ptr], [1 x ptr] }
  35 // CK0-64: [[KMP_PRIVATES_T_2]] = type { [1 x ptr], [1 x ptr], [1 x i64], [1 x ptr] }
  36 // CK0: [[KMP_TASK_T_WITH_PRIVATES_4:%.+]] = type { [[KMP_TASK_T]], [[KMP_PRIVATES_T_5:%.+]] }
  37 // CK0-32: [[KMP_PRIVATES_T_5]] = type { [1 x i64], [1 x ptr], [1 x ptr], [1 x ptr] }
  38 // CK0-64: [[KMP_PRIVATES_T_5]] = type { [1 x ptr], [1 x ptr], [1 x i64], [1 x ptr] }
  39 // CK0: [[KMP_TASK_T_WITH_PRIVATES_7:%.+]] = type { [[KMP_TASK_T]], [[KMP_PRIVATES_T_8:%.+]] }
  40 // CK0-32: [[KMP_PRIVATES_T_8]] = type { [1 x i64], [1 x ptr], [1 x ptr], [1 x ptr] }
  41 // CK0-64: [[KMP_PRIVATES_T_8]] = type { [1 x ptr], [1 x ptr], [1 x i64], [1 x ptr] }
  42 // CK0: [[KMP_TASK_T_WITH_PRIVATES_10:%.+]] = type { [[KMP_TASK_T]], [[KMP_PRIVATES_T_11:%.+]] }
  43 // CK0-32: [[KMP_PRIVATES_T_11]] = type { [1 x i64], [1 x ptr], [1 x ptr], [1 x ptr] }
  44 // CK0-64: [[KMP_PRIVATES_T_11]] = type { [1 x ptr], [1 x ptr], [1 x i64], [1 x ptr] }
  45
  46 // CK0-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}.region_id = weak constant i8 0
  47 // CK0-64: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
  48 // CK0-32: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
  49 // CK0: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 35]
  50 // CK0-64: [[NWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
  51 // CK0-32: [[NWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
  52 // CK0: [[NWTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 35]
  53 // CK0-64: [[TEAMSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
  54 // CK0-32: [[TEAMSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
  55 // CK0: [[TEAMTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 33]
  56 // CK0-64: [[TEAMNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
  57 // CK0-32: [[TEAMNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
  58 // CK0: [[TEAMNWTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 33]
  59 // CK0-64: [[EDSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
  60 // CK0-32: [[EDSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
  61 // CK0: [[EDTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 1]
  62 // CK0-64: [[EDNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
  63 // CK0-32: [[EDNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
  64 // CK0: [[EDNWTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 1]
  65 // CK0-64: [[EXDSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
  66 // CK0-32: [[EXDSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
  67 // CK0: [[EXDTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 2]
  68 // CK0-64: [[EXDNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
  69 // CK0-32: [[EXDNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
  70 // CK0: [[EXDNWTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 2]
  71 // CK0-64: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
  72 // CK0-32: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
  73 // CK0: [[TTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 1]
  74 // CK0-64: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
  75 // CK0-32: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
  76 // CK0: [[FTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 2]
  77 // CK0-64: [[FNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
  78 // CK0-32: [[FNWSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
  79 // CK0: [[FNWTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 2]
  80
  81 class C {
  82 public:
  83   int a;
  84   double *b;
  85 };
  86
  87 #pragma omp declare mapper(id: C s) map(s.a, s.b[0:2])
  88
  89 // CK0: define {{.*}}void [[MPRFUNC:@[.]omp_mapper[.].*C[.]id]](ptr{{.*}}, ptr{{.*}}, ptr{{.*}}, i64{{.*}}, i64{{.*}}, ptr{{.*}})
  90 // CK0: store ptr %{{[^,]+}}, ptr [[HANDLEADDR:%[^,]+]]
  91 // CK0: store ptr %{{[^,]+}}, ptr [[BPTRADDR:%[^,]+]]
  92 // CK0: store ptr %{{[^,]+}}, ptr [[VPTRADDR:%[^,]+]]
  93 // CK0: store i64 %{{[^,]+}}, ptr [[SIZEADDR:%[^,]+]]
  94 // CK0: store i64 %{{[^,]+}}, ptr [[TYPEADDR:%[^,]+]]
  95 // CK0-DAG: [[BYTESIZE:%.+]] = load i64, ptr [[SIZEADDR]]
  96 // CK0-64-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 16
  97 // CK0-32-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 8
  98 // CK0-DAG: [[TYPE:%.+]] = load i64, ptr [[TYPEADDR]]
  99 // CK0-DAG: [[HANDLE:%.+]] = load ptr, ptr [[HANDLEADDR]]
 100 // CK0-DAG: [[BPTR:%.+]] = load ptr, ptr [[BPTRADDR]]
 101 // CK0-DAG: [[BEGIN:%.+]] = load ptr, ptr [[VPTRADDR]]
 102 // CK0-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1
 103 // CK0-DAG: [[PTREND:%.+]] = getelementptr %class.C, ptr [[BEGIN]], i64 [[SIZE]]
 104 // CK0-DAG: [[PTRSNE:%.+]] = icmp ne ptr [[BPTR]], [[BEGIN]]
 105 // CK0-DAG: [[PTRANDOBJ:%.+]] = and i64 [[TYPE]], 16
 106 // CK0-DAG: [[ISPTRANDOBJ:%.+]] = icmp ne i64 [[PTRANDOBJ]], 0
 107 // CK0-DAG: [[CMPA:%.+]] = and i1 [[PTRSNE]], [[ISPTRANDOBJ]]
 108 // CK0-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[CMPA]]
 109 // CK0-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
 110 // CK0-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0
 111 // CK0-DAG: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]]
 112 // CK0: br i1 [[CMP1]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]]
 113 // CK0: [[INIT]]
 114 // CK0-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
 115 // CK0-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
 116
 117 // Remove movement mappings and mark as implicit
 118 // CK0-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
 119 // CK0-DAG: [[ITYPE1:%.+]] = or i64 [[ITYPE]], 512
 120 // CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BPTR]], ptr [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE1]], {{.*}})
 121 // CK0: br label %[[LHEAD:[^,]+]]
 122
 123 // CK0: [[LHEAD]]
 124 // CK0: [[ISEMPTY:%.+]] = icmp eq ptr [[BEGIN]], [[PTREND]]
 125 // CK0: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]]
 126 // CK0: [[LBODY]]
 127 // CK0: [[PTR:%.+]] = phi ptr [ [[BEGIN]], %{{.+}} ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
 128 // CK0-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, ptr [[PTR]], i32 0, i32 0
 129 // CK0-DAG: [[BBEGIN:%.+]] = getelementptr inbounds %class.C, ptr [[PTR]], i32 0, i32 1
 130 // CK0-DAG: [[BBEGIN2:%.+]] = getelementptr inbounds %class.C, ptr [[PTR]], i32 0, i32 1
 131 // CK0-DAG: [[BARRBEGIN:%.+]] = load ptr, ptr [[BBEGIN2]]
 132 // CK0-DAG: [[BARRBEGINGEP:%.+]] = getelementptr inbounds double, ptr [[BARRBEGIN]], i[[sz:64|32]] 0
 133 // CK0-DAG: [[BEND:%.+]] = getelementptr ptr, ptr [[BBEGIN]], i32 1
 134 // CK0-DAG: [[ABEGINI:%.+]] = ptrtoint ptr [[ABEGIN]] to i64
 135 // CK0-DAG: [[BENDI:%.+]] = ptrtoint ptr [[BEND]] to i64
 136 // CK0-DAG: [[CSIZE:%.+]] = sub i64 [[BENDI]], [[ABEGINI]]
 137 // CK0-DAG: [[CUSIZE:%.+]] = sdiv exact i64 [[CSIZE]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
 138 // CK0-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(ptr [[HANDLE]])
 139 // CK0-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48
 140 // CK0-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 0, [[SHIPRESIZE]]
 141 // CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
 142 // CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
 143 // CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
 144 // CK0-DAG: [[ALLOC]]
 145 // CK0-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
 146 // CK0-DAG: br label %[[TYEND:[^,]+]]
 147 // CK0-DAG: [[ALLOCELSE]]
 148 // CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
 149 // CK0-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
 150 // CK0-DAG: [[TO]]
 151 // CK0-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
 152 // CK0-DAG: br label %[[TYEND]]
 153 // CK0-DAG: [[TOELSE]]
 154 // CK0-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
 155 // CK0-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
 156 // CK0-DAG: [[FROM]]
 157 // CK0-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
 158 // CK0-DAG: br label %[[TYEND]]
 159 // CK0-DAG: [[TYEND]]
 160 // CK0-DAG: [[PHITYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
 161 // CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 [[CUSIZE]], i64 [[PHITYPE0]], {{.*}})
 162 // 281474976710659 == 0x1,000,000,003
 163 // CK0-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]]
 164 // CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
 165 // CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
 166 // CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
 167 // CK0-DAG: [[ALLOC]]
 168 // CK0-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
 169 // CK0-DAG: br label %[[TYEND:[^,]+]]
 170 // CK0-DAG: [[ALLOCELSE]]
 171 // CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
 172 // CK0-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
 173 // CK0-DAG: [[TO]]
 174 // CK0-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
 175 // CK0-DAG: br label %[[TYEND]]
 176 // CK0-DAG: [[TOELSE]]
 177 // CK0-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
 178 // CK0-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
 179 // CK0-DAG: [[FROM]]
 180 // CK0-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
 181 // CK0-DAG: br label %[[TYEND]]
 182 // CK0-DAG: [[TYEND]]
 183 // CK0-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
 184 // CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 4, i64 [[TYPE1]], {{.*}})
 185 // 281474976710675 == 0x1,000,000,013
 186 // CK0-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710675, [[SHIPRESIZE]]
 187 // CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
 188 // CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
 189 // CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
 190 // CK0-DAG: [[ALLOC]]
 191 // CK0-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
 192 // CK0-DAG: br label %[[TYEND:[^,]+]]
 193 // CK0-DAG: [[ALLOCELSE]]
 194 // CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
 195 // CK0-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
 196 // CK0-DAG: [[TO]]
 197 // CK0-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
 198 // CK0-DAG: br label %[[TYEND]]
 199 // CK0-DAG: [[TOELSE]]
 200 // CK0-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
 201 // CK0-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
 202 // CK0-DAG: [[FROM]]
 203 // CK0-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
 204 // CK0-DAG: br label %[[TYEND]]
 205 // CK0-DAG: [[TYEND]]
 206 // CK0-DAG: [[TYPE2:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
 207 // CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BBEGIN]], ptr [[BARRBEGINGEP]], i64 16, i64 [[TYPE2]], {{.*}})
 208 // CK0: [[PTRNEXT]] = getelementptr %class.C, ptr [[PTR]], i32 1
 209 // CK0: [[ISDONE:%.+]] = icmp eq ptr [[PTRNEXT]], [[PTREND]]
 210 // CK0: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
 211
 212 // CK0: [[LEXIT]]
 213 // CK0: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1
 214 // CK0: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
 215 // CK0: [[ISNOTDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
 216 // CK0: [[CMP1:%.+]] = and i1 [[ISARRAY]], [[ISNOTDEL]]
 217 // CK0: br i1 [[CMP1]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
 218 // CK0: [[EVALDEL]]
 219 // CK0-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
 220 // CK0-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
 221
 222 // Remove movement mappings and mark as implicit
 223 // CK0-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
 224 // CK0-DAG: [[DTYPE1:%.+]] = or i64 [[DTYPE]], 512
 225 // CK0: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BPTR]], ptr [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE1]], {{.*}})
 226 // CK0: br label %[[DONE]]
 227 // CK0: [[DONE]]
 228 // CK0: ret void
 229
 230
 231 // CK0-LABEL: define {{.*}}void @{{.*}}foo{{.*}}
 232 void foo(int a){
 233   int i = a;
 234   C c;
 235   c.a = a;
 236
 237 // CK0-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 238 // CK0-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 239 // CK0-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 240 // CK0-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 241 // CK0-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 242 // CK0-DAG: [[MARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 7
 243 // CK0-DAG: store ptr [[MPR:%.+]], ptr [[MARG]]
 244 // CK0-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
 245 // CK0-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
 246 // CK0-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
 247 // CK0-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
 248 // CK0-DAG: [[MPR1:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i[[sz]] 0, i[[sz]] 0
 249 // CK0-DAG: store ptr [[VAL:%[^,]+]], ptr [[BP1]]
 250 // CK0-DAG: store ptr [[VAL]], ptr [[P1]]
 251 // CK0-DAG: store ptr [[MPRFUNC]], ptr [[MPR1]]
 252 // CK0: call void [[KERNEL_1:@.+]](ptr [[VAL]])
 253 #pragma omp target map(mapper(id), tofrom \
 254                        : c)
 255   {
 256     ++c.a;
 257   }
 258
 259   // CK0: [[BP2GEP:%.+]] = getelementptr inbounds [1 x ptr], ptr [[OFFLOAD_BP2:%[^,]+]], i32 0, i32 0
 260   // CK0: store ptr [[CADDR:%[^,]+]], ptr [[BP2GEP]], align
 261   // CK0: [[P2GEP:%.+]] = getelementptr inbounds [1 x ptr], ptr [[OFFLOAD_P2:%[^,]+]], i32 0, i32 0
 262   // CK0: store ptr [[CADDR]], ptr [[P2GEP]], align
 263   // CK0: [[MAPPER2GEP:%.+]] = getelementptr inbounds [1 x ptr], ptr [[OFFLOAD_MAPPER2:%[^,]+]], i[[SZ]] 0, i[[SZ]] 0
 264   // CK0: store ptr [[MPRFUNC]], ptr [[MAPPER2GEP]], align
 265   // CK0: [[BP2:%.+]] = getelementptr inbounds [1 x ptr], ptr [[OFFLOAD_BP2]], i32 0, i32 0
 266   // CK0: [[P2:%.+]] = getelementptr inbounds [1 x ptr], ptr [[OFFLOAD_P2]], i32 0, i32 0
 267   // CK0-32: [[TASK:%.+]] = call ptr @__kmpc_omp_target_task_alloc(ptr {{@.+}}, i32 {{%.+}}, i32 1, i32 40, i32 4, ptr [[TASK_ENTRY:@.+]], i64 -1)
 268   // CK0-64: [[TASK:%.+]] = call ptr @__kmpc_omp_target_task_alloc(ptr {{@.+}}, i32 {{%.+}}, i32 1, i64 72, i64 8, ptr [[TASK_ENTRY:@.+]], i64 -1)
 269   // CK0: [[TASK_WITH_PRIVATES:%.+]] = getelementptr inbounds [[KMP_TASK_T_WITH_PRIVATES]], ptr [[TASK]], i32 0, i32 1
 270   // CK0: {{.+}} = call i32 @__kmpc_omp_task(ptr @1, i32 {{.+}}, ptr [[TASK]])
 271   #pragma omp target map(mapper(id),tofrom: c) nowait
 272   {
 273     ++c.a;
 274   }
 275
 276 // CK0-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 0, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 277 // CK0-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 278 // CK0-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 279 // CK0-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 280 // CK0-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 281 // CK0-DAG: [[MARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 7
 282 // CK0-DAG: store ptr [[MPRGEP:%.+]], ptr [[MARG]]
 283 // CK0-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
 284 // CK0-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
 285 // CK0-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
 286 // CK0-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
 287 // CK0-DAG: [[MPR1:%.+]] = getelementptr inbounds {{.+}}[[MPRGEP]], i[[sz]] 0, i[[sz]] 0
 288 // CK0-DAG: store ptr [[VAL:%[^,]+]], ptr [[BP1]]
 289 // CK0-DAG: store ptr [[VAL]], ptr [[P1]]
 290 // CK0-DAG: store ptr [[MPRFUNC]], ptr [[MPR1]]
 291 // CK0: call void [[KERNEL_3:@.+]](ptr [[VAL]])
 292 #pragma omp target teams map(mapper(id), to \
 293                              : c)
 294   {
 295     ++c.a;
 296   }
 297
 298   // CK0-32: [[TASK_1:%.+]] = call ptr @__kmpc_omp_target_task_alloc(ptr {{@.+}}, i32 {{%.+}}, i32 1, i32 40, i32 4, ptr [[TASK_ENTRY_1:@.+]], i64 -1)
 299   // CK0-64: [[TASK_1:%.+]] = call ptr @__kmpc_omp_target_task_alloc(ptr {{@.+}}, i32 {{%.+}}, i32 1, i64 72, i64 8, ptr [[TASK_ENTRY_1:@.+]], i64 -1)
 300   // CK0: [[TASK_CAST_GET_1:%.+]] = getelementptr inbounds [[KMP_TASK_T_WITH_PRIVATES_1]], ptr [[TASK_1]], i32 0, i32 0
 301   // CK0: {{.+}} = getelementptr inbounds [[KMP_TASK_T]], ptr [[TASK_CAST_GET_1]], i32 0, i32 0
 302   // CK0: {{.+}} = call i32 @__kmpc_omp_task(ptr @1, i32 {{.+}}, ptr [[TASK_1]])
 303   #pragma omp target teams map(mapper(id),to: c) nowait
 304   {
 305     ++c.a;
 306   }
 307
 308   // CK0-DAG: call void @__tgt_target_data_begin_mapper(ptr @{{.+}}, i64 {{.+}}, i32 1, ptr [[BPGEP:%[0-9]+]], ptr [[PGEP:%[0-9]+]], {{.+}}[[EDSIZES]], {{.+}}[[EDTYPES]], ptr null, ptr [[MPR:%.+]])
 309   // CK0-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
 310   // CK0-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
 311   // CK0-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
 312   // CK0-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
 313   // CK0-DAG: [[MPR1:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i[[sz]] 0, i[[sz]] 0
 314   // CK0-DAG: store ptr [[VAL:%[^,]+]], ptr [[BP1]]
 315   // CK0-DAG: store ptr [[VAL]], ptr [[P1]]
 316   // CK0-DAG: store ptr [[MPRFUNC]], ptr [[MPR1]]
 317   #pragma omp target enter data map(mapper(id),to: c)
 318
 319   // CK0-DAG: call i32 @__kmpc_omp_task(ptr @{{[^,]+}}, i32 %{{[^,]+}}, ptr [[TASK_2:%.+]])
 320   // CK0-DAG: [[TASK_2]] = call ptr @__kmpc_omp_target_task_alloc(ptr @{{[^,]+}}, i32 %{{[^,]+}}, i32 1, i[[sz]] {{40|72}}, i[[sz]] 1, ptr [[OMP_TASK_ENTRY_18:@[^,]+]], i64 -1)
 321   // CK0-DAG: [[PRIVATES:%.+]] = getelementptr inbounds [[KMP_TASK_T_WITH_PRIVATES_4]], ptr [[TASK_2]], i32 0, i32 1
 322   // CK0-32-DAG: [[FPBPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_5]], ptr [[PRIVATES]], i32 0, i32 1
 323   // CK0-64-DAG: [[FPBPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_5]], ptr [[PRIVATES]], i32 0, i32 0
 324   // CK0-DAG: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{4|8}} [[FPBPGEP]], ptr align {{4|8}} [[BPGEP:%.+]], i[[sz]] {{4|8}}, i1 false)
 325   // CK0-DAG: [[BPGEP]] = getelementptr inbounds [1 x ptr], ptr [[BP:%.+]], i32 0, i32 0
 326   // CK0-DAG: [[BPGEP:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP]], i32 0, i32 0
 327   // CK0-DAG: store ptr [[C:%[^,]+]], ptr [[BPGEP]], align
 328   // CK0-32-DAG: [[FPPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_5]], ptr [[PRIVATES]], i32 0, i32 2
 329   // CK0-64-DAG: [[FPPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_5]], ptr [[PRIVATES]], i32 0, i32 1
 330   // CK0-DAG: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{4|8}} [[FPPGEP]], ptr align {{4|8}} [[PGEP:%.+]], i[[sz]] {{4|8}}, i1 false)
 331   // CK0-DAG: [[PGEP]] = getelementptr inbounds [1 x ptr], ptr [[P:%.+]], i32 0, i32 0
 332   // CK0-DAG: [[PGEP:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P]], i32 0, i32 0
 333   // CK0-DAG: store ptr [[C]], ptr [[PGEP]], align
 334   // CK0-32-DAG: [[FPSZGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_5]], ptr [[PRIVATES]], i32 0, i32 0
 335   // CK0-64-DAG: [[FPSZGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_5]], ptr [[PRIVATES]], i32 0, i32 2
 336   // CK0-DAG: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{4|8}} [[FPSZGEP]], ptr align {{4|8}} [[EDNWSIZES]], i[[sz]] {{4|8}}, i1 false)
 337   // CK0-DAG: [[FPMPRGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_5]], ptr [[PRIVATES]], i32 0, i32 3
 338   // CK0-DAG: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{4|8}} [[FPMPRGEP]], ptr align {{4|8}} [[MPR:%.+]], i[[sz]] {{4|8}}, i1 false)
 339   // CK0-DAG: [[MPRGEP:%.+]] = getelementptr inbounds [1 x ptr], ptr [[MPR]], i[[sz]] 0, i[[sz]] 0
 340   // CK0-DAG: store ptr [[MPRFUNC]], ptr [[MPRGEP]], align
 341   #pragma omp target enter data map(mapper(id),to: c) nowait
 342
 343   // CK0-DAG: call void @__tgt_target_data_end_mapper(ptr @{{.+}}, i64 {{.+}}, i32 1, ptr [[BPGEP:%[0-9]+]], ptr [[PGEP:%[0-9]+]], {{.+}}[[EXDSIZES]], {{.+}}[[EXDTYPES]], ptr null, ptr [[MPR:%.+]])
 344   // CK0-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
 345   // CK0-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
 346   // CK0-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
 347   // CK0-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
 348   // CK0-DAG: [[MPR1:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i[[sz]] 0, i[[sz]] 0
 349   // CK0-DAG: store ptr [[VAL:%[^,]+]], ptr [[BP1]]
 350   // CK0-DAG: store ptr [[VAL]], ptr [[P1]]
 351   // CK0-DAG: store ptr [[MPRFUNC]], ptr [[MPR1]]
 352   #pragma omp target exit data map(mapper(id),from: c)
 353
 354   // CK0-DAG: call i32 @__kmpc_omp_task(ptr @{{[^,]+}}, i32 %{{[^,]+}}, ptr [[TASK_3:%.+]])
 355   // CK0-DAG: [[TASK_3]] = call ptr @__kmpc_omp_target_task_alloc(ptr @{{[^,]+}}, i32 %{{[^,]+}}, i32 1, i[[sz]] {{40|72}}, i[[sz]] 1, ptr [[OMP_TASK_ENTRY_25:@[^,]+]], i64 -1)
 356   // CK0-DAG: [[PRIVATES:%.+]] = getelementptr inbounds [[KMP_TASK_T_WITH_PRIVATES_7]], ptr [[TASK_3]], i32 0, i32 1
 357   // CK0-32-DAG: [[FPBPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_8]], ptr [[PRIVATES]], i32 0, i32 1
 358   // CK0-64-DAG: [[FPBPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_8]], ptr [[PRIVATES]], i32 0, i32 0
 359   // CK0-DAG: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{4|8}} [[FPBPGEP]], ptr align {{4|8}} [[BPGEP:%.+]], i[[sz]] {{4|8}}, i1 false)
 360   // CK0-DAG: [[BPGEP]] = getelementptr inbounds [1 x ptr], ptr [[BP:%.+]], i32 0, i32 0
 361   // CK0-DAG: [[BPGEP:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP]], i32 0, i32 0
 362   // CK0-DAG: store ptr [[C:%[^,]+]], ptr [[BPGEP]], align
 363   // CK0-32-DAG: [[FPPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_8]], ptr [[PRIVATES]], i32 0, i32 2
 364   // CK0-64-DAG: [[FPPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_8]], ptr [[PRIVATES]], i32 0, i32 1
 365   // CK0-DAG: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{4|8}} [[FPPGEP]], ptr align {{4|8}} [[PGEP:%.+]], i[[sz]] {{4|8}}, i1 false)
 366   // CK0-DAG: [[PGEP]] = getelementptr inbounds [1 x ptr], ptr [[P:%.+]], i32 0, i32 0
 367   // CK0-DAG: [[PGEP:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P]], i32 0, i32 0
 368   // CK0-DAG: store ptr [[C]], ptr [[PGEP]], align
 369   // CK0-32-DAG: [[FPSZGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_8]], ptr [[PRIVATES]], i32 0, i32 0
 370   // CK0-64-DAG: [[FPSZGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_8]], ptr [[PRIVATES]], i32 0, i32 2
 371   // CK0-DAG: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{4|8}} [[FPSZGEP]], ptr align {{4|8}} [[EXDNWSIZES]], i[[sz]] {{4|8}}, i1 false)
 372   // CK0-DAG: [[FPMPRGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_8]], ptr [[PRIVATES]], i32 0, i32 3
 373   // CK0-DAG: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{4|8}} [[FPMPRGEP]], ptr align {{4|8}} [[MPR:%.+]], i[[sz]] {{4|8}}, i1 false)
 374   // CK0-DAG: [[MPRGEP:%.+]] = getelementptr inbounds [1 x ptr], ptr [[MPR]], i[[sz]] 0, i[[sz]] 0
 375   // CK0-DAG: store ptr [[MPRFUNC]], ptr [[MPRGEP]], align
 376   #pragma omp target exit data map(mapper(id),from: c) nowait
 377
 378   // CK0-DAG: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[TGEPBP:%.+]], ptr [[TGEPP:%.+]], ptr [[TSIZES]], ptr [[TTYPES]], ptr null, ptr [[TMPR:%.+]])
 379   // CK0-DAG: [[TGEPBP]] = getelementptr inbounds {{.+}}[[TBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
 380   // CK0-DAG: [[TGEPP]] = getelementptr inbounds {{.+}}[[TP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
 381   // CK0-DAG: [[TBP0:%.+]] = getelementptr inbounds {{.+}}[[TBP]], i{{.+}} 0, i{{.+}} 0
 382   // CK0-DAG: [[TP0:%.+]] = getelementptr inbounds {{.+}}[[TP]], i{{.+}} 0, i{{.+}} 0
 383   // CK0-DAG: [[TMPR1:%.+]] = getelementptr inbounds {{.+}}[[TMPR]], i[[sz]] 0, i[[sz]] 0
 384   // CK0-DAG: store ptr [[VAL]], ptr [[TBP0]]
 385   // CK0-DAG: store ptr [[VAL]], ptr [[TP0]]
 386   // CK0-DAG: store ptr [[MPRFUNC]], ptr [[TMPR1]]
 387   #pragma omp target update to(mapper(id): c)
 388
 389   // CK0-DAG: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[FGEPBP:%.+]], ptr [[FGEPP:%.+]], ptr [[FSIZES]], ptr [[FTYPES]], ptr null, ptr [[FMPR:%.+]])
 390   // CK0-DAG: [[FGEPBP]] = getelementptr inbounds {{.+}}[[FBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
 391   // CK0-DAG: [[FGEPP]] = getelementptr inbounds {{.+}}[[FP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
 392   // CK0-DAG: [[FBP0:%.+]] = getelementptr inbounds {{.+}}[[FBP]], i{{.+}} 0, i{{.+}} 0
 393   // CK0-DAG: [[FP0:%.+]] = getelementptr inbounds {{.+}}[[FP]], i{{.+}} 0, i{{.+}} 0
 394   // CK0-DAG: [[FMPR1:%.+]] = getelementptr inbounds {{.+}}[[FMPR]], i[[sz]] 0, i[[sz]] 0
 395   // CK0-DAG: store ptr [[VAL]], ptr [[FBP0]]
 396   // CK0-DAG: store ptr [[VAL]], ptr [[FP0]]
 397   // CK0-DAG: store ptr [[MPRFUNC]], ptr [[FMPR1]]
 398   #pragma omp target update from(mapper(id): c)
 399
 400   // CK0-DAG: call i32 @__kmpc_omp_task(ptr @{{[^,]+}}, i32 %{{[^,]+}}, ptr [[TASK_4:%.+]])
 401   // CK0-DAG: [[TASK_4]] = call ptr @__kmpc_omp_target_task_alloc(ptr @{{[^,]+}}, i32 %{{[^,]+}}, i32 1, i[[sz]] {{40|72}}, i[[sz]] 1, ptr [[OMP_TASK_ENTRY_34:@[^,]+]], i64 -1)
 402   // CK0-DAG: [[PRIVATES:%.+]] = getelementptr inbounds [[KMP_TASK_T_WITH_PRIVATES_10]], ptr [[TASK_4]], i32 0, i32 1
 403   // CK0-32-DAG: [[FPBPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_11]], ptr [[PRIVATES]], i32 0, i32 1
 404   // CK0-64-DAG: [[FPBPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_11]], ptr [[PRIVATES]], i32 0, i32 0
 405   // CK0-DAG: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{4|8}} [[FPBPGEP]], ptr align {{4|8}} [[BPGEP:%.+]], i[[sz]] {{4|8}}, i1 false)
 406   // CK0-DAG: [[BPGEP]] = getelementptr inbounds [1 x ptr], ptr [[BP:%.+]], i32 0, i32 0
 407   // CK0-DAG: [[BPGEP:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP]], i32 0, i32 0
 408   // CK0-DAG: store ptr [[C:%[^,]+]], ptr [[BPGEP]], align
 409   // CK0-32-DAG: [[FPPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_11]], ptr [[PRIVATES]], i32 0, i32 2
 410   // CK0-64-DAG: [[FPPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_11]], ptr [[PRIVATES]], i32 0, i32 1
 411   // CK0-DAG: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{4|8}} [[FPPGEP]], ptr align {{4|8}} [[PGEP:%.+]], i[[sz]] {{4|8}}, i1 false)
 412   // CK0-DAG: [[PGEP]] = getelementptr inbounds [1 x ptr], ptr [[P:%.+]], i32 0, i32 0
 413   // CK0-DAG: [[PGEP:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P]], i32 0, i32 0
 414   // CK0-DAG: store ptr [[C]], ptr [[PGEP]], align
 415   // CK0-32-DAG: [[FPSZGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_11]], ptr [[PRIVATES]], i32 0, i32 0
 416   // CK0-64-DAG: [[FPSZGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_11]], ptr [[PRIVATES]], i32 0, i32 2
 417   // CK0-DAG: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{4|8}} [[FPSZGEP]], ptr align {{4|8}} [[FNWSIZES]], i[[sz]] {{4|8}}, i1 false)
 418   // CK0-DAG: [[FPMPRGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_11]], ptr [[PRIVATES]], i32 0, i32 3
 419   // CK0-DAG: call void @llvm.memcpy.p0.p0.i[[sz]](ptr align {{4|8}} [[FPMPRGEP]], ptr align {{4|8}} [[MPR:%.+]], i[[sz]] {{4|8}}, i1 false)
 420   // CK0-DAG: [[MPRGEP:%.+]] = getelementptr inbounds [1 x ptr], ptr [[MPR]], i[[sz]] 0, i[[sz]] 0
 421   // CK0-DAG: store ptr [[MPRFUNC]], ptr [[MPRGEP]], align
 422   #pragma omp target update from(mapper(id): c) nowait
 423 }
 424
 425
 426 // CK0: define internal void [[KERNEL_1]](ptr {{.+}}[[ARG:%.+]])
 427 // CK0: [[ADDR:%.+]] = alloca ptr,
 428 // CK0: store ptr [[ARG]], ptr [[ADDR]]
 429 // CK0: [[CADDR:%.+]] = load ptr, ptr [[ADDR]]
 430 // CK0: [[CAADDR:%.+]] = getelementptr inbounds %class.C, ptr [[CADDR]], i32 0, i32 0
 431 // CK0: [[VAL:%[^,]+]] = load i32, ptr [[CAADDR]]
 432 // CK0: {{.+}} = add nsw i32 [[VAL]], 1
 433 // CK0: }
 434
 435 // CK0: define internal void [[KERNEL_2:@.+]](ptr {{.+}}[[ARG:%.+]])
 436 // CK0: [[ADDR:%.+]] = alloca ptr,
 437 // CK0: store ptr [[ARG]], ptr [[ADDR]]
 438 // CK0: [[CADDR:%.+]] = load ptr, ptr [[ADDR]]
 439 // CK0: [[CAADDR:%.+]] = getelementptr inbounds %class.C, ptr [[CADDR]], i32 0, i32 0
 440 // CK0: [[VAL:%[^,]+]] = load i32, ptr [[CAADDR]]
 441 // CK0: {{.+}} = add nsw i32 [[VAL]], 1
 442 // CK0: }
 443
 444 // CK0: define internal void [[OUTLINED:@.+]](i32 {{.*}}{{[^,]+}}, ptr noalias noundef [[CTXARG:%.+]])
 445 // CK0-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 446 // CK0-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 447 // CK0-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 448 // CK0-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 449 // CK0-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 450 // CK0-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
 451 // CK0-DAG: store ptr [[SIZEGEP:%.+]], ptr [[SARG]]
 452 // CK0-DAG: [[MARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 7
 453 // CK0-DAG: store ptr [[MPRGEP:%.+]], ptr [[MARG]]
 454 // CK0-DAG: [[BPGEP]] = getelementptr inbounds [1 x ptr], ptr [[BPFPADDR:%.+]], i[[SZ]] 0, i[[SZ]] 0
 455 // CK0-DAG: [[PGEP]] = getelementptr inbounds [1 x ptr], ptr [[PFPADDR:%.+]], i[[SZ]] 0, i[[SZ]] 0
 456 // CK0-DAG: [[SIZEGEP]] = getelementptr inbounds [1 x i64], ptr [[SIZEFPADDR:%.+]], i[[SZ]] 0, i[[SZ]] 0
 457 // CK0-DAG: [[MPRGEP]] = getelementptr inbounds [1 x ptr], ptr [[MPRFPADDR:%.+]], i[[SZ]] 0, i[[SZ]] 0
 458 // CK0-DAG: [[BPFPADDR]] = load ptr, ptr [[FPPTRADDR_BP:%.+]], align
 459 // CK0-DAG: [[PFPADDR]] = load ptr, ptr [[FPPTRADDR_P:%.+]], align
 460 // CK0-DAG: [[SIZEFPADDR]] = load ptr, ptr [[FPPTRADDR_SIZE:%.+]], align
 461 // CK0-DAG: [[MPRFPADDR]] = load ptr, ptr [[FPPTRADDR_MPR:%.+]], align
 462 // CK0-DAG: call void %1(ptr %2, {{.+}}[[FPPTRADDR_BP]], {{.+}}[[FPPTRADDR_P]], {{.+}}[[FPPTRADDR_SIZE]], {{.+}}[[FPPTRADDR_MPR]])
 463 // CK0-DAG: call void [[KERNEL_2:@.+]](ptr [[KERNELARG:%.+]])
 464 // CK0-DAG: [[KERNELARG]] = load ptr, ptr [[KERNELARGGEP:%.+]], align
 465 // CK0-DAG: [[KERNELARGGEP]] = getelementptr inbounds [[ANON_T]], ptr [[CTX:%.+]], i32 0, i32 0
 466 // CK0-DAG: [[CTX]] = load ptr, ptr [[CTXADDR:%.+]], align
 467 // CK0-DAG: store ptr [[CTXARG]], ptr [[CTXADDR]], align
 468 // CK0: }
 469
 470 // CK0: define internal {{.*}}i32 [[TASK_ENTRY]](i32 {{.*}}%0, ptr noalias noundef %1)
 471 // CK0: store ptr %1, ptr [[ADDR:%.+]], align
 472 // CK0: [[TASK_T_WITH_PRIVATES:%.+]] = load ptr, ptr [[ADDR]], align
 473 // CK0: [[TASKGEP:%.+]] = getelementptr inbounds [[KMP_TASK_T_WITH_PRIVATES]], ptr [[TASK_T_WITH_PRIVATES]], i32 0, i32 0
 474 // CK0: [[SHAREDSGEP:%.+]] = getelementptr inbounds [[KMP_TASK_T]], ptr [[TASKGEP]], i32 0, i32 0
 475 // CK0: [[SHAREDS:%.+]] = load ptr, ptr [[SHAREDSGEP]], align
 476 // CK0: [[PRIVATESGEP:%.+]] = getelementptr inbounds [[KMP_TASK_T_WITH_PRIVATES]], ptr [[TASK_T_WITH_PRIVATES]], i32 0, i32 1
 477 // CK0: call void [[OUTLINED]](i32 {{%.+}}, ptr {{%.+}}, ptr [[PRIVATESGEP]], {{.+}}, ptr [[TASK_T_WITH_PRIVATES]], ptr [[SHAREDS]])
 478 // CK0: }
 479
 480 // CK0: define internal void [[OUTLINE_1:@.+]](i32 {{.*}}%.global_tid.{{.+}}, ptr noalias noundef [[CTXARG:%.+]])
 481 // CK0-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 0, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 482 // CK0-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 483 // CK0-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 484 // CK0-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 485 // CK0-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 486 // CK0-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4
 487 // CK0-DAG: store ptr [[SIZEGEP:%.+]], ptr [[SARG]]
 488 // CK0-DAG: [[MARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 7
 489 // CK0-DAG: store ptr [[MPRGEP:%.+]], ptr [[MARG]]
 490 // CK0-DAG: [[BPGEP]] = getelementptr inbounds [1 x ptr], ptr [[BPFPADDR:%.+]], i[[SZ]] 0, i[[SZ]] 0
 491 // CK0-DAG: [[PGEP]] = getelementptr inbounds [1 x ptr], ptr [[PFPADDR:%.+]], i[[SZ]] 0, i[[SZ]] 0
 492 // CK0-DAG: [[SIZEGEP]] = getelementptr inbounds [1 x i64], ptr [[SIZEFPADDR:%.+]], i[[SZ]] 0, i[[SZ]] 0
 493 // CK0-DAG: [[MPRGEP]] = getelementptr inbounds [1 x ptr], ptr [[MPRFPADDR:%.+]], i[[SZ]] 0, i[[SZ]] 0
 494 // CK0-DAG: [[BPFPADDR]] = load ptr, ptr [[FPPTRADDR_BP:%.+]], align
 495 // CK0-DAG: [[PFPADDR]] = load ptr, ptr [[FPPTRADDR_P:%.+]], align
 496 // CK0-DAG: [[SIZEFPADDR]] = load ptr, ptr [[FPPTRADDR_SIZE:%.+]], align
 497 // CK0-DAG: [[MPRFPADDR]] = load ptr, ptr [[FPPTRADDR_MPR:%.+]], align
 498 // CK0-DAG: call void %1(ptr %2, {{.+}}[[FPPTRADDR_BP]], {{.+}}[[FPPTRADDR_P]], {{.+}}[[FPPTRADDR_SIZE]], {{.+}}[[FPPTRADDR_MPR]])
 499 // CK0-DAG: call void [[KERNEL_2:@.+]](ptr [[KERNELARG:%.+]])
 500 // CK0-DAG: [[KERNELARG]] = load ptr, ptr [[KERNELARGGEP:%.+]], align
 501 // CK0-DAG: [[KERNELARGGEP]] = getelementptr inbounds [[ANON_T_0]], ptr [[CTX:%.+]], i32 0, i32 0
 502 // CK0-DAG: [[CTX]] = load ptr, ptr [[CTXADDR:%.+]], align
 503 // CK0-DAG: store ptr [[CTXARG]], ptr [[CTXADDR]], align
 504 // CK0: }
 505
 506 // CK0: define internal {{.*}}i32 [[TASK_ENTRY_1]](i32 {{.*}}%0, ptr noalias noundef %1)
 507 // CK0: store ptr %1, ptr [[ADDR:%.+]], align
 508 // CK0: [[TASK_T_WITH_PRIVATES:%.+]] = load ptr, ptr [[ADDR]], align
 509 // CK0: [[TASKGEP:%.+]] = getelementptr inbounds [[KMP_TASK_T_WITH_PRIVATES_1]], ptr [[TASK_T_WITH_PRIVATES]], i32 0, i32 0
 510 // CK0: [[SHAREDSGEP:%.+]] = getelementptr inbounds [[KMP_TASK_T]], ptr [[TASKGEP]], i32 0, i32 0
 511 // CK0: [[SHAREDS:%.+]] = load ptr, ptr [[SHAREDSGEP]], align
 512 // CK0: [[PRIVATESGEP:%.+]] = getelementptr inbounds [[KMP_TASK_T_WITH_PRIVATES_1]], ptr [[TASK_T_WITH_PRIVATES]], i32 0, i32 1
 513 // CK0: call void [[OUTLINE_1]](i32 {{%.+}}, ptr {{%.+}}, ptr [[PRIVATESGEP]], {{.+}}, ptr [[TASK_T_WITH_PRIVATES]], ptr [[SHAREDS]])
 514 // CK0: }
 515
 516 // CK0: define internal void [[OMP_OUTLINED_16:@.+]](i32{{.*}} %{{[^,]+}}, ptr noalias noundef %{{[^,]+}}, ptr noalias noundef %{{[^,]+}}
 517 // CK0-DAG: call void @__tgt_target_data_begin_nowait_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[BP:%[^,]+]], ptr [[P:%[^,]+]], ptr [[SZ:%[^,]+]], ptr [[EDNWTYPES]], ptr null, ptr [[MPR:%.+]])
 518 // CK0-DAG: [[BP]] = getelementptr inbounds [1 x ptr], ptr [[BPADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
 519 // CK0-DAG: [[P]] = getelementptr inbounds [1 x ptr], ptr [[PADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
 520 // CK0-DAG: [[SZ]] = getelementptr inbounds [1 x i64], ptr [[SZADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
 521 // CK0-DAG: [[MPR]] = getelementptr inbounds [1 x ptr], ptr [[MPRADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
 522 // CK0-DAG: [[BPADDR]] = load ptr, ptr [[FPBPADDR:%[^,]+]], align
 523 // CK0-DAG: [[PADDR]] = load ptr, ptr [[FPPADDR:%[^,]+]], align
 524 // CK0-DAG: [[SZADDR]] = load ptr, ptr [[FPSZADDR:%[^,]+]], align
 525 // CK0-DAG: [[MPRADDR]] = load ptr, ptr [[FPMPRADDR:%[^,]+]], align
 526 // CK0-DAG: call void %{{.+}}(ptr %{{[^,]+}}, ptr [[FPBPADDR]], ptr [[FPPADDR]], ptr [[FPSZADDR]], ptr [[FPMPRADDR]])
 527 // CK0: ret void
 528 // CK0: }
 529
 530 // CK0: define internal {{.*}}i32 [[OMP_TASK_ENTRY_18]](i32 {{.*}}%{{[^,]+}}, ptr noalias noundef %{{[^,]+}})
 531 // CK0:   call void [[OMP_OUTLINED_16]]
 532 // CK0:   ret i32 0
 533 // CK0: }
 534
 535 // CK0: define internal void [[OMP_OUTLINED_23:@.+]](i32{{.*}} %{{[^,]+}}, ptr noalias noundef %{{[^,]+}}, ptr noalias noundef %{{[^,]+}}
 536 // CK0-DAG: call void @__tgt_target_data_end_nowait_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[BP:%[^,]+]], ptr [[P:%[^,]+]], ptr [[SZ:%[^,]+]], ptr [[EXDNWTYPES]], ptr null, ptr [[MPR:%.+]])
 537 // CK0-DAG: [[BP]] = getelementptr inbounds [1 x ptr], ptr [[BPADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
 538 // CK0-DAG: [[P]] = getelementptr inbounds [1 x ptr], ptr [[PADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
 539 // CK0-DAG: [[SZ]] = getelementptr inbounds [1 x i64], ptr [[SZADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
 540 // CK0-DAG: [[MPR]] = getelementptr inbounds [1 x ptr], ptr [[MPRADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
 541 // CK0-DAG: [[BPADDR]] = load ptr, ptr [[FPBPADDR:%[^,]+]], align
 542 // CK0-DAG: [[PADDR]] = load ptr, ptr [[FPPADDR:%[^,]+]], align
 543 // CK0-DAG: [[SZADDR]] = load ptr, ptr [[FPSZADDR:%[^,]+]], align
 544 // CK0-DAG: [[MPRADDR]] = load ptr, ptr [[FPMPRADDR:%[^,]+]], align
 545 // CK0-DAG: call void %{{.+}}(ptr %{{[^,]+}}, ptr [[FPBPADDR]], ptr [[FPPADDR]], ptr [[FPSZADDR]], ptr [[FPMPRADDR]])
 546 // CK0: }
 547
 548 // CK0: define internal {{.*}}i32 [[OMP_TASK_ENTRY_25]](i32 {{.*}}%{{[^,]+}}, ptr noalias noundef %{{[^,]+}})
 549 // CK0:   call void [[OMP_OUTLINED_23]]
 550 // CK0:   ret i32 0
 551 // CK0: }
 552
 553 // CK0: define internal void [[OMP_OUTLINED_32:@.+]](i32{{.*}} %{{[^,]+}}, ptr noalias noundef %{{[^,]+}}, ptr noalias noundef %{{[^,]+}}
 554 // CK0-DAG: call void @__tgt_target_data_update_nowait_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[BP:%[^,]+]], ptr [[P:%[^,]+]], ptr [[SZ:%[^,]+]], ptr [[FNWTYPES]], ptr null, ptr [[MPR:%.+]])
 555 // CK0-DAG: [[BP]] = getelementptr inbounds [1 x ptr], ptr [[BPADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
 556 // CK0-DAG: [[P]] = getelementptr inbounds [1 x ptr], ptr [[PADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
 557 // CK0-DAG: [[SZ]] = getelementptr inbounds [1 x i64], ptr [[SZADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
 558 // CK0-DAG: [[MPR]] = getelementptr inbounds [1 x ptr], ptr [[MPRADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
 559 // CK0-DAG: [[BPADDR]] = load ptr, ptr [[FPBPADDR:%[^,]+]], align
 560 // CK0-DAG: [[PADDR]] = load ptr, ptr [[FPPADDR:%[^,]+]], align
 561 // CK0-DAG: [[SZADDR]] = load ptr, ptr [[FPSZADDR:%[^,]+]], align
 562 // CK0-DAG: [[MPRADDR]] = load ptr, ptr [[FPMPRADDR:%[^,]+]], align
 563 // CK0-DAG: call void %{{.+}}(ptr %{{[^,]+}}, ptr [[FPBPADDR]], ptr [[FPPADDR]], ptr [[FPSZADDR]], ptr [[FPMPRADDR]])
 564 // CK0: }
 565
 566 // CK0: define internal {{.*}}i32 [[OMP_TASK_ENTRY_34]](i32 {{.*}}%{{[^,]+}}, ptr noalias noundef %{{[^,]+}})
 567 // CK0:   call void [[OMP_OUTLINED_32]]
 568 // CK0:   ret i32 0
 569 // CK0: }
 570
 571 #endif // CK0
 572
 573
 574 ///==========================================================================///
 575 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK1 --check-prefix CK1-64 %s
 576 // RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
 577 // RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK1 --check-prefix CK1-64 %s
 578 // RUN: %clang_cc1 -DCK1 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK1 --check-prefix CK1-32 %s
 579 // RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
 580 // RUN: %clang_cc1 -DCK1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK1 --check-prefix CK1-32 %s
 581
 582 // RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 583 // RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
 584 // RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 585 // RUN: %clang_cc1 -DCK1 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 586 // RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
 587 // RUN: %clang_cc1 -DCK1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 588
 589 #ifdef CK1
 590 // C++ template
 591
 592 template <class T>
 593 class C {
 594 public:
 595   T a;
 596 };
 597
 598 #pragma omp declare mapper(id: C<int> s) map(s.a)
 599
 600 // CK1-LABEL: define {{.*}}void @.omp_mapper.{{.*}}C{{.*}}.id{{.*}}(ptr{{.*}}, ptr{{.*}}, ptr{{.*}}, i64{{.*}}, i64{{.*}}, ptr{{.*}})
 601 // CK1: store ptr %{{[^,]+}}, ptr [[HANDLEADDR:%[^,]+]]
 602 // CK1: store ptr %{{[^,]+}}, ptr [[BPTRADDR:%[^,]+]]
 603 // CK1: store ptr %{{[^,]+}}, ptr [[VPTRADDR:%[^,]+]]
 604 // CK1: store i64 %{{[^,]+}}, ptr [[SIZEADDR:%[^,]+]]
 605 // CK1: store i64 %{{[^,]+}}, ptr [[TYPEADDR:%[^,]+]]
 606 // CK1-DAG: [[BYTESIZE:%.+]] = load i64, ptr [[SIZEADDR]]
 607 // CK1-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 4
 608 // CK1-DAG: [[TYPE:%.+]] = load i64, ptr [[TYPEADDR]]
 609 // CK1-DAG: [[HANDLE:%.+]] = load ptr, ptr [[HANDLEADDR]]
 610 // CK1-DAG: [[BPTR:%.+]] = load ptr, ptr [[BPTRADDR]]
 611 // CK1-DAG: [[BEGIN:%.+]] = load ptr, ptr [[VPTRADDR]]
 612 // CK1-DAG: [[PTREND:%.+]] = getelementptr %class.C, ptr [[BEGIN]], i64 [[SIZE]]
 613 // CK1-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1
 614 // CK1-DAG: [[PTRSNE:%.+]] = icmp ne ptr [[BPTR]], [[BEGIN]]
 615 // CK1-DAG: [[PTRANDOBJ:%.+]] = and i64 [[TYPE]], 16
 616 // CK1-DAG: [[ISPTRANDOBJ:%.+]] = icmp ne i64 [[PTRANDOBJ]], 0
 617 // CK1-DAG: [[CMPA:%.+]] = and i1 [[PTRSNE]], [[ISPTRANDOBJ]]
 618 // CK1-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[CMPA]]
 619 // CK1-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
 620 // CK1-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0
 621 // CK1-DAG: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]]
 622 // CK1: br i1 [[CMP1]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]
 623
 624 // CK1: [[INITEVALDEL]]
 625 // CK1-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 4
 626
 627 // Remove movement mappings and mark as implicit
 628 // CK1-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
 629 // CK1-DAG: [[ITYPE1:%.+]] = or i64 [[ITYPE]], 512
 630 // CK1: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BPTR]], ptr [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE1]], {{.*}})
 631 // CK1: br label %[[LHEAD:[^,]+]]
 632
 633 // CK1: [[LHEAD]]
 634 // CK1: [[ISEMPTY:%.+]] = icmp eq ptr [[BEGIN]], [[PTREND]]
 635 // CK1: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]]
 636 // CK1: [[LBODY]]
 637 // CK1: [[PTR:%.+]] = phi ptr [ [[BEGIN]], %{{.+}} ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
 638 // CK1-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, ptr [[PTR]], i32 0, i32 0
 639 // CK1-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(ptr [[HANDLE]])
 640 // CK1-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48
 641 // CK1-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 3, [[SHIPRESIZE]]
 642 // CK1-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
 643 // CK1-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
 644 // CK1-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
 645 // CK1-DAG: [[ALLOC]]
 646 // CK1-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
 647 // CK1-DAG: br label %[[TYEND:[^,]+]]
 648 // CK1-DAG: [[ALLOCELSE]]
 649 // CK1-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
 650 // CK1-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
 651 // CK1-DAG: [[TO]]
 652 // CK1-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
 653 // CK1-DAG: br label %[[TYEND]]
 654 // CK1-DAG: [[TOELSE]]
 655 // CK1-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
 656 // CK1-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
 657 // CK1-DAG: [[FROM]]
 658 // CK1-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
 659 // CK1-DAG: br label %[[TYEND]]
 660 // CK1-DAG: [[TYEND]]
 661 // CK1-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
 662 // CK1: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 4, i64 [[TYPE1]], {{.*}})
 663 // CK1: [[PTRNEXT]] = getelementptr %class.C, ptr [[PTR]], i32 1
 664 // CK1: [[ISDONE:%.+]] = icmp eq ptr [[PTRNEXT]], [[PTREND]]
 665 // CK1: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
 666
 667 // CK1: [[LEXIT]]
 668 // CK1: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1
 669 // CK1: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
 670 // CK1: [[ISNOTDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
 671 // CK1: [[CMP1:%.+]] = and i1 [[ISARRAY]], [[ISNOTDEL]]
 672 // CK1-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 4
 673
 674 // Remove movement mappings and mark as implicit
 675 // CK1-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
 676 // CK1-DAG: [[DTYPE1:%.+]] = or i64 [[DTYPE]], 512
 677 // CK1: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BPTR]], ptr [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE1]], {{.*}})
 678 // CK1: br label %[[DONE]]
 679 // CK1: [[DONE]]
 680 // CK1: ret void
 681
 682 #endif // CK1
 683
 684
 685 ///==========================================================================///
 686 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK2 --check-prefix CK2-64 %s
 687 // RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
 688 // RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK2 --check-prefix CK2-64 %s
 689 // RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK2 --check-prefix CK2-32 %s
 690 // RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
 691 // RUN: %clang_cc1 -DCK2 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK2 --check-prefix CK2-32 %s
 692
 693 // RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 694 // RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
 695 // RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 696 // RUN: %clang_cc1 -DCK2 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 697 // RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
 698 // RUN: %clang_cc1 -DCK2 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 699
 700 #ifdef CK2
 701 // Nested mappers.
 702
 703 class B {
 704 public:
 705   double a;
 706 };
 707
 708 class C {
 709 public:
 710   double a;
 711   B b;
 712 };
 713
 714 #pragma omp declare mapper(B s) map(s.a)
 715
 716 #pragma omp declare mapper(id: C s) map(s.b)
 717
 718 // CK2: define {{.*}}void [[BMPRFUNC:@[.]omp_mapper[.].*B[.]default]](ptr{{.*}}, ptr{{.*}}, ptr{{.*}}, i64{{.*}}, i64{{.*}}, ptr{{.*}})
 719
 720 // CK2-LABEL: define {{.*}}void @.omp_mapper.{{.*}}C{{.*}}.id(ptr{{.*}}, ptr{{.*}}, ptr{{.*}}, i64{{.*}}, i64{{.*}}, ptr{{.*}})
 721 // CK2: store ptr %{{[^,]+}}, ptr [[HANDLEADDR:%[^,]+]]
 722 // CK2: store ptr %{{[^,]+}}, ptr [[BPTRADDR:%[^,]+]]
 723 // CK2: store ptr %{{[^,]+}}, ptr [[VPTRADDR:%[^,]+]]
 724 // CK2: store i64 %{{[^,]+}}, ptr [[SIZEADDR:%[^,]+]]
 725 // CK2: store i64 %{{[^,]+}}, ptr [[TYPEADDR:%[^,]+]]
 726 // CK2-DAG: [[BYTESIZE:%.+]] = load i64, ptr [[SIZEADDR]]
 727 // CK2-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 16
 728 // CK2-DAG: [[TYPE:%.+]] = load i64, ptr [[TYPEADDR]]
 729 // CK2-DAG: [[HANDLE:%.+]] = load ptr, ptr [[HANDLEADDR]]
 730 // CK2-DAG: [[BPTR:%.+]] = load ptr, ptr [[BPTRADDR]]
 731 // CK2-DAG: [[BEGIN:%.+]] = load ptr, ptr [[VPTRADDR]]
 732 // CK2-DAG: [[PTREND:%.+]] = getelementptr %class.C, ptr [[BEGIN]], i64 [[SIZE]]
 733 // CK2-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1
 734 // CK2-DAG: [[PTRSNE:%.+]] = icmp ne ptr [[BPTR]], [[BEGIN]]
 735 // CK2-DAG: [[PTRANDOBJ:%.+]] = and i64 [[TYPE]], 16
 736 // CK2-DAG: [[ISPTRANDOBJ:%.+]] = icmp ne i64 [[PTRANDOBJ]], 0
 737 // CK2-DAG: [[CMPA:%.+]] = and i1 [[PTRSNE]], [[ISPTRANDOBJ]]
 738 // CK2-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[CMPA]]
 739 // CK2-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
 740 // CK2-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0
 741 // CK2-DAG: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]]
 742 // CK2: br i1 [[CMP1]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]
 743
 744 // CK2: [[INITEVALDEL]]
 745 // CK2-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
 746
 747 // Remove movement mappings and mark as implicit
 748 // CK2-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
 749 // CK2-DAG: [[ITYPE1:%.+]] = or i64 [[ITYPE]], 512
 750 // CK2: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BPTR]], ptr [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE1]], {{.*}})
 751 // CK2: br label %[[LHEAD:[^,]+]]
 752
 753 // CK2: [[LHEAD]]
 754 // CK2: [[ISEMPTY:%.+]] = icmp eq ptr [[BEGIN]], [[PTREND]]
 755 // CK2: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]]
 756 // CK2: [[LBODY]]
 757 // CK2: [[PTR:%.+]] = phi ptr [ [[BEGIN]], %{{.+}} ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
 758 // CK2-DAG: [[BBEGIN:%.+]] = getelementptr inbounds %class.C, ptr [[PTR]], i32 0, i32 1
 759 // CK2-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(ptr [[HANDLE]])
 760 // CK2-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48
 761 // CK2-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 3, [[SHIPRESIZE]]
 762 // CK2-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
 763 // CK2-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
 764 // CK2-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
 765 // CK2-DAG: [[ALLOC]]
 766 // CK2-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
 767 // CK2-DAG: br label %[[TYEND:[^,]+]]
 768 // CK2-DAG: [[ALLOCELSE]]
 769 // CK2-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
 770 // CK2-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
 771 // CK2-DAG: [[TO]]
 772 // CK2-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
 773 // CK2-DAG: br label %[[TYEND]]
 774 // CK2-DAG: [[TOELSE]]
 775 // CK2-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
 776 // CK2-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
 777 // CK2-DAG: [[FROM]]
 778 // CK2-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
 779 // CK2-DAG: br label %[[TYEND]]
 780 // CK2-DAG: [[TYEND]]
 781 // CK2-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
 782 // CK2: call void [[BMPRFUNC]](ptr [[HANDLE]], ptr [[PTR]], ptr [[BBEGIN]], i64 8, i64 [[TYPE1]], {{.*}})
 783 // CK2: [[PTRNEXT]] = getelementptr %class.C, ptr [[PTR]], i32 1
 784 // CK2: [[ISDONE:%.+]] = icmp eq ptr [[PTRNEXT]], [[PTREND]]
 785 // CK2: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
 786
 787 // CK2: [[LEXIT]]
 788 // CK2: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1
 789 // CK2: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
 790 // CK2: [[ISNOTDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
 791 // CK2: [[CMP1:%.+]] = and i1 [[ISARRAY]], [[ISNOTDEL]]
 792 // CK2: br i1 [[CMP1]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
 793 // CK2: [[EVALDEL]]
 794 // CK2-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
 795
 796 // Remove movement mappings and mark as implicit
 797 // CK2-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
 798 // CK2-DAG: [[DTYPE1:%.+]] = or i64 [[DTYPE]], 512
 799 // CK2: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BPTR]], ptr [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE1]], {{.*}})
 800 // CK2: br label %[[DONE]]
 801 // CK2: [[DONE]]
 802 // CK2: ret void
 803
 804 #endif // CK2
 805
 806
 807 ///==========================================================================///
 808 // RUN: %clang_cc1 -DCK3 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK3 %s
 809 // RUN: %clang_cc1 -DCK3 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
 810 // RUN: %clang_cc1 -DCK3 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK3 %s
 811 // RUN: %clang_cc1 -DCK3 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK3 %s
 812 // RUN: %clang_cc1 -DCK3 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
 813 // RUN: %clang_cc1 -DCK3 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK3 %s
 814
 815 // RUN: %clang_cc1 -DCK3 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 816 // RUN: %clang_cc1 -DCK3 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
 817 // RUN: %clang_cc1 -DCK3 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 818 // RUN: %clang_cc1 -DCK3 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 819 // RUN: %clang_cc1 -DCK3 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
 820 // RUN: %clang_cc1 -DCK3 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 821
 822 #ifdef CK3
 823 // map of array sections and nested components.
 824
 825 // CK3-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}.region_id = weak constant i8 0
 826 // CK3-DAG: [[SIZES:@.+]] = {{.+}}constant [2 x i64] [i64 {{8|16}}, i64 {{80|160}}]
 827 // CK3-DAG: [[TYPES:@.+]] = {{.+}}constant [2 x i64] [i64 35, i64 35]
 828
 829 class C {
 830 public:
 831   int a;
 832   double *b;
 833 };
 834
 835 class B {
 836 public:
 837   C c;
 838 };
 839
 840 #pragma omp declare mapper(id: C s) map(s.a, s.b[0:2])
 841
 842 // CK3: define {{.*}}void [[MPRFUNC:@[.]omp_mapper[.].*C[.]id]](ptr{{.*}}, ptr{{.*}}, ptr{{.*}}, i64{{.*}}, i64{{.*}}, ptr{{.*}})
 843
 844 // CK3-LABEL: define {{.*}}void @{{.*}}foo{{.*}}
 845 void foo(int a){
 846   // CK3-DAG: [[CVAL:%.+]] = alloca [10 x %class.C]
 847   // CK3-DAG: [[BVAL:%.+]] = alloca %class.B
 848   C c[10];
 849   B b;
 850
 851   // CK3-DAG: [[BC:%.+]] = getelementptr inbounds %class.B, ptr [[BVAL]], i32 0, i32 0
 852
 853 // CK3-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])
 854 // CK3-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2
 855 // CK3-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]
 856 // CK3-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3
 857 // CK3-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]
 858 // CK3-DAG: [[MARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 7
 859 // CK3-DAG: store ptr [[MPR:%.+]], ptr [[MARG]]
 860 // CK3-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
 861 // CK3-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
 862 // CK3-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
 863 // CK3-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
 864 // CK3-DAG: [[MPR1:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i{{64|32}} 0, i{{64|32}} 0
 865 // CK3-DAG: store ptr [[BVAL]], ptr [[BP1]]
 866 // CK3-DAG: store ptr [[BC]], ptr [[P1]]
 867 // CK3-DAG: store ptr [[MPRFUNC]], ptr [[MPR1]]
 868 // CK3-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1
 869 // CK3-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1
 870 // CK3-DAG: [[MPR2:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i{{64|32}} 0, i{{64|32}} 1
 871 // CK3-DAG: store ptr [[CVAL]], ptr [[BP2]]
 872 // CK3-DAG: [[CVALGEP:%.+]] = getelementptr inbounds {{.+}}[[CVAL]], i{{64|32}} 0, i{{64|32}} 0
 873 // CK3-DAG: store ptr [[CVALGEP]], ptr [[P2]]
 874 // CK3-DAG: store ptr [[MPRFUNC]], ptr [[MPR2]]
 875 // CK3: call void [[KERNEL:@.+]](ptr [[BVAL]], ptr [[CVAL]])
 876 #pragma omp target map(mapper(id), tofrom \
 877                        : c [0:10], b.c)
 878   for (int i = 0; i < 10; i++) {
 879     b.c.a += ++c[i].a;
 880   }
 881 }
 882
 883
 884 // CK3: define internal void [[KERNEL]](ptr {{[^,]+}}, ptr {{[^,]+}})
 885
 886 #endif // CK3
 887
 888 ///==========================================================================///
 889 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK4 --check-prefix CK4-64 %s
 890 // RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
 891 // RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK4 --check-prefix CK4-64 %s
 892 // RUN: %clang_cc1 -DCK4 -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK4 --check-prefix CK4-32 %s
 893 // RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
 894 // RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK4 --check-prefix CK4-32 %s
 895
 896 // RUN: %clang_cc1 -DCK4 -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 897 // RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
 898 // RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 899 // RUN: %clang_cc1 -DCK4 -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 900 // RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s
 901 // RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
 902
 903 #ifdef CK4
 904 // Mapper function code generation and runtime interface.
 905
 906 // CK4-64: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
 907 // CK4-32: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
 908 // PRESENT=0x1000 | TO=0x1 = 0x1001
 909 // CK4: [[TTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1001]]]
 910
 911 // CK4-64: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
 912 // CK4-32: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]
 913 // PRESENT=0x1000 | FROM=0x2 = 0x1002
 914 // CK4: [[FTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1002]]]
 915
 916 class C {
 917 public:
 918   int a;
 919   double *b;
 920 };
 921
 922 #pragma omp declare mapper(id: C s) map(s.a, s.b[0:2])
 923
 924 // CK4: define {{.*}}void [[MPRFUNC:@[.]omp_mapper[.].*C[.]id]](ptr{{.*}}, ptr{{.*}}, ptr{{.*}}, i64{{.*}}, i64{{.*}}, ptr{{.*}})
 925 // CK4: store ptr %{{[^,]+}}, ptr [[HANDLEADDR:%[^,]+]]
 926 // CK4: store ptr %{{[^,]+}}, ptr [[BPTRADDR:%[^,]+]]
 927 // CK4: store ptr %{{[^,]+}}, ptr [[VPTRADDR:%[^,]+]]
 928 // CK4: store i64 %{{[^,]+}}, ptr [[SIZEADDR:%[^,]+]]
 929 // CK4: store i64 %{{[^,]+}}, ptr [[TYPEADDR:%[^,]+]]
 930 // CK4-DAG: [[BYTESIZE:%.+]] = load i64, ptr [[SIZEADDR]]
 931 // CK4-64-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 16
 932 // CK4-32-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 8
 933 // CK4-DAG: [[TYPE:%.+]] = load i64, ptr [[TYPEADDR]]
 934 // CK4-DAG: [[HANDLE:%.+]] = load ptr, ptr [[HANDLEADDR]]
 935 // CK4-DAG: [[BPTR:%.+]] = load ptr, ptr [[BPTRADDR]]
 936 // CK4-DAG: [[BEGIN:%.+]] = load ptr, ptr [[VPTRADDR]]
 937 // CK4-DAG: [[PTREND:%.+]] = getelementptr %class.C, ptr [[BEGIN]], i64 [[SIZE]]
 938 // CK4-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1
 939 // CK4-DAG: [[PTRSNE:%.+]] = icmp ne ptr [[BPTR]], [[BEGIN]]
 940 // CK4-DAG: [[PTRANDOBJ:%.+]] = and i64 [[TYPE]], 16
 941 // CK4-DAG: [[ISPTRANDOBJ:%.+]] = icmp ne i64 [[PTRANDOBJ]], 0
 942 // CK4-DAG: [[CMPA:%.+]] = and i1 [[PTRSNE]], [[ISPTRANDOBJ]]
 943 // CK4-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[CMPA]]
 944 // CK4-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
 945 // CK4-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0
 946 // CK4-DAG: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]]
 947 // CK4: br i1 [[CMP1]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]
 948
 949 // CK4: [[INITEVALDEL]]
 950 // CK4-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
 951 // CK4-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
 952
 953 // Remove movement mappings and mark as implicit
 954 // CK4-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4
 955 // CK4-DAG: [[ITYPE1:%.+]] = or i64 [[ITYPE]], 512
 956 // CK4: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BPTR]], ptr [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE1]], {{.*}})
 957 // CK4: br label %[[LHEAD:[^,]+]]
 958
 959 // CK4: [[LHEAD]]
 960 // CK4: [[ISEMPTY:%.+]] = icmp eq ptr [[BEGIN]], [[PTREND]]
 961 // CK4: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]]
 962 // CK4: [[LBODY]]
 963 // CK4: [[PTR:%.+]] = phi ptr [ [[BEGIN]], %{{.+}} ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]
 964 // CK4-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, ptr [[PTR]], i32 0, i32 0
 965 // CK4-DAG: [[BBEGIN:%.+]] = getelementptr inbounds %class.C, ptr [[PTR]], i32 0, i32 1
 966 // CK4-DAG: [[BBEGIN2:%.+]] = getelementptr inbounds %class.C, ptr [[PTR]], i32 0, i32 1
 967 // CK4-DAG: [[BARRBEGIN:%.+]] = load ptr, ptr [[BBEGIN2]]
 968 // CK4-DAG: [[BARRBEGINGEP:%.+]] = getelementptr inbounds double, ptr [[BARRBEGIN]], i[[sz:64|32]] 0
 969 // CK4-DAG: [[BEND:%.+]] = getelementptr ptr, ptr [[BBEGIN]], i32 1
 970 // CK4-DAG: [[ABEGINI:%.+]] = ptrtoint ptr [[ABEGIN]] to i64
 971 // CK4-DAG: [[BENDI:%.+]] = ptrtoint ptr [[BEND]] to i64
 972 // CK4-DAG: [[CSIZE:%.+]] = sub i64 [[BENDI]], [[ABEGINI]]
 973 // CK4-DAG: [[CUSIZE:%.+]] = sdiv exact i64 [[CSIZE]], ptrtoint (ptr getelementptr (i8, ptr null, i32 1) to i64)
 974 // CK4-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(ptr [[HANDLE]])
 975 // CK4-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48
 976 // CK4-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 0, [[SHIPRESIZE]]
 977 // CK4-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
 978 // CK4-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
 979 // CK4-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
 980 // CK4-DAG: [[ALLOC]]
 981 // CK4-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
 982 // CK4-DAG: br label %[[TYEND:[^,]+]]
 983 // CK4-DAG: [[ALLOCELSE]]
 984 // CK4-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
 985 // CK4-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
 986 // CK4-DAG: [[TO]]
 987 // CK4-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
 988 // CK4-DAG: br label %[[TYEND]]
 989 // CK4-DAG: [[TOELSE]]
 990 // CK4-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
 991 // CK4-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
 992 // CK4-DAG: [[FROM]]
 993 // CK4-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
 994 // CK4-DAG: br label %[[TYEND]]
 995 // CK4-DAG: [[TYEND]]
 996 // CK4-DAG: [[PHITYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
 997 // CK4: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 [[CUSIZE]], i64 [[PHITYPE0]], {{.*}})
 998 // 281474976710659 == 0x1,000,000,003
 999 // CK4-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]]
1000 // CK4-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
1001 // CK4-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
1002 // CK4-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
1003 // CK4-DAG: [[ALLOC]]
1004 // CK4-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
1005 // CK4-DAG: br label %[[TYEND:[^,]+]]
1006 // CK4-DAG: [[ALLOCELSE]]
1007 // CK4-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
1008 // CK4-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
1009 // CK4-DAG: [[TO]]
1010 // CK4-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
1011 // CK4-DAG: br label %[[TYEND]]
1012 // CK4-DAG: [[TOELSE]]
1013 // CK4-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
1014 // CK4-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
1015 // CK4-DAG: [[FROM]]
1016 // CK4-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
1017 // CK4-DAG: br label %[[TYEND]]
1018 // CK4-DAG: [[TYEND]]
1019 // CK4-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
1020 // CK4: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 4, i64 [[TYPE1]], {{.*}})
1021 // 281474976710675 == 0x1,000,000,013
1022 // CK4-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710675, [[SHIPRESIZE]]
1023 // CK4-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3
1024 // CK4-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0
1025 // CK4-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]
1026 // CK4-DAG: [[ALLOC]]
1027 // CK4-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4
1028 // CK4-DAG: br label %[[TYEND:[^,]+]]
1029 // CK4-DAG: [[ALLOCELSE]]
1030 // CK4-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1
1031 // CK4-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]
1032 // CK4-DAG: [[TO]]
1033 // CK4-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3
1034 // CK4-DAG: br label %[[TYEND]]
1035 // CK4-DAG: [[TOELSE]]
1036 // CK4-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2
1037 // CK4-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]
1038 // CK4-DAG: [[FROM]]
1039 // CK4-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2
1040 // CK4-DAG: br label %[[TYEND]]
1041 // CK4-DAG: [[TYEND]]
1042 // CK4-DAG: [[TYPE2:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]
1043 // CK4: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BBEGIN]], ptr [[BARRBEGINGEP]], i64 16, i64 [[TYPE2]], {{.*}})
1044 // CK4: [[PTRNEXT]] = getelementptr %class.C, ptr [[PTR]], i32 1
1045 // CK4: [[ISDONE:%.+]] = icmp eq ptr [[PTRNEXT]], [[PTREND]]
1046 // CK4: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]
1047
1048 // CK4: [[LEXIT]]
1049 // CK4: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1
1050 // CK4: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8
1051 // CK4: [[ISNOTDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0
1052 // CK4: [[CMP1:%.+]] = and i1 [[ISARRAY]], [[ISNOTDEL]]
1053 // CK4: br i1 [[CMP1]], label %[[EVALDEL:[^,]+]], label %[[DONE]]
1054 // CK4: [[EVALDEL]]
1055 // CK4-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16
1056 // CK4-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8
1057
1058 // Remove movement mappings and mark as implicit
1059 // CK4-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4
1060 // CK4-DAG: [[DTYPE1:%.+]] = or i64 [[DTYPE]], 512
1061 // CK4: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BPTR]], ptr [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE1]], {{.*}})
1062 // CK4: br label %[[DONE]]
1063 // CK4: [[DONE]]
1064 // CK4: ret void
1065
1066
1067 // CK4-LABEL: define {{.*}}void @{{.*}}foo{{.*}}
1068 void foo(int a){
1069   int i = a;
1070   C c;
1071   c.a = a;
1072
1073   // CK4-DAG: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[TGEPBP:%.+]], ptr [[TGEPP:%.+]], ptr [[TSIZES]], ptr [[TTYPES]], ptr null, ptr [[TMPR:%.+]])
1074   // CK4-DAG: [[TGEPBP]] = getelementptr inbounds {{.+}}[[TBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
1075   // CK4-DAG: [[TGEPP]] = getelementptr inbounds {{.+}}[[TP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
1076   // CK4-DAG: [[TBP0:%.+]] = getelementptr inbounds {{.+}}[[TBP]], i{{.+}} 0, i{{.+}} 0
1077   // CK4-DAG: [[TP0:%.+]] = getelementptr inbounds {{.+}}[[TP]], i{{.+}} 0, i{{.+}} 0
1078   // CK4-DAG: [[TMPR1:%.+]] = getelementptr inbounds {{.+}}[[TMPR]], i[[sz]] 0, i[[sz]] 0
1079   // CK4-DAG: store ptr [[VAL:%[^,]+]], ptr [[TBP0]]
1080   // CK4-DAG: store ptr [[VAL]], ptr [[TP0]]
1081   // CK4-DAG: store ptr [[MPRFUNC]], ptr [[TMPR1]]
1082   #pragma omp target update to(present, mapper(id): c)
1083
1084   // CK4-DAG: call void @__tgt_target_data_update_mapper(ptr @{{.+}}, i64 -1, i32 1, ptr [[FGEPBP:%.+]], ptr [[FGEPP:%.+]], ptr [[FSIZES]], ptr [[FTYPES]], ptr null, ptr [[FMPR:%.+]])
1085   // CK4-DAG: [[FGEPBP]] = getelementptr inbounds {{.+}}[[FBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
1086   // CK4-DAG: [[FGEPP]] = getelementptr inbounds {{.+}}[[FP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
1087   // CK4-DAG: [[FBP0:%.+]] = getelementptr inbounds {{.+}}[[FBP]], i{{.+}} 0, i{{.+}} 0
1088   // CK4-DAG: [[FP0:%.+]] = getelementptr inbounds {{.+}}[[FP]], i{{.+}} 0, i{{.+}} 0
1089   // CK4-DAG: [[FMPR1:%.+]] = getelementptr inbounds {{.+}}[[FMPR]], i[[sz]] 0, i[[sz]] 0
1090   // CK4-DAG: store ptr [[VAL]], ptr [[FBP0]]
1091   // CK4-DAG: store ptr [[VAL]], ptr [[FP0]]
1092   // CK4-DAG: store ptr [[MPRFUNC]], ptr [[FMPR1]]
1093   #pragma omp target update from(mapper(id), present: c)
1094 }
1095
1096 #endif // CK4
1097
1098 #endif // HEADER