diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -7384,10 +7384,9 @@ // &p, &p, sizeof(float*), TARGET_PARAM | TO | FROM // // map(p[1:24]) + // &p, &p[1], 24*sizeof(float), TARGET_PARAM | TO | FROM | PTR_AND_OBJ + // in unified shared memory mode // p, &p[1], 24*sizeof(float), TARGET_PARAM | TO | FROM - // for data directives - // p, p, sizeof(float*), TARGET_PARAM | TO | FROM - // p, &p[1], 24*sizeof(float), PTR_AND_OBJ | TO | FROM // // map(s) // &s, &s, sizeof(S2), TARGET_PARAM | TO | FROM @@ -7522,6 +7521,7 @@ // Track if the map information being generated is the first for a list of // components. bool IsExpressionFirstInfo = true; + bool FirstPointerInComplexData = false; Address BP = Address::invalid(); const Expr *AssocExpr = I->getAssociatedExpression(); const auto *AE = dyn_cast(AssocExpr); @@ -7564,17 +7564,14 @@ QualType Ty = I->getAssociatedDeclaration()->getType().getNonReferenceType(); if (Ty->isAnyPointerType() && std::next(I) != CE) { - BP = CGF.EmitLoadOfPointer(BP, Ty->castAs()); - - // For non-data directives, we do not need to generate individual map - // information for the pointer, it can be associated with the combined - // storage. - if (CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory() || - !CurDir.is() || - !isOpenMPTargetDataManagementDirective( - CurDir.get() - ->getDirectiveKind())) - ++I; + // No need to generate individual map information for the pointer, it + // can be associated with the combined storage if shared memory mode is + // active. + if (CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory()) + BP = CGF.EmitLoadOfPointer(BP, Ty->castAs()); + else + FirstPointerInComplexData = true; + ++I; } } @@ -7609,8 +7606,19 @@ EncounteredME = dyn_cast(I->getAssociatedExpression()); // If we encounter a PTR_AND_OBJ entry from now on it should be marked // as MEMBER_OF the parent struct. - if (EncounteredME) + if (EncounteredME) { ShouldBeMemberOf = true; + // Do not emit as complex pointer if this is actually not array-like + // expression. + if (FirstPointerInComplexData) { + QualType Ty = std::prev(I) + ->getAssociatedDeclaration() + ->getType() + .getNonReferenceType(); + BP = CGF.EmitLoadOfPointer(BP, Ty->castAs()); + FirstPointerInComplexData = false; + } + } } auto Next = std::next(I); @@ -7750,10 +7758,11 @@ // same expression except for the first one. We also need to signal // this map is the first one that relates with the current capture // (there is a set of entries for each capture). - OpenMPOffloadMappingFlags Flags = getMapTypeBits( - MapType, MapModifiers, IsImplicit, - !IsExpressionFirstInfo || RequiresReference, - IsCaptureFirstInfo && !RequiresReference); + OpenMPOffloadMappingFlags Flags = + getMapTypeBits(MapType, MapModifiers, IsImplicit, + !IsExpressionFirstInfo || RequiresReference || + FirstPointerInComplexData, + IsCaptureFirstInfo && !RequiresReference); if (!IsExpressionFirstInfo) { // If we have a PTR_AND_OBJ pair where the OBJ is a pointer as well, @@ -7811,6 +7820,7 @@ IsExpressionFirstInfo = false; IsCaptureFirstInfo = false; + FirstPointerInComplexData = false; } } } @@ -8055,6 +8065,7 @@ // emission of that entry until the whole struct has been processed. llvm::MapVector> DeferredInfo; + MapCombinedInfoTy UseDevicePtrCombinedInfo; for (const auto *C : CurExecDir->getClausesOfKind()) { @@ -8074,13 +8085,22 @@ // We potentially have map information for this declaration already. // Look for the first set of components that refer to it. if (It != Info.end()) { - auto CI = std::find_if( - It->second.begin(), It->second.end(), [VD](const MapInfo &MI) { - return MI.Components.back().getAssociatedDeclaration() == VD; - }); + auto *CI = llvm::find_if(It->second, [VD](const MapInfo &MI) { + return MI.Components.back().getAssociatedDeclaration() == VD; + }); // If we found a map entry, signal that the pointer has to be returned // and move on to the next declaration. - if (CI != It->second.end()) { + // Exclude cases where the base pointer is mapped as array subscript, + // array section or array shaping. The base address is passed as a + // pointer to base in this case and cannot be used as a base for + // use_device_ptr list item. + auto PrevCI = std::next(CI->Components.rbegin()); + if (CI != It->second.end() && + (CGF.CGM.getOpenMPRuntime().hasRequiresUnifiedSharedMemory() || + isa(IE) || + !VD->getType().getNonReferenceType()->isPointerType() || + PrevCI == CI->Components.rend() || + isa(PrevCI->getAssociatedExpression()))) { CI->ReturnDevicePointer = true; continue; } @@ -8103,13 +8123,13 @@ } else { llvm::Value *Ptr = CGF.EmitLoadOfScalar(CGF.EmitLValue(IE), IE->getExprLoc()); - CombinedInfo.BasePointers.emplace_back(Ptr, VD); - CombinedInfo.Pointers.push_back(Ptr); - CombinedInfo.Sizes.push_back( + UseDevicePtrCombinedInfo.BasePointers.emplace_back(Ptr, VD); + UseDevicePtrCombinedInfo.Pointers.push_back(Ptr); + UseDevicePtrCombinedInfo.Sizes.push_back( llvm::Constant::getNullValue(CGF.Int64Ty)); - CombinedInfo.Types.push_back(OMP_MAP_RETURN_PARAM | - OMP_MAP_TARGET_PARAM); - CombinedInfo.Mappers.push_back(nullptr); + UseDevicePtrCombinedInfo.Types.push_back(OMP_MAP_RETURN_PARAM | + OMP_MAP_TARGET_PARAM); + UseDevicePtrCombinedInfo.Mappers.push_back(nullptr); } } } @@ -8260,6 +8280,8 @@ // We need to append the results of this capture to what we already have. CombinedInfo.append(CurInfo); } + // Append data for use_device_ptr clauses. + CombinedInfo.append(UseDevicePtrCombinedInfo); } /// Generate all the base pointers, section pointers, sizes, map types, and diff --git a/clang/test/OpenMP/target_data_codegen.cpp b/clang/test/OpenMP/target_data_codegen.cpp --- a/clang/test/OpenMP/target_data_codegen.cpp +++ b/clang/test/OpenMP/target_data_codegen.cpp @@ -555,7 +555,7 @@ void test_close_modifier(int arg) { S2 *ps; - // CK5: private unnamed_addr constant [6 x i64] [i64 1059, i64 32, i64 562949953422339, i64 562949953421328, i64 16, i64 1043] + // CK5: private unnamed_addr constant [5 x i64] [i64 1059, i64 32, i64 562949953421328, i64 16, i64 1043] #pragma omp target data map(close,tofrom: arg, ps->ps->ps->ps->s) { ++(arg); @@ -634,20 +634,17 @@ // Make sure the struct picks up present even if another element of the struct // doesn't have present. - // CK8: private unnamed_addr constant [15 x i64] + // CK8: private unnamed_addr constant [11 x i64] // ps1 // // PRESENT=0x1000 | TARGET_PARAM=0x20 = 0x1020 // MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x1000000000003 - // MEMBER_OF_1=0x1000000000000 | PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x1000000000013 - // MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x1000000001003 // MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1000000001010 // PRESENT=0x1000 | PTR_AND_OBJ=0x10 = 0x1010 // PRESENT=0x1000 | PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x1013 // // CK8-SAME: {{^}} [i64 [[#0x1020]], i64 [[#0x1000000000003]], - // CK8-SAME: {{^}} i64 [[#0x1000000000013]], i64 [[#0x1000000001003]], // CK8-SAME: {{^}} i64 [[#0x1000000001010]], i64 [[#0x1010]], i64 [[#0x1013]], // arg @@ -659,16 +656,13 @@ // ps2 // // PRESENT=0x1000 | TARGET_PARAM=0x20 = 0x1020 - // MEMBER_OF_9=0x9000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x9000000001003 - // MEMBER_OF_9=0x9000000000000 | PRESENT=0x1000 | PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x9000000001013 - // MEMBER_OF_9=0x9000000000000 | FROM=0x2 | TO=0x1 = 0x9000000000003 - // MEMBER_OF_9=0x9000000000000 | PTR_AND_OBJ=0x10 = 0x9000000000010 + // MEMBER_OF_7=0x9000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x7000000001003 + // MEMBER_OF_7=0x9000000000000 | PTR_AND_OBJ=0x10 = 0x7000000000010 // PTR_AND_OBJ=0x10 = 0x10 // PTR_AND_OBJ=0x10 | FROM=0x2 | TO=0x1 = 0x13 // - // CK8-SAME: {{^}} i64 [[#0x1020]], i64 [[#0x9000000001003]], - // CK8-SAME: {{^}} i64 [[#0x9000000001013]], i64 [[#0x9000000000003]], - // CK8-SAME: {{^}} i64 [[#0x9000000000010]], i64 [[#0x10]], i64 [[#0x13]]] + // CK8-SAME: {{^}} i64 [[#0x1020]], i64 [[#0x7000000001003]], + // CK8-SAME: {{^}} i64 [[#0x7000000000010]], i64 [[#0x10]], i64 [[#0x13]]] #pragma omp target data map(tofrom: ps1->s) \ map(present,tofrom: arg, ps1->ps->ps->ps->s, ps2->s) \ map(tofrom: ps2->ps->ps->ps->s) diff --git a/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp b/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp --- a/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp +++ b/clang/test/OpenMP/target_data_use_device_ptr_codegen.cpp @@ -22,18 +22,18 @@ double *g; // CK1: @g = global double* -// CK1: [[MTYPE00:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19] -// CK1: [[MTYPE01:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19] -// CK1: [[MTYPE03:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19] -// CK1: [[MTYPE04:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19] -// CK1: [[MTYPE05:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19] -// CK1: [[MTYPE06:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19] -// CK1: [[MTYPE07:@.+]] = {{.*}}constant [2 x i64] [i64 99, i64 19] -// CK1: [[MTYPE08:@.+]] = {{.*}}constant [4 x i64] [i64 99, i64 19, i64 35, i64 19] -// CK1: [[MTYPE09:@.+]] = {{.*}}constant [4 x i64] [i64 99, i64 19, i64 99, i64 19] -// CK1: [[MTYPE10:@.+]] = {{.*}}constant [4 x i64] [i64 99, i64 19, i64 99, i64 19] -// CK1: [[MTYPE11:@.+]] = {{.*}}constant [3 x i64] [i64 96, i64 35, i64 19] -// CK1: [[MTYPE12:@.+]] = {{.*}}constant [3 x i64] [i64 96, i64 35, i64 19] +// CK1: [[MTYPE00:@.+]] = {{.*}}constant [2 x i64] [i64 51, i64 96] +// CK1: [[MTYPE01:@.+]] = {{.*}}constant [2 x i64] [i64 51, i64 96] +// CK1: [[MTYPE03:@.+]] = {{.*}}constant [2 x i64] [i64 51, i64 96] +// CK1: [[MTYPE04:@.+]] = {{.*}}constant [2 x i64] [i64 51, i64 96] +// CK1: [[MTYPE05:@.+]] = {{.*}}constant [2 x i64] [i64 51, i64 96] +// CK1: [[MTYPE06:@.+]] = {{.*}}constant [2 x i64] [i64 51, i64 96] +// CK1: [[MTYPE07:@.+]] = {{.*}}constant [2 x i64] [i64 51, i64 96] +// CK1: [[MTYPE08:@.+]] = {{.*}}constant [3 x i64] [i64 51, i64 51, i64 96] +// CK1: [[MTYPE09:@.+]] = {{.*}}constant [4 x i64] [i64 51, i64 51, i64 96, i64 96] +// CK1: [[MTYPE10:@.+]] = {{.*}}constant [4 x i64] [i64 51, i64 51, i64 96, i64 96] +// CK1: [[MTYPE11:@.+]] = {{.*}}constant [2 x i64] [i64 51, i64 96] +// CK1: [[MTYPE12:@.+]] = {{.*}}constant [2 x i64] [i64 51, i64 96] // CK1-LABEL: @_Z3foo template @@ -41,13 +41,14 @@ float *l; T *t; - // CK1: [[T:%.+]] = load double*, double** [[DECL:@g]], // CK1: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0 + // CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to double*** + // CK1: store double** @g, double*** [[CBP]], + // CK1: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 1 // CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to double** - // CK1: store double* [[T]], double** [[CBP]], // CK1: call void @__tgt_target_data_begin{{.+}}[[MTYPE00]] // CK1: [[VAL:%.+]] = load double*, double** [[CBP]], - // CK1-NOT: store double* [[VAL]], double** [[DECL]], + // CK1-NOT: store double* [[VAL]], double** @g, // CK1: store double* [[VAL]], double** [[PVT:%.+]], // CK1: [[TT:%.+]] = load double*, double** [[PVT]], // CK1: getelementptr inbounds double, double* [[TT]], i32 1 @@ -56,14 +57,15 @@ ++g; } // CK1: call void @__tgt_target_data_end{{.+}}[[MTYPE00]] - // CK1: [[TTT:%.+]] = load double*, double** [[DECL]], + // CK1: [[TTT:%.+]] = load double*, double** @g, // CK1: getelementptr inbounds double, double* [[TTT]], i32 1 ++g; - // CK1: [[T1:%.+]] = load float*, float** [[DECL:%.+]], // CK1: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0 - // CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to float** - // CK1: store float* [[T1]], float** [[CBP]], + // CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to float*** + // CK1: store float** [[DECL:%.+]], float*** [[CBP]], + // CK1: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 1 + // CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to float* // CK1: call void @__tgt_target_data_begin{{.+}}[[MTYPE01]] // CK1: [[VAL:%.+]] = load float*, float** [[CBP]], // CK1-NOT: store float* [[VAL]], float** [[DECL]], @@ -91,10 +93,11 @@ // CK1: getelementptr inbounds float, float* [[TTT]], i32 1 ++l; - // CK1: [[T1:%.+]] = load float*, float** [[DECL:%.+]], // CK1: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0 + // CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to float*** + // CK1: store float** [[DECL:%.+]], float*** [[CBP]], + // CK1: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 1 // CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to float** - // CK1: store float* [[T1]], float** [[CBP]], // CK1: call void @__tgt_target_data_begin{{.+}}[[MTYPE03]] // CK1: [[VAL:%.+]] = load float*, float** [[CBP]], // CK1-NOT: store float* [[VAL]], float** [[DECL]], @@ -114,10 +117,11 @@ // CK1: br i1 [[CMP]], label %[[BTHEN:.+]], label %[[BELSE:.+]] // CK1: [[BTHEN]]: - // CK1: [[T1:%.+]] = load float*, float** [[DECL:%.+]], // CK1: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0 + // CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to float*** + // CK1: store float** [[DECL:%.+]], float*** [[CBP]], + // CK1: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 1 // CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to float** - // CK1: store float* [[T1]], float** [[CBP]], // CK1: call void @__tgt_target_data_begin{{.+}}[[MTYPE04]] // CK1: [[VAL:%.+]] = load float*, float** [[CBP]], // CK1-NOT: store float* [[VAL]], float** [[DECL]], @@ -151,10 +155,11 @@ ++l; // CK1: [[T2:%.+]] = load float**, float*** [[DECL:%.+]], - // CK1: [[T1:%.+]] = load float*, float** [[T2]], // CK1: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0 + // CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to float*** + // CK1: store float** [[T2:%.+]], float*** [[CBP]], + // CK1: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 1 // CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to float** - // CK1: store float* [[T1]], float** [[CBP]], // CK1: call void @__tgt_target_data_begin{{.+}}[[MTYPE05]] // CK1: [[VAL:%.+]] = load float*, float** [[CBP]], // CK1: store float* [[VAL]], float** [[PVTV:%.+]], @@ -173,10 +178,11 @@ // CK1: getelementptr inbounds float, float* [[TTTT]], i32 1 ++lr; - // CK1: [[T1:%.+]] = load i32*, i32** [[DECL:%.+]], // CK1: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0 + // CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to i32*** + // CK1: store i32** [[DECL:%.+]], i32*** [[CBP]], + // CK1: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 1 // CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to i32** - // CK1: store i32* [[T1]], i32** [[CBP]], // CK1: call void @__tgt_target_data_begin{{.+}}[[MTYPE06]] // CK1: [[VAL:%.+]] = load i32*, i32** [[CBP]], // CK1-NOT: store i32* [[VAL]], i32** [[DECL]], @@ -192,11 +198,13 @@ // CK1: getelementptr inbounds i32, i32* [[TTT]], i32 1 ++t; + // CK1: load i32**, i32*** % // CK1: [[T2:%.+]] = load i32**, i32*** [[DECL:%.+]], - // CK1: [[T1:%.+]] = load i32*, i32** [[T2]], // CK1: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 0 + // CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to i32*** + // CK1: store i32** [[T2]], i32*** [[CBP]], + // CK1: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 1 // CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to i32** - // CK1: store i32* [[T1]], i32** [[CBP]], // CK1: call void @__tgt_target_data_begin{{.+}}[[MTYPE07]] // CK1: [[VAL:%.+]] = load i32*, i32** [[CBP]], // CK1: store i32* [[VAL]], i32** [[PVTV:%.+]], @@ -215,10 +223,11 @@ // CK1: getelementptr inbounds i32, i32* [[TTTT]], i32 1 ++tr; - // CK1: [[T1:%.+]] = load float*, float** [[DECL:%.+]], - // CK1: [[BP:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* %{{.+}}, i32 0, i32 + // CK1: [[BP:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %{{.+}}, i32 0, i32 0 + // CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to float*** + // CK1: store float** [[DECL:%.+]], float*** [[CBP]], + // CK1: [[BP:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %{{.+}}, i32 0, i32 2 // CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to float** - // CK1: store float* [[T1]], float** [[CBP]], // CK1: call void @__tgt_target_data_begin{{.+}}[[MTYPE08]] // CK1: [[VAL:%.+]] = load float*, float** [[CBP]], // CK1-NOT: store float* [[VAL]], float** [[DECL]], @@ -235,6 +244,8 @@ ++l; ++t; + // CK1: bitcast i8** {{%.+}} to float*** + // CK1: bitcast i8** {{%.+}} to i32*** // CK1: [[_CBP:%.+]] = bitcast i8** {{%.+}} to float** // CK1: [[CBP:%.+]] = bitcast i8** {{%.+}} to i32** // CK1: call void @__tgt_target_data_begin{{.+}}[[MTYPE09]] @@ -257,6 +268,8 @@ // CK1: getelementptr inbounds i32, i32* [[TTT]], i32 1 ++l; ++t; + // CK1: bitcast i8** {{%.+}} to float*** + // CK1: bitcast i8** {{%.+}} to i32*** // CK1: [[_CBP:%.+]] = bitcast i8** {{%.+}} to float** // CK1: [[CBP:%.+]] = bitcast i8** {{%.+}} to i32** // CK1: call void @__tgt_target_data_begin{{.+}}[[MTYPE10]] @@ -280,7 +293,7 @@ ++l; ++t; // CK1: [[T1:%.+]] = load i32*, i32** [[DECL:%.+]], - // CK1: [[BP:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %{{.+}}, i32 0, i32 0 + // CK1: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 1 // CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to i32** // CK1: store i32* [[T1]], i32** [[CBP]], // CK1: call void @__tgt_target_data_begin{{.+}}[[MTYPE11]] @@ -300,7 +313,7 @@ // CK1: [[T2:%.+]] = load i32**, i32*** [[DECL:%.+]], // CK1: [[T1:%.+]] = load i32*, i32** [[T2]], - // CK1: [[BP:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %{{.+}}, i32 0, i32 0 + // CK1: [[BP:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* %{{.+}}, i32 0, i32 1 // CK1: [[CBP:%.+]] = bitcast i8** [[BP]] to i32** // CK1: store i32* [[T1]], i32** [[CBP]], // CK1: call void @__tgt_target_data_begin{{.+}}[[MTYPE12]] @@ -348,7 +361,7 @@ // CK2: [[ST:%.+]] = type { double*, double** } // CK2: [[MTYPE00:@.+]] = {{.*}}constant [2 x i64] [i64 32, i64 281474976710739] // CK2: [[MTYPE01:@.+]] = {{.*}}constant [2 x i64] [i64 32, i64 281474976710739] -// CK2: [[MTYPE02:@.+]] = {{.*}}constant [4 x i64] [i64 35, i64 19, i64 32, i64 844424930132048] +// CK2: [[MTYPE02:@.+]] = {{.*}}constant [3 x i64] [i64 51, i64 32, i64 562949953421392] // CK2: [[MTYPE03:@.+]] = {{.*}}constant [3 x i64] [i64 32, i64 281474976710739, i64 281474976710736] template @@ -404,7 +417,7 @@ // CK2: getelementptr inbounds double, double* [[TTTT]], i32 1 b++; - // CK2: [[BP:%.+]] = getelementptr inbounds [4 x i8*], [4 x i8*]* %{{.+}}, i32 0, i32 3 + // CK2: [[BP:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* %{{.+}}, i32 0, i32 2 // CK2: [[CBP:%.+]] = bitcast i8** [[BP]] to double*** // CK2: store double** [[RVAL:%.+]], double*** [[CBP]], // CK2: call void @__tgt_target_data_begin{{.+}}[[MTYPE02]] diff --git a/clang/test/OpenMP/target_map_codegen.cpp b/clang/test/OpenMP/target_map_codegen.cpp --- a/clang/test/OpenMP/target_map_codegen.cpp +++ b/clang/test/OpenMP/target_map_codegen.cpp @@ -1657,25 +1657,25 @@ // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK19: [[SIZE10:@.+]] = private {{.*}}constant [1 x i64] [i64 240] -// CK19: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i64] [i64 35] +// CK19: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i64] [i64 51] // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK19: [[SIZE11:@.+]] = private {{.*}}constant [1 x i64] [i64 240] -// CK19: [[MTYPE11:@.+]] = private {{.*}}constant [1 x i64] [i64 32] +// CK19: [[MTYPE11:@.+]] = private {{.*}}constant [1 x i64] [i64 48] // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK19: [[SIZE12:@.+]] = private {{.*}}constant [1 x i64] [i64 4] -// CK19: [[MTYPE12:@.+]] = private {{.*}}constant [1 x i64] [i64 33] +// CK19: [[MTYPE12:@.+]] = private {{.*}}constant [1 x i64] [i64 49] // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 -// CK19: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i64] [i64 32] +// CK19: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i64] [i64 48] // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 -// CK19: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i64] [i64 33] +// CK19: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i64] [i64 49] // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK19: [[SIZE15:@.+]] = private {{.*}}constant [1 x i64] [i64 4] -// CK19: [[MTYPE15:@.+]] = private {{.*}}constant [1 x i64] [i64 34] +// CK19: [[MTYPE15:@.+]] = private {{.*}}constant [1 x i64] [i64 50] // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK19-USE: [[MTYPE16:@.+]] = private {{.*}}constant [2 x i64] [i64 800, i64 33] @@ -1735,11 +1735,11 @@ // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK19: [[SIZE28:@.+]] = private {{.*}}constant [3 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 16] -// CK19: [[MTYPE28:@.+]] = private {{.*}}constant [3 x i64] [i64 35, i64 16, i64 19] +// CK19: [[MTYPE28:@.+]] = private {{.*}}constant [3 x i64] [i64 51, i64 16, i64 19] // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK19: [[SIZE29:@.+]] = private {{.*}}constant [3 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 4] -// CK19: [[MTYPE29:@.+]] = private {{.*}}constant [3 x i64] [i64 35, i64 16, i64 19] +// CK19: [[MTYPE29:@.+]] = private {{.*}}constant [3 x i64] [i64 51, i64 16, i64 19] // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK19-USE: [[MTYPE30:@.+]] = private {{.*}}constant [4 x i64] [i64 800, i64 800, i64 800, i64 35] @@ -1794,7 +1794,7 @@ // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK19: [[SIZE42:@.+]] = private {{.*}}constant [3 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 104] -// CK19: [[MTYPE42:@.+]] = private {{.*}}constant [3 x i64] [i64 35, i64 16, i64 19] +// CK19: [[MTYPE42:@.+]] = private {{.*}}constant [3 x i64] [i64 51, i64 16, i64 19] // CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK19: [[MTYPE43:@.+]] = private {{.*}}constant [1 x i64] [i64 35] @@ -2071,11 +2071,10 @@ // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 - // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** + // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32*** // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** - // CK19-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]] + // CK19-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]] // CK19-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] - // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 20 // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] @@ -2095,11 +2094,10 @@ // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 - // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** + // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32*** // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** - // CK19-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]] + // CK19-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]] // CK19-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] - // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 0 // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] @@ -2119,11 +2117,10 @@ // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 - // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** + // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32*** // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** - // CK19-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]] + // CK19-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]] // CK19-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] - // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 15 // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] @@ -2145,13 +2142,12 @@ // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 - // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** + // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32*** // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** - // CK19-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]] + // CK19-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]] // CK19-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], i{{.+}}* [[S0]] // CK19-DAG: [[CSVAL0]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}} - // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} %{{.*}} // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] @@ -2173,13 +2169,12 @@ // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 - // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** + // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32*** // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** - // CK19-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]] + // CK19-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]] // CK19-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], i{{.+}}* [[S0]] // CK19-DAG: [[CSVAL0]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}} - // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 0 // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] @@ -2199,11 +2194,10 @@ // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 - // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** + // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32*** // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** - // CK19-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]] + // CK19-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]] // CK19-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] - // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} %{{.*}} // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] @@ -2629,11 +2623,10 @@ // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 - // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**** + // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32***** // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**** - // CK19-DAG: store i32*** [[VAR0:%.+]], i32**** [[CBP0]] + // CK19-DAG: store i32**** [[PTR:%.+]], i32***** [[CBP0]] // CK19-DAG: store i32*** [[SEC0:%.+]], i32**** [[CP0]] - // CK19-DAG: [[VAR0]] = load i32***, i32**** [[PTR:%[^,]+]], // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32*** [[SEC00:[^,]+]], i{{.+}} 1 // CK19-DAG: [[SEC00]] = load i32***, i32**** [[PTR]], @@ -2677,11 +2670,10 @@ // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 - // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**** + // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32***** // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**** - // CK19-DAG: store i32*** [[VAR0:%.+]], i32**** [[CBP0]] + // CK19-DAG: store i32**** [[PTR:%.+]], i32***** [[CBP0]] // CK19-DAG: store i32*** [[SEC0:%.+]], i32**** [[CP0]] - // CK19-DAG: [[VAR0]] = load i32***, i32**** [[PTR:%[^,]+]], // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32*** [[SEC00:[^,]+]], i{{.+}} 1 // CK19-DAG: [[SEC00]] = load i32***, i32**** [[PTR]], @@ -3247,11 +3239,10 @@ // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 - // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to double**** + // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to double***** // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double**** - // CK19-DAG: store double*** [[VAR0:%.+]], double**** [[CBP0]] + // CK19-DAG: store double**** [[PTR:%.+]], double***** [[CBP0]] // CK19-DAG: store double*** [[SEC0:%.+]], double**** [[CP0]] - // CK19-DAG: [[VAR0]] = load double***, double**** [[PTR:%[^,]+]], // CK19-DAG: [[SEC0]] = getelementptr {{.*}}double*** [[SEC00:[^,]+]], i{{.+}} 0 // CK19-DAG: [[SEC00]] = load double***, double**** [[PTR]], @@ -3432,7 +3423,7 @@ // CK20-LABEL: @.__omp_offloading_{{.*}}explicit_maps_references_and_function_args{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK20: [[SIZE03:@.+]] = private {{.*}}constant [1 x i64] [i64 12] -// CK20: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i64] [i64 34] +// CK20: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i64] [i64 50] // CK20-LABEL: explicit_maps_references_and_function_args{{.*}}( void explicit_maps_references_and_function_args (int a, float b, int (&c)[10], float *d){ @@ -3508,11 +3499,10 @@ // CK20-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK20-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 - // CK20-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float** + // CK20-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float*** // CK20-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float** - // CK20-DAG: store float* [[RVAR0:%.+]], float** [[CBP0]] + // CK20-DAG: store float** [[VAR0:%.+]], float*** [[CBP0]] // CK20-DAG: store float* [[SEC0:%.+]], float** [[CP0]] - // CK20-DAG: [[RVAR0]] = load float*, float** [[VAR0:%[^,]+]] // CK20-DAG: [[SEC0]] = getelementptr {{.*}}float* [[RVAR00:%.+]], i{{.+}} 2 // CK20-DAG: [[RVAR00]] = load float*, float** [[VAR0]] @@ -3595,7 +3585,7 @@ // CK21-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK21: [[SIZE01:@.+]] = private {{.*}}constant [1 x i64] [i64 492] -// CK21: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 35] +// CK21: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 51] // CK21-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK21: [[MTYPE02:@.+]] = private {{.*}}constant [2 x i64] [i64 32, i64 281474976710674] @@ -3665,11 +3655,10 @@ // CK21-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK21-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 - // CK21-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** + // CK21-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32*** // CK21-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** - // CK21-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]] + // CK21-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]] // CK21-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] - // CK21-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] // CK21-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 0 // CK21-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] @@ -3874,7 +3863,7 @@ // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK22: [[SIZE04:@.+]] = private {{.*}}constant [1 x i64] [i64 20] -// CK22: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i64] [i64 35] +// CK22: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i64] [i64 51] // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK22: [[SIZE05:@.+]] = private {{.*}}constant [1 x i64] [i64 4] @@ -3894,7 +3883,7 @@ // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK22: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] [i64 20] -// CK22: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 35] +// CK22: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 51] // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK22: [[SIZE10:@.+]] = private {{.*}}constant [1 x i64] [i64 4] @@ -3914,7 +3903,7 @@ // CK22-LABEL: @.__omp_offloading_{{.*}}explicit_maps_globals{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK22: [[SIZE14:@.+]] = private {{.*}}constant [1 x i64] [i64 20] -// CK22: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i64] [i64 35] +// CK22: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i64] [i64 51] int a; int c[100]; @@ -4010,11 +3999,10 @@ // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 - // CK22-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** + // CK22-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32*** // CK22-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** - // CK22-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]] + // CK22-DAG: store i32** @d, i32*** [[CBP0]] // CK22-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] - // CK22-DAG: [[RVAR0]] = load i32*, i32** @d // CK22-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 2 // CK22-DAG: [[RVAR00]] = load i32*, i32** @d @@ -4093,11 +4081,10 @@ // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 - // CK22-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]** + // CK22-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]*** // CK22-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[ST]]** - // CK22-DAG: store [[ST]]* [[RVAR0:%.+]], [[ST]]** [[CBP0]] + // CK22-DAG: store [[ST]]** @sd, [[ST]]*** [[CBP0]] // CK22-DAG: store [[ST]]* [[SEC0:%.+]], [[ST]]** [[CP0]] - // CK22-DAG: [[RVAR0]] = load [[ST]]*, [[ST]]** @sd // CK22-DAG: [[SEC0]] = getelementptr {{.*}}[[ST]]* [[RVAR00:%.+]], i{{.+}} 2 // CK22-DAG: [[RVAR00]] = load [[ST]]*, [[ST]]** @sd @@ -4176,11 +4163,10 @@ // CK22-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK22-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 - // CK22-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[STT]]** + // CK22-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[STT]]*** // CK22-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[STT]]** - // CK22-DAG: store [[STT]]* [[RVAR0:%.+]], [[STT]]** [[CBP0]] + // CK22-DAG: store [[STT]]** @std, [[STT]]*** [[CBP0]] // CK22-DAG: store [[STT]]* [[SEC0:%.+]], [[STT]]** [[CP0]] - // CK22-DAG: [[RVAR0]] = load [[STT]]*, [[STT]]** @std // CK22-DAG: [[SEC0]] = getelementptr {{.*}}[[STT]]* [[RVAR00:%.+]], i{{.+}} 2 // CK22-DAG: [[RVAR00]] = load [[STT]]*, [[STT]]** @std @@ -4259,7 +4245,7 @@ // CK23-LABEL: @.__omp_offloading_{{.*}}explicit_maps_inside_captured{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK23: [[SIZE05:@.+]] = private {{.*}}constant [1 x i64] [i64 16] -// CK23: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i64] [i64 35] +// CK23: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i64] [i64 51] // CK23-LABEL: explicit_maps_inside_captured{{.*}}( int explicit_maps_inside_captured(int a){ @@ -4375,11 +4361,10 @@ // CK23-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK23-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 - // CK23-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float** + // CK23-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float*** // CK23-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float** - // CK23-DAG: store float* [[RVAR0:%.+]], float** [[CBP0]] + // CK23-DAG: store float** [[VAR0:%.+]], float*** [[CBP0]] // CK23-DAG: store float* [[SEC0:%.+]], float** [[CP0]] - // CK23-DAG: [[RVAR0]] = load float*, float** [[VAR0:%[^,]+]] // CK23-DAG: [[SEC0]] = getelementptr {{.*}}float* [[RVAR00:%.+]], i{{.+}} 2 // CK23-DAG: [[RVAR00]] = load float*, float** [[VAR00:%[^,]+]] // CK23-DAG: [[VAR0]] = load float**, float*** [[CAP0:%[^,]+]] @@ -5414,15 +5399,15 @@ // CK27-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK27: [[SIZE01:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer -// CK27: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 35] +// CK27: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 51] // CK27-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK27: [[SIZE02:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer -// CK27: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i64] [i64 35] +// CK27: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i64] [i64 51] // CK27-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK27: [[SIZE03:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer -// CK27: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i64] [i64 35] +// CK27: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i64] [i64 51] // CK27-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK27-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 @@ -5470,11 +5455,10 @@ // CK27-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK27-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 - // CK27-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** + // CK27-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32*** // CK27-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** - // CK27-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]] + // CK27-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]] // CK27-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] - // CK27-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] // CK27-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 0 // CK27-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] @@ -5491,11 +5475,10 @@ // CK27-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK27-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 - // CK27-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** + // CK27-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32*** // CK27-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** - // CK27-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]] + // CK27-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]] // CK27-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] - // CK27-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] // CK27-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 0 // CK27-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] @@ -5512,11 +5495,10 @@ // CK27-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK27-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 - // CK27-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** + // CK27-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32*** // CK27-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** - // CK27-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]] + // CK27-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]] // CK27-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] - // CK27-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] // CK27-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} %{{.+}} // CK27-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] @@ -5657,7 +5639,7 @@ // CK28-LABEL: @.__omp_offloading_{{.*}}explicit_maps_pointer_references{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 // CK28: [[SIZE01:@.+]] = private {{.*}}constant [1 x i64] [i64 400] -// CK28: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 35] +// CK28: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 51] // CK28-LABEL: explicit_maps_pointer_references{{.*}}( void explicit_maps_pointer_references (int *p){ @@ -5690,11 +5672,10 @@ // CK28-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK28-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 - // CK28-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** + // CK28-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32*** // CK28-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** - // CK28-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]] + // CK28-DAG: store i32** [[VAR00:%.+]], i32*** [[CBP0]] // CK28-DAG: store i32* [[VAR1:%.+]], i32** [[CP0]] - // CK28-DAG: [[VAR0]] = load i32*, i32** [[VAR00:%.+]], // CK28-DAG: [[VAR00]] = load i32**, i32*** [[VAR000:%.+]], // CK28-DAG: [[VAR1]] = getelementptr inbounds i32, i32* [[VAR11:%.+]], i{{64|32}} 2 // CK28-DAG: [[VAR11]] = load i32*, i32** [[VAR111:%.+]], @@ -6474,8 +6455,8 @@ // SIMD-ONLY32-NOT: {{__kmpc|__tgt}} #ifdef CK32 -// CK32-DAG: [[MTYPE_TO:@.+]] = {{.+}}constant [1 x i64] [i64 33] -// CK32-DAG: [[MTYPE_FROM:@.+]] = {{.+}}constant [1 x i64] [i64 34] +// CK32-DAG: [[MTYPE_TO:@.+]] = {{.+}}constant [1 x i64] [i64 49] +// CK32-DAG: [[MTYPE_FROM:@.+]] = {{.+}}constant [1 x i64] [i64 50] void array_shaping(float *f, int sa) { @@ -6488,14 +6469,13 @@ // CK32-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 // CK32-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 - // CK32-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float** + // CK32-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float*** // CK32-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to float** - // CK32-DAG: store float* [[F1:%.+]], float** [[BPC0]], + // CK32-DAG: store float** [[F_ADDR:%.+]], float*** [[BPC0]], // CK32-DAG: store float* [[F2:%.+]], float** [[PC0]], // CK32-DAG: store i64 [[SIZE:%.+]], i64* [[S0]], - // CK32-DAG: [[F1]] = load float*, float** [[F_ADDR:%.+]], // CK32-DAG: [[F2]] = load float*, float** [[F_ADDR]], // CK32-64-DAG: [[SIZE]] = mul nuw i64 [[SZ1:%.+]], 4 // CK32-64-DAG: [[SZ1]] = mul nuw i64 12, %{{.+}} @@ -6514,14 +6494,13 @@ // CK32-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 // CK32-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 - // CK32-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float** + // CK32-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float*** // CK32-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to float** - // CK32-DAG: store float* [[F1:%.+]], float** [[BPC0]], + // CK32-DAG: store float** [[F_ADDR:%.+]], float*** [[BPC0]], // CK32-DAG: store float* [[F2:%.+]], float** [[PC0]], // CK32-DAG: store i64 [[SIZE:%.+]], i64* [[S0]], - // CK32-DAG: [[F1]] = load float*, float** [[F_ADDR:%.+]], // CK32-DAG: [[F2]] = load float*, float** [[F_ADDR]], // CK32-64-DAG: [[SIZE]] = mul nuw i64 [[SZ1:%.+]], 5 // CK32-64-DAG: [[SZ1]] = mul nuw i64 4, %{{.+}} diff --git a/clang/test/OpenMP/target_update_codegen.cpp b/clang/test/OpenMP/target_update_codegen.cpp --- a/clang/test/OpenMP/target_update_codegen.cpp +++ b/clang/test/OpenMP/target_update_codegen.cpp @@ -310,18 +310,18 @@ #ifdef CK5 -// CK5: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}, i{{64|32}} 4] -// CK5: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 33, i64 17] +// CK5: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4] +// CK5: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 49] // CK5-LABEL: lvalue void lvalue(int *B, int l, int e) { - // CK5-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) + // CK5-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) // CK5-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK5-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] - // CK5-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 - // CK5-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK5-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK5-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 // CK5-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32*** // CK5-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to i32** // CK5-DAG: store i32** [[B_ADDR:%.+]], i32*** [[BPC0]] @@ -351,18 +351,18 @@ #ifdef CK6 -// CK6: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}, i{{64|32}} 4] -// CK6: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 33, i64 17] +// CK6: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4] +// CK6: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 49] // CK6-LABEL: lvalue void lvalue(int *B, int l, int e) { - // CK6-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) + // CK6-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) // CK6-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK6-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] - // CK6-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 - // CK6-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK6-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK6-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 // CK6-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32*** // CK6-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to i32** // CK6-DAG: store i32** [[B_ADDR:%.+]], i32*** [[BPC0]] @@ -397,18 +397,18 @@ #ifdef CK7 -// CK7: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}, i{{64|32}} 4] -// CK7: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 33, i64 17] +// CK7: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 4] +// CK7: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 49] // CK7-LABEL: lvalue void lvalue(int *B, int l, int e) { - // CK7-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) + // CK7-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) // CK7-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK7-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] - // CK7-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 - // CK7-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK7-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK7-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 // CK7-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32*** // CK7-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to i32** // CK7-DAG: store i32** [[B_ADDR:%.+]], i32*** [[BPC0]] @@ -446,18 +446,18 @@ // SIMD-ONLY0-NOT: {{__kmpc|__tgt}} #ifdef CK8 -// CK8: [[SIZE00:@.+]] = {{.+}}constant [3 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}, i{{64|32}} {{8|4}}, i{{64|32}} 4] -// CK8: [[MTYPE00:@.+]] = {{.+}}constant [3 x i64] [i64 33, i64 16, i64 17] +// CK8: [[SIZE00:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}, i{{64|32}} 4] +// CK8: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 49, i64 17] // CK8-LABEL: lvalue void lvalue(int **B, int l, int e) { - // CK8-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}], [3 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}, [3 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) + // CK8-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}], [2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}, [2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) // CK8-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK8-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] - // CK8-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 - // CK8-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 + // CK8-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK8-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 // CK8-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32*** // CK8-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to i32** // CK8-DAG: store i32** [[ARRAY_IDX_1:%.+]], i32*** [[BPC0]] @@ -466,13 +466,11 @@ // CK8-DAG: [[ARRAY_IDX_1]] = getelementptr inbounds i32*, i32** [[ADD_PTR:%.+]], i{{.+}} 1 // CK8-64-DAG: [[ADD_PTR]] = getelementptr inbounds i32*, i32** [[B_VAL:%.+]], i{{.+}} [[IDX_EXT:%.+]] // CK8-64-DAG: [[IDX_EXT]] = sext i32 [[L_VAL:%.+]] to i64 - // CK8-DAG: [[ARRAY_IDX_4]] = getelementptr inbounds i32, i32* [[FIVE:%.+]], i{{.+}} 2 - // CK8-64-DAG: [[ARRAY_IDX_3:%.+]] = getelementptr inbounds i32*, i32** [[ADD_PTR_2:%.+]], i{{.+}} 1 - // CK8-64-DAG: [[ADD_PTR_2]] = getelementptr inbounds i32*, i32** %3, i{{.+}} [[IDX_EXT_1:%.+]] - // CK8-64-DAG: [[IDX_EXT_1]] = sext i32 [[L_VAL:%.+]] to i{{.+}} + // CK8-64-DAG: [[ARRAY_IDX_4]] = getelementptr inbounds i32, i32* [[FIVE:%.+]], i{{.+}} 2 // CK8-32-DAG: [[ADD_PTR]] = getelementptr inbounds i32*, i32** [[B_VAL:%.+]], i{{.+}} [[L_VAL:%.+]] - // CK8-32-DAG: [[ARRAY_IDX_4:%.+]] = getelementptr inbounds i32*, i32** [[ADD_PTR_2:%.+]], i{{.+}} 1 - // CK8-32-DAG: [[ADD_PTR_2]] = getelementptr inbounds i32*, i32** %3, i{{.+}} [[L_VAL:%.+]] + // CK8-32-DAG: [[ARRAY_IDX_4:%.+]] = getelementptr inbounds i32, i32* [[ADD_PTR_2:%.+]], i{{.+}} 2 + // CK8-32-DAG: [[ADD_PTR_2]] = load i32*, i32** [[ARRAY_IDX_2:%.+]], + // CK8-32-DAG: [[ARRAY_IDX_2]] = getelementptr inbounds i32*, i32** %{{.+}}, i32 1 #pragma omp target update to((B+l)[1][2]) (B+l)[1][2] += e; #pragma omp target update from((B+l)[1][2]) @@ -501,19 +499,19 @@ double *p; }; -// CK9: [[MTYPE00:@.+]] = {{.+}}constant [3 x i64] [i64 32, i64 281474976710657, i64 281474976710673] +// CK9: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673] // CK9-LABEL: lvalue void lvalue(struct S *s, int l, int e) { - // CK9-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}, [3 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) + // CK9-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}, [2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) // CK9-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK9-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] // CK9-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]] // - // CK9-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 - // CK9-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 - // CK9-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 2 + // CK9-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK9-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK9-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1 // CK9-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to double*** // CK9-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to double** // CK9-DAG: store double** [[P:%.+]], double*** [[BPC0]] @@ -551,19 +549,19 @@ double *p; }; -// CK10: [[MTYPE00:@.+]] = {{.+}}constant [3 x i64] [i64 32, i64 281474976710657, i64 281474976710673] +// CK10: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673] // CK10-LABEL: lvalue void lvalue(struct S *s, int l, int e) { - // CK10-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}, [3 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) + // CK10-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}, [2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) // CK10-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK10-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] // CK10-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]] // - // CK10-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 - // CK10-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 - // CK10-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 2 + // CK10-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK10-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK10-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1 // CK10-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to double*** // CK10-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to double** // CK10-DAG: store double** [[P_VAL:%.+]], double*** [[BPC0]] @@ -601,19 +599,19 @@ struct S { double *p; }; -// CK11: [[MTYPE00:@.+]] = {{.+}}constant [3 x i64] [i64 32, i64 281474976710657, i64 281474976710673] +// CK11: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673] // CK11-LABEL: lvalue void lvalue(struct S *s, int l, int e) { - // CK11-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}, [3 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) + // CK11-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}, [2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) // CK11-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK11-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] // CK11-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]] // - // CK11-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 - // CK11-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 - // CK11-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 2 + // CK11-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK11-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK11-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1 // CK11-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to double*** // CK11-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to double** // CK11-DAG: store double** [[P:%.+]], double*** [[BPC0]] @@ -653,41 +651,41 @@ double *p; struct S *sp; }; -// CK12: [[MTYPE00:@.+]] = {{.+}}constant [4 x i64] [i64 32, i64 281474976710657, i64 281474976710672, i64 17] +// CK12: [[MTYPE00:@.+]] = {{.+}}constant [3 x i64] [i64 32, i64 281474976710672, i64 17] // CK12-LABEL: lvalue void lvalue(struct S *s, int l, int e) { - // CK12-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 4, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[4 x i{{.+}}, [4 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) + // CK12-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i{{.+}} [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}, [3 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) // CK12-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK12-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] // CK12-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]] // - // CK12-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3 - // CK12-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3 - // CK12-DAG: [[SIZE2:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 3 + // CK12-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 + // CK12-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 + // CK12-DAG: [[SIZE2:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 2 // CK12-DAG: [[BPC2:%.+]] = bitcast i8** [[BP2]] to double*** // CK12-DAG: [[PC2:%.+]] = bitcast i8** [[P2]] to double** // CK12-DAG: store double** [[P_VAL:%.+]], double*** [[BPC2]] // CK12-DAG: store double* [[SIX:%.+]], double** [[PC2]] // CK12-DAG: store i{{.+}} 8, i{{.+}}* [[SIZE2]] - // CK12-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 - // CK12-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 - // CK12-DAG: [[SIZE1:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 2 + // CK12-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK12-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK12-DAG: [[SIZE1:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1 // CK12-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to [[STRUCT_S:%.+]]*** // CK12-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to double*** // CK12-DAG: store [[STRUCT_S]]** [[SP:%.+]], [[STRUCT_S]]*** [[BPC1]] // CK12-DAG: store double** [[P_VAL:%.+]], double*** [[PC1]] // CK12-DAG: store i{{.+}} {{4|8}}, i{{.+}}* [[SIZE1]] - // CK12-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 - // CK12-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 - // CK12-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1 + // CK12-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK12-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK12-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 0 // CK12-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to [[STRUCT_S:%.+]]** // CK12-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to [[STRUCT_S]]*** - // CK12-DAG: store [[STRUCT_S]]** [[S:%.+]], [[STRUCT_S]]*** [[S_VAL:%.+]] - // CK12-DAG: store i{{.+}} {{.+}}, i{{.+}}* [[SIZE0]] + // CK12-DAG: store [[STRUCT_S]]* [[ZERO:%.+]], [[STRUCT_S]]** [[BPC0]] // CK12-DAG: [[SP]] = getelementptr inbounds [[STRUCT_S]], [[STRUCT_S]]* [[ONE:%.+]], i32 0, i32 1 - // CK12-DAG: [[ONE]] = load %struct.S*, %struct.S** [[S]], + // CK12-DAG: [[ONE]] = load %struct.S*, %struct.S** [[S:%.+]], + // CK12-DAG: [[ZERO]] = load %struct.S*, %struct.S** [[S]], #pragma omp target update to(*(s->sp->p)) *(s->sp->p) = e; #pragma omp target update from(*(s->sp->p)) @@ -711,18 +709,18 @@ // SIMD-ONLY0-NOT: {{__kmpc|__tgt}} #ifdef CK13 -// CK13: [[SIZE00:@.+]] = {{.+}}constant [2 x i64] [i64 {{8|4}}, i64 4] -// CK13: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 33, i64 17] +// CK13: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 4] +// CK13: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 49] // CK13-LABEL: lvalue void lvalue(int **BB, int a, int b) { - // CK13-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) + // CK13-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) // CK13-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK13-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] - // CK13-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 - // CK13-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK13-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK13-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 // CK13-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to i32**** // CK13-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to i32** // CK13-DAG: store i32*** [[BB_ADDR:%.+]], i32**** [[BPC0]] @@ -831,7 +829,7 @@ // SIMD-ONLY0-NOT: {{__kmpc|__tgt}} #ifdef CK15 -// CK15: [[MTYPE00:@.+]] = {{.+}}constant [3 x i64] [i64 32, i64 281474976710657, i64 281474976710673] +// CK15: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 32, i64 281474976710673] struct SSA { double *p; @@ -842,14 +840,14 @@ //CK-15-LABEL: lvalue_member void lvalue_member(SSA *sap) { - // CK15-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 3, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[3 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) + // CK15-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GSIZE:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) // CK15-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK15-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] // CK15-DAG: [[GSIZE]] = getelementptr inbounds {{.+}}[[SIZE:%[^,]+]] - // CK15-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 - // CK15-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 - // CK15-DAG: [[SIZE2:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 2 + // CK15-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 + // CK15-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK15-DAG: [[SIZE2:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1 // CK15-DAG: [[BPC2:%.+]] = bitcast i8** [[BP2]] to double*** // CK15-DAG: [[PC2:%.+]] = bitcast i8** [[P2]] to double** // CK15-DAG: store double** [[P_VAL:%.+]], double*** [[BPC2]] @@ -858,20 +856,11 @@ // CK15-DAG: [[ADD_PTR]] = getelementptr inbounds double, double* [[THREE:%.+]], i{{.+}} 3 // CK15-DAG: [[THREE]] = load double*, double** [[P_VAL_1:%.+]] // CK15-DAG: [[P_VAL]] = getelementptr inbounds [[SSA:%.+]], [[SSA:%.+]]* [[THIS:%.+]], i32 0, i32 0 - // CK15-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 - // CK15-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 - // CK15-DAG: [[SIZE1:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 1 - // CK15-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to [[SSA]]** - // CK15-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to [[SSA]]*** - // CK15-DAG: store [[SSA]]* [[SAP_VAL:%.+]], [[SSA]]** [[BPC1]], - // CK15-DAG: store [[SSA]]** [[SAP_ADDR:%.+]], [[SSA]]*** [[PC1]] - // CK15-DAG: store i{{.+}} {{8|4}}, i{{.+}}* [[SIZE1]] - // CK15-DAG: [[SAP_VAL]] = load [[SSA]]*, [[SSA]]** [[SAP_ADDR]], // CK15-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK15-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 // CK15-DAG: [[SIZE0:%.+]] = getelementptr inbounds {{.+}}[[SIZE]], i{{.+}} 0, i{{.+}} 0 - // CK15-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to [[SSA]]*** - // CK15-DAG: store [[SSA]]** [[SAP_ADDR]], [[SSA]]*** [[BPC0]], + // CK15-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to [[SSA]]** + // CK15-DAG: store [[SSA]]* [[SAP_ADDR:%.+]], [[SSA]]** [[BPC0]], // CK15-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to double*** // CK15-DAG: store double** [[P_VAL]], double*** [[PC0]], // CK15-DAG: store i{{.+}} [[COMPUTE_SIZE:%.+]], i{{.+}}* [[SIZE0]] @@ -904,18 +893,18 @@ // SIMD-ONLY0-NOT: {{__kmpc|__tgt}} #ifdef CK16 -// CK16: [[SIZE00:@.+]] = {{.+}}constant [2 x i64] [i64 {{8|4}}, i64 4] -// CK16: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 33, i64 17] +// CK16: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 4] +// CK16: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 49] //CK16-LABEL: lvalue_find_base void lvalue_find_base(float *f, int *i) { - // CK16-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) + // CK16-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) // CK16-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK16-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] - // CK16-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 - // CK16-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK16-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK16-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 // CK16-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float*** // CK16-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to float** // CK16-DAG: store float** [[F_ADDR:%.+]], float*** [[BPC0]] @@ -948,8 +937,8 @@ // SIMD-ONLY0-NOT: {{__kmpc|__tgt}} #ifdef CK17 -// CK17: [[SIZE00:@.+]] = {{.+}}constant [2 x i64] [i64 {{4|8}}, i64 4] -// CK17: [[MTYPE00:@.+]] = {{.+}}constant [2 x i64] [i64 33, i64 17] +// CK17: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 4] +// CK17: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 49] struct SSA { int i; @@ -959,12 +948,12 @@ //CK17-LABEL: lvalue_find_base void lvalue_find_base(float **f, SSA *sa) { - // CK17-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) + // CK17-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) // CK17-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK17-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] - // CK17-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 - // CK17-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 + // CK17-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK17-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 // CK17-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float**** // CK17-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to float** // CK17-DAG: store float*** [[F_ADDR:%.+]], float**** [[BPC0]], @@ -1005,37 +994,25 @@ // SIMD-ONLY18-NOT: {{__kmpc|__tgt}} #ifdef CK18 -// CK18-DAG: [[MTYPE_TO:@.+]] = {{.+}}constant [2 x i64] [i64 33, i64 16] -// CK18-DAG: [[MTYPE_FROM:@.+]] = {{.+}}constant [2 x i64] [i64 34, i64 16] +// CK18-DAG: [[MTYPE_TO:@.+]] = {{.+}}constant [1 x i64] [i64 49] +// CK18-DAG: [[MTYPE_FROM:@.+]] = {{.+}}constant [1 x i64] [i64 50] //CK18-LABEL: array_shaping void array_shaping(float *f, int sa) { - // CK18-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE_TO]]{{.+}}, i8** null) + // CK18-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE_TO]]{{.+}}, i8** null) // CK18-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK18-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] // CK18-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] - // CK18-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 - // CK18-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 - // CK18-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 - - // CK18-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float** - // CK18-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to float*** - - // CK18-DAG: store float* [[F1:%.+]], float** [[BPC0]], - // CK18-DAG: store float** [[F_ADDR:%.+]], float*** [[PC0]], - // CK18-DAG: store i64 {{8|4}}, i64* [[S0]], - // CK18-DAG: [[F1]] = load float*, float** [[F_ADDR]], - - // CK18-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 - // CK18-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 - // CK18-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 + // CK18-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK18-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK18-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 // CK18-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to float*** // CK18-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to float** - // CK18-DAG: store float** [[F_ADDR]], float*** [[BPC1]], + // CK18-DAG: store float** [[F_ADDR:%.+]], float*** [[BPC1]], // CK18-DAG: store float* [[F2:%.+]], float** [[PC1]], // CK18-DAG: store i64 [[SIZE:%.+]], i64* [[S1]], // CK18-DAG: [[F2]] = load float*, float** [[F_ADDR]], @@ -1047,31 +1024,19 @@ // CK18-32-DAG: [[SZ2]] = mul nuw i32 12, %{{.+}} #pragma omp target update to(([3][sa][4])f) sa = 1; - // CK18-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE_FROM]]{{.+}}, i8** null) + // CK18-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE_FROM]]{{.+}}, i8** null) // CK18-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK18-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] // CK18-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] - // CK18-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 - // CK18-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 - // CK18-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 - - // CK18-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to float** - // CK18-DAG: [[PC0:%.+]] = bitcast i8** [[P0]] to float*** - - // CK18-DAG: store float* [[F1:%.+]], float** [[BPC0]], - // CK18-DAG: store float** [[F_ADDR:%.+]], float*** [[PC0]], - // CK18-DAG: store i64 {{8|4}}, i64* [[S0]], - // CK18-DAG: [[F1]] = load float*, float** [[F_ADDR]], - - // CK18-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 - // CK18-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 - // CK18-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 + // CK18-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 + // CK18-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 + // CK18-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 // CK18-DAG: [[BPC1:%.+]] = bitcast i8** [[BP1]] to float*** // CK18-DAG: [[PC1:%.+]] = bitcast i8** [[P1]] to float** - // CK18-DAG: store float** [[F_ADDR]], float*** [[BPC1]], + // CK18-DAG: store float** [[F_ADDR:%.+]], float*** [[BPC1]], // CK18-DAG: store float* [[F2:%.+]], float** [[PC1]], // CK18-DAG: store i64 [[SIZE:%.+]], i64* [[S1]], // CK18-DAG: [[F2]] = load float*, float** [[F_ADDR]], diff --git a/openmp/libomptarget/src/omptarget.cpp b/openmp/libomptarget/src/omptarget.cpp --- a/openmp/libomptarget/src/omptarget.cpp +++ b/openmp/libomptarget/src/omptarget.cpp @@ -316,6 +316,11 @@ // may be considered a hack, we could revise the scheme in the future. bool UpdateRef = !(arg_types[i] & OMP_TGT_MAPTYPE_MEMBER_OF); if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) { + // Without OMP_TGT_MAPTYPE_MEMBER_OF, OMP_TGT_MAPTYPE_PTR_AND_OBJ can be + // used only if complex data with base pointer is mapped. No need to + // update ref counter for the base pointer in this case. + UpdateRef = + UpdateRef && (RTLs->RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY); DP("Has a pointer entry: \n"); // Base is address of pointer. // @@ -872,14 +877,9 @@ return OFFLOAD_FAIL; } } - } else if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) { - TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBase, sizeof(void *), IsLast, - false, IsHostPtr); - TgtBaseOffset = 0; // no offset for ptrs. - DP("Obtained target argument " DPxMOD " from host pointer " DPxMOD " to " - "object " DPxMOD "\n", DPxPTR(TgtPtrBegin), DPxPTR(HstPtrBase), - DPxPTR(HstPtrBase)); } else { + if (arg_types[i] & OMP_TGT_MAPTYPE_PTR_AND_OBJ) + HstPtrBase = *reinterpret_cast(HstPtrBase); TgtPtrBegin = Device.getTgtPtrBegin(HstPtrBegin, arg_sizes[i], IsLast, false, IsHostPtr); TgtBaseOffset = (intptr_t)HstPtrBase - (intptr_t)HstPtrBegin; diff --git a/openmp/libomptarget/test/env/base_ptr_ref_count.c b/openmp/libomptarget/test/env/base_ptr_ref_count.c new file mode 100644 --- /dev/null +++ b/openmp/libomptarget/test/env/base_ptr_ref_count.c @@ -0,0 +1,51 @@ +// RUN: %libomptarget-compile-aarch64-unknown-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-aarch64-unknown-linux-gnu 2>&1 | %fcheck-aarch64-unknown-linux-gnu +// RUN: %libomptarget-compile-powerpc64-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-powerpc64-ibm-linux-gnu 2>&1 | %fcheck-powerpc64-ibm-linux-gnu +// RUN: %libomptarget-compile-powerpc64le-ibm-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-powerpc64le-ibm-linux-gnu 2>&1 | %fcheck-powerpc64le-ibm-linux-gnu +// RUN: %libomptarget-compile-x86_64-pc-linux-gnu && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-x86_64-pc-linux-gnu 2>&1 | %fcheck-x86_64-pc-linux-gnu +// RUN: %libomptarget-compile-nvptx64-nvidia-cuda && env LIBOMPTARGET_DEBUG=1 %libomptarget-run-nvptx64-nvidia-cuda 2>&1 | %fcheck-nvptx64-nvidia-cuda +// REQUIRES: libomptarget-debug + +#include + +int *allocate(size_t n) { + int *ptr = malloc(sizeof(int) * n); +#pragma omp target enter data map(to : ptr[:n]) + return ptr; +} + +void deallocate(int *ptr, size_t n) { +#pragma omp target exit data map(delete : ptr[:n]) + free(ptr); +} + +#pragma omp declare target +int *cnt; +void foo() { + ++(*cnt); +} +#pragma omp end declare target + +int main(void) { + int *A = allocate(10); + int *V = allocate(10); + deallocate(A, 10); + deallocate(V, 10); +// CHECK-NOT: RefCount=2 + cnt = malloc(sizeof(int)); + *cnt = 0; +#pragma omp target map(cnt[:1]) + foo(); + printf("Cnt = %d.\n", *cnt); +// CHECK: Cnt = 1. + *cnt = 0; +#pragma omp target data map(cnt[:1]) +#pragma omp target + foo(); + printf("Cnt = %d.\n", *cnt); +// CHECK: Cnt = 1. + free(cnt); + + return 0; +} + +