diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index ce8846140d41eb..c20094f5bc9933 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -7126,6 +7126,7 @@ class MappableExprsHandler { std::pair HighestElem = { 0, Address::invalid()}; Address Base = Address::invalid(); + bool IsArraySection = false; }; private: @@ -7785,7 +7786,8 @@ class MappableExprsHandler { break; } llvm::Value *Size = getExprTypeSize(I->getAssociatedExpression()); - if (!IsMemberPointerOrAddr) { + if (!IsMemberPointerOrAddr || + (Next == CE && MapType != OMPC_MAP_unknown)) { CombinedInfo.BasePointers.push_back(BP.getPointer()); CombinedInfo.Pointers.push_back(LB.getPointer()); CombinedInfo.Sizes.push_back( @@ -7853,6 +7855,10 @@ class MappableExprsHandler { } } + // Need to emit combined struct for array sections. + if (IsFinalArraySection || IsNonContiguous) + PartialStruct.IsArraySection = true; + // If we have a final array section, we are done with this expression. if (IsFinalArraySection) break; @@ -8188,6 +8194,10 @@ class MappableExprsHandler { MapFlagsArrayTy &CurTypes, const StructRangeInfoTy &PartialStruct, bool NotTargetParams = false) const { + if (CurTypes.size() == 1 && + ((CurTypes.back() & OMP_MAP_MEMBER_OF) != OMP_MAP_MEMBER_OF) && + !PartialStruct.IsArraySection) + return; // Base is the base of the struct CombinedInfo.BasePointers.push_back(PartialStruct.Base.getPointer()); // Pointer is the address of the lowest element @@ -9938,7 +9948,7 @@ void CGOpenMPRuntime::emitTargetCall( MappedVarSet.insert(CI->getCapturedVar()); else MappedVarSet.insert(nullptr); - if (CurInfo.BasePointers.empty()) + if (CurInfo.BasePointers.empty() && !PartialStruct.Base.isValid()) MEHandler.generateDefaultMapInfo(*CI, **RI, *CV, CurInfo); // Generate correct mapping for variables captured by reference in // lambdas. @@ -9947,7 +9957,7 @@ void CGOpenMPRuntime::emitTargetCall( CurInfo, LambdaPointers); } // We expect to have at least an element of information for this capture. - assert(!CurInfo.BasePointers.empty() && + assert((!CurInfo.BasePointers.empty() || PartialStruct.Base.isValid()) && "Non-existing map pointer for capture!"); assert(CurInfo.BasePointers.size() == CurInfo.Pointers.size() && CurInfo.BasePointers.size() == CurInfo.Sizes.size() && diff --git a/clang/test/OpenMP/declare_mapper_codegen.cpp b/clang/test/OpenMP/declare_mapper_codegen.cpp index 115179dbabfbad..95a8e86a0e7377 100644 --- a/clang/test/OpenMP/declare_mapper_codegen.cpp +++ b/clang/test/OpenMP/declare_mapper_codegen.cpp @@ -686,57 +686,18 @@ class C { // CK1: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ] // CK1: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]] // CK1-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 0 -// CK1-DAG: [[AEND:%.+]] = getelementptr i32, i32* [[ABEGIN]], i32 1 -// CK1-DAG: [[ABEGINV:%.+]] = bitcast i32* [[ABEGIN]] to i8* -// CK1-DAG: [[AENDV:%.+]] = bitcast i32* [[AEND]] to i8* -// CK1-DAG: [[ABEGINI:%.+]] = ptrtoint i8* [[ABEGINV]] to i64 -// CK1-DAG: [[AENDI:%.+]] = ptrtoint i8* [[AENDV]] to i64 -// CK1-DAG: [[CSIZE:%.+]] = sub i64 [[AENDI]], [[ABEGINI]] -// CK1-DAG: [[CUSIZE:%.+]] = sdiv exact i64 [[CSIZE]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) -// CK1-DAG: [[BPTRADDR0BC:%.+]] = bitcast %class.C* [[OBJ]] to i8* -// CK1-DAG: [[PTRADDR0BC:%.+]] = bitcast i32* [[ABEGIN]] to i8* // CK1-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(i8* [[HANDLE]]) // CK1-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48 -// CK1-DAG: br label %[[MEMBER:[^,]+]] -// CK1-DAG: [[MEMBER]] -// CK1-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]] -// CK1-DAG: [[MEMBERCOM]] -// CK1-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 32, [[SHIPRESIZE]] -// CK1-DAG: br label %[[LTYPE]] -// CK1-DAG: [[LTYPE]] -// CK1-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 32, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ] -// CK1-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3 -// CK1-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0 -// CK1-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]] -// CK1-DAG: [[ALLOC]] -// CK1-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4 -// CK1-DAG: br label %[[TYEND:[^,]+]] -// CK1-DAG: [[ALLOCELSE]] -// CK1-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1 -// CK1-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]] -// CK1-DAG: [[TO]] -// CK1-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3 -// CK1-DAG: br label %[[TYEND]] -// CK1-DAG: [[TOELSE]] -// CK1-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2 -// CK1-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]] -// CK1-DAG: [[FROM]] -// CK1-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2 -// CK1-DAG: br label %[[TYEND]] -// CK1-DAG: [[TYEND]] -// CK1-DAG: [[TYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ] -// CK1-64: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR0BC]], i8* [[PTRADDR0BC]], i64 [[CUSIZE]], i64 [[TYPE0]]) // CK1-DAG: [[BPTRADDR1BC:%.+]] = bitcast %class.C* [[OBJ]] to i8* // CK1-DAG: [[PTRADDR1BC:%.+]] = bitcast i32* [[ABEGIN]] to i8* // CK1-DAG: br label %[[MEMBER:[^,]+]] // CK1-DAG: [[MEMBER]] -// CK1-DAG: br i1 false, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]] +// CK1-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]] // CK1-DAG: [[MEMBERCOM]] -// 281474976710659 == 0x1,000,000,003 -// CK1-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]] +// CK1-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 35, [[SHIPRESIZE]] // CK1-DAG: br label %[[LTYPE]] // CK1-DAG: [[LTYPE]] -// CK1-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 281474976710659, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ] +// CK1-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 35, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ] // CK1-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3 // CK1-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0 // CK1-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]] @@ -849,57 +810,18 @@ class C { // CK2: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ] // CK2: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]] // CK2-DAG: [[BBEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1 -// CK2-DAG: [[BEND:%.+]] = getelementptr %class.B, %class.B* [[BBEGIN]], i32 1 -// CK2-DAG: [[BBEGINV:%.+]] = bitcast %class.B* [[BBEGIN]] to i8* -// CK2-DAG: [[BENDV:%.+]] = bitcast %class.B* [[BEND]] to i8* -// CK2-DAG: [[BBEGINI:%.+]] = ptrtoint i8* [[BBEGINV]] to i64 -// CK2-DAG: [[BENDI:%.+]] = ptrtoint i8* [[BENDV]] to i64 -// CK2-DAG: [[BSIZE:%.+]] = sub i64 [[BENDI]], [[BBEGINI]] -// CK2-DAG: [[BUSIZE:%.+]] = sdiv exact i64 [[BSIZE]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) -// CK2-DAG: [[BPTRADDR0BC:%.+]] = bitcast %class.C* [[OBJ]] to i8* -// CK2-DAG: [[PTRADDR0BC:%.+]] = bitcast %class.B* [[BBEGIN]] to i8* // CK2-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(i8* [[HANDLE]]) // CK2-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48 -// CK2-DAG: br label %[[MEMBER:[^,]+]] -// CK2-DAG: [[MEMBER]] -// CK2-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]] -// CK2-DAG: [[MEMBERCOM]] -// CK2-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 32, [[SHIPRESIZE]] -// CK2-DAG: br label %[[LTYPE]] -// CK2-DAG: [[LTYPE]] -// CK2-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 32, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ] -// CK2-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3 -// CK2-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0 -// CK2-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]] -// CK2-DAG: [[ALLOC]] -// CK2-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4 -// CK2-DAG: br label %[[TYEND:[^,]+]] -// CK2-DAG: [[ALLOCELSE]] -// CK2-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1 -// CK2-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]] -// CK2-DAG: [[TO]] -// CK2-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3 -// CK2-DAG: br label %[[TYEND]] -// CK2-DAG: [[TOELSE]] -// CK2-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2 -// CK2-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]] -// CK2-DAG: [[FROM]] -// CK2-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2 -// CK2-DAG: br label %[[TYEND]] -// CK2-DAG: [[TYEND]] -// CK2-DAG: [[TYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ] -// CK2-64: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR0BC]], i8* [[PTRADDR0BC]], i64 [[BUSIZE]], i64 [[TYPE0]]) // CK2-DAG: [[BPTRADDR1BC:%.+]] = bitcast %class.C* [[OBJ]] to i8* // CK2-DAG: [[PTRADDR1BC:%.+]] = bitcast %class.B* [[BBEGIN]] to i8* // CK2-DAG: br label %[[MEMBER:[^,]+]] // CK2-DAG: [[MEMBER]] -// CK2-DAG: br i1 false, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]] +// CK2-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]] // CK2-DAG: [[MEMBERCOM]] -// 281474976710659 == 0x1,000,000,003 -// CK2-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]] +// CK2-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 35, [[SHIPRESIZE]] // CK2-DAG: br label %[[LTYPE]] // CK2-DAG: [[LTYPE]] -// CK2-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 281474976710659, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ] +// CK2-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 35, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ] // CK2-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3 // CK2-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0 // CK2-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]] @@ -962,7 +884,8 @@ class C { // map of array sections and nested components. // CK3-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}.region_id = weak constant i8 0 -// CK3: [[TYPES:@.+]] = {{.+}}constant [3 x i64] [i64 32, i64 281474976710659, i64 35] +// CK3-DAG: [[SIZES:@.+]] = {{.+}}constant [2 x i64] [i64 {{8|16}}, i64 {{80|160}}] +// CK3-DAG: [[TYPES:@.+]] = {{.+}}constant [2 x i64] [i64 35, i64 35] class C { public: @@ -987,52 +910,28 @@ void foo(int a){ B b; // CK3-DAG: [[BC:%.+]] = getelementptr inbounds %class.B, %class.B* [[BVAL]], i32 0, i32 0 - // CK3-DAG: [[BCEND:%.+]] = getelementptr %class.C, %class.C* [[BC]], i32 1 - // CK3-DAG: [[BCC:%.+]] = bitcast %class.C* [[BC]] to i8* - // CK3-DAG: [[BCENDC:%.+]] = bitcast %class.C* [[BCEND]] to i8* - // CK3-DAG: [[BCI:%.+]] = ptrtoint i8* [[BCC]] to i64 - // CK3-DAG: [[BCENDI:%.+]] = ptrtoint i8* [[BCENDC]] to i64 - // CK3-DAG: [[BSIZE:%.+]] = sub i64 [[BCENDI]], [[BCI]] - // CK3-DAG: [[BSIZED:%.+]] = sdiv exact i64 [[BSIZE]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) - - // CK3-DAG: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 3, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], i64* [[SGEP:%[^,]+]], {{.+}}[[TYPES]]{{.+}}, i8** [[MPRGEP:%.+]]) + + // CK3-DAG: call i32 @__tgt_target_mapper(i64 {{.+}}, i8* {{.+}}, i32 2, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}, i8** [[MPRGEP:%.+]]) // CK3-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 // CK3-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 - // CK3-DAG: [[SGEP]] = getelementptr inbounds {{.+}}[[SIZES:%[^,]+]], i32 0, i32 0 - // CK3-DAG: [[MPRGEP]] = bitcast [3 x i8*]* [[MPR:%[^,]+]] to i8** + // CK3-DAG: [[MPRGEP]] = bitcast [2 x i8*]* [[MPR:%[^,]+]] to i8** // CK3-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 // CK3-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 - // CK3-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[SIZES]], i32 0, i32 0 // CK3-DAG: [[MPR1:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i{{64|32}} 0, i{{64|32}} 0 // CK3-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to %class.B** // CK3-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C** // CK3-DAG: store %class.B* [[BVAL]], %class.B** [[CBP1]] // CK3-DAG: store %class.C* [[BC]], %class.C** [[CP1]] - // CK3-DAG: store i64 [[BSIZED]], i64* [[S1]] - // CK3-DAG: store i8* null, i8** [[MPR1]] + // CK3-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR1]] // CK3-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1 // CK3-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1 - // CK3-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[SIZES]], i32 0, i32 1 // CK3-DAG: [[MPR2:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i{{64|32}} 0, i{{64|32}} 1 - // CK3-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to %class.B** + // CK3-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [10 x %class.C]** // CK3-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to %class.C** - // CK3-DAG: store %class.B* [[BVAL]], %class.B** [[CBP2]] - // CK3-DAG: store %class.C* [[BC]], %class.C** [[CP2]] - // CK3-64-DAG: store i64 16, i64* [[S2]] - // CK3-32-DAG: store i64 8, i64* [[S2]] - // CK3-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR2]] - // CK3-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 2 - // CK3-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 2 - // CK3-DAG: [[S3:%.+]] = getelementptr inbounds {{.+}}[[SIZES]], i32 0, i32 2 - // CK3-DAG: [[MPR3:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i{{64|32}} 0, i{{64|32}} 2 - // CK3-DAG: [[CBP3:%.+]] = bitcast i8** [[BP3]] to [10 x %class.C]** - // CK3-DAG: [[CP3:%.+]] = bitcast i8** [[P3]] to %class.C** - // CK3-DAG: store [10 x %class.C]* [[CVAL]], [10 x %class.C]** [[CBP3]] + // CK3-DAG: store [10 x %class.C]* [[CVAL]], [10 x %class.C]** [[CBP2]] // CK3-DAG: [[CVALGEP:%.+]] = getelementptr inbounds {{.+}}[[CVAL]], i{{64|32}} 0, i{{64|32}} 0 - // CK3-DAG: store %class.C* [[CVALGEP]], %class.C** [[CP3]] - // CK3-64-DAG: store i64 160, i64* [[S3]] - // CK3-32-DAG: store i64 80, i64* [[S3]] - // CK3-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR3]] + // CK3-DAG: store %class.C* [[CVALGEP]], %class.C** [[CP2]] + // CK3-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR2]] // CK3: call void [[KERNEL:@.+]](%class.B* [[BVAL]], [10 x %class.C]* [[CVAL]]) #pragma omp target map(mapper(id),tofrom: c[0:10], b.c) for (int i = 0; i < 10; i++) { diff --git a/clang/test/OpenMP/target_codegen.cpp b/clang/test/OpenMP/target_codegen.cpp index 31f342057f7866..dff010c84b33ab 100644 --- a/clang/test/OpenMP/target_codegen.cpp +++ b/clang/test/OpenMP/target_codegen.cpp @@ -81,7 +81,7 @@ // CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i64] [i64 800, i64 800, i64 547] // CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [4 x i64] [i64 4, i64 2, i64 1, i64 40] // CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i64] [i64 800, i64 800, i64 800, i64 547] -// CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [6 x i64] [i64 32, i64 281474976711171, i64 800, i64 800, i64 800, i64 547] +// CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [5 x i64] [i64 547, i64 800, i64 800, i64 800, i64 547] // CHECK-DAG: [[SIZET9:@.+]] = private unnamed_addr constant [1 x i64] [i64 12] // CHECK-DAG: [[MAPT10:@.+]] = private unnamed_addr constant [1 x i64] [i64 35] // CHECK-DAG: @{{.*}} = weak constant i8 0 @@ -579,66 +579,57 @@ int bar(int n){ // CHECK-32: [[CSZSIZE:%.+]] = mul nuw i32 [[CELEMSIZE2]], 2 // CHECK-32: [[CSIZE:%.+]] = sext i32 [[CSZSIZE]] to i64 -// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_mapper(i64 -1, i8* @{{[^,]+}}, i32 6, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* [[SR:%[^,]+]], i64* getelementptr inbounds ([6 x i64], [6 x i64]* [[MAPT7]], i32 0, i32 0), i8** null) -// CHECK-DAG: [[BPR]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP:%.+]], i32 0, i32 0 -// CHECK-DAG: [[PR]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P:%.+]], i32 0, i32 0 -// CHECK-DAG: [[SR]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S:%.+]], i32 0, i32 0 -// CHECK-DAG: [[SADDR0:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 0, i32 [[IDX0:0]] -// CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 0, i32 [[IDX0]] -// CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 0, i32 [[IDX0]] -// CHECK-DAG: [[SADDR1:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 0, i32 [[IDX1:1]] -// CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 0, i32 [[IDX1]] -// CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 0, i32 [[IDX1]] -// CHECK-DAG: [[SADDR2:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 0, i32 [[IDX2:2]] -// CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 0, i32 [[IDX2]] -// CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 0, i32 [[IDX2]] -// CHECK-DAG: [[SADDR3:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 0, i32 [[IDX3:3]] -// CHECK-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 0, i32 [[IDX3]] -// CHECK-DAG: [[PADDR3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 0, i32 [[IDX3]] -// CHECK-DAG: [[SADDR4:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 0, i32 [[IDX4:4]] -// CHECK-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 0, i32 [[IDX4]] -// CHECK-DAG: [[PADDR4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 0, i32 [[IDX4]] -// CHECK-DAG: [[SADDR5:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 0, i32 [[IDX5:5]] -// CHECK-DAG: [[BPADDR5:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 0, i32 [[IDX5]] -// CHECK-DAG: [[PADDR5:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 0, i32 [[IDX5]] +// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_mapper(i64 -1, i8* @{{[^,]+}}, i32 5, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* [[SR:%[^,]+]], i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[MAPT7]], i32 0, i32 0), i8** null) +// CHECK-DAG: [[BPR]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP:%.+]], i32 0, i32 0 +// CHECK-DAG: [[PR]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P:%.+]], i32 0, i32 0 +// CHECK-DAG: [[SR]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S:%.+]], i32 0, i32 0 +// CHECK-DAG: [[SADDR0:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 0, i32 [[IDX0:0]] +// CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 0, i32 [[IDX0]] +// CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 0, i32 [[IDX0]] +// CHECK-DAG: [[SADDR1:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 0, i32 [[IDX1:1]] +// CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 0, i32 [[IDX1]] +// CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 0, i32 [[IDX1]] +// CHECK-DAG: [[SADDR2:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 0, i32 [[IDX2:2]] +// CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 0, i32 [[IDX2]] +// CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 0, i32 [[IDX2]] +// CHECK-DAG: [[SADDR3:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 0, i32 [[IDX3:3]] +// CHECK-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 0, i32 [[IDX3]] +// CHECK-DAG: [[PADDR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 0, i32 [[IDX3]] +// CHECK-DAG: [[SADDR4:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 0, i32 [[IDX4:4]] +// CHECK-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 0, i32 [[IDX4]] +// CHECK-DAG: [[PADDR4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 0, i32 [[IDX4]] // The names below are not necessarily consistent with the names used for the // addresses above as some are repeated. -// CHECK-DAG: [[CBPADDR5:%.+]] = bitcast i8** [[BPADDR5]] to i16** -// CHECK-DAG: [[CPADDR5:%.+]] = bitcast i8** [[PADDR5]] to i16** -// CHECK-DAG: store i16* %{{.+}}, i16** [[CBPADDR5]] -// CHECK-DAG: store i16* %{{.+}}, i16** [[CPADDR5]] -// CHECK-DAG: store i64 [[CSIZE]], i64* [[SADDR5]] - -// CHECK-DAG: [[CBPADDR4:%.+]] = bitcast i8** [[BPADDR4]] to i[[SZ]]* -// CHECK-DAG: [[CPADDR4:%.+]] = bitcast i8** [[PADDR4]] to i[[SZ]]* -// CHECK-DAG: store i[[SZ]] [[VLA0]], i[[SZ]]* [[CBPADDR4]] -// CHECK-DAG: store i[[SZ]] [[VLA0]], i[[SZ]]* [[CPADDR4]] -// CHECK-DAG: store i64 {{4|8}}, i64* [[SADDR4]] +// CHECK-DAG: [[CBPADDR4:%.+]] = bitcast i8** [[BPADDR4]] to i16** +// CHECK-DAG: [[CPADDR4:%.+]] = bitcast i8** [[PADDR4]] to i16** +// CHECK-DAG: store i16* %{{.+}}, i16** [[CBPADDR4]] +// CHECK-DAG: store i16* %{{.+}}, i16** [[CPADDR4]] +// CHECK-DAG: store i64 [[CSIZE]], i64* [[SADDR4]] // CHECK-DAG: [[CBPADDR3:%.+]] = bitcast i8** [[BPADDR3]] to i[[SZ]]* // CHECK-DAG: [[CPADDR3:%.+]] = bitcast i8** [[PADDR3]] to i[[SZ]]* -// CHECK-DAG: store i[[SZ]] 2, i[[SZ]]* [[CBPADDR3]] -// CHECK-DAG: store i[[SZ]] 2, i[[SZ]]* [[CPADDR3]] +// CHECK-DAG: store i[[SZ]] [[VLA0]], i[[SZ]]* [[CBPADDR3]] +// CHECK-DAG: store i[[SZ]] [[VLA0]], i[[SZ]]* [[CPADDR3]] // CHECK-DAG: store i64 {{4|8}}, i64* [[SADDR3]] // CHECK-DAG: [[CBPADDR2:%.+]] = bitcast i8** [[BPADDR2]] to i[[SZ]]* // CHECK-DAG: [[CPADDR2:%.+]] = bitcast i8** [[PADDR2]] to i[[SZ]]* -// CHECK-DAG: store i[[SZ]] [[B_CVAL]], i[[SZ]]* [[CBPADDR2]] -// CHECK-DAG: store i[[SZ]] [[B_CVAL]], i[[SZ]]* [[CPADDR2]] -// CHECK-DAG: store i64 4, i64* [[SADDR2]] +// CHECK-DAG: store i[[SZ]] 2, i[[SZ]]* [[CBPADDR2]] +// CHECK-DAG: store i[[SZ]] 2, i[[SZ]]* [[CPADDR2]] +// CHECK-DAG: store i64 {{4|8}}, i64* [[SADDR2]] + +// CHECK-DAG: [[CBPADDR1:%.+]] = bitcast i8** [[BPADDR1]] to i[[SZ]]* +// CHECK-DAG: [[CPADDR1:%.+]] = bitcast i8** [[PADDR1]] to i[[SZ]]* +// CHECK-DAG: store i[[SZ]] [[B_CVAL]], i[[SZ]]* [[CBPADDR1]] +// CHECK-DAG: store i[[SZ]] [[B_CVAL]], i[[SZ]]* [[CPADDR1]] +// CHECK-DAG: store i64 4, i64* [[SADDR1]] // CHECK-DAG: [[CBPADDR0:%.+]] = bitcast i8** [[BPADDR0]] to [[S1]]** // CHECK-DAG: [[CPADDR0:%.+]] = bitcast i8** [[PADDR0]] to double** // CHECK-DAG: store [[S1]]* [[THIS:%.+]], [[S1]]** [[CBPADDR0]] // CHECK-DAG: store double* [[A:%.+]], double** [[CPADDR0]] -// CHECK-DAG: store i64 %{{.+}}, i64* [[SADDR0]] - -// CHECK-DAG: [[CBPADDR1:%.+]] = bitcast i8** [[BPADDR1]] to [[S1]]** -// CHECK-DAG: [[CPADDR1:%.+]] = bitcast i8** [[PADDR1]] to double** -// CHECK-DAG: store [[S1]]* [[THIS]], [[S1]]** [[CBPADDR1]] -// CHECK-DAG: store double* [[A]], double** [[CPADDR1]] -// CHECK-DAG: store i64 8, i64* [[SADDR1]] +// CHECK-DAG: store i64 8, i64* [[SADDR0]] // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] diff --git a/clang/test/OpenMP/target_data_member_codegen.cpp b/clang/test/OpenMP/target_data_member_codegen.cpp index af21fb2863658c..21730cf8939f31 100644 --- a/clang/test/OpenMP/target_data_member_codegen.cpp +++ b/clang/test/OpenMP/target_data_member_codegen.cpp @@ -25,14 +25,15 @@ #ifndef HEADER #define HEADER -// CHECK: [[MAPTYPES:@.+]] = private unnamed_addr constant [1 x i64] zeroinitializer +// CHECK-DAG: [[SIZES:@.+]] = private unnamed_addr constant [1 x i64] [i64 {{8|4}}] +// CHECK-DAG: [[MAPTYPES:@.+]] = private unnamed_addr constant [1 x i64] [i64 3] struct S { double *p; }; // CHECK-LABEL: @main int main() { - // CHECK: call i32 @__tgt_target_mapper(i64 -1, i8* @.{{.+}}, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* %{{.+}}, i64* getelementptr inbounds ([1 x i64], [1 x i64]* [[MAPTYPES]], i32 0, i32 0), i8** null) + // CHECK: call i32 @__tgt_target_mapper(i64 -1, i8* @.{{.+}}, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* getelementptr inbounds ([1 x i64], [1 x i64]* [[SIZES]], i32 0, i32 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* [[MAPTYPES]], i32 0, i32 0), i8** null) S s; #pragma omp target map(s.p) { } diff --git a/clang/test/OpenMP/target_firstprivate_codegen.cpp b/clang/test/OpenMP/target_firstprivate_codegen.cpp index e4a1a0302411d6..f329b2d641f997 100644 --- a/clang/test/OpenMP/target_firstprivate_codegen.cpp +++ b/clang/test/OpenMP/target_firstprivate_codegen.cpp @@ -58,7 +58,7 @@ struct TT { // CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [9 x i64] [i64 288, i64 161, i64 800, i64 161, i64 161, i64 800, i64 800, i64 161, i64 161] // CHECK-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [2 x i{{32|64}}] [i{{32|64}} 0, i{{32|64}} 8] // CHECK-DAG: [[MAPT3:@.+]] = private unnamed_addr constant [2 x i64] [i64 32, i64 37] -// CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [6 x i64] [i64 32, i64 281474976711171, i64 288, i64 800, i64 800, i64 161] +// CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [5 x i64] [i64 547, i64 288, i64 800, i64 800, i64 161] // CHECK-DAG: [[SIZET5:@.+]] = private unnamed_addr constant [3 x i{{32|64}}] [i[[SZ]] 4, i[[SZ]] 1, i[[SZ]] 40] // CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i64] [i64 288, i64 288, i64 161] // CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [2 x i{{32|64}}] [i[[SZ]] 4, i[[SZ]] 40] @@ -438,69 +438,60 @@ struct S1 { // on the host side, we first generate r1, then the static function and the template above // CHECK: define{{.+}} i32 {{.+}}([[S1]]* {{.+}}, i{{[0-9]+}} {{.+}}) - // CHECK: [[BASE_PTRS4:%.+]] = alloca [6 x i8*], - // CHECK: [[PTRS4:%.+]] = alloca [6 x i8*], - // CHECK: [[SIZET4:%.+]] = alloca [6 x i{{[0-9]+}}], + // CHECK: [[BASE_PTRS4:%.+]] = alloca [5 x i8*], + // CHECK: [[PTRS4:%.+]] = alloca [5 x i8*], + // CHECK: [[SIZET4:%.+]] = alloca [5 x i{{[0-9]+}}], // map(this: this ptr is implicitly captured (not firstprivate matter) - // CHECK: [[BP0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASE_PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 + // CHECK: [[BP0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BASE_PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 // CHECK: [[CBP0:%.+]] = bitcast i8** [[BP0]] to %struct.S1** // CHECK: store %struct.S1* [[THIS:%.+]], %struct.S1** [[CBP0]], - // CHECK: [[P0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 + // CHECK: [[P0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 // CHECK: [[CP0:%.+]] = bitcast i8** [[P0]] to double** // CHECK: store double* [[A:%.+]], double** [[CP0]], - // CHECK: [[SZ0:%.+]] = getelementptr inbounds [6 x i{{[0-9]+}}], [6 x i{{[0-9]+}}]* [[SIZET4]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 - // CHECK: store i{{[0-9]+}} %{{.+}}, i{{[0-9]+}}* [[SZ0]], - - // CHECK: [[BP1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASE_PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 - // CHECK: [[CBP1:%.+]] = bitcast i8** [[BP1]] to %struct.S1** - // CHECK: store %struct.S1* [[THIS]], %struct.S1** [[CBP1]], - // CHECK: [[P1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 - // CHECK: [[CP1:%.+]] = bitcast i8** [[P1]] to double** - // CHECK: store double* [[A]], double** [[CP1]], - // CHECK: [[SZ1:%.+]] = getelementptr inbounds [6 x i{{[0-9]+}}], [6 x i{{[0-9]+}}]* [[SIZET4]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 - // CHECK: store i{{[0-9]+}} 8, i{{[0-9]+}}* [[SZ1]], + // CHECK: [[SZ0:%.+]] = getelementptr inbounds [5 x i{{[0-9]+}}], [5 x i{{[0-9]+}}]* [[SIZET4]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 + // CHECK: store i{{[0-9]+}} 8, i{{[0-9]+}}* [[SZ0]], // firstprivate(b): base_ptr = b, ptr = b, size = 4 (pass by-value) - // CHECK: [[BASE_PTRS_GEP4_1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASE_PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 + // CHECK: [[BASE_PTRS_GEP4_1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BASE_PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 // CHECK: [[BCAST_TOPTR:%.+]] = bitcast i8** [[BASE_PTRS_GEP4_1]] to i{{[0-9]+}}* // CHECK: store i{{[0-9]+}} [[B_CAST:%.+]], i{{[0-9]+}}* [[BCAST_TOPTR]], - // CHECK: [[PTRS_GEP4_1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 + // CHECK: [[PTRS_GEP4_1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 // CHECK: [[BCAST_TOPTR:%.+]] = bitcast i8** [[PTRS_GEP4_1]] to i{{[0-9]+}}* // CHECK: store i{{[0-9]+}} [[B_CAST]], i{{[0-9]+}}* [[BCAST_TOPTR]], - // CHECK: [[SIZES_GEP4_1:%.+]] = getelementptr inbounds [6 x i{{[0-9]+}}], [6 x i{{[0-9]+}}]* [[SIZET4]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 + // CHECK: [[SIZES_GEP4_1:%.+]] = getelementptr inbounds [5 x i{{[0-9]+}}], [5 x i{{[0-9]+}}]* [[SIZET4]], i{{[0-9]+}} 0, i{{[0-9]+}} 1 // CHECK: store i{{[0-9]+}} 4, i{{[0-9]+}}* [[SIZES_GEP4_1]], // firstprivate(c), 3 entries: 2, n, c - // CHECK: [[BASE_PTRS_GEP4_2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASE_PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 3 + // CHECK: [[BASE_PTRS_GEP4_2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BASE_PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 // CHECK: [[BCAST_TOPTR:%.+]] = bitcast i8** [[BASE_PTRS_GEP4_2]] to i{{[0-9]+}}* // CHECK: store i{{[0-9]+}} 2, i{{[0-9]+}}* [[BCAST_TOPTR]], - // CHECK: [[PTRS_GEP4_2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 3 + // CHECK: [[PTRS_GEP4_2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 // CHECK: [[BCAST_TOPTR:%.+]] = bitcast i8** [[PTRS_GEP4_2]] to i{{[0-9]+}}* // CHECK: store i{{[0-9]+}} 2, i{{[0-9]+}}* [[BCAST_TOPTR]], - // CHECK: [[SIZES_GEP4_2:%.+]] = getelementptr inbounds [6 x i{{[0-9]+}}], [6 x i{{[0-9]+}}]* [[SIZET4]], i{{[0-9]+}} 0, i{{[0-9]+}} 3 + // CHECK: [[SIZES_GEP4_2:%.+]] = getelementptr inbounds [5 x i{{[0-9]+}}], [5 x i{{[0-9]+}}]* [[SIZET4]], i{{[0-9]+}} 0, i{{[0-9]+}} 2 // CHECK-64: store i{{[0-9]+}} 8, i{{[0-9]+}}* [[SIZES_GEP4_2]], // CHECK-32: store i{{[0-9]+}} 4, i{{[0-9]+}}* [[SIZES_GEP4_2]], - // CHECK: [[BASE_PTRS_GEP4_3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASE_PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 4 + // CHECK: [[BASE_PTRS_GEP4_3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BASE_PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 3 // CHECK: [[BCAST_TOPTR:%.+]] = bitcast i8** [[BASE_PTRS_GEP4_3]] to i{{[0-9]+}}* // CHECK: store i{{[0-9]+}} [[N:%.+]], i{{[0-9]+}}* [[BCAST_TOPTR]], - // CHECK: [[PTRS_GEP4_3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 4 + // CHECK: [[PTRS_GEP4_3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 3 // CHECK: [[BCAST_TOPTR:%.+]] = bitcast i8** [[PTRS_GEP4_3]] to i{{[0-9]+}}* // CHECK: store i{{[0-9]+}} [[N]], i{{[0-9]+}}* [[BCAST_TOPTR]], - // CHECK: [[SIZES_GEP4_3:%.+]] = getelementptr inbounds [6 x i{{[0-9]+}}], [6 x i{{[0-9]+}}]* [[SIZET4]], i{{[0-9]+}} 0, i{{[0-9]+}} 4 + // CHECK: [[SIZES_GEP4_3:%.+]] = getelementptr inbounds [5 x i{{[0-9]+}}], [5 x i{{[0-9]+}}]* [[SIZET4]], i{{[0-9]+}} 0, i{{[0-9]+}} 3 // CHECK-64: store i{{[0-9]+}} 8, i{{[0-9]+}}* [[SIZES_GEP4_3]], // CHECK-32: store i{{[0-9]+}} 4, i{{[0-9]+}}* [[SIZES_GEP4_3]], - // CHECK: [[BASE_PTRS_GEP4_4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BASE_PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 5 + // CHECK: [[BASE_PTRS_GEP4_4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BASE_PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 4 // CHECK: [[BCAST_TOPTR:%.+]] = bitcast i8** [[BASE_PTRS_GEP4_4]] to i{{[0-9]+}}** // CHECK: store i{{[0-9]+}}* [[B:%.+]], i{{[0-9]+}}** [[BCAST_TOPTR]], - // CHECK: [[PTRS_GEP4_4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 5 + // CHECK: [[PTRS_GEP4_4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[PTRS4]], i{{[0-9]+}} 0, i{{[0-9]+}} 4 // CHECK: [[BCAST_TOPTR:%.+]] = bitcast i8** [[PTRS_GEP4_4]] to i{{[0-9]+}}** // CHECK: store i{{[0-9]+}}* [[B]], i{{[0-9]+}}** [[BCAST_TOPTR]], - // CHECK: [[SIZES_GEP4_4:%.+]] = getelementptr inbounds [6 x i{{[0-9]+}}], [6 x i{{[0-9]+}}]* [[SIZET4]], i{{[0-9]+}} 0, i{{[0-9]+}} 5 + // CHECK: [[SIZES_GEP4_4:%.+]] = getelementptr inbounds [5 x i{{[0-9]+}}], [5 x i{{[0-9]+}}]* [[SIZET4]], i{{[0-9]+}} 0, i{{[0-9]+}} 4 // CHECK: store i{{[0-9]+}} [[B_SIZE:%.+]], i{{[0-9]+}}* [[SIZES_GEP4_4]], // only check that we use the map types stored in the global variable - // CHECK: call i32 @__tgt_target_mapper(i64 -1, {{.+}}, i32 6, i8** {{.+}}, i8** {{.+}}, i{{[0-9]+}}* {{.+}}, i64* getelementptr inbounds ([6 x i64], [6 x i64]* [[MAPT4]], i32 0, i32 0), i8** null) + // CHECK: call i32 @__tgt_target_mapper(i64 -1, {{.+}}, i32 5, i8** {{.+}}, i8** {{.+}}, i{{[0-9]+}}* {{.+}}, i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[MAPT4]], i32 0, i32 0), i8** null) // TCHECK: define weak void @__omp_offloading_{{.+}}([[S1]]* [[TH:%.+]], i{{[0-9]+}} [[B_IN:%.+]], i{{[0-9]+}} [[VLA:%.+]], i{{[0-9]+}} [[VLA1:%.+]], i{{[0-9]+}}{{.+}} [[C_IN:%.+]]) // TCHECK: [[TH_ADDR:%.+]] = alloca [[S1]]*, diff --git a/clang/test/OpenMP/target_map_codegen_20.cpp b/clang/test/OpenMP/target_map_codegen_20.cpp index b3080e447ed8e3..639a06252a29d1 100644 --- a/clang/test/OpenMP/target_map_codegen_20.cpp +++ b/clang/test/OpenMP/target_map_codegen_20.cpp @@ -64,8 +64,9 @@ // CK21: [[ST:%.+]] = type { i32, i32, float* } // CK21-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 -// CK21-USE: [[MTYPE00:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710659] -// CK21-NOUSE: [[MTYPE00:@.+]] = private {{.*}}constant [2 x i64] [i64 0, i64 281474976710659] +// CK21: [[SIZE00:@.+]] = private {{.*}}constant [1 x i64] [i64 4] +// CK21-USE: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i64] [i64 35] +// CK21-NOUSE: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i64] [i64 3] // CK21-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK21: [[SIZE01:@.+]] = private {{.*}}constant [1 x i64] [i64 492] @@ -103,31 +104,18 @@ struct CC { T *lb; // Region 00 - // CK21-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) + // CK21-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]]{{.+}}, {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) // CK21-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK21-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] - // CK21-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] // CK21-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK21-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 - // CK21-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 // CK21-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]** // CK21-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** // CK21-DAG: store [[ST]]* [[VAR0:%.+]], [[ST]]** [[CBP0]] // CK21-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] - // CK21-DAG: store i64 {{%.+}}, i64* [[S0]] // CK21-DAG: [[SEC0]] = getelementptr {{.*}}[[ST]]* [[VAR0:%.+]], i{{.+}} 0, i{{.+}} 0 - // CK21-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 - // CK21-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 - // CK21-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 - // CK21-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]** - // CK21-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** - // CK21-DAG: store [[ST]]* [[VAR1:%.+]], [[ST]]** [[CBP1]] - // CK21-DAG: store i32* [[SEC1:%.+]], i32** [[CP1]] - // CK21-DAG: store i64 {{.+}}, i64* [[S1]] - // CK21-DAG: [[SEC1]] = getelementptr {{.*}}[[ST]]* [[VAR1:%.+]], i{{.+}} 0, i{{.+}} 0 - // CK21-USE: call void [[CALL00:@.+]]([[ST]]* {{[^,]+}}) // CK21-NOUSE: call void [[CALL00:@.+]]() #pragma omp target map(A) diff --git a/clang/test/OpenMP/target_map_codegen_23.cpp b/clang/test/OpenMP/target_map_codegen_23.cpp index 5bc469772266ad..d941aede15b7d5 100644 --- a/clang/test/OpenMP/target_map_codegen_23.cpp +++ b/clang/test/OpenMP/target_map_codegen_23.cpp @@ -57,16 +57,20 @@ struct SC{ }; // CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 -// CK24: [[MTYPE01:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710659] +// CK24: [[SIZE01:@.+]] = private {{.*}}constant [1 x i64] [i64 4] +// CK24: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 35] // CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 -// CK24: [[MTYPE13:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710659] +// CK24: [[SIZE13:@.+]] = private {{.*}}constant [1 x i64] [i64 4] +// CK24: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i64] [i64 35] // CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 -// CK24: [[MTYPE14:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710659] +// CK24: [[SIZE14:@.+]] = private {{.*}}constant [1 x i64] [i64 {{48|56}}] +// CK24: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i64] [i64 35] // CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 -// CK24: [[MTYPE15:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710659] +// CK24: [[SIZE15:@.+]] = private {{.*}}constant [1 x i64] [i64 4] +// CK24: [[MTYPE15:@.+]] = private {{.*}}constant [1 x i64] [i64 35] // CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK24: [[MTYPE16:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710659] @@ -75,7 +79,8 @@ struct SC{ // CK24: [[MTYPE17:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710675] // CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 -// CK24: [[MTYPE18:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710659] +// CK24: [[SIZE18:@.+]] = private {{.*}}constant [1 x i64] [i64 4] +// CK24: [[MTYPE18:@.+]] = private {{.*}}constant [1 x i64] [i64 35] // CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_struct_fields{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK24: [[MTYPE19:@.+]] = private {{.*}}constant [3 x i64] [i64 32, i64 281474976710659, i64 281474976710675] @@ -101,30 +106,18 @@ int explicit_maps_struct_fields(int a){ SC *p; // Region 01 -// CK24-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE01]]{{.+}}, i8** null) +// CK24-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE01]]{{.+}}, {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}}, i8** null) // CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] -// CK24-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] // CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 -// CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 // CK24-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SC]]** // CK24-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** // CK24-DAG: store [[SC]]* [[VAR0:%.+]], [[SC]]** [[CBP0]] // CK24-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] -// CK24-DAG: store i64 {{%.+}}, i64* [[S0]] // CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 0 -// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SC]]** -// CK24-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** -// CK24-DAG: store [[SC]]* [[VAR0]], [[SC]]** [[CBP1]] -// CK24-DAG: store i32* [[SEC0]], i32** [[CP1]] -// CK24-DAG: store i64 {{.+}}, i64* [[S1]] - // CK24: call void [[CALL01:@.+]]([[SC]]* {{[^,]+}}) #pragma omp target map(s.a) { s.a++; } @@ -133,30 +126,18 @@ int explicit_maps_struct_fields(int a){ // Same thing but starting from a pointer. // // Region 13 -// CK24-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE13]]{{.+}}, i8** null) +// CK24-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE13]]{{.+}}, {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE13]]{{.+}}, i8** null) // CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] -// CK24-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] // CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 -// CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 // CK24-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SC]]** // CK24-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** // CK24-DAG: store [[SC]]* [[VAR0:%.+]], [[SC]]** [[CBP0]] // CK24-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] -// CK24-DAG: store i64 {{%.+}}, i64* [[S0]] // CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 0 -// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SC]]** -// CK24-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** -// CK24-DAG: store [[SC]]* [[VAR0]], [[SC]]** [[CBP1]] -// CK24-DAG: store i32* [[SEC0]], i32** [[CP1]] -// CK24-DAG: store i64 {{.+}}, i64* [[S1]] - // CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} // CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} @@ -165,31 +146,19 @@ int explicit_maps_struct_fields(int a){ { p->a++; } // Region 14 -// CK24-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE14]]{{.+}}, i8** null) +// CK24-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE14]]{{.+}}, {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE14]]{{.+}}, i8** null) // CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] -// CK24-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] // CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 -// CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 // CK24-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SC]]** // CK24-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SA]]** // CK24-DAG: store [[SC]]* [[VAR0:%.+]], [[SC]]** [[CBP0]] // CK24-DAG: store [[SA]]* [[SEC0:%.+]], [[SA]]** [[CP0]] -// CK24-DAG: store i64 {{%.+}}, i64* [[S0]] // CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SB]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 1 // CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SC]]** -// CK24-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to [[SA]]** -// CK24-DAG: store [[SC]]* [[VAR0]], [[SC]]** [[CBP1]] -// CK24-DAG: store [[SA]]* [[SEC0]], [[SA]]** [[CP1]] -// CK24-DAG: store i64 {{.+}}, i64* [[S1]] - // CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} // CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} @@ -198,32 +167,20 @@ int explicit_maps_struct_fields(int a){ { p->a++; } // Region 15 -// CK24-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:64|32]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE15]]{{.+}}, i8** null) +// CK24-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE15]]{{.+}}, {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE15]]{{.+}}, i8** null) // CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] -// CK24-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] // CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 -// CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 // CK24-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SC]]** // CK24-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** // CK24-DAG: store [[SC]]* [[VAR0:%.+]], [[SC]]** [[CBP0]] // CK24-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] -// CK24-DAG: store i64 {{%.+}}, i64* [[S0]] // CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SA]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0 // CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SB]]* [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 1 // CK24-DAG: [[SEC000]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SC]]** -// CK24-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** -// CK24-DAG: store [[SC]]* [[VAR0]], [[SC]]** [[CBP1]] -// CK24-DAG: store i32* [[SEC0]], i32** [[CP1]] -// CK24-DAG: store i64 {{.+}}, i64* [[S1]] - // CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} // CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} @@ -301,33 +258,21 @@ int explicit_maps_struct_fields(int a){ { p->a++; } // Region 18 -// CK24-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE18]]{{.+}}, i8** null) +// CK24-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE18]]{{.+}}, {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE18]]{{.+}}, i8** null) // CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] -// CK24-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] // CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 -// CK24-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 // CK24-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SC]]** // CK24-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** // CK24-DAG: store [[SC]]* [[VAR0:%.+]], [[SC]]** [[CBP0]] // CK24-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] -// CK24-DAG: store i64 {{%.+}}, i64* [[S0]] // CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SA]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0 // CK24-DAG: [[SEC00]] = getelementptr {{.*}}[10 x [[SA]]]* [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 3 // CK24-DAG: [[SEC000]] = getelementptr {{.*}}[[SB]]* [[SEC0000:%[^,]+]], i{{.+}} 0, i{{.+}} 2 // CK24-DAG: [[SEC0000]] = getelementptr {{.*}}[[SC]]* [[VAR00:%.+]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SC]]** -// CK24-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** -// CK24-DAG: store [[SC]]* [[VAR0]], [[SC]]** [[CBP1]] -// CK24-DAG: store i32* [[SEC0]], i32** [[CP1]] -// CK24-DAG: store i64 {{.+}}, i64* [[S1]] - // CK24-DAG: [[VAR0]] = load [[SC]]*, [[SC]]** %{{.+}} // CK24-DAG: [[VAR00]] = load [[SC]]*, [[SC]]** %{{.+}} diff --git a/clang/test/OpenMP/target_map_codegen_24.cpp b/clang/test/OpenMP/target_map_codegen_24.cpp index 1f22eb31de53a9..15cfc7d8ca2b98 100644 --- a/clang/test/OpenMP/target_map_codegen_24.cpp +++ b/clang/test/OpenMP/target_map_codegen_24.cpp @@ -37,7 +37,8 @@ // CK25: [[CA01:%.+]] = type { i32* } // CK25-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 -// CK25: [[MTYPE00:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710657] +// CK25: [[SIZE00:@.+]] = private {{.*}}constant [1 x i64] [i64 4] +// CK25: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i64] [i64 33] // CK25-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK25: [[SIZE01:@.+]] = private {{.*}}constant [1 x i64] [i64 4] @@ -52,30 +53,18 @@ struct CC { int foo(T arg) { // Region 00 - // CK25-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) + // CK25-DAG: call i32 @__tgt_target_mapper(i64 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]]{{.+}}, {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) // CK25-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK25-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] - // CK25-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] // CK25-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK25-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 - // CK25-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 // CK25-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]** // CK25-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** // CK25-DAG: store [[ST]]* [[VAR0:%.+]], [[ST]]** [[CBP0]] // CK25-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] - // CK25-DAG: store i64 {{%.+}}, i64* [[S0]] // CK25-DAG: [[SEC0]] = getelementptr {{.*}}[[ST]]* [[VAR0:%.+]], i{{.+}} 0, i{{.+}} 0 - // CK25-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 - // CK25-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 - // CK25-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 - // CK25-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]** - // CK25-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** - // CK25-DAG: store [[ST]]* [[VAR0]], [[ST]]** [[CBP1]] - // CK25-DAG: store i32* [[SEC0]], i32** [[CP1]] - // CK25-DAG: store i64 {{.+}}, i64* [[S1]] - // CK25: call void [[CALL00:@.+]]([[ST]]* {{[^,]+}}) #pragma omp target map(to:A) { diff --git a/clang/test/OpenMP/target_parallel_codegen.cpp b/clang/test/OpenMP/target_parallel_codegen.cpp index 0d7e86caae6d7e..9ccd46703b69c1 100644 --- a/clang/test/OpenMP/target_parallel_codegen.cpp +++ b/clang/test/OpenMP/target_parallel_codegen.cpp @@ -99,7 +99,7 @@ // CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i64] [i64 800, i64 800, i64 547] // CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [4 x i64] [i64 4, i64 2, i64 1, i64 40] // CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i64] [i64 800, i64 800, i64 800, i64 547] -// CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [6 x i64] [i64 32, i64 281474976711171, i64 800, i64 800, i64 800, i64 547] +// CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [5 x i64] [i64 547, i64 800, i64 800, i64 800, i64 547] // CHECK-DAG: @{{.*}} = weak constant i8 0 // CHECK-DAG: @{{.*}} = weak constant i8 0 // CHECK-DAG: @{{.*}} = weak constant i8 0 @@ -583,25 +583,25 @@ int bar(int n){ // CHECK-32: [[CSZSIZE:%.+]] = mul nuw i32 [[CELEMSIZE2]], 2 // CHECK-32: [[CSIZE:%.+]] = sext i32 [[CSZSIZE]] to i64 -// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 6, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* [[SR:%[^,]+]], i64* getelementptr inbounds ([6 x i64], [6 x i64]* [[MAPT7]], i32 0, i32 0), i8** null, i32 1, i32 0) -// CHECK-DAG: [[BPR]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP:%.+]], i32 0, i32 0 -// CHECK-DAG: [[PR]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P:%.+]], i32 0, i32 0 -// CHECK-DAG: [[SR]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S:%.+]], i32 0, i32 0 -// CHECK-DAG: [[SADDR0:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX0:[0-9]+]] -// CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX0]] -// CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX0]] -// CHECK-DAG: [[SADDR1:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX1:[0-9]+]] -// CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX1]] -// CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX1]] -// CHECK-DAG: [[SADDR2:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX2:[0-9]+]] -// CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX2]] -// CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX2]] -// CHECK-DAG: [[SADDR3:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX3:[0-9]+]] -// CHECK-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX3]] -// CHECK-DAG: [[PADDR3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX3]] -// CHECK-DAG: [[SADDR4:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX4:[0-9]+]] -// CHECK-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX4]] -// CHECK-DAG: [[PADDR4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX4]] +// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 5, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* [[SR:%[^,]+]], i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[MAPT7]], i32 0, i32 0), i8** null, i32 1, i32 0) +// CHECK-DAG: [[BPR]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP:%.+]], i32 0, i32 0 +// CHECK-DAG: [[PR]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P:%.+]], i32 0, i32 0 +// CHECK-DAG: [[SR]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S:%.+]], i32 0, i32 0 +// CHECK-DAG: [[SADDR0:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX0:[0-9]+]] +// CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX0]] +// CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX0]] +// CHECK-DAG: [[SADDR1:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX1:[0-9]+]] +// CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX1]] +// CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX1]] +// CHECK-DAG: [[SADDR2:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX2:[0-9]+]] +// CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX2]] +// CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX2]] +// CHECK-DAG: [[SADDR3:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX3:[0-9]+]] +// CHECK-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX3]] +// CHECK-DAG: [[PADDR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX3]] +// CHECK-DAG: [[SADDR4:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX4:[0-9]+]] +// CHECK-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX4]] +// CHECK-DAG: [[PADDR4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX4]] // The names below are not necessarily consistent with the names used for the // addresses above as some are repeated. @@ -621,9 +621,9 @@ int bar(int n){ // CHECK-DAG: [[CPADDR2]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* // CHECK-DAG: store [[S1]]* %{{.+}}, [[S1]]** [[CBPADDR3:%.+]], -// CHECK-DAG: store [[S1]]* %{{.+}}, [[S1]]** [[CPADDR3:%.+]], +// CHECK-DAG: store double* %{{.+}}, double** [[CPADDR3:%.+]], // CHECK-DAG: [[CBPADDR3]] = bitcast i8** {{%[^,]+}} to [[S1]]** -// CHECK-DAG: [[CPADDR3]] = bitcast i8** {{%[^,]+}} to [[S1]]** +// CHECK-DAG: [[CPADDR3]] = bitcast i8** {{%[^,]+}} to double* // CHECK-DAG: store i16* %{{.+}}, i16** [[CBPADDR4:%.+]], // CHECK-DAG: store i16* %{{.+}}, i16** [[CPADDR4:%.+]], @@ -634,7 +634,7 @@ int bar(int n){ // CHECK-DAG: store i64 [[CSIZE]], i64* {{%[^,]+}} // CHECK-DAG: store i64 {{4|8}}, i64* {{%[^,]+}} // CHECK-DAG: store i64 {{4|8}}, i64* {{%[^,]+}} -// CHECK-DAG: store i64 %{{.+}}, i64* {{%[^,]+}} +// CHECK-DAG: store i64 {{4|8}}, i64* {{%[^,]+}} // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:[^,]+]], label %[[END:[^,]+]] diff --git a/clang/test/OpenMP/target_parallel_for_codegen.cpp b/clang/test/OpenMP/target_parallel_for_codegen.cpp index 5e2c865c944d5e..0da49e7eedfaca 100644 --- a/clang/test/OpenMP/target_parallel_for_codegen.cpp +++ b/clang/test/OpenMP/target_parallel_for_codegen.cpp @@ -101,7 +101,7 @@ // CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i64] [i64 800, i64 800, i64 547] // CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [4 x i64] [i64 4, i64 2, i64 1, i64 40] // CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i64] [i64 800, i64 800, i64 800, i64 547] -// CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [6 x i64] [i64 32, i64 281474976711171, i64 800, i64 800, i64 800, i64 547] +// CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [5 x i64] [i64 547, i64 800, i64 800, i64 800, i64 547] // CHECK-DAG: @{{.*}} = weak constant i8 0 // CHECK-DAG: @{{.*}} = weak constant i8 0 // CHECK-DAG: @{{.*}} = weak constant i8 0 @@ -634,28 +634,25 @@ int bar(int n){ // CHECK-32: [[CSZSIZE:%.+]] = mul nuw i32 [[CELEMSIZE2]], 2 // CHECK-32: [[CSIZE:%.+]] = sext i32 [[CSZSIZE]] to i64 -// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 6, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* [[SR:%[^,]+]], i64* getelementptr inbounds ([6 x i64], [6 x i64]* [[MAPT7]], i32 0, i32 0), i8** null, i32 1, i32 0) -// CHECK-DAG: [[BPR]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP:%.+]], i32 0, i32 0 -// CHECK-DAG: [[PR]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P:%.+]], i32 0, i32 0 -// CHECK-DAG: [[SR]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S:%.+]], i32 0, i32 0 -// CHECK-DAG: [[SADDR0:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX0:[0-9]+]] -// CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX0]] -// CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX0]] -// CHECK-DAG: [[SADDR1:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX1:[0-9]+]] -// CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX1]] -// CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX1]] -// CHECK-DAG: [[SADDR2:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX2:[0-9]+]] -// CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX2]] -// CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX2]] -// CHECK-DAG: [[SADDR3:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX3:[0-9]+]] -// CHECK-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX3]] -// CHECK-DAG: [[PADDR3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX3]] -// CHECK-DAG: [[SADDR4:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX4:[0-9]+]] -// CHECK-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX4]] -// CHECK-DAG: [[PADDR4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX4]] -// CHECK-DAG: [[SADDR5:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX5:[0-9]+]] -// CHECK-DAG: [[BPADDR5:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX5]] -// CHECK-DAG: [[PADDR5:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX5]] +// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 5, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* [[SR:%[^,]+]], i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[MAPT7]], i32 0, i32 0), i8** null, i32 1, i32 0) +// CHECK-DAG: [[BPR]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP:%.+]], i32 0, i32 0 +// CHECK-DAG: [[PR]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P:%.+]], i32 0, i32 0 +// CHECK-DAG: [[SR]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S:%.+]], i32 0, i32 0 +// CHECK-DAG: [[SADDR0:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX0:[0-9]+]] +// CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX0]] +// CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX0]] +// CHECK-DAG: [[SADDR1:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX1:[0-9]+]] +// CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX1]] +// CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX1]] +// CHECK-DAG: [[SADDR2:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX2:[0-9]+]] +// CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX2]] +// CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX2]] +// CHECK-DAG: [[SADDR3:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX3:[0-9]+]] +// CHECK-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX3]] +// CHECK-DAG: [[PADDR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX3]] +// CHECK-DAG: [[SADDR4:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX4:[0-9]+]] +// CHECK-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX4]] +// CHECK-DAG: [[PADDR4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX4]] // The names below are not necessarily consistent with the names used for the // addresses above as some are repeated. @@ -683,16 +680,10 @@ int bar(int n){ // CHECK-DAG: [[CPADDR3]] = bitcast i8** {{%[^,]+}} to double** // CHECK-DAG: store i64 8, i64* {{%[^,]+}} -// CHECK-DAG: store [[S1]]* %{{.+}}, [[S1]]** [[CBPADDR4:%.+]], -// CHECK-DAG: store double* %{{.+}}, double** [[CPADDR4:%.+]], -// CHECK-DAG: [[CBPADDR4]] = bitcast i8** {{%[^,]+}} to [[S1]]** -// CHECK-DAG: [[CPADDR4]] = bitcast i8** {{%[^,]+}} to double** -// CHECK-DAG: store i64 8, i64* {{%[^,]+}} - -// CHECK-DAG: store i16* %{{.+}}, i16** [[CBPADDR5:%.+]], -// CHECK-DAG: store i16* %{{.+}}, i16** [[CPADDR5:%.+]], -// CHECK-DAG: [[CBPADDR5]] = bitcast i8** {{%[^,]+}} to i16** -// CHECK-DAG: [[CPADDR5]] = bitcast i8** {{%[^,]+}} to i16** +// CHECK-DAG: store i16* %{{.+}}, i16** [[CBPADDR4:%.+]], +// CHECK-DAG: store i16* %{{.+}}, i16** [[CPADDR4:%.+]], +// CHECK-DAG: [[CBPADDR4]] = bitcast i8** {{%[^,]+}} to i16** +// CHECK-DAG: [[CPADDR4]] = bitcast i8** {{%[^,]+}} to i16** // CHECK-DAG: store i64 [[CSIZE]], i64* {{%[^,]+}} // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 diff --git a/clang/test/OpenMP/target_parallel_for_simd_codegen.cpp b/clang/test/OpenMP/target_parallel_for_simd_codegen.cpp index 802eb95b132440..18cb489719b286 100644 --- a/clang/test/OpenMP/target_parallel_for_simd_codegen.cpp +++ b/clang/test/OpenMP/target_parallel_for_simd_codegen.cpp @@ -91,8 +91,8 @@ // CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i64] [i64 800, i64 800, i64 547] // CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [4 x i64] [i64 4, i64 2, i64 1, i64 40] // CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i64] [i64 800, i64 800, i64 800, i64 547] -// OMP45-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [6 x i64] [i64 32, i64 281474976711171, i64 800, i64 800, i64 800, i64 547] -// OMP50-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [7 x i64] [i64 32, i64 281474976711171, i64 800, i64 800, i64 800, i64 547, i64 800] +// OMP45-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [5 x i64] [i64 547, i64 800, i64 800, i64 800, i64 547] +// OMP50-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [6 x i64] [i64 547, i64 800, i64 800, i64 800, i64 547, i64 800] // CHECK-DAG: @{{.*}} = weak constant i8 0 // CHECK-DAG: @{{.*}} = weak constant i8 0 // CHECK-DAG: @{{.*}} = weak constant i8 0 @@ -608,53 +608,47 @@ int bar(int n){ // CHECK-32: [[CSZSIZE:%.+]] = mul nuw i32 [[CELEMSIZE2]], 2 // CHECK-32: [[CSIZE:%.+]] = sext i32 [[CSZSIZE]] to i64 -// OMP45-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 6, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* [[SR:%[^,]+]], i64* getelementptr inbounds ([6 x i64], [6 x i64]* [[MAPT7]], i32 0, i32 0), i8** null, i32 1, i32 0) -// OMP50-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 7, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* [[SR:%[^,]+]], i64* getelementptr inbounds ([7 x i64], [7 x i64]* [[MAPT7]], i32 0, i32 0), i8** null, i32 1, i32 %{{.+}}) -// OMP45-DAG: [[BPR]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP:%.+]], i32 0, i32 0 -// OMP45-DAG: [[PR]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P:%.+]], i32 0, i32 0 -// OMP45-DAG: [[SR]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S:%.+]], i32 0, i32 0 -// OMP45-DAG: [[SADDR0:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX0:[0-9]+]] -// OMP45-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX0]] -// OMP45-DAG: [[PADDR0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX0]] -// OMP45-DAG: [[SADDR1:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX1:[0-9]+]] -// OMP45-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX1]] -// OMP45-DAG: [[PADDR1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX1]] -// OMP45-DAG: [[SADDR2:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX2:[0-9]+]] -// OMP45-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX2]] -// OMP45-DAG: [[PADDR2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX2]] -// OMP45-DAG: [[SADDR3:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX3:[0-9]+]] -// OMP45-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX3]] -// OMP45-DAG: [[PADDR3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX3]] -// OMP45-DAG: [[SADDR4:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX4:[0-9]+]] -// OMP45-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX4]] -// OMP45-DAG: [[PADDR4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX4]] -// OMP45-DAG: [[SADDR5:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX5:[0-9]+]] -// OMP45-DAG: [[BPADDR5:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX5]] -// OMP45-DAG: [[PADDR5:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX5]] -// OMP50-DAG: [[BPR]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[BP:%.+]], i32 0, i32 0 -// OMP50-DAG: [[PR]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[P:%.+]], i32 0, i32 0 -// OMP50-DAG: [[SR]] = getelementptr inbounds [7 x i64], [7 x i64]* [[S:%.+]], i32 0, i32 0 -// OMP50-DAG: [[SADDR0:%.+]] = getelementptr inbounds [7 x i64], [7 x i64]* [[S]], i32 [[IDX0:[0-9]+]] -// OMP50-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[BP]], i32 [[IDX0]] -// OMP50-DAG: [[PADDR0:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[P]], i32 [[IDX0]] -// OMP50-DAG: [[SADDR1:%.+]] = getelementptr inbounds [7 x i64], [7 x i64]* [[S]], i32 [[IDX1:[0-9]+]] -// OMP50-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[BP]], i32 [[IDX1]] -// OMP50-DAG: [[PADDR1:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[P]], i32 [[IDX1]] -// OMP50-DAG: [[SADDR2:%.+]] = getelementptr inbounds [7 x i64], [7 x i64]* [[S]], i32 [[IDX2:[0-9]+]] -// OMP50-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[BP]], i32 [[IDX2]] -// OMP50-DAG: [[PADDR2:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[P]], i32 [[IDX2]] -// OMP50-DAG: [[SADDR3:%.+]] = getelementptr inbounds [7 x i64], [7 x i64]* [[S]], i32 [[IDX3:[0-9]+]] -// OMP50-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[BP]], i32 [[IDX3]] -// OMP50-DAG: [[PADDR3:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[P]], i32 [[IDX3]] -// OMP50-DAG: [[SADDR4:%.+]] = getelementptr inbounds [7 x i64], [7 x i64]* [[S]], i32 [[IDX4:[0-9]+]] -// OMP50-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[BP]], i32 [[IDX4]] -// OMP50-DAG: [[PADDR4:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[P]], i32 [[IDX4]] -// OMP50-DAG: [[SADDR5:%.+]] = getelementptr inbounds [7 x i64], [7 x i64]* [[S]], i32 [[IDX5:[0-9]+]] -// OMP50-DAG: [[BPADDR5:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[BP]], i32 [[IDX5]] -// OMP50-DAG: [[PADDR5:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[P]], i32 [[IDX5]] -// OMP50-DAG: [[SADDR6:%.+]] = getelementptr inbounds [7 x i64], [7 x i64]* [[S]], i32 [[IDX6:[0-9]+]] -// OMP50-DAG: [[BPADDR6:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[BP]], i32 [[IDX6]] -// OMP50-DAG: [[PADDR6:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[P]], i32 [[IDX6]] +// OMP45-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 5, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* [[SR:%[^,]+]], i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[MAPT7]], i32 0, i32 0), i8** null, i32 1, i32 0) +// OMP50-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 6, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* [[SR:%[^,]+]], i64* getelementptr inbounds ([6 x i64], [6 x i64]* [[MAPT7]], i32 0, i32 0), i8** null, i32 1, i32 %{{.+}}) +// OMP45-DAG: [[BPR]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP:%.+]], i32 0, i32 0 +// OMP45-DAG: [[PR]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P:%.+]], i32 0, i32 0 +// OMP45-DAG: [[SR]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S:%.+]], i32 0, i32 0 +// OMP45-DAG: [[SADDR0:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX0:[0-9]+]] +// OMP45-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX0]] +// OMP45-DAG: [[PADDR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX0]] +// OMP45-DAG: [[SADDR1:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX1:[0-9]+]] +// OMP45-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX1]] +// OMP45-DAG: [[PADDR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX1]] +// OMP45-DAG: [[SADDR2:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX2:[0-9]+]] +// OMP45-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX2]] +// OMP45-DAG: [[PADDR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX2]] +// OMP45-DAG: [[SADDR3:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX3:[0-9]+]] +// OMP45-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX3]] +// OMP45-DAG: [[PADDR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX3]] +// OMP45-DAG: [[SADDR4:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX4:[0-9]+]] +// OMP45-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX4]] +// OMP45-DAG: [[PADDR4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX4]] +// OMP50-DAG: [[BPR]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP:%.+]], i32 0, i32 0 +// OMP50-DAG: [[PR]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P:%.+]], i32 0, i32 0 +// OMP50-DAG: [[SR]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S:%.+]], i32 0, i32 0 +// OMP50-DAG: [[SADDR0:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX0:[0-9]+]] +// OMP50-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX0]] +// OMP50-DAG: [[PADDR0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX0]] +// OMP50-DAG: [[SADDR1:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX1:[0-9]+]] +// OMP50-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX1]] +// OMP50-DAG: [[PADDR1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX1]] +// OMP50-DAG: [[SADDR2:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX2:[0-9]+]] +// OMP50-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX2]] +// OMP50-DAG: [[PADDR2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX2]] +// OMP50-DAG: [[SADDR3:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX3:[0-9]+]] +// OMP50-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX3]] +// OMP50-DAG: [[PADDR3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX3]] +// OMP50-DAG: [[SADDR4:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX4:[0-9]+]] +// OMP50-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX4]] +// OMP50-DAG: [[PADDR4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX4]] +// OMP50-DAG: [[SADDR5:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX5:[0-9]+]] +// OMP50-DAG: [[BPADDR5:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX5]] +// OMP50-DAG: [[PADDR5:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX5]] // The names below are not necessarily consistent with the names used for the // addresses above as some are repeated. @@ -682,22 +676,16 @@ int bar(int n){ // CHECK-DAG: [[CPADDR3]] = bitcast i8** {{%[^,]+}} to double** // CHECK-DAG: store i64 8, i64* {{%[^,]+}} -// CHECK-DAG: store [[S1]]* %{{.+}}, [[S1]]** [[CBPADDR4:%.+]], -// CHECK-DAG: store double* %{{.+}}, double** [[CPADDR4:%.+]], -// CHECK-DAG: [[CBPADDR4]] = bitcast i8** {{%[^,]+}} to [[S1]]** -// CHECK-DAG: [[CPADDR4]] = bitcast i8** {{%[^,]+}} to double** -// CHECK-DAG: store i64 8, i64* {{%[^,]+}} - -// CHECK-DAG: store i16* %{{.+}}, i16** [[CBPADDR5:%.+]], -// CHECK-DAG: store i16* %{{.+}}, i16** [[CPADDR5:%.+]], -// CHECK-DAG: [[CBPADDR5]] = bitcast i8** {{%[^,]+}} to i16** -// CHECK-DAG: [[CPADDR5]] = bitcast i8** {{%[^,]+}} to i16** +// CHECK-DAG: store i16* %{{.+}}, i16** [[CBPADDR4:%.+]], +// CHECK-DAG: store i16* %{{.+}}, i16** [[CPADDR4:%.+]], +// CHECK-DAG: [[CBPADDR4]] = bitcast i8** {{%[^,]+}} to i16** +// CHECK-DAG: [[CPADDR4]] = bitcast i8** {{%[^,]+}} to i16** // CHECK-DAG: store i64 [[CSIZE]], i64* {{%[^,]+}} -// OMP50-DAG: store i[[SZ]] [[SIMD_COND]], i[[SZ]]* [[CBPADDR6:%.+]], -// OMP50-DAG: store i[[SZ]] [[SIMD_COND]], i[[SZ]]* [[CPADDR6:%.+]], -// OMP50-DAG: [[CBPADDR6]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* -// OMP50-DAG: [[CPADDR6]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* +// OMP50-DAG: store i[[SZ]] [[SIMD_COND]], i[[SZ]]* [[CBPADDR5:%.+]], +// OMP50-DAG: store i[[SZ]] [[SIMD_COND]], i[[SZ]]* [[CPADDR5:%.+]], +// OMP50-DAG: [[CBPADDR5]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* +// OMP50-DAG: [[CPADDR5]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* // OMP50-DAG: store i64 1, i64* {{%[^,]+}} // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 diff --git a/clang/test/OpenMP/target_parallel_if_codegen.cpp b/clang/test/OpenMP/target_parallel_if_codegen.cpp index 862bbe3c080690..09c96d1c4fc521 100644 --- a/clang/test/OpenMP/target_parallel_if_codegen.cpp +++ b/clang/test/OpenMP/target_parallel_if_codegen.cpp @@ -186,7 +186,7 @@ int bar(int n){ // CHECK: store i8 [[FB]], i8* [[CONV]], align // CHECK: [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align // -// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 4, {{.*}}, i8** null, i32 1, i32 [[NT:%.+]]) +// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 3, {{.*}}, i8** null, i32 1, i32 [[NT:%.+]]) // CHECK-DAG: [[NT]] = select i1 %{{.+}}, i32 0, i32 1 // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] @@ -213,7 +213,7 @@ int bar(int n){ // CHECK: br i1 [[CMP]], label {{%?}}[[IF_THEN:.+]], label {{%?}}[[IF_ELSE:.+]] // // CHECK: [[IF_THEN]] -// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 3, {{.*}}, i8** null, i32 1, i32 [[NT:%.+]]) +// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 2, {{.*}}, i8** null, i32 1, i32 [[NT:%.+]]) // CHECK-DAG: [[NT]] = select i1 %{{.+}}, i32 0, i32 1 // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK-NEXT: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] diff --git a/clang/test/OpenMP/target_parallel_num_threads_codegen.cpp b/clang/test/OpenMP/target_parallel_num_threads_codegen.cpp index e9a55c20298bc9..4b7886009770f7 100644 --- a/clang/test/OpenMP/target_parallel_num_threads_codegen.cpp +++ b/clang/test/OpenMP/target_parallel_num_threads_codegen.cpp @@ -189,7 +189,7 @@ int bar(int n){ // CHECK: [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align // CHECK: [[THREADS:%.+]] = load i32, i32* [[CAPE_ADDR]], align // -// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 4, {{.*}}, i8** null, i32 1, i32 [[THREADS]]) +// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 3, {{.*}}, i8** null, i32 1, i32 [[THREADS]]) // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // @@ -200,7 +200,7 @@ int bar(int n){ // // // -// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 2, {{.+}}, i8** null, i32 1, i32 1024) +// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 1, {{.+}}, i8** null, i32 1, i32 1024) // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // diff --git a/clang/test/OpenMP/target_teams_codegen.cpp b/clang/test/OpenMP/target_teams_codegen.cpp index e3e90c46de56d9..c40a4527c68ee6 100644 --- a/clang/test/OpenMP/target_teams_codegen.cpp +++ b/clang/test/OpenMP/target_teams_codegen.cpp @@ -101,7 +101,7 @@ // CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i64] [i64 800, i64 800, i64 547] // CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [4 x i64] [i64 4, i64 2, i64 1, i64 40] // CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i64] [i64 800, i64 800, i64 800, i64 547] -// CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [6 x i64] [i64 32, i64 281474976711171, i64 800, i64 800, i64 800, i64 547] +// CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [5 x i64] [i64 547, i64 800, i64 800, i64 800, i64 547] // CHECK-DAG: @{{.*}} = weak constant i8 0 // CHECK-DAG: @{{.*}} = weak constant i8 0 // CHECK-DAG: @{{.*}} = weak constant i8 0 @@ -629,25 +629,25 @@ int bar(int n){ // CHECK-32: [[CSZSIZE:%.+]] = mul nuw i32 [[CELEMSIZE2]], 2 // CHECK-32: [[CSIZE:%.+]] = sext i32 [[CSZSIZE]] to i64 -// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 6, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* [[SR:%[^,]+]], i64* getelementptr inbounds ([6 x i64], [6 x i64]* [[MAPT7]], i32 0, i32 0), i8** null, i32 0, i32 0) -// CHECK-DAG: [[BPR]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP:%.+]], i32 0, i32 0 -// CHECK-DAG: [[PR]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P:%.+]], i32 0, i32 0 -// CHECK-DAG: [[SR]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S:%.+]], i32 0, i32 0 -// CHECK-DAG: [[SADDR0:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX0:[0-9]+]] -// CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX0]] -// CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX0]] -// CHECK-DAG: [[SADDR1:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX1:[0-9]+]] -// CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX1]] -// CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX1]] -// CHECK-DAG: [[SADDR2:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX2:[0-9]+]] -// CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX2]] -// CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX2]] -// CHECK-DAG: [[SADDR3:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX3:[0-9]+]] -// CHECK-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX3]] -// CHECK-DAG: [[PADDR3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX3]] -// CHECK-DAG: [[SADDR4:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX4:[0-9]+]] -// CHECK-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX4]] -// CHECK-DAG: [[PADDR4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX4]] +// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 5, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* [[SR:%[^,]+]], i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[MAPT7]], i32 0, i32 0), i8** null, i32 0, i32 0) +// CHECK-DAG: [[BPR]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP:%.+]], i32 0, i32 0 +// CHECK-DAG: [[PR]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P:%.+]], i32 0, i32 0 +// CHECK-DAG: [[SR]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S:%.+]], i32 0, i32 0 +// CHECK-DAG: [[SADDR0:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX0:[0-9]+]] +// CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX0]] +// CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX0]] +// CHECK-DAG: [[SADDR1:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX1:[0-9]+]] +// CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX1]] +// CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX1]] +// CHECK-DAG: [[SADDR2:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX2:[0-9]+]] +// CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX2]] +// CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX2]] +// CHECK-DAG: [[SADDR3:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX3:[0-9]+]] +// CHECK-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX3]] +// CHECK-DAG: [[PADDR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX3]] +// CHECK-DAG: [[SADDR4:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX4:[0-9]+]] +// CHECK-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX4]] +// CHECK-DAG: [[PADDR4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX4]] // The names below are not necessarily consistent with the names used for the // addresses above as some are repeated. @@ -670,21 +670,15 @@ int bar(int n){ // CHECK-DAG: store i64 4, i64* {{%[^,]+}} // CHECK-DAG: store [[S1]]* %{{.+}}, [[S1]]** [[CBPADDR3:%.+]], -// CHECK-DAG: store [[S1]]* %{{.+}}, [[S1]]** [[CPADDR3:%.+]], +// CHECK-DAG: store double* %{{.+}}, double** [[CPADDR3:%.+]], // CHECK-DAG: [[CBPADDR3]] = bitcast i8** {{%[^,]+}} to [[S1]]** -// CHECK-DAG: [[CPADDR3]] = bitcast i8** {{%[^,]+}} to [[S1]]** +// CHECK-DAG: [[CPADDR3]] = bitcast i8** {{%[^,]+}} to double** // CHECK-DAG: store i64 8, i64* {{%[^,]+}} -// CHECK-DAG: store [[S1]]* %{{.+}}, [[S1]]** [[CBPADDR4:%.+]], -// CHECK-DAG: store [[S1]]* %{{.+}}, [[S1]]** [[CPADDR4:%.+]], -// CHECK-DAG: [[CBPADDR4]] = bitcast i8** {{%[^,]+}} to [[S1]]** -// CHECK-DAG: [[CPADDR4]] = bitcast i8** {{%[^,]+}} to [[S1]]** -// CHECK-DAG: store i64 8, i64* {{%[^,]+}} - -// CHECK-DAG: store i16* %{{.+}}, i16** [[CBPADDR5:%.+]], -// CHECK-DAG: store i16* %{{.+}}, i16** [[CPADDR5:%.+]], -// CHECK-DAG: [[CBPADDR5]] = bitcast i8** {{%[^,]+}} to i16** -// CHECK-DAG: [[CPADDR5]] = bitcast i8** {{%[^,]+}} to i16** +// CHECK-DAG: store i16* %{{.+}}, i16** [[CBPADDR4:%.+]], +// CHECK-DAG: store i16* %{{.+}}, i16** [[CPADDR4:%.+]], +// CHECK-DAG: [[CBPADDR4]] = bitcast i8** {{%[^,]+}} to i16** +// CHECK-DAG: [[CPADDR4]] = bitcast i8** {{%[^,]+}} to i16** // CHECK-DAG: store i64 [[CSIZE]], i64* {{%[^,]+}} // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 diff --git a/clang/test/OpenMP/target_teams_distribute_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_codegen.cpp index 1be4866695a19f..dacdcffda0419f 100644 --- a/clang/test/OpenMP/target_teams_distribute_codegen.cpp +++ b/clang/test/OpenMP/target_teams_distribute_codegen.cpp @@ -99,7 +99,7 @@ // CHECK-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [2 x i64] [i64 4, i64 2] // CHECK-DAG: [[MAPT3:@.+]] = private unnamed_addr constant [2 x i64] [i64 800, i64 800] // CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [10 x i64] [i64 800, i64 547, i64 800, i64 547, i64 547, i64 800, i64 800, i64 547, i64 547, i64 800] -// CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [6 x i64] [i64 32, i64 281474976711171, i64 800, i64 800, i64 800, i64 547] +// CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [5 x i64] [i64 547, i64 800, i64 800, i64 800, i64 547] // CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [5 x i64] [i64 4, i64 4, i64 2, i64 1, i64 40] // CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [5 x i64] [i64 800, i64 800, i64 800, i64 800, i64 547] // CHECK-DAG: [[SIZET7:@.+]] = private unnamed_addr constant [3 x i64] [i64 4, i64 2, i64 40] @@ -618,28 +618,25 @@ int bar(int n){ // CHECK-32: [[CSSZSIZE:%.+]] = mul nuw i[[SZ]] [[CELEMSIZE2]], 2 // CHECK-32: [[CSIZE:%.+]] = sext i32 [[CSSZSIZE]] to i64 -// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 6, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* [[SR:%[^,]+]], i64* getelementptr inbounds ([6 x i64], [6 x i64]* [[MAPT5]], i32 0, i32 0), i8** null, i32 0, i32 0) -// CHECK-DAG: [[BPR]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP:%.+]], i32 0, i32 0 -// CHECK-DAG: [[PR]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P:%.+]], i32 0, i32 0 -// CHECK-DAG: [[SR]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S:%.+]], i32 0, i32 0 -// CHECK-DAG: [[SADDR0:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX0:[0-9]+]] -// CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX0]] -// CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX0]] -// CHECK-DAG: [[SADDR1:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX1:[0-9]+]] -// CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX1]] -// CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX1]] -// CHECK-DAG: [[SADDR2:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX2:[0-9]+]] -// CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX2]] -// CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX2]] -// CHECK-DAG: [[SADDR3:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX3:[0-9]+]] -// CHECK-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX3]] -// CHECK-DAG: [[PADDR3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX3]] -// CHECK-DAG: [[SADDR4:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX3:[0-9]+]] -// CHECK-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX3]] -// CHECK-DAG: [[PADDR4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX3]] -// CHECK-DAG: [[SADDR5:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX3:[0-9]+]] -// CHECK-DAG: [[BPADDR5:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX3]] -// CHECK-DAG: [[PADDR5:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX3]] +// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 5, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* [[SR:%[^,]+]], i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[MAPT5]], i32 0, i32 0), i8** null, i32 0, i32 0) +// CHECK-DAG: [[BPR]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP:%.+]], i32 0, i32 0 +// CHECK-DAG: [[PR]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P:%.+]], i32 0, i32 0 +// CHECK-DAG: [[SR]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S:%.+]], i32 0, i32 0 +// CHECK-DAG: [[SADDR0:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX0:[0-9]+]] +// CHECK-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX0]] +// CHECK-DAG: [[PADDR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX0]] +// CHECK-DAG: [[SADDR1:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX1:[0-9]+]] +// CHECK-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX1]] +// CHECK-DAG: [[PADDR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX1]] +// CHECK-DAG: [[SADDR2:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX2:[0-9]+]] +// CHECK-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX2]] +// CHECK-DAG: [[PADDR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX2]] +// CHECK-DAG: [[SADDR3:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX3:[0-9]+]] +// CHECK-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX3]] +// CHECK-DAG: [[PADDR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX3]] +// CHECK-DAG: [[SADDR4:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX3:[0-9]+]] +// CHECK-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX3]] +// CHECK-DAG: [[PADDR4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX3]] // The names below are not necessarily consistent with the names used for the // addresses above as some are repeated. @@ -667,16 +664,10 @@ int bar(int n){ // CHECK-DAG: [[CPADDR3]] = bitcast i8** {{%[^,]+}} to double** // CHECK-DAG: store i64 8, i64* {{%[^,]+}} -// CHECK-DAG: store [[S1]]* %{{.+}}, [[S1]]** [[CBPADDR4:%.+]], -// CHECK-DAG: store double* %{{.+}}, double** [[CPADDR4:%.+]], -// CHECK-DAG: [[CBPADDR4]] = bitcast i8** {{%[^,]+}} to [[S1]]** -// CHECK-DAG: [[CPADDR4]] = bitcast i8** {{%[^,]+}} to double** -// CHECK-DAG: store i64 8, i64* {{%[^,]+}} - -// CHECK-DAG: store i16* %{{.+}}, i16** [[CBPADDR5:%.+]], -// CHECK-DAG: store i16* %{{.+}}, i16** [[CPADDR5:%.+]], -// CHECK-DAG: [[CBPADDR5]] = bitcast i8** {{%[^,]+}} to i16** -// CHECK-DAG: [[CPADDR5]] = bitcast i8** {{%[^,]+}} to i16** +// CHECK-DAG: store i16* %{{.+}}, i16** [[CBPADDR4:%.+]], +// CHECK-DAG: store i16* %{{.+}}, i16** [[CPADDR4:%.+]], +// CHECK-DAG: [[CBPADDR4]] = bitcast i8** {{%[^,]+}} to i16** +// CHECK-DAG: [[CPADDR4]] = bitcast i8** {{%[^,]+}} to i16** // CHECK-DAG: store i64 [[CSIZE]], i64* {{%[^,]+}} // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 diff --git a/clang/test/OpenMP/target_teams_distribute_simd_codegen.cpp b/clang/test/OpenMP/target_teams_distribute_simd_codegen.cpp index e4500cb7f30a55..1e9544f0ab8c62 100644 --- a/clang/test/OpenMP/target_teams_distribute_simd_codegen.cpp +++ b/clang/test/OpenMP/target_teams_distribute_simd_codegen.cpp @@ -91,8 +91,8 @@ // CHECK-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [2 x i64] [i64 4, i64 2] // CHECK-DAG: [[MAPT3:@.+]] = private unnamed_addr constant [2 x i64] [i64 800, i64 800] // CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [9 x i64] [i64 800, i64 547, i64 800, i64 547, i64 547, i64 800, i64 800, i64 547, i64 547] -// OMP45-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [6 x i64] [i64 32, i64 281474976711171, i64 800, i64 800, i64 800, i64 547] -// OMP50-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [7 x i64] [i64 32, i64 281474976711171, i64 800, i64 800, i64 800, i64 547, i64 800] +// OMP45-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [5 x i64] [i64 547, i64 800, i64 800, i64 800, i64 547] +// OMP50-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [6 x i64] [i64 547, i64 800, i64 800, i64 800, i64 547, i64 800] // CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [5 x i64] [i64 4, i64 4, i64 2, i64 1, i64 40] // CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [5 x i64] [i64 800, i64 800, i64 800, i64 800, i64 547] // CHECK-DAG: [[SIZET7:@.+]] = private unnamed_addr constant [3 x i64] [i64 4, i64 2, i64 40] @@ -608,53 +608,47 @@ int bar(int n){ // CHECK-32: [[CSZSIZE:%.+]] = mul nuw i32 [[CELEMSIZE2]], 2 // CHECK-32: [[CSIZE:%.+]] = sext i32 [[CSZSIZE]] to i64 -// OMP45-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 6, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* [[SR:%[^,]+]], i64* getelementptr inbounds ([6 x i64], [6 x i64]* [[MAPT5]], i32 0, i32 0), i8** null, i32 0, i32 1) -// OMP45-DAG: [[BPR]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP:%.+]], i32 0, i32 0 -// OMP45-DAG: [[PR]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P:%.+]], i32 0, i32 0 -// OMP45-DAG: [[SR]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S:%.+]], i32 0, i32 0 -// OMP45-DAG: [[SADDR0:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX0:[0-9]+]] -// OMP45-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX0]] -// OMP45-DAG: [[PADDR0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX0]] -// OMP45-DAG: [[SADDR1:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX1:[0-9]+]] -// OMP45-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX1]] -// OMP45-DAG: [[PADDR1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX1]] -// OMP45-DAG: [[SADDR2:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX2:[0-9]+]] -// OMP45-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX2]] -// OMP45-DAG: [[PADDR2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX2]] -// OMP45-DAG: [[SADDR3:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX3:[0-9]+]] -// OMP45-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX3]] -// OMP45-DAG: [[PADDR3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX3]] -// OMP45-DAG: [[SADDR4:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX4:[0-9]+]] -// OMP45-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX4]] -// OMP45-DAG: [[PADDR4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX4]] -// OMP45-DAG: [[SADDR5:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX5:[0-9]+]] -// OMP45-DAG: [[BPADDR5:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX5]] -// OMP45-DAG: [[PADDR5:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX5]] -// OMP50-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 7, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* [[SR:%[^,]+]], i64* getelementptr inbounds ([7 x i64], [7 x i64]* [[MAPT5]], i32 0, i32 0), i8** null, i32 0, i32 1) -// OMP50-DAG: [[BPR]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[BP:%.+]], i32 0, i32 0 -// OMP50-DAG: [[PR]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[P:%.+]], i32 0, i32 0 -// OMP50-DAG: [[SR]] = getelementptr inbounds [7 x i64], [7 x i64]* [[S:%.+]], i32 0, i32 0 -// OMP50-DAG: [[SADDR0:%.+]] = getelementptr inbounds [7 x i64], [7 x i64]* [[S]], i32 [[IDX0:[0-9]+]] -// OMP50-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[BP]], i32 [[IDX0]] -// OMP50-DAG: [[PADDR0:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[P]], i32 [[IDX0]] -// OMP50-DAG: [[SADDR1:%.+]] = getelementptr inbounds [7 x i64], [7 x i64]* [[S]], i32 [[IDX1:[0-9]+]] -// OMP50-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[BP]], i32 [[IDX1]] -// OMP50-DAG: [[PADDR1:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[P]], i32 [[IDX1]] -// OMP50-DAG: [[SADDR2:%.+]] = getelementptr inbounds [7 x i64], [7 x i64]* [[S]], i32 [[IDX2:[0-9]+]] -// OMP50-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[BP]], i32 [[IDX2]] -// OMP50-DAG: [[PADDR2:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[P]], i32 [[IDX2]] -// OMP50-DAG: [[SADDR3:%.+]] = getelementptr inbounds [7 x i64], [7 x i64]* [[S]], i32 [[IDX3:[0-9]+]] -// OMP50-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[BP]], i32 [[IDX3]] -// OMP50-DAG: [[PADDR3:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[P]], i32 [[IDX3]] -// OMP50-DAG: [[SADDR4:%.+]] = getelementptr inbounds [7 x i64], [7 x i64]* [[S]], i32 [[IDX4:[0-9]+]] -// OMP50-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[BP]], i32 [[IDX4]] -// OMP50-DAG: [[PADDR4:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[P]], i32 [[IDX4]] -// OMP50-DAG: [[SADDR5:%.+]] = getelementptr inbounds [7 x i64], [7 x i64]* [[S]], i32 [[IDX5:[0-9]+]] -// OMP50-DAG: [[BPADDR5:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[BP]], i32 [[IDX5]] -// OMP50-DAG: [[PADDR5:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[P]], i32 [[IDX5]] -// OMP50-DAG: [[SADDR6:%.+]] = getelementptr inbounds [7 x i64], [7 x i64]* [[S]], i32 [[IDX6:[0-9]+]] -// OMP50-DAG: [[BPADDR6:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[BP]], i32 [[IDX6]] -// OMP50-DAG: [[PADDR6:%.+]] = getelementptr inbounds [7 x i8*], [7 x i8*]* [[P]], i32 [[IDX6]] +// OMP45-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 5, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* [[SR:%[^,]+]], i64* getelementptr inbounds ([5 x i64], [5 x i64]* [[MAPT5]], i32 0, i32 0), i8** null, i32 0, i32 1) +// OMP45-DAG: [[BPR]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP:%.+]], i32 0, i32 0 +// OMP45-DAG: [[PR]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P:%.+]], i32 0, i32 0 +// OMP45-DAG: [[SR]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S:%.+]], i32 0, i32 0 +// OMP45-DAG: [[SADDR0:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX0:[0-9]+]] +// OMP45-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX0]] +// OMP45-DAG: [[PADDR0:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX0]] +// OMP45-DAG: [[SADDR1:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX1:[0-9]+]] +// OMP45-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX1]] +// OMP45-DAG: [[PADDR1:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX1]] +// OMP45-DAG: [[SADDR2:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX2:[0-9]+]] +// OMP45-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX2]] +// OMP45-DAG: [[PADDR2:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX2]] +// OMP45-DAG: [[SADDR3:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX3:[0-9]+]] +// OMP45-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX3]] +// OMP45-DAG: [[PADDR3:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX3]] +// OMP45-DAG: [[SADDR4:%.+]] = getelementptr inbounds [5 x i64], [5 x i64]* [[S]], i32 [[IDX4:[0-9]+]] +// OMP45-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[BP]], i32 [[IDX4]] +// OMP45-DAG: [[PADDR4:%.+]] = getelementptr inbounds [5 x i8*], [5 x i8*]* [[P]], i32 [[IDX4]] +// OMP50-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 6, i8** [[BPR:%[^,]+]], i8** [[PR:%[^,]+]], i64* [[SR:%[^,]+]], i64* getelementptr inbounds ([6 x i64], [6 x i64]* [[MAPT5]], i32 0, i32 0), i8** null, i32 0, i32 1) +// OMP50-DAG: [[BPR]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP:%.+]], i32 0, i32 0 +// OMP50-DAG: [[PR]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P:%.+]], i32 0, i32 0 +// OMP50-DAG: [[SR]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S:%.+]], i32 0, i32 0 +// OMP50-DAG: [[SADDR0:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX0:[0-9]+]] +// OMP50-DAG: [[BPADDR0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX0]] +// OMP50-DAG: [[PADDR0:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX0]] +// OMP50-DAG: [[SADDR1:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX1:[0-9]+]] +// OMP50-DAG: [[BPADDR1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX1]] +// OMP50-DAG: [[PADDR1:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX1]] +// OMP50-DAG: [[SADDR2:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX2:[0-9]+]] +// OMP50-DAG: [[BPADDR2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX2]] +// OMP50-DAG: [[PADDR2:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX2]] +// OMP50-DAG: [[SADDR3:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX3:[0-9]+]] +// OMP50-DAG: [[BPADDR3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX3]] +// OMP50-DAG: [[PADDR3:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX3]] +// OMP50-DAG: [[SADDR4:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX4:[0-9]+]] +// OMP50-DAG: [[BPADDR4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX4]] +// OMP50-DAG: [[PADDR4:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX4]] +// OMP50-DAG: [[SADDR5:%.+]] = getelementptr inbounds [6 x i64], [6 x i64]* [[S]], i32 [[IDX5:[0-9]+]] +// OMP50-DAG: [[BPADDR5:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[BP]], i32 [[IDX5]] +// OMP50-DAG: [[PADDR5:%.+]] = getelementptr inbounds [6 x i8*], [6 x i8*]* [[P]], i32 [[IDX5]] // The names below are not necessarily consistent with the names used for the // addresses above as some are repeated. @@ -662,42 +656,36 @@ int bar(int n){ // CHECK-DAG: store double* %{{.+}}, double** [[CPADDR0:%.+]], // CHECK-DAG: [[CBPADDR0]] = bitcast i8** {{%[^,]+}} to [[S1]]** // CHECK-DAG: [[CPADDR0]] = bitcast i8** {{%[^,]+}} to double** -// CHECK-DAG: store i64 %{{.+}}, i64* {{%[^,]+}} - -// CHECK-DAG: store [[S1]]* %{{.+}}, [[S1]]** [[CBPADDR1:%.+]], -// CHECK-DAG: store double* %{{.+}}, double** [[CPADDR1:%.+]], -// CHECK-DAG: [[CBPADDR1]] = bitcast i8** {{%[^,]+}} to [[S1]]** -// CHECK-DAG: [[CPADDR1]] = bitcast i8** {{%[^,]+}} to double** // CHECK-DAG: store i64 {{8|4}}, i64* {{%[^,]+}} -// CHECK-DAG: store i[[SZ]] [[B_CVAL]], i[[SZ]]* [[CBPADDR2:%.+]], -// CHECK-DAG: store i[[SZ]] [[B_CVAL]], i[[SZ]]* [[CPADDR2:%.+]], +// CHECK-DAG: store i[[SZ]] [[B_CVAL]], i[[SZ]]* [[CBPADDR1:%.+]], +// CHECK-DAG: store i[[SZ]] [[B_CVAL]], i[[SZ]]* [[CPADDR1:%.+]], +// CHECK-DAG: [[CBPADDR1]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* +// CHECK-DAG: [[CPADDR1]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* +// CHECK-DAG: store i64 4, i64* {{%[^,]+}} + +// CHECK-DAG: store i[[SZ]] 2, i[[SZ]]* [[CBPADDR2:%.+]], +// CHECK-DAG: store i[[SZ]] 2, i[[SZ]]* [[CPADDR2:%.+]], // CHECK-DAG: [[CBPADDR2]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* // CHECK-DAG: [[CPADDR2]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* -// CHECK-DAG: store i64 4, i64* {{%[^,]+}} +// CHECK-DAG: store i64 {{4|8}}, i64* {{%[^,]+}} -// CHECK-DAG: store i[[SZ]] 2, i[[SZ]]* [[CBPADDR3:%.+]], -// CHECK-DAG: store i[[SZ]] 2, i[[SZ]]* [[CPADDR3:%.+]], +// CHECK-DAG: store i[[SZ]] [[VLA0]], i[[SZ]]* [[CBPADDR3:%.+]], +// CHECK-DAG: store i[[SZ]] [[VLA0]], i[[SZ]]* [[CPADDR3:%.+]], // CHECK-DAG: [[CBPADDR3]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* // CHECK-DAG: [[CPADDR3]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* // CHECK-DAG: store i64 {{4|8}}, i64* {{%[^,]+}} -// CHECK-DAG: store i[[SZ]] [[VLA0]], i[[SZ]]* [[CBPADDR4:%.+]], -// CHECK-DAG: store i[[SZ]] [[VLA0]], i[[SZ]]* [[CPADDR4:%.+]], -// CHECK-DAG: [[CBPADDR4]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* -// CHECK-DAG: [[CPADDR4]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* -// CHECK-DAG: store i64 {{4|8}}, i64* {{%[^,]+}} - -// CHECK-DAG: store i16* %{{.+}}, i16** [[CBPADDR5:%.+]], -// CHECK-DAG: store i16* %{{.+}}, i16** [[CPADDR5:%.+]], -// CHECK-DAG: [[CBPADDR5]] = bitcast i8** {{%[^,]+}} to i16** -// CHECK-DAG: [[CPADDR5]] = bitcast i8** {{%[^,]+}} to i16** +// CHECK-DAG: store i16* %{{.+}}, i16** [[CBPADDR4:%.+]], +// CHECK-DAG: store i16* %{{.+}}, i16** [[CPADDR4:%.+]], +// CHECK-DAG: [[CBPADDR4]] = bitcast i8** {{%[^,]+}} to i16** +// CHECK-DAG: [[CPADDR4]] = bitcast i8** {{%[^,]+}} to i16** // CHECK-DAG: store i64 [[CSIZE]], i64* {{%[^,]+}} -// OMP50-DAG: store i[[SZ]] [[IF_C:%.+]], i[[SZ]]* [[CBPADDR6:%.+]], -// OMP50-DAG: store i[[SZ]] [[IF_C]], i[[SZ]]* [[CPADDR6:%.+]], -// OMP50-DAG: [[CBPADDR6]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* -// OMP50-DAG: [[CPADDR6]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* +// OMP50-DAG: store i[[SZ]] [[IF_C:%.+]], i[[SZ]]* [[CBPADDR5:%.+]], +// OMP50-DAG: store i[[SZ]] [[IF_C]], i[[SZ]]* [[CPADDR5:%.+]], +// OMP50-DAG: [[CBPADDR5]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* +// OMP50-DAG: [[CPADDR5]] = bitcast i8** {{%[^,]+}} to i[[SZ]]* // OMP50-DAG: store i64 1, i64* {{%[^,]+}} // CHECK-NEXT: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 diff --git a/clang/test/OpenMP/target_teams_num_teams_codegen.cpp b/clang/test/OpenMP/target_teams_num_teams_codegen.cpp index db67a45127997b..9cff89b49f867d 100644 --- a/clang/test/OpenMP/target_teams_num_teams_codegen.cpp +++ b/clang/test/OpenMP/target_teams_num_teams_codegen.cpp @@ -189,7 +189,7 @@ int bar(int n){ // CHECK: [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align // CHECK: [[TEAMS:%.+]] = load i32, i32* [[CAPE_ADDR]], align // -// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 4, {{.*}}, i8** null, i32 [[TEAMS]], i32 0) +// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 3, {{.*}}, i8** null, i32 [[TEAMS]], i32 0) // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // @@ -200,7 +200,7 @@ int bar(int n){ // // // -// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 2, {{.+}}, i8** null, i32 1024, i32 0) +// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 1, {{.+}}, i8** null, i32 1024, i32 0) // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // diff --git a/clang/test/OpenMP/target_teams_thread_limit_codegen.cpp b/clang/test/OpenMP/target_teams_thread_limit_codegen.cpp index 637e2ede86d5d8..288579c4f13baa 100644 --- a/clang/test/OpenMP/target_teams_thread_limit_codegen.cpp +++ b/clang/test/OpenMP/target_teams_thread_limit_codegen.cpp @@ -189,7 +189,7 @@ int bar(int n){ // CHECK: [[ARG:%.+]] = load i[[SZ]], i[[SZ]]* [[CAPEC_ADDR]], align // CHECK: [[TL:%.+]] = load i32, i32* [[CAPE_ADDR]], align // -// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 4, {{.*}}, i8** null, i32 0, i32 [[TL]]) +// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 3, {{.*}}, i8** null, i32 0, i32 [[TL]]) // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // @@ -200,7 +200,7 @@ int bar(int n){ // // // -// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 2, {{.+}}, i8** null, i32 0, i32 1024) +// CHECK-DAG: [[RET:%.+]] = call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 1, {{.+}}, i8** null, i32 0, i32 1024) // CHECK: [[ERROR:%.+]] = icmp ne i32 [[RET]], 0 // CHECK: br i1 [[ERROR]], label %[[FAIL:.+]], label %[[END:[^,]+]] // diff --git a/clang/test/OpenMP/teams_codegen.cpp b/clang/test/OpenMP/teams_codegen.cpp index 09d02025a56b4a..cea9ee75ef2a53 100644 --- a/clang/test/OpenMP/teams_codegen.cpp +++ b/clang/test/OpenMP/teams_codegen.cpp @@ -205,7 +205,7 @@ struct SS{ int foo(void) { int comp = 1; - // CK3-DAG: call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* %{{[^,]+}}, i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i8** null, i32 [[NT:%[^,]+]], i32 123) + // CK3-DAG: call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 2, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i8** null, i32 [[NT:%[^,]+]], i32 123) // CK3-DAG: [[NT]] = load i32, i32* [[NTA:%[^,]+]], // CK3-DAG: [[NTA]] = getelementptr inbounds [[SSI]], [[SSI]]* [[NTB:%[^,]+]], i32 0, i32 0 @@ -218,7 +218,7 @@ struct SS{ ++comp; } - // CK3-DAG: call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* %{{[^,]+}}, i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i8** null, i32 456, i32 [[TL:%[^,]+]]) + // CK3-DAG: call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 2, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i8** null, i32 456, i32 [[TL:%[^,]+]]) // CK3-DAG: [[TL]] = add nsw i32 [[TLA:%[^,]+]], 123 // CK3-DAG: [[TLA]] = fptosi float [[TLB:%[^,]+]] to i32 diff --git a/clang/test/OpenMP/teams_distribute_codegen.cpp b/clang/test/OpenMP/teams_distribute_codegen.cpp index fa12fc09f81444..7acd70bd0ab7c8 100644 --- a/clang/test/OpenMP/teams_distribute_codegen.cpp +++ b/clang/test/OpenMP/teams_distribute_codegen.cpp @@ -158,7 +158,7 @@ struct SS{ // CK3: define {{.*}}i32 @{{.+}}foo{{.+}}( int foo(void) { - // CK3: call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 2, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* %{{.+}}, i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i8** null, i32 0, i32 0) + // CK3: call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i8** null, i32 0, i32 0) // CK3: call void @[[OFFL1:.+]]([[SSI]]* %{{.+}}) #pragma omp target #pragma omp teams distribute diff --git a/clang/test/OpenMP/teams_distribute_parallel_for_codegen.cpp b/clang/test/OpenMP/teams_distribute_parallel_for_codegen.cpp index 679fd549baf855..e69b28176568bf 100644 --- a/clang/test/OpenMP/teams_distribute_parallel_for_codegen.cpp +++ b/clang/test/OpenMP/teams_distribute_parallel_for_codegen.cpp @@ -161,7 +161,7 @@ struct SS{ // CK3: define {{.*}}i32 @{{.+}}foo{{.+}}( int foo(void) { - // CK3: call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 2, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* %{{.+}}, i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i8** null, i32 0, i32 0) + // CK3: call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 1, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i8** null, i32 0, i32 0) // CK3: call void @[[OFFL1:.+]]([[SSI]]* %{{.+}}) #pragma omp target #pragma omp teams distribute parallel for diff --git a/clang/test/OpenMP/teams_distribute_parallel_for_simd_codegen.cpp b/clang/test/OpenMP/teams_distribute_parallel_for_simd_codegen.cpp index 4118739226293f..e392571f2f27a7 100644 --- a/clang/test/OpenMP/teams_distribute_parallel_for_simd_codegen.cpp +++ b/clang/test/OpenMP/teams_distribute_parallel_for_simd_codegen.cpp @@ -169,7 +169,7 @@ struct SS{ // CK3: define {{.*}}i32 @{{.+}}foo{{.+}}( int foo(void) { int i; - // CK3: call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 3, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i{{64|32}}* %{{[^,]+}}, i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i8** null + // CK3: call i32 @__tgt_target_teams_mapper(i64 -1, i8* @{{[^,]+}}, i32 2, i8** %{{[^,]+}}, i8** %{{[^,]+}}, i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i64* {{.+}}@{{[^,]+}}, i32 0, i32 0), i8** null // CK3: call void @[[OFFL1:.+]]([[SSI]]* %{{.+}}) #pragma omp target #pragma omp teams distribute parallel for simd safelen(4) aligned(a) linear(i)