1161 lines · cpp
1// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}2 3// expected-no-diagnostics4#ifndef HEADER5#define HEADER6 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 %s9// 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 %s10// 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 %s11// 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 %s12// 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 %s13// 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 %s14 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 %s16// 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 %s17// 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 %s18// 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 %s19// 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 %s20// 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 %s21 22#ifdef CK023// 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 047// 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 81class C {82public: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]], 1691// CK0-32-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 892// CK0-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 193// 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]], 1696// CK0-DAG: [[ISPTRANDOBJ:%.+]] = icmp ne i64 [[PTRANDOBJ]], 097// CK0-DAG: [[CMPA:%.+]] = and i1 [[PTRSNE]], [[ISPTRANDOBJ]]98// CK0-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[CMPA]]99// CK0-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8100// CK0-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0101// 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]], 16105// CK0-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8106 107// Remove movement mappings and mark as implicit108// CK0-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4109// CK0-DAG: [[ITYPE1:%.+]] = or i64 [[ITYPE]], 512110// 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 0119// CK0-DAG: [[BBEGIN:%.+]] = getelementptr inbounds nuw %class.C, ptr [[PTR]], i32 0, i32 1120// CK0-DAG: [[BBEGIN2:%.+]] = getelementptr inbounds nuw %class.C, ptr [[PTR]], i32 0, i32 1121// CK0-DAG: [[BARRBEGIN:%.+]] = load ptr, ptr [[BBEGIN2]]122// CK0-DAG: [[BARRBEGINGEP:%.+]] = getelementptr inbounds nuw double, ptr [[BARRBEGIN]], i[[sz:64|32]] 0123// CK0-DAG: [[BEND:%.+]] = getelementptr ptr, ptr [[BBEGIN]], i32 1124// CK0-DAG: [[ABEGINI:%.+]] = ptrtoint ptr [[ABEGIN]] to i64125// CK0-DAG: [[BENDI:%.+]] = ptrtoint ptr [[BEND]] to i64126// 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]], 48130// CK0-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 0, [[SHIPRESIZE]]131// CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3132// CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0133// CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]134// CK0-DAG: [[ALLOC]]135// CK0-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4136// CK0-DAG: br label %[[TYEND:[^,]+]]137// CK0-DAG: [[ALLOCELSE]]138// CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1139// CK0-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]140// CK0-DAG: [[TO]]141// CK0-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3142// CK0-DAG: br label %[[TYEND]]143// CK0-DAG: [[TOELSE]]144// CK0-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2145// CK0-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]146// CK0-DAG: [[FROM]]147// CK0-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2148// 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,003153// CK0-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]]154// CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3155// CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0156// CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]157// CK0-DAG: [[ALLOC]]158// CK0-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4159// CK0-DAG: br label %[[TYEND:[^,]+]]160// CK0-DAG: [[ALLOCELSE]]161// CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1162// CK0-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]163// CK0-DAG: [[TO]]164// CK0-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3165// CK0-DAG: br label %[[TYEND]]166// CK0-DAG: [[TOELSE]]167// CK0-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2168// CK0-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]169// CK0-DAG: [[FROM]]170// CK0-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2171// 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,013176// CK0-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710675, [[SHIPRESIZE]]177// CK0-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3178// CK0-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0179// CK0-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]180// CK0-DAG: [[ALLOC]]181// CK0-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4182// CK0-DAG: br label %[[TYEND:[^,]+]]183// CK0-DAG: [[ALLOCELSE]]184// CK0-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1185// CK0-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]186// CK0-DAG: [[TO]]187// CK0-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3188// CK0-DAG: br label %[[TYEND]]189// CK0-DAG: [[TOELSE]]190// CK0-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2191// CK0-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]192// CK0-DAG: [[FROM]]193// CK0-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2194// 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 1199// 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]], 1204// CK0: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8205// CK0: [[ISNOTDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0206// 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]], 16210// CK0-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8211 212// Remove movement mappings and mark as implicit213// CK0-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4214// CK0-DAG: [[DTYPE1:%.+]] = or i64 [[DTYPE]], 512215// 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 void219 220 221// CK0-LABEL: define {{.*}}void @{{.*}}foo{{.*}}222void 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 2229// CK0-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]230// CK0-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3231// CK0-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]232// CK0-DAG: [[MARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 7233// CK0-DAG: store ptr [[MPR:%.+]], ptr [[MARG]]234// CK0-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0235// CK0-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0236// CK0-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0237// CK0-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0238// CK0-DAG: [[MPR1:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i[[sz:32|64]] 0, i[[sz]] 0239// 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 0250 // CK0: store ptr [[CADDR:%[^,]+]], ptr [[BP2GEP]], align251 // CK0: [[P2GEP:%.+]] = getelementptr inbounds [1 x ptr], ptr [[OFFLOAD_P2:%[^,]+]], i32 0, i32 0252 // CK0: store ptr [[CADDR]], ptr [[P2GEP]], align253 // CK0: [[MAPPER2GEP:%.+]] = getelementptr inbounds [1 x ptr], ptr [[OFFLOAD_MAPPER2:%[^,]+]], i[[SZ:32|64]] 0, i[[SZ]] 0254 // CK0: store ptr [[MPRFUNC]], ptr [[MAPPER2GEP]], align255 // CK0: [[BP2:%.+]] = getelementptr inbounds [1 x ptr], ptr [[OFFLOAD_BP2]], i32 0, i32 0256 // CK0: [[P2:%.+]] = getelementptr inbounds [1 x ptr], ptr [[OFFLOAD_P2]], i32 0, i32 0257 // 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 1260 // CK0: {{.+}} = call i32 @__kmpc_omp_task(ptr @1, i32 {{.+}}, ptr [[TASK]])261 #pragma omp target map(mapper(id),tofrom: c) nowait262 {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 2268// CK0-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]269// CK0-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3270// CK0-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]271// CK0-DAG: [[MARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 7272// CK0-DAG: store ptr [[MPRGEP:%.+]], ptr [[MARG]]273// CK0-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0274// CK0-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0275// CK0-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0276// CK0-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0277// CK0-DAG: [[MPR1:%.+]] = getelementptr inbounds {{.+}}[[MPRGEP]], i[[sz]] 0, i[[sz]] 0278// 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 0291 // CK0: {{.+}} = getelementptr inbounds nuw [[KMP_TASK_T]], ptr [[TASK_CAST_GET_1]], i32 0, i32 0292 // CK0: {{.+}} = call i32 @__kmpc_omp_task(ptr @1, i32 {{.+}}, ptr [[TASK_1]])293 #pragma omp target teams map(mapper(id),to: c) nowait294 {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 0300 // CK0-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0301 // CK0-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0302 // CK0-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0303 // CK0-DAG: [[MPR1:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i[[sz]] 0, i[[sz]] 0304 // 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 1312 // CK0-32-DAG: [[FPBPGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T_5]], ptr [[PRIVATES]], i32 0, i32 1313 // CK0-64-DAG: [[FPBPGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T_5]], ptr [[PRIVATES]], i32 0, i32 0314 // 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 0316 // CK0-DAG: [[BPGEP:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP]], i32 0, i32 0317 // CK0-DAG: store ptr [[C:%[^,]+]], ptr [[BPGEP]], align318 // CK0-32-DAG: [[FPPGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T_5]], ptr [[PRIVATES]], i32 0, i32 2319 // CK0-64-DAG: [[FPPGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T_5]], ptr [[PRIVATES]], i32 0, i32 1320 // 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 0322 // CK0-DAG: [[PGEP:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P]], i32 0, i32 0323 // CK0-DAG: store ptr [[C]], ptr [[PGEP]], align324 // CK0-32-DAG: [[FPSZGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T_5]], ptr [[PRIVATES]], i32 0, i32 0325 // CK0-64-DAG: [[FPSZGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T_5]], ptr [[PRIVATES]], i32 0, i32 2326 // 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 3328 // 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]] 0330 // CK0-DAG: store ptr [[MPRFUNC]], ptr [[MPRGEP]], align331 #pragma omp target enter data map(mapper(id),to: c) nowait332 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 0335 // CK0-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0336 // CK0-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0337 // CK0-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0338 // CK0-DAG: [[MPR1:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i[[sz]] 0, i[[sz]] 0339 // 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 1347 // CK0-32-DAG: [[FPBPGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T_8]], ptr [[PRIVATES]], i32 0, i32 1348 // CK0-64-DAG: [[FPBPGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T_8]], ptr [[PRIVATES]], i32 0, i32 0349 // 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 0351 // CK0-DAG: [[BPGEP:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP]], i32 0, i32 0352 // CK0-DAG: store ptr [[C:%[^,]+]], ptr [[BPGEP]], align353 // CK0-32-DAG: [[FPPGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T_8]], ptr [[PRIVATES]], i32 0, i32 2354 // CK0-64-DAG: [[FPPGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T_8]], ptr [[PRIVATES]], i32 0, i32 1355 // 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 0357 // CK0-DAG: [[PGEP:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P]], i32 0, i32 0358 // CK0-DAG: store ptr [[C]], ptr [[PGEP]], align359 // CK0-32-DAG: [[FPSZGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T_8]], ptr [[PRIVATES]], i32 0, i32 0360 // CK0-64-DAG: [[FPSZGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T_8]], ptr [[PRIVATES]], i32 0, i32 2361 // 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 3363 // 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]] 0365 // CK0-DAG: store ptr [[MPRFUNC]], ptr [[MPRGEP]], align366 #pragma omp target exit data map(mapper(id),from: c) nowait367 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{{.+}} 0370 // CK0-DAG: [[TGEPP]] = getelementptr inbounds {{.+}}[[TP:%[^,]+]], i{{.+}} 0, i{{.+}} 0371 // CK0-DAG: [[TBP0:%.+]] = getelementptr inbounds {{.+}}[[TBP]], i{{.+}} 0, i{{.+}} 0372 // CK0-DAG: [[TP0:%.+]] = getelementptr inbounds {{.+}}[[TP]], i{{.+}} 0, i{{.+}} 0373 // CK0-DAG: [[TMPR1:%.+]] = getelementptr inbounds {{.+}}[[TMPR]], i[[sz]] 0, i[[sz]] 0374 // 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{{.+}} 0381 // CK0-DAG: [[FGEPP]] = getelementptr inbounds {{.+}}[[FP:%[^,]+]], i{{.+}} 0, i{{.+}} 0382 // CK0-DAG: [[FBP0:%.+]] = getelementptr inbounds {{.+}}[[FBP]], i{{.+}} 0, i{{.+}} 0383 // CK0-DAG: [[FP0:%.+]] = getelementptr inbounds {{.+}}[[FP]], i{{.+}} 0, i{{.+}} 0384 // CK0-DAG: [[FMPR1:%.+]] = getelementptr inbounds {{.+}}[[FMPR]], i[[sz]] 0, i[[sz]] 0385 // 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 1393 // CK0-32-DAG: [[FPBPGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T_11]], ptr [[PRIVATES]], i32 0, i32 1394 // CK0-64-DAG: [[FPBPGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T_11]], ptr [[PRIVATES]], i32 0, i32 0395 // 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 0397 // CK0-DAG: [[BPGEP:%.+]] = getelementptr inbounds [1 x ptr], ptr [[BP]], i32 0, i32 0398 // CK0-DAG: store ptr [[C:%[^,]+]], ptr [[BPGEP]], align399 // CK0-32-DAG: [[FPPGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T_11]], ptr [[PRIVATES]], i32 0, i32 2400 // CK0-64-DAG: [[FPPGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T_11]], ptr [[PRIVATES]], i32 0, i32 1401 // 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 0403 // CK0-DAG: [[PGEP:%.+]] = getelementptr inbounds [1 x ptr], ptr [[P]], i32 0, i32 0404 // CK0-DAG: store ptr [[C]], ptr [[PGEP]], align405 // CK0-32-DAG: [[FPSZGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T_11]], ptr [[PRIVATES]], i32 0, i32 0406 // CK0-64-DAG: [[FPSZGEP:%.+]] = getelementptr inbounds nuw [[KMP_PRIVATES_T_11]], ptr [[PRIVATES]], i32 0, i32 2407 // 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 3409 // 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]] 0411 // CK0-DAG: store ptr [[MPRFUNC]], ptr [[MPRGEP]], align412 #pragma omp target update from(mapper(id): c) nowait413}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 0421// CK0: [[VAL:%[^,]+]] = load i32, ptr [[CAADDR]]422// CK0: {{.+}} = add nsw i32 [[VAL]], 1423// 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 0430// CK0: [[VAL:%[^,]+]] = load i32, ptr [[CAADDR]]431// CK0: {{.+}} = add nsw i32 [[VAL]], 1432// 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 2437// CK0-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]438// CK0-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3439// CK0-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]440// CK0-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4441// CK0-DAG: store ptr [[SIZEGEP:%.+]], ptr [[SARG]]442// CK0-DAG: [[MARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 7443// CK0-DAG: store ptr [[MPRGEP:%.+]], ptr [[MARG]]444// CK0-DAG: [[BPGEP]] = getelementptr inbounds [1 x ptr], ptr [[BPFPADDR:%.+]], i[[SZ]] 0, i[[SZ]] 0445// CK0-DAG: [[PGEP]] = getelementptr inbounds [1 x ptr], ptr [[PFPADDR:%.+]], i[[SZ]] 0, i[[SZ]] 0446// CK0-DAG: [[SIZEGEP]] = getelementptr inbounds [1 x i64], ptr [[SIZEFPADDR:%.+]], i[[SZ]] 0, i[[SZ]] 0447// CK0-DAG: [[MPRGEP]] = getelementptr inbounds [1 x ptr], ptr [[MPRFPADDR:%.+]], i[[SZ]] 0, i[[SZ]] 0448// CK0-DAG: [[BPFPADDR]] = load ptr, ptr [[FPPTRADDR_BP:%.+]], align449// CK0-DAG: [[PFPADDR]] = load ptr, ptr [[FPPTRADDR_P:%.+]], align450// CK0-DAG: [[SIZEFPADDR]] = load ptr, ptr [[FPPTRADDR_SIZE:%.+]], align451// CK0-DAG: [[MPRFPADDR]] = load ptr, ptr [[FPPTRADDR_MPR:%.+]], align452// 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:%.+]], align455// CK0-DAG: [[KERNELARGGEP]] = getelementptr inbounds nuw [[ANON_T]], ptr [[CTX:%.+]], i32 0, i32 0456// CK0-DAG: [[CTX]] = load ptr, ptr [[CTXADDR:%.+]], align457// CK0-DAG: store ptr [[CTXARG]], ptr [[CTXADDR]], align458// CK0: }459 460// CK0: define internal {{.*}}i32 [[TASK_ENTRY]](i32 {{.*}}%0, ptr noalias noundef %1)461// CK0: store ptr %1, ptr [[ADDR:%.+]], align462// CK0: [[TASK_T_WITH_PRIVATES:%.+]] = load ptr, ptr [[ADDR]], align463// CK0: [[TASKGEP:%.+]] = getelementptr inbounds nuw [[KMP_TASK_T_WITH_PRIVATES]], ptr [[TASK_T_WITH_PRIVATES]], i32 0, i32 0464// CK0: [[SHAREDSGEP:%.+]] = getelementptr inbounds nuw [[KMP_TASK_T]], ptr [[TASKGEP]], i32 0, i32 0465// CK0: [[SHAREDS:%.+]] = load ptr, ptr [[SHAREDSGEP]], align466// CK0: [[PRIVATESGEP:%.+]] = getelementptr inbounds nuw [[KMP_TASK_T_WITH_PRIVATES]], ptr [[TASK_T_WITH_PRIVATES]], i32 0, i32 1467// 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 2473// CK0-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]474// CK0-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3475// CK0-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]476// CK0-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4477// CK0-DAG: store ptr [[SIZEGEP:%.+]], ptr [[SARG]]478// CK0-DAG: [[MARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 7479// CK0-DAG: store ptr [[MPRGEP:%.+]], ptr [[MARG]]480// CK0-DAG: [[BPGEP]] = getelementptr inbounds [1 x ptr], ptr [[BPFPADDR:%.+]], i[[SZ]] 0, i[[SZ]] 0481// CK0-DAG: [[PGEP]] = getelementptr inbounds [1 x ptr], ptr [[PFPADDR:%.+]], i[[SZ]] 0, i[[SZ]] 0482// CK0-DAG: [[SIZEGEP]] = getelementptr inbounds [1 x i64], ptr [[SIZEFPADDR:%.+]], i[[SZ]] 0, i[[SZ]] 0483// CK0-DAG: [[MPRGEP]] = getelementptr inbounds [1 x ptr], ptr [[MPRFPADDR:%.+]], i[[SZ]] 0, i[[SZ]] 0484// CK0-DAG: [[BPFPADDR]] = load ptr, ptr [[FPPTRADDR_BP:%.+]], align485// CK0-DAG: [[PFPADDR]] = load ptr, ptr [[FPPTRADDR_P:%.+]], align486// CK0-DAG: [[SIZEFPADDR]] = load ptr, ptr [[FPPTRADDR_SIZE:%.+]], align487// CK0-DAG: [[MPRFPADDR]] = load ptr, ptr [[FPPTRADDR_MPR:%.+]], align488// 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:%.+]], align491// CK0-DAG: [[KERNELARGGEP]] = getelementptr inbounds nuw [[ANON_T_0]], ptr [[CTX:%.+]], i32 0, i32 0492// CK0-DAG: [[CTX]] = load ptr, ptr [[CTXADDR:%.+]], align493// CK0-DAG: store ptr [[CTXARG]], ptr [[CTXADDR]], align494// CK0: }495 496// CK0: define internal {{.*}}i32 [[TASK_ENTRY_1]](i32 {{.*}}%0, ptr noalias noundef %1)497// CK0: store ptr %1, ptr [[ADDR:%.+]], align498// CK0: [[TASK_T_WITH_PRIVATES:%.+]] = load ptr, ptr [[ADDR]], align499// CK0: [[TASKGEP:%.+]] = getelementptr inbounds nuw [[KMP_TASK_T_WITH_PRIVATES_1]], ptr [[TASK_T_WITH_PRIVATES]], i32 0, i32 0500// CK0: [[SHAREDSGEP:%.+]] = getelementptr inbounds nuw [[KMP_TASK_T]], ptr [[TASKGEP]], i32 0, i32 0501// CK0: [[SHAREDS:%.+]] = load ptr, ptr [[SHAREDSGEP]], align502// CK0: [[PRIVATESGEP:%.+]] = getelementptr inbounds nuw [[KMP_TASK_T_WITH_PRIVATES_1]], ptr [[TASK_T_WITH_PRIVATES]], i32 0, i32 1503// 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]] 0509// CK0-DAG: [[P]] = getelementptr inbounds [1 x ptr], ptr [[PADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0510// CK0-DAG: [[SZ]] = getelementptr inbounds [1 x i64], ptr [[SZADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0511// CK0-DAG: [[MPR]] = getelementptr inbounds [1 x ptr], ptr [[MPRADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0512// CK0-DAG: [[BPADDR]] = load ptr, ptr [[FPBPADDR:%[^,]+]], align513// CK0-DAG: [[PADDR]] = load ptr, ptr [[FPPADDR:%[^,]+]], align514// CK0-DAG: [[SZADDR]] = load ptr, ptr [[FPSZADDR:%[^,]+]], align515// CK0-DAG: [[MPRADDR]] = load ptr, ptr [[FPMPRADDR:%[^,]+]], align516// CK0-DAG: call void %{{.+}}(ptr %{{[^,]+}}, ptr [[FPBPADDR]], ptr [[FPPADDR]], ptr [[FPSZADDR]], ptr [[FPMPRADDR]])517// CK0: ret void518// 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 0523// 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]] 0528// CK0-DAG: [[P]] = getelementptr inbounds [1 x ptr], ptr [[PADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0529// CK0-DAG: [[SZ]] = getelementptr inbounds [1 x i64], ptr [[SZADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0530// CK0-DAG: [[MPR]] = getelementptr inbounds [1 x ptr], ptr [[MPRADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0531// CK0-DAG: [[BPADDR]] = load ptr, ptr [[FPBPADDR:%[^,]+]], align532// CK0-DAG: [[PADDR]] = load ptr, ptr [[FPPADDR:%[^,]+]], align533// CK0-DAG: [[SZADDR]] = load ptr, ptr [[FPSZADDR:%[^,]+]], align534// CK0-DAG: [[MPRADDR]] = load ptr, ptr [[FPMPRADDR:%[^,]+]], align535// 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 0541// 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]] 0546// CK0-DAG: [[P]] = getelementptr inbounds [1 x ptr], ptr [[PADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0547// CK0-DAG: [[SZ]] = getelementptr inbounds [1 x i64], ptr [[SZADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0548// CK0-DAG: [[MPR]] = getelementptr inbounds [1 x ptr], ptr [[MPRADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0549// CK0-DAG: [[BPADDR]] = load ptr, ptr [[FPBPADDR:%[^,]+]], align550// CK0-DAG: [[PADDR]] = load ptr, ptr [[FPPADDR:%[^,]+]], align551// CK0-DAG: [[SZADDR]] = load ptr, ptr [[FPSZADDR:%[^,]+]], align552// CK0-DAG: [[MPRADDR]] = load ptr, ptr [[FPMPRADDR:%[^,]+]], align553// 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 0559// CK0: }560 561#endif // CK0562 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 %s566// 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 %s567// 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 %s568// 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 %s569// 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 %s570// 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 %s571 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 %s573// 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 %s574// 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 %s575// 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 %s576// 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 %s577// 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 %s578 579#ifdef CK1580// C++ template581 582template <class T>583class C {584public: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]], 4592// CK1-DAG: [[PTREND:%.+]] = getelementptr %class.C, ptr [[BEGIN]], i64 [[SIZE]]593// CK1-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1594// CK1-DAG: [[PTRSNE:%.+]] = icmp ne ptr [[BPTR]], [[BEGIN]]595// CK1-DAG: [[PTRANDOBJ:%.+]] = and i64 [[TYPE]], 16596// CK1-DAG: [[ISPTRANDOBJ:%.+]] = icmp ne i64 [[PTRANDOBJ]], 0597// CK1-DAG: [[CMPA:%.+]] = and i1 [[PTRSNE]], [[ISPTRANDOBJ]]598// CK1-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[CMPA]]599// CK1-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8600// CK1-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0601// 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]], 4606 607// Remove movement mappings and mark as implicit608// CK1-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4609// CK1-DAG: [[ITYPE1:%.+]] = or i64 [[ITYPE]], 512610// 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 0619// CK1-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(ptr [[HANDLE]])620// CK1-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48621// CK1-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 3, [[SHIPRESIZE]]622// CK1-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3623// CK1-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0624// CK1-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]625// CK1-DAG: [[ALLOC]]626// CK1-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4627// CK1-DAG: br label %[[TYEND:[^,]+]]628// CK1-DAG: [[ALLOCELSE]]629// CK1-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1630// CK1-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]631// CK1-DAG: [[TO]]632// CK1-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3633// CK1-DAG: br label %[[TYEND]]634// CK1-DAG: [[TOELSE]]635// CK1-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2636// CK1-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]637// CK1-DAG: [[FROM]]638// CK1-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2639// 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 1644// 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]], 1649// CK1: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8650// CK1: [[ISNOTDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0651// CK1: [[CMP1:%.+]] = and i1 [[ISARRAY]], [[ISNOTDEL]]652// CK1-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 4653 654// Remove movement mappings and mark as implicit655// CK1-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4656// CK1-DAG: [[DTYPE1:%.+]] = or i64 [[DTYPE]], 512657// 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 void661 662#endif // CK1663 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 %s667// 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 %s668// 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 %s669// 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 %s670// 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 %s671// 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 %s672 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 %s674// 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 %s675// 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 %s676// 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 %s677// 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 %s678// 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 %s679 680#ifdef CK2681// Nested mappers.682 683class B {684public:685 double a;686};687 688class C {689public: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]], 16702// CK2-DAG: [[PTREND:%.+]] = getelementptr %class.C, ptr [[BEGIN]], i64 [[SIZE]]703// CK2-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1704// CK2-DAG: [[PTRSNE:%.+]] = icmp ne ptr [[BPTR]], [[BEGIN]]705// CK2-DAG: [[PTRANDOBJ:%.+]] = and i64 [[TYPE]], 16706// CK2-DAG: [[ISPTRANDOBJ:%.+]] = icmp ne i64 [[PTRANDOBJ]], 0707// CK2-DAG: [[CMPA:%.+]] = and i1 [[PTRSNE]], [[ISPTRANDOBJ]]708// CK2-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[CMPA]]709// CK2-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8710// CK2-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0711// 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]], 16716 717// Remove movement mappings and mark as implicit718// CK2-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4719// CK2-DAG: [[ITYPE1:%.+]] = or i64 [[ITYPE]], 512720// 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 1729// CK2-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(ptr [[HANDLE]])730// CK2-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48731// CK2-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 3, [[SHIPRESIZE]]732// CK2-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3733// CK2-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0734// CK2-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]735// CK2-DAG: [[ALLOC]]736// CK2-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4737// CK2-DAG: br label %[[TYEND:[^,]+]]738// CK2-DAG: [[ALLOCELSE]]739// CK2-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1740// CK2-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]741// CK2-DAG: [[TO]]742// CK2-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3743// CK2-DAG: br label %[[TYEND]]744// CK2-DAG: [[TOELSE]]745// CK2-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2746// CK2-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]747// CK2-DAG: [[FROM]]748// CK2-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2749// 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 1754// 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]], 1759// CK2: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8760// CK2: [[ISNOTDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0761// 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]], 16765 766// Remove movement mappings and mark as implicit767// CK2-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4768// CK2-DAG: [[DTYPE1:%.+]] = or i64 [[DTYPE]], 512769// 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 void773 774#endif // CK2775 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 %s779// 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 %s780// 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 %s781// 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 %s782// 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 %s783// 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 %s784 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 %s786// 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 %s787// 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 %s788// 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 %s789// 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 %s790// 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 %s791 792#ifdef CK3793// map of array sections and nested components.794 795// CK3-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}.region_id = weak constant i8 0796// 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 799class C {800public:801 int a;802 double *b;803};804 805class B {806public: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{{.*}}815void foo(int a){816 // CK3-DAG: [[CVAL:%.+]] = alloca [10 x %class.C]817 // CK3-DAG: [[BVAL:%.+]] = alloca %class.B818 C c[10];819 B b;820 821 // CK3-DAG: [[BC:%.+]] = getelementptr inbounds nuw %class.B, ptr [[BVAL]], i32 0, i32 0822 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 2825// CK3-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]826// CK3-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3827// CK3-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]828// CK3-DAG: [[MARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 7829// CK3-DAG: store ptr [[MPR:%.+]], ptr [[MARG]]830// CK3-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0831// CK3-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0832// CK3-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0833// CK3-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0834// CK3-DAG: [[MPR1:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i{{64|32}} 0, i{{64|32}} 0835// 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 1839// CK3-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1840// CK3-DAG: [[MPR2:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i{{64|32}} 0, i{{64|32}} 1841// CK3-DAG: store ptr [[CVAL]], ptr [[BP2]]842// CK3-DAG: [[CVALGEP:%.+]] = getelementptr inbounds {{.+}}[[CVAL]], i{{64|32}} 0, i{{64|32}} 0843// 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 // CK3857 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 %s860// 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 %s861// 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 %s862// 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 %s863// 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 %s864// 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 %s865 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 %s867// 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 %s868// 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 %s869// 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 %s870// 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 %s871// 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 %s872 873#ifdef CK4874// 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 = 0x1001879// 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 = 0x1002884// CK4: [[FTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1002]]]885 886class C {887public: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]], 16896// CK4-32-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 8897// CK4-DAG: [[PTREND:%.+]] = getelementptr %class.C, ptr [[BEGIN]], i64 [[SIZE]]898// CK4-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 1899// CK4-DAG: [[PTRSNE:%.+]] = icmp ne ptr [[BPTR]], [[BEGIN]]900// CK4-DAG: [[PTRANDOBJ:%.+]] = and i64 [[TYPE]], 16901// CK4-DAG: [[ISPTRANDOBJ:%.+]] = icmp ne i64 [[PTRANDOBJ]], 0902// CK4-DAG: [[CMPA:%.+]] = and i1 [[PTRSNE]], [[ISPTRANDOBJ]]903// CK4-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[CMPA]]904// CK4-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8905// CK4-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0906// 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]], 16911// CK4-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8912 913// Remove movement mappings and mark as implicit914// CK4-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4915// CK4-DAG: [[ITYPE1:%.+]] = or i64 [[ITYPE]], 512916// 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 0925// CK4-DAG: [[BBEGIN:%.+]] = getelementptr inbounds nuw %class.C, ptr [[PTR]], i32 0, i32 1926// CK4-DAG: [[BBEGIN2:%.+]] = getelementptr inbounds nuw %class.C, ptr [[PTR]], i32 0, i32 1927// CK4-DAG: [[BARRBEGIN:%.+]] = load ptr, ptr [[BBEGIN2]]928// CK4-DAG: [[BARRBEGINGEP:%.+]] = getelementptr inbounds nuw double, ptr [[BARRBEGIN]], i[[sz:64|32]] 0929// CK4-DAG: [[BEND:%.+]] = getelementptr ptr, ptr [[BBEGIN]], i32 1930// CK4-DAG: [[ABEGINI:%.+]] = ptrtoint ptr [[ABEGIN]] to i64931// CK4-DAG: [[BENDI:%.+]] = ptrtoint ptr [[BEND]] to i64932// 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]], 48936// CK4-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 0, [[SHIPRESIZE]]937// CK4-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3938// CK4-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0939// CK4-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]940// CK4-DAG: [[ALLOC]]941// CK4-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4942// CK4-DAG: br label %[[TYEND:[^,]+]]943// CK4-DAG: [[ALLOCELSE]]944// CK4-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1945// CK4-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]946// CK4-DAG: [[TO]]947// CK4-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3948// CK4-DAG: br label %[[TYEND]]949// CK4-DAG: [[TOELSE]]950// CK4-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2951// CK4-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]952// CK4-DAG: [[FROM]]953// CK4-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2954// 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,003959// CK4-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]]960// CK4-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3961// CK4-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0962// CK4-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]963// CK4-DAG: [[ALLOC]]964// CK4-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4965// CK4-DAG: br label %[[TYEND:[^,]+]]966// CK4-DAG: [[ALLOCELSE]]967// CK4-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1968// CK4-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]969// CK4-DAG: [[TO]]970// CK4-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3971// CK4-DAG: br label %[[TYEND]]972// CK4-DAG: [[TOELSE]]973// CK4-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2974// CK4-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]975// CK4-DAG: [[FROM]]976// CK4-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2977// 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,013982// CK4-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 281474976710675, [[SHIPRESIZE]]983// CK4-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3984// CK4-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0985// CK4-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]986// CK4-DAG: [[ALLOC]]987// CK4-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4988// CK4-DAG: br label %[[TYEND:[^,]+]]989// CK4-DAG: [[ALLOCELSE]]990// CK4-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1991// CK4-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]992// CK4-DAG: [[TO]]993// CK4-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3994// CK4-DAG: br label %[[TYEND]]995// CK4-DAG: [[TOELSE]]996// CK4-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2997// CK4-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]998// CK4-DAG: [[FROM]]999// CK4-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -21000// 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 11005// 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]], 11010// CK4: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 81011// CK4: [[ISNOTDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 01012// 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]], 161016// CK4-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 81017 1018// Remove movement mappings and mark as implicit1019// CK4-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -41020// CK4-DAG: [[DTYPE1:%.+]] = or i64 [[DTYPE]], 5121021// 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 void1025 1026 1027// CK4-LABEL: define {{.*}}void @{{.*}}foo{{.*}}1028void 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{{.+}} 01035 // CK4-DAG: [[TGEPP]] = getelementptr inbounds {{.+}}[[TP:%[^,]+]], i{{.+}} 0, i{{.+}} 01036 // CK4-DAG: [[TBP0:%.+]] = getelementptr inbounds {{.+}}[[TBP]], i{{.+}} 0, i{{.+}} 01037 // CK4-DAG: [[TP0:%.+]] = getelementptr inbounds {{.+}}[[TP]], i{{.+}} 0, i{{.+}} 01038 // CK4-DAG: [[TMPR1:%.+]] = getelementptr inbounds {{.+}}[[TMPR]], i[[sz]] 0, i[[sz]] 01039 // 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{{.+}} 01046 // CK4-DAG: [[FGEPP]] = getelementptr inbounds {{.+}}[[FP:%[^,]+]], i{{.+}} 0, i{{.+}} 01047 // CK4-DAG: [[FBP0:%.+]] = getelementptr inbounds {{.+}}[[FBP]], i{{.+}} 0, i{{.+}} 01048 // CK4-DAG: [[FP0:%.+]] = getelementptr inbounds {{.+}}[[FP]], i{{.+}} 0, i{{.+}} 01049 // CK4-DAG: [[FMPR1:%.+]] = getelementptr inbounds {{.+}}[[FMPR]], i[[sz]] 0, i[[sz]] 01050 // 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 // CK41057 1058///==========================================================================///1059// RUN: %clang_cc1 -DCK5 -verify -fopenmp -fopenmp-version=52 -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 CK5 %s1060// RUN: %clang_cc1 -DCK5 -fopenmp -fopenmp-version=52 -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 %s1061// RUN: %clang_cc1 -DCK5 -verify -fopenmp-version=52 -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 CK5 %s1062// RUN: %clang_cc1 -DCK5 -fopenmp -fopenmp-version=52 -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 %s1063 1064// RUN: %clang_cc1 -DCK5 -verify -fopenmp-simd -fopenmp-version=52 -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 %s1065// RUN: %clang_cc1 -DCK5 -fopenmp-simd -fopenmp-version=52 -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 %s1066// RUN: %clang_cc1 -DCK5 -verify -fopenmp-simd -fopenmp-version=52 -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 %s1067// RUN: %clang_cc1 -DCK5 -fopenmp-simd -fopenmp-version=52 -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 %s1068 1069#ifdef CK51070typedef struct myvec {1071 int a;1072 double *b;1073} myvec_t;1074 1075#pragma omp declare mapper(id: myvec_t v) map(iterator(it=0:v.a), tofrom: v.b[it])1076// CK5: @[[ITER:[a-zA-Z0-9_]+]] = global i32 0, align 41077 1078void foo(){ 1079 myvec_t s;1080 #pragma omp target map(mapper(id), to:s)1081 {1082 }1083}1084 1085// CK5: define {{.*}}void [[MPRFUNC:@[.]omp_mapper[.].*myvec[.]id]](ptr noundef [[HANDLE:%.+]], ptr noundef [[BPTR:%.+]], ptr noundef [[BEGIN:%.+]], i64 noundef [[BYTESIZE:%.+]], i64 noundef [[TYPE:%.+]], ptr{{.*}})1086// CK5-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], {{.*}}1087// CK5-DAG: [[PTREND:%.+]] = getelementptr %struct.myvec, ptr [[BEGIN]], i64 [[SIZE]]1088// CK5-DAG: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 11089// CK5-DAG: [[PTRSNE:%.+]] = icmp ne ptr [[BPTR]], [[BEGIN]]1090// CK5-DAG: [[PTRANDOBJ:%.+]] = and i64 [[TYPE]], 161091// CK5-DAG: [[ISPTRANDOBJ:%.+]] = icmp ne i64 [[PTRANDOBJ]], 01092// CK5-DAG: [[CMPA:%.+]] = and i1 [[PTRSNE]], [[ISPTRANDOBJ]]1093// CK5-DAG: [[CMP:%.+]] = or i1 [[ISARRAY]], [[CMPA]]1094// CK5-DAG: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 81095// CK5-DAG: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 01096// CK5-DAG: [[CMP1:%.+]] = and i1 [[CMP]], [[ISNOTDEL]]1097// CK5: br i1 [[CMP1]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]]1098 1099// CK5: [[INITEVALDEL]]1100// CK5-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], {{.*}}1101 1102// Remove movement mappings and mark as implicit1103// CK5-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -41104// CK5-DAG: [[ITYPE1:%.+]] = or i64 [[ITYPE]], 5121105// CK5: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BPTR]], ptr [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE1]], {{.*}})1106// CK5: br label %[[LHEAD:[^,]+]]1107 1108// CK5: [[LHEAD]]1109// CK5: [[ISEMPTY:%.+]] = icmp eq ptr [[BEGIN]], [[PTREND]]1110// CK5: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]]1111// CK5: [[LBODY]]1112// CK5: [[PTR:%.+]] = phi ptr [ [[BEGIN]], %{{.+}} ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ]1113// CK5-DAG: [[ABEGIN:%.+]] = getelementptr inbounds nuw %struct.myvec, ptr [[PTR]], i32 0, i32 11114// CK5-DAG: load i32, ptr @[[ITER]], align 41115// CK5-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(ptr [[HANDLE]])1116// CK5-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 481117// CK5-DAG: [[MEMBERTYPE:%.+]] = add nuw i64 0, [[SHIPRESIZE]]1118// CK5-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 31119// CK5-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 01120// CK5-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]]1121// CK5-DAG: [[ALLOC]]1122// CK5-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -41123// CK5-DAG: br label %[[TYEND:[^,]+]]1124// CK5-DAG: [[ALLOCELSE]]1125// CK5-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 11126// CK5-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]]1127// CK5-DAG: [[TO]]1128// CK5-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -31129// CK5-DAG: br label %[[TYEND]]1130// CK5-DAG: [[TOELSE]]1131// CK5-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 21132// CK5-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]]1133// CK5-DAG: [[FROM]]1134// CK5-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -21135// CK5-DAG: br label %[[TYEND]]1136// CK5-DAG: [[TYEND]]1137// CK5-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ]1138// CK5: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[PTR]], ptr [[ABEGIN]], i64 {{.*}}, i64 [[TYPE1]], {{.*}})1139// CK5: [[PTRNEXT]] = getelementptr %struct.myvec, ptr [[PTR]], i32 11140// CK5: [[ISDONE:%.+]] = icmp eq ptr [[PTRNEXT]], [[PTREND]]1141// CK5: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]]1142 1143// CK5: [[LEXIT]]1144// CK5: [[ISARRAY:%.+]] = icmp sgt i64 [[SIZE]], 11145// CK5: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 81146// CK5: [[ISNOTDEL:%.+]] = icmp ne i64 [[TYPEDEL]], {{.*}}1147// CK5: [[CMP1:%.+]] = and i1 [[ISARRAY]], [[ISNOTDEL]]1148// CK5-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], {{.*}}1149 1150// Remove movement mappings and mark as implicit1151// CK5-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -41152// CK5-DAG: [[DTYPE1:%.+]] = or i64 [[DTYPE]], 5121153// CK5: call void @__tgt_push_mapper_component(ptr [[HANDLE]], ptr [[BPTR]], ptr [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE1]], {{.*}})1154// CK5: br label %[[DONE]]1155// CK5: [[DONE]]1156// CK5: ret void1157 1158#endif // CK51159 1160#endif // HEADER1161