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