From: Alexey Bataev Date: Tue, 26 Sep 2017 13:47:31 +0000 (+0000) Subject: [OPENMP] Generate implicit map|firstprivate clauses for target-based X-Git-Tag: llvmorg-6.0.0-rc1~7180 X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=f47c4b41840975f33153d6a817f8f45c4169bc74;p=platform%2Fupstream%2Fllvm.git [OPENMP] Generate implicit map|firstprivate clauses for target-based directives. If the variable is used in the target-based region but is not found in any private|mapping clause, then generate implicit firstprivate|map clauses for these implicitly mapped variables. llvm-svn: 314205 --- diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index 42b7d16..fb2ce6c 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -862,18 +862,7 @@ static void EmitOMPAggregateInit(CodeGenFunction &CGF, Address DestAddr, } LValue ReductionCodeGen::emitSharedLValue(CodeGenFunction &CGF, const Expr *E) { - if (const auto *OASE = dyn_cast(E)) - return CGF.EmitOMPArraySectionExpr(OASE); - if (const auto *ASE = dyn_cast(E)) - return CGF.EmitLValue(ASE); - auto *OrigVD = cast(cast(E)->getDecl()); - DeclRefExpr DRE(const_cast(OrigVD), - CGF.CapturedStmtInfo && - CGF.CapturedStmtInfo->lookup(OrigVD) != nullptr, - E->getType(), VK_LValue, E->getExprLoc()); - // Store the address of the original variable associated with the LHS - // implicit variable. - return CGF.EmitLValue(&DRE); + return CGF.EmitOMPSharedLValue(E); } LValue ReductionCodeGen::emitSharedLValueUB(CodeGenFunction &CGF, @@ -5978,6 +5967,8 @@ public: OMP_MAP_PRIVATE_PTR = 0x80, /// \brief Pass the element to the device by value. OMP_MAP_PRIVATE_VAL = 0x100, + /// Implicit map + OMP_MAP_IMPLICIT = 0x200, }; /// Class that associates information with a base pointer to be passed to the @@ -6148,7 +6139,7 @@ private: OMPClauseMappableExprCommon::MappableExprComponentListRef Components, MapBaseValuesArrayTy &BasePointers, MapValuesArrayTy &Pointers, MapValuesArrayTy &Sizes, MapFlagsArrayTy &Types, - bool IsFirstComponentList) const { + bool IsFirstComponentList, bool IsImplicit) const { // The following summarizes what has to be generated for each map and the // types bellow. The generated information is expressed in this order: @@ -6283,8 +6274,7 @@ private: } else { // The base is the reference to the variable. // BP = &Var. - BP = CGF.EmitLValue(cast(I->getAssociatedExpression())) - .getPointer(); + BP = CGF.EmitOMPSharedLValue(I->getAssociatedExpression()).getPointer(); // If the variable is a pointer and is being dereferenced (i.e. is not // the last component), the base has to be the pointer itself, not its @@ -6303,6 +6293,7 @@ private: } } + unsigned DefaultFlags = IsImplicit ? OMP_MAP_IMPLICIT : 0; for (; I != CE; ++I) { auto Next = std::next(I); @@ -6337,7 +6328,8 @@ private: isa(Next->getAssociatedExpression())) && "Unexpected expression"); - auto *LB = CGF.EmitLValue(I->getAssociatedExpression()).getPointer(); + llvm::Value *LB = + CGF.EmitOMPSharedLValue(I->getAssociatedExpression()).getPointer(); auto *Size = getExprTypeSize(I->getAssociatedExpression()); // If we have a member expression and the current component is a @@ -6352,9 +6344,11 @@ private: BasePointers.push_back(BP); Pointers.push_back(RefAddr); Sizes.push_back(CGF.getTypeSize(CGF.getContext().VoidPtrTy)); - Types.push_back(getMapTypeBits( - /*MapType*/ OMPC_MAP_alloc, /*MapTypeModifier=*/OMPC_MAP_unknown, - !IsExpressionFirstInfo, IsCaptureFirstInfo)); + Types.push_back(DefaultFlags | + getMapTypeBits( + /*MapType*/ OMPC_MAP_alloc, + /*MapTypeModifier=*/OMPC_MAP_unknown, + !IsExpressionFirstInfo, IsCaptureFirstInfo)); IsExpressionFirstInfo = false; IsCaptureFirstInfo = false; // The reference will be the next base address. @@ -6369,9 +6363,9 @@ private: // 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). - Types.push_back(getMapTypeBits(MapType, MapTypeModifier, - !IsExpressionFirstInfo, - IsCaptureFirstInfo)); + Types.push_back(DefaultFlags | getMapTypeBits(MapType, MapTypeModifier, + !IsExpressionFirstInfo, + IsCaptureFirstInfo)); // If we have a final array section, we are done with this expression. if (IsFinalArraySection) @@ -6383,7 +6377,6 @@ private: IsExpressionFirstInfo = false; IsCaptureFirstInfo = false; - continue; } } } @@ -6445,20 +6438,19 @@ public: RPK_MemberReference, }; OMPClauseMappableExprCommon::MappableExprComponentListRef Components; - OpenMPMapClauseKind MapType; - OpenMPMapClauseKind MapTypeModifier; - ReturnPointerKind ReturnDevicePointer; + OpenMPMapClauseKind MapType = OMPC_MAP_unknown; + OpenMPMapClauseKind MapTypeModifier = OMPC_MAP_unknown; + ReturnPointerKind ReturnDevicePointer = RPK_None; + bool IsImplicit = false; - MapInfo() - : MapType(OMPC_MAP_unknown), MapTypeModifier(OMPC_MAP_unknown), - ReturnDevicePointer(RPK_None) {} + MapInfo() = default; MapInfo( OMPClauseMappableExprCommon::MappableExprComponentListRef Components, OpenMPMapClauseKind MapType, OpenMPMapClauseKind MapTypeModifier, - ReturnPointerKind ReturnDevicePointer) + ReturnPointerKind ReturnDevicePointer, bool IsImplicit) : Components(Components), MapType(MapType), MapTypeModifier(MapTypeModifier), - ReturnDevicePointer(ReturnDevicePointer) {} + ReturnDevicePointer(ReturnDevicePointer), IsImplicit(IsImplicit) {} }; // We have to process the component lists that relate with the same @@ -6472,25 +6464,29 @@ public: const ValueDecl *D, OMPClauseMappableExprCommon::MappableExprComponentListRef L, OpenMPMapClauseKind MapType, OpenMPMapClauseKind MapModifier, - MapInfo::ReturnPointerKind ReturnDevicePointer) { + MapInfo::ReturnPointerKind ReturnDevicePointer, bool IsImplicit) { const ValueDecl *VD = D ? cast(D->getCanonicalDecl()) : nullptr; - Info[VD].push_back({L, MapType, MapModifier, ReturnDevicePointer}); + Info[VD].emplace_back(L, MapType, MapModifier, ReturnDevicePointer, + IsImplicit); }; // FIXME: MSVC 2013 seems to require this-> to find member CurDir. for (auto *C : this->CurDir.getClausesOfKind()) - for (auto L : C->component_lists()) + for (auto L : C->component_lists()) { InfoGen(L.first, L.second, C->getMapType(), C->getMapTypeModifier(), - MapInfo::RPK_None); + MapInfo::RPK_None, C->isImplicit()); + } for (auto *C : this->CurDir.getClausesOfKind()) - for (auto L : C->component_lists()) + for (auto L : C->component_lists()) { InfoGen(L.first, L.second, OMPC_MAP_to, OMPC_MAP_unknown, - MapInfo::RPK_None); + MapInfo::RPK_None, C->isImplicit()); + } for (auto *C : this->CurDir.getClausesOfKind()) - for (auto L : C->component_lists()) + for (auto L : C->component_lists()) { InfoGen(L.first, L.second, OMPC_MAP_from, OMPC_MAP_unknown, - MapInfo::RPK_None); + MapInfo::RPK_None, C->isImplicit()); + } // Look at the use_device_ptr clause information and mark the existing map // entries as such. If there is no map information for an entry in the @@ -6551,9 +6547,9 @@ public: // Remember the current base pointer index. unsigned CurrentBasePointersIdx = BasePointers.size(); // FIXME: MSVC 2013 seems to require this-> to find the member method. - this->generateInfoForComponentList(L.MapType, L.MapTypeModifier, - L.Components, BasePointers, Pointers, - Sizes, Types, IsFirstComponentList); + this->generateInfoForComponentList( + L.MapType, L.MapTypeModifier, L.Components, BasePointers, Pointers, + Sizes, Types, IsFirstComponentList, L.IsImplicit); // If this entry relates with a device pointer, set the relevant // declaration and add the 'return pointer' flag. @@ -6617,7 +6613,8 @@ public: for (auto L : It->second) { generateInfoForComponentList( /*MapType=*/OMPC_MAP_to, /*MapTypeModifier=*/OMPC_MAP_unknown, L, - BasePointers, Pointers, Sizes, Types, IsFirstComponentList); + BasePointers, Pointers, Sizes, Types, IsFirstComponentList, + /*IsImplicit=*/false); IsFirstComponentList = false; } return; @@ -6637,9 +6634,9 @@ public: "We got information for the wrong declaration??"); assert(!L.second.empty() && "Not expecting declaration with no component lists."); - generateInfoForComponentList(C->getMapType(), C->getMapTypeModifier(), - L.second, BasePointers, Pointers, Sizes, - Types, IsFirstComponentList); + generateInfoForComponentList( + C->getMapType(), C->getMapTypeModifier(), L.second, BasePointers, + Pointers, Sizes, Types, IsFirstComponentList, C->isImplicit()); IsFirstComponentList = false; } diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp index 4dcdaa2..20f1535 100644 --- a/clang/lib/CodeGen/CGStmtOpenMP.cpp +++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp @@ -138,6 +138,22 @@ public: } // namespace +LValue CodeGenFunction::EmitOMPSharedLValue(const Expr *E) { + if (auto *OrigDRE = dyn_cast(E)) { + if (auto *OrigVD = dyn_cast(OrigDRE->getDecl())) { + OrigVD = OrigVD->getCanonicalDecl(); + bool IsCaptured = + LambdaCaptureFields.lookup(OrigVD) || + (CapturedStmtInfo && CapturedStmtInfo->lookup(OrigVD)) || + (CurCodeDecl && isa(CurCodeDecl)); + DeclRefExpr DRE(const_cast(OrigVD), IsCaptured, + OrigDRE->getType(), VK_LValue, OrigDRE->getExprLoc()); + return EmitLValue(&DRE); + } + } + return EmitLValue(E); +} + llvm::Value *CodeGenFunction::getTypeSize(QualType Ty) { auto &C = getContext(); llvm::Value *Size = nullptr; diff --git a/clang/lib/CodeGen/CodeGenFunction.h b/clang/lib/CodeGen/CodeGenFunction.h index d77cea4..635df4bc 100644 --- a/clang/lib/CodeGen/CodeGenFunction.h +++ b/clang/lib/CodeGen/CodeGenFunction.h @@ -2886,6 +2886,9 @@ public: const CodeGenLoopBoundsTy &CodeGenLoopBounds, const CodeGenDispatchBoundsTy &CGDispatchBounds); + /// Emits the lvalue for the expression with possibly captured variable. + LValue EmitOMPSharedLValue(const Expr *E); + private: /// Helpers for blocks llvm::Value *EmitBlockLiteral(const CGBlockInfo &Info); diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp index e4a57db..735e827 100644 --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -35,6 +35,11 @@ using namespace clang; // Stack of data-sharing attributes for variables //===----------------------------------------------------------------------===// +static Expr *CheckMapClauseExpressionBase( + Sema &SemaRef, Expr *E, + OMPClauseMappableExprCommon::MappableExprComponentList &CurComponents, + OpenMPClauseKind CKind); + namespace { /// \brief Default data sharing attributes, which can be applied to directive. enum DefaultDataSharingAttributes { @@ -1787,6 +1792,7 @@ class DSAAttrChecker : public StmtVisitor { bool ErrorFound; CapturedStmt *CS; llvm::SmallVector ImplicitFirstprivate; + llvm::SmallVector ImplicitMap; llvm::DenseMap VarsWithInheritedDSA; llvm::DenseSet ImplicitDeclarations; @@ -1819,6 +1825,50 @@ public: return; } + if (isOpenMPTargetExecutionDirective(DKind) && + !Stack->isLoopControlVariable(VD).first) { + if (!Stack->checkMappableExprComponentListsForDecl( + VD, /*CurrentRegionOnly=*/true, + [](OMPClauseMappableExprCommon::MappableExprComponentListRef + StackComponents, + OpenMPClauseKind) { + // Variable is used if it has been marked as an array, array + // section or the variable iself. + return StackComponents.size() == 1 || + std::all_of( + std::next(StackComponents.rbegin()), + StackComponents.rend(), + [](const OMPClauseMappableExprCommon:: + MappableComponent &MC) { + return MC.getAssociatedDeclaration() == + nullptr && + (isa( + MC.getAssociatedExpression()) || + isa( + MC.getAssociatedExpression())); + }); + })) { + bool CapturedByCopy = false; + // By default lambdas are captured as firstprivates. + if (const auto *RD = + VD->getType().getNonReferenceType()->getAsCXXRecordDecl()) + if (RD->isLambda()) + CapturedByCopy = true; + CapturedByCopy = + CapturedByCopy || + llvm::any_of( + CS->captures(), [VD](const CapturedStmt::Capture &I) { + return I.capturesVariableByCopy() && + I.getCapturedVar()->getCanonicalDecl() == VD; + }); + if (CapturedByCopy) + ImplicitFirstprivate.emplace_back(E); + else + ImplicitMap.emplace_back(E); + return; + } + } + // OpenMP [2.9.3.6, Restrictions, p.2] // A list item that appears in a reduction clause of the innermost // enclosing worksharing or parallel construct may not be accessed in an @@ -1848,40 +1898,104 @@ public: if (E->isTypeDependent() || E->isValueDependent() || E->containsUnexpandedParameterPack() || E->isInstantiationDependent()) return; + auto *FD = dyn_cast(E->getMemberDecl()); + if (!FD) + return; + OpenMPDirectiveKind DKind = Stack->getCurrentDirective(); if (isa(E->getBase()->IgnoreParens())) { - if (auto *FD = dyn_cast(E->getMemberDecl())) { - auto DVar = Stack->getTopDSA(FD, false); - // Check if the variable has explicit DSA set and stop analysis if it - // so. - if (DVar.RefExpr || !ImplicitDeclarations.insert(FD).second) - return; + auto DVar = Stack->getTopDSA(FD, false); + // Check if the variable has explicit DSA set and stop analysis if it + // so. + if (DVar.RefExpr || !ImplicitDeclarations.insert(FD).second) + return; - auto ELoc = E->getExprLoc(); - auto DKind = Stack->getCurrentDirective(); - // OpenMP [2.9.3.6, Restrictions, p.2] - // A list item that appears in a reduction clause of the innermost - // enclosing worksharing or parallel construct may not be accessed in - // an explicit task. - DVar = Stack->hasInnermostDSA( - FD, [](OpenMPClauseKind C) -> bool { return C == OMPC_reduction; }, - [](OpenMPDirectiveKind K) -> bool { - return isOpenMPParallelDirective(K) || - isOpenMPWorksharingDirective(K) || - isOpenMPTeamsDirective(K); - }, - /*FromParent=*/true); - if (isOpenMPTaskingDirective(DKind) && DVar.CKind == OMPC_reduction) { - ErrorFound = true; - SemaRef.Diag(ELoc, diag::err_omp_reduction_in_task); - ReportOriginalDSA(SemaRef, Stack, FD, DVar); + if (isOpenMPTargetExecutionDirective(DKind) && + !Stack->isLoopControlVariable(FD).first && + !Stack->checkMappableExprComponentListsForDecl( + FD, /*CurrentRegionOnly=*/true, + [](OMPClauseMappableExprCommon::MappableExprComponentListRef + StackComponents, + OpenMPClauseKind) { + return isa( + cast( + StackComponents.back().getAssociatedExpression()) + ->getBase() + ->IgnoreParens()); + })) { + // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C/C++, p.3] + // A bit-field cannot appear in a map clause. + // + if (FD->isBitField()) { + SemaRef.Diag(E->getMemberLoc(), + diag::err_omp_bit_fields_forbidden_in_clause) + << E->getSourceRange() << getOpenMPClauseName(OMPC_map); return; } + ImplicitMap.emplace_back(E); + return; + } - // Define implicit data-sharing attributes for task. - DVar = Stack->getImplicitDSA(FD, false); - if (isOpenMPTaskingDirective(DKind) && DVar.CKind != OMPC_shared && - !Stack->isLoopControlVariable(FD).first) - ImplicitFirstprivate.push_back(E); + auto ELoc = E->getExprLoc(); + // OpenMP [2.9.3.6, Restrictions, p.2] + // A list item that appears in a reduction clause of the innermost + // enclosing worksharing or parallel construct may not be accessed in + // an explicit task. + DVar = Stack->hasInnermostDSA( + FD, [](OpenMPClauseKind C) -> bool { return C == OMPC_reduction; }, + [](OpenMPDirectiveKind K) -> bool { + return isOpenMPParallelDirective(K) || + isOpenMPWorksharingDirective(K) || isOpenMPTeamsDirective(K); + }, + /*FromParent=*/true); + if (isOpenMPTaskingDirective(DKind) && DVar.CKind == OMPC_reduction) { + ErrorFound = true; + SemaRef.Diag(ELoc, diag::err_omp_reduction_in_task); + ReportOriginalDSA(SemaRef, Stack, FD, DVar); + return; + } + + // Define implicit data-sharing attributes for task. + DVar = Stack->getImplicitDSA(FD, false); + if (isOpenMPTaskingDirective(DKind) && DVar.CKind != OMPC_shared && + !Stack->isLoopControlVariable(FD).first) + ImplicitFirstprivate.push_back(E); + return; + } + if (isOpenMPTargetExecutionDirective(DKind) && !FD->isBitField()) { + OMPClauseMappableExprCommon::MappableExprComponentList CurComponents; + CheckMapClauseExpressionBase(SemaRef, E, CurComponents, OMPC_map); + auto *VD = cast( + CurComponents.back().getAssociatedDeclaration()->getCanonicalDecl()); + if (!Stack->checkMappableExprComponentListsForDecl( + VD, /*CurrentRegionOnly=*/true, + [&CurComponents]( + OMPClauseMappableExprCommon::MappableExprComponentListRef + StackComponents, + OpenMPClauseKind) { + if (CurComponents.size() < StackComponents.size()) + return false; + auto CCI = CurComponents.rbegin(); + for (const auto &SC : llvm::reverse(StackComponents)) { + // Do both expressions have the same kind? + if (CCI->getAssociatedExpression()->getStmtClass() != + SC.getAssociatedExpression()->getStmtClass()) + if (!(isa( + SC.getAssociatedExpression()) && + isa( + CCI->getAssociatedExpression()))) + return false; + + Decl *CCD = CCI->getAssociatedDeclaration(); + Decl *SCD = SC.getAssociatedDeclaration(); + CCD = CCD ? CCD->getCanonicalDecl() : nullptr; + SCD = SCD ? SCD->getCanonicalDecl() : nullptr; + if (SCD != CCD) + return false; + std::advance(CCI, 1); + } + return true; + })) { + Visit(E->getBase()); } } else Visit(E->getBase()); @@ -1889,12 +2003,16 @@ public: void VisitOMPExecutableDirective(OMPExecutableDirective *S) { for (auto *C : S->clauses()) { // Skip analysis of arguments of implicitly defined firstprivate clause - // for task directives. - if (C && (!isa(C) || C->getLocStart().isValid())) + // for task|target directives. + // Skip analysis of arguments of implicitly defined map clause for target + // directives. + if (C && !((isa(C) || isa(C)) && + C->isImplicit())) { for (auto *CC : C->children()) { if (CC) Visit(CC); } + } } } void VisitStmt(Stmt *S) { @@ -1905,7 +2023,10 @@ public: } bool isErrorFound() { return ErrorFound; } - ArrayRef getImplicitFirstprivate() { return ImplicitFirstprivate; } + ArrayRef getImplicitFirstprivate() const { + return ImplicitFirstprivate; + } + ArrayRef getImplicitMap() const { return ImplicitMap; } llvm::DenseMap &getVarsWithInheritedDSA() { return VarsWithInheritedDSA; } @@ -2638,6 +2759,8 @@ StmtResult Sema::ActOnOpenMPExecutableDirective( SmallVector ImplicitFirstprivates( DSAChecker.getImplicitFirstprivate().begin(), DSAChecker.getImplicitFirstprivate().end()); + SmallVector ImplicitMaps(DSAChecker.getImplicitMap().begin(), + DSAChecker.getImplicitMap().end()); // Mark taskgroup task_reduction descriptors as implicitly firstprivate. for (auto *C : Clauses) { if (auto *IRC = dyn_cast(C)) { @@ -2656,6 +2779,17 @@ StmtResult Sema::ActOnOpenMPExecutableDirective( } else ErrorFound = true; } + if (!ImplicitMaps.empty()) { + if (OMPClause *Implicit = ActOnOpenMPMapClause( + OMPC_MAP_unknown, OMPC_MAP_tofrom, /*IsMapTypeImplicit=*/true, + SourceLocation(), SourceLocation(), ImplicitMaps, + SourceLocation(), SourceLocation(), SourceLocation())) { + ClausesWithImplicit.emplace_back(Implicit); + ErrorFound |= + cast(Implicit)->varlist_size() != ImplicitMaps.size(); + } else + ErrorFound = true; + } } llvm::SmallVector AllowedNameModifiers; diff --git a/clang/test/OpenMP/target_codegen.cpp b/clang/test/OpenMP/target_codegen.cpp index 21c82d1..6ed4e59 100644 --- a/clang/test/OpenMP/target_codegen.cpp +++ b/clang/test/OpenMP/target_codegen.cpp @@ -38,12 +38,12 @@ // CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [1 x i32] [i32 288] // CHECK-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [2 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2] // CHECK-DAG: [[MAPT3:@.+]] = private unnamed_addr constant [2 x i32] [i32 288, i32 288] -// CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [9 x i32] [i32 288, i32 35, i32 288, i32 35, i32 35, i32 288, i32 288, i32 35, i32 35] +// CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [9 x i32] [i32 288, i32 547, i32 288, i32 547, i32 547, i32 288, i32 288, i32 547, i32 547] // CHECK-DAG: [[SIZET5:@.+]] = private unnamed_addr constant [3 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2, i[[SZ]] 40] -// CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i32] [i32 288, i32 288, i32 35] +// CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i32] [i32 288, i32 288, i32 547] // CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [4 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2, i[[SZ]] 1, i[[SZ]] 40] -// CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i32] [i32 288, i32 288, i32 288, i32 35] -// CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [5 x i32] [i32 35, i32 288, i32 288, i32 288, i32 35] +// CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i32] [i32 288, i32 288, i32 288, i32 547] +// CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [5 x i32] [i32 547, i32 288, i32 288, i32 288, i32 547] // CHECK-DAG: @{{.*}} = private constant i8 0 // CHECK-DAG: @{{.*}} = private constant i8 0 // CHECK-DAG: @{{.*}} = private constant i8 0 @@ -510,9 +510,9 @@ int bar(int n){ // CHECK-DAG: store i[[SZ]] 4, i[[SZ]]* [[SADDR1]] // CHECK-DAG: [[CBPADDR0:%.+]] = bitcast i8** [[BPADDR0]] to [[S1]]** -// CHECK-DAG: [[CPADDR0:%.+]] = bitcast i8** [[PADDR0]] to [[S1]]** +// CHECK-DAG: [[CPADDR0:%.+]] = bitcast i8** [[PADDR0]] to double** // CHECK-DAG: store [[S1]]* %{{.+}}, [[S1]]** [[CBPADDR0]] -// CHECK-DAG: store [[S1]]* %{{.+}}, [[S1]]** [[CPADDR0]] +// CHECK-DAG: store double* %{{.+}}, double** [[CPADDR0]] // CHECK-DAG: store i[[SZ]] 8, i[[SZ]]* [[SADDR0]] // CHECK-DAG: [[CBPADDR4:%.+]] = bitcast i8** [[BPADDR4]] to i16** diff --git a/clang/test/OpenMP/target_firstprivate_codegen.cpp b/clang/test/OpenMP/target_firstprivate_codegen.cpp index 1fb0ef9..64f5225 100644 --- a/clang/test/OpenMP/target_firstprivate_codegen.cpp +++ b/clang/test/OpenMP/target_firstprivate_codegen.cpp @@ -37,7 +37,7 @@ struct TT{ // CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [9 x i32] [i32 288, i32 161, i32 288, i32 161, i32 161, i32 288, i32 288, i32 161, i32 161] // CHECK-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [1 x i{{32|64}}] zeroinitializer // CHECK-DAG: [[MAPT3:@.+]] = private unnamed_addr constant [1 x i32] [i32 32] -// CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [5 x i32] [i32 35, i32 288, i32 288, i32 288, i32 161] +// CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [5 x i32] [i32 547, i32 288, i32 288, i32 288, i32 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 i32] [i32 288, i32 288, i32 161] // CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [2 x i{{32|64}}] [i[[SZ]] 4, i[[SZ]] 40] diff --git a/clang/test/OpenMP/target_map_codegen.cpp b/clang/test/OpenMP/target_map_codegen.cpp index 820815c..460a02e 100644 --- a/clang/test/OpenMP/target_map_codegen.cpp +++ b/clang/test/OpenMP/target_map_codegen.cpp @@ -360,8 +360,8 @@ void implicit_maps_host_global (int a){ // CK7-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 8] // Map types: OMP_MAP_PRIVATE_VAL | OMP_MAP_IS_FIRST = 288 // CK7-64-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288] -// Map types: OMP_MAP_TO | OMP_MAP_IS_FIRST = 33 -// CK7-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 33] +// Map types: OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_IS_FIRST | OMP_MAP_IMPLICIT = 547 +// CK7-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 547] // CK7-LABEL: implicit_maps_double void implicit_maps_double (int a){ @@ -461,8 +461,8 @@ void implicit_maps_float (int a){ #ifdef CK9 // CK9-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 16] -// Map types: OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST = 35 -// CK9-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 35] +// Map types: OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST + OMP_MAP_IMPLICIT = 547 +// CK9-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 547] // CK9-LABEL: implicit_maps_array void implicit_maps_array (int a){ @@ -544,8 +544,8 @@ void implicit_maps_pointer (){ #ifdef CK11 // CK11-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 16] -// Map types: OMP_MAP_TO + OMP_MAP_IS_FIRST = 33 -// CK11-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 33] +// Map types: OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST + OMP_MAP_IMPLICIT = 547 +// CK11-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 547] // CK11-LABEL: implicit_maps_double_complex void implicit_maps_double_complex (int a){ @@ -589,8 +589,8 @@ void implicit_maps_double_complex (int a){ // CK12-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 8] // Map types: OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288 // CK12-64-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 288] -// Map types: OMP_MAP_TO + OMP_MAP_IS_FIRST = 33 -// CK12-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 33] +// Map types: OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST + OMP_MAP_IMPLICIT = 547 +// CK12-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 547] // CK12-LABEL: implicit_maps_float_complex void implicit_maps_float_complex (int a){ @@ -648,8 +648,8 @@ void implicit_maps_float_complex (int a){ // Map types: // - OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288 (vla size) // - OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288 (vla size) -// - OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST = 35 -// CK13-DAG: [[TYPES:@.+]] = {{.+}}constant [3 x i32] [i32 288, i32 288, i32 35] +// - OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST + OMP_IMPICIT_MAP = 547 +// CK13-DAG: [[TYPES:@.+]] = {{.+}}constant [3 x i32] [i32 288, i32 288, i32 547] // CK13-LABEL: implicit_maps_variable_length_array void implicit_maps_variable_length_array (int a){ @@ -717,11 +717,12 @@ void implicit_maps_variable_length_array (int a){ #ifdef CK14 // CK14-DAG: [[ST:%.+]] = type { i32, double } -// CK14-DAG: [[SIZES:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{16|12}}, i{{64|32}} 4] +// CK14-DAG: [[SIZES:@.+]] = {{.+}}constant [3 x i[[sz:64|32]]] [i{{64|32}} 4, i{{64|32}} 8, i{{64|32}} 4] // Map types: -// - OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST = 35 +// - OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST + OMP_IMPLICIT_MAP = 547 +// - OMP_MAP_TO + OMP_MAP_FROM + OMP_IMPLICIT_MAP = 515 // - OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288 -// CK14-DAG: [[TYPES:@.+]] = {{.+}}constant [2 x i32] [i32 35, i32 288] +// CK14-DAG: [[TYPES:@.+]] = {{.+}}constant [3 x i32] [i32 547, i32 515, i32 288] class SSS { public: @@ -744,19 +745,26 @@ void implicit_maps_class (int a){ SSS sss(a, (double)a); // CK14: define {{.*}}void @{{.+}}foo{{.+}}([[ST]]* {{[^,]+}}, i32 {{[^,]+}}) - // CK14-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 2, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}) + // CK14-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 3, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}) // CK14-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 // CK14-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 // CK14-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 // CK14-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 // CK14-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]** - // CK14-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[ST]]** + // CK14-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** // CK14-DAG: store [[ST]]* [[DECL:%.+]], [[ST]]** [[CBP0]] - // CK14-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CP0]] + // CK14-DAG: store i32* %{{.+}}, i32** [[CP0]] // CK14-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1 // CK14-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1 + // CK14-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]** + // CK14-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double** + // CK14-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CBP1]] + // CK14-DAG: store double* %{{.+}}, double** [[CP1]] + + // CK14-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 2 + // CK14-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 2 // CK14-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i[[sz]]* // CK14-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i[[sz]]* // CK14-DAG: store i[[sz]] [[VAL:%.+]], i[[sz]]* [[CBP1]] @@ -791,17 +799,17 @@ void implicit_maps_class (int a){ #ifdef CK15 // CK15: [[ST:%.+]] = type { i32, double, i32* } -// CK15: [[SIZES:@.+]] = {{.+}}constant [2 x i[[sz:64|32]]] [i{{64|32}} {{24|16}}, i{{64|32}} 4] +// CK15: [[SIZES:@.+]] = {{.+}}constant [5 x i[[sz:64|32]]] [i{{64|32}} 4, i{{64|32}} 8, i{{64|32}} {{8|4}}, i{{64|32}} 4, i{{64|32}} 4] // Map types: -// - OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST = 35 +// - OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST + OMP_MAP_IMPLICIT = 547 // - OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288 -// CK15: [[TYPES:@.+]] = {{.+}}constant [2 x i32] [i32 35, i32 288] +// CK15: [[TYPES:@.+]] = {{.+}}constant [5 x i32] [i32 547, i32 515, i32 512, i32 531, i32 288] -// CK15: [[SIZES2:@.+]] = {{.+}}constant [2 x i[[sz]]] [i{{64|32}} {{24|16}}, i{{64|32}} 4] +// CK15: [[SIZES2:@.+]] = {{.+}}constant [5 x i[[sz]]] [i{{64|32}} 4, i{{64|32}} 8, i{{64|32}} {{8|4}}, i{{64|32}} 4, i{{64|32}} 4] // Map types: // - OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST = 35 // - OMP_MAP_PRIVATE_VAL + OMP_MAP_IS_FIRST = 288 -// CK15: [[TYPES2:@.+]] = {{.+}}constant [2 x i32] [i32 35, i32 288] +// CK15: [[TYPES2:@.+]] = {{.+}}constant [5 x i32] [i32 547, i32 515, i32 512, i32 531, i32 288] template class SSST { @@ -836,23 +844,44 @@ void implicit_maps_templated_class (int a){ SSST<123> ssst(a, (double)a, a); // CK15: define {{.*}}void @{{.+}}foo{{.+}}([[ST]]* {{[^,]+}}, i32 {{[^,]+}}) - // CK15-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 2, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}) + // CK15-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 5, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES]]{{.+}}, {{.+}}[[TYPES]]{{.+}}) // CK15-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 // CK15-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 // CK15-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 // CK15-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 // CK15-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]** - // CK15-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[ST]]** + // CK15-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** // CK15-DAG: store [[ST]]* [[DECL:%.+]], [[ST]]** [[CBP0]] - // CK15-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CP0]] + // CK15-DAG: store i32* %{{.+}}, i32** [[CP0]] // CK15-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1 // CK15-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1 - // CK15-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i[[sz]]* - // CK15-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i[[sz]]* - // CK15-DAG: store i[[sz]] [[VAL:%.+]], i[[sz]]* [[CBP1]] - // CK15-DAG: store i[[sz]] [[VAL]], i[[sz]]* [[CP1]] + // CK15-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]** + // CK15-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double** + // CK15-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CBP1]] + // CK15-DAG: store double* %{{.+}}, double** [[CP1]] + + // CK15-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 2 + // CK15-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 2 + // CK15-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [[ST]]** + // CK15-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to i32*** + // CK15-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CBP2]] + // CK15-DAG: store i32** %{{.+}}, i32*** [[CP2]] + + // CK15-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 3 + // CK15-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 3 + // CK15-DAG: [[CBP3:%.+]] = bitcast i8** [[BP3]] to i32*** + // CK15-DAG: [[CP3:%.+]] = bitcast i8** [[P3]] to i32** + // CK15-DAG: store i32** %{{.+}}, i32*** [[CBP3]] + // CK15-DAG: store i32* %{{.+}}, i32** [[CP3]] + + // CK15-DAG: [[BP4:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 4 + // CK15-DAG: [[P4:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 4 + // CK15-DAG: [[CBP4:%.+]] = bitcast i8** [[BP4]] to i[[sz]]* + // CK15-DAG: [[CP4:%.+]] = bitcast i8** [[P4]] to i[[sz]]* + // CK15-DAG: store i[[sz]] [[VAL:%.+]], i[[sz]]* [[CBP4]] + // CK15-DAG: store i[[sz]] [[VAL]], i[[sz]]* [[CP4]] // CK15-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]], // CK15-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32* // CK15-64-DAG: store i32 {{.+}}, i32* [[CADDR]], @@ -861,23 +890,44 @@ void implicit_maps_templated_class (int a){ ssst.foo(456); // CK15: define {{.*}}void @{{.+}}bar{{.+}}([[ST]]* {{[^,]+}}, i32 {{[^,]+}}) - // CK15-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 2, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES2]]{{.+}}, {{.+}}[[TYPES2]]{{.+}}) + // CK15-DAG: call i32 @__tgt_target(i32 {{.+}}, i8* {{.+}}, i32 5, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[SIZES2]]{{.+}}, {{.+}}[[TYPES2]]{{.+}}) // CK15-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 // CK15-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 // CK15-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 // CK15-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 // CK15-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]** - // CK15-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[ST]]** - // CK15-DAG: store [[ST]]* [[DECL:%.+]], [[ST]]** [[CBP0]] - // CK15-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CP0]] + // CK15-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** + // CK15-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CBP0]] + // CK15-DAG: store i32* %{{.+}}, i32** [[CP0]] // CK15-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1 // CK15-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1 - // CK15-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i[[sz]]* - // CK15-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i[[sz]]* - // CK15-DAG: store i[[sz]] [[VAL:%.+]], i[[sz]]* [[CBP1]] - // CK15-DAG: store i[[sz]] [[VAL]], i[[sz]]* [[CP1]] + // CK15-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]** + // CK15-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double** + // CK15-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CBP1]] + // CK15-DAG: store double* %{{.+}}, double** [[CP1]] + + // CK15-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 2 + // CK15-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 2 + // CK15-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [[ST]]** + // CK15-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to i32*** + // CK15-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CBP2]] + // CK15-DAG: store i32** %{{.+}}, i32*** [[CP2]] + + // CK15-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 3 + // CK15-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 3 + // CK15-DAG: [[CBP3:%.+]] = bitcast i8** [[BP3]] to i32*** + // CK15-DAG: [[CP3:%.+]] = bitcast i8** [[P3]] to i32** + // CK15-DAG: store i32** %{{.+}}, i32*** [[CBP3]] + // CK15-DAG: store i32* %{{.+}}, i32** [[CP3]] + + // CK15-DAG: [[BP4:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 4 + // CK15-DAG: [[P4:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 4 + // CK15-DAG: [[CBP4:%.+]] = bitcast i8** [[BP4]] to i[[sz]]* + // CK15-DAG: [[CP4:%.+]] = bitcast i8** [[P4]] to i[[sz]]* + // CK15-DAG: store i[[sz]] [[VAL:%.+]], i[[sz]]* [[CBP4]] + // CK15-DAG: store i[[sz]] [[VAL]], i[[sz]]* [[CP4]] // CK15-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]], // CK15-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to i32* // CK15-64-DAG: store i32 {{.+}}, i32* [[CADDR]], @@ -973,8 +1023,8 @@ void implicit_maps_templated_function (int a){ // CK17-DAG: [[ST:%.+]] = type { i32, double } // CK17-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} {{16|12}}] -// Map types: OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST = 35 -// CK17-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 35] +// Map types: OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_IS_FIRST + OMP_MAP_IMPLICIT = 547 +// CK17-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i32] [i32 547] class SSS { public: @@ -3331,39 +3381,6 @@ struct SC{ // CK24: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4] // CK24: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i32] [i32 35] -// CK24: [[SIZE02:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] {{56|48}}] -// CK24: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i32] [i32 35] - -// CK24: [[SIZE03:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] -// CK24: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i32] [i32 35] - -// CK24: [[SIZE04:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 20] -// CK24: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i32] [i32 35] - -// CK24: [[SIZE05:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{3560|2880}}] -// CK24: [[MTYPE05:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19] - -// CK24: [[SIZE06:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] -// CK24: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i32] [i32 35] - -// CK24: [[SIZE07:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4] -// CK24: [[MTYPE07:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19] - -// CK24: [[SIZE08:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4] -// CK24: [[MTYPE08:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19] - -// CK24: [[SIZE09:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] 4] -// CK24: [[MTYPE09:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19] - -// CK24: [[SIZE10:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 8] -// CK24: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i32] [i32 35] - -// CK24: [[SIZE11:@.+]] = private {{.*}}constant [2 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}] -// CK24: [[MTYPE11:@.+]] = private {{.*}}constant [2 x i32] [i32 35, i32 19] - -// CK24: [[SIZE12:@.+]] = private {{.*}}constant [4 x i[[Z]]] [i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] {{8|4}}, i[[Z]] 4] -// CK24: [[MTYPE12:@.+]] = private {{.*}}constant [4 x i32] [i32 35, i32 19, i32 19, i32 19] - // CK24: [[SIZE13:@.+]] = private {{.*}}constant [1 x i[[Z]]] [i[[Z]] 4] // CK24: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i32] [i32 35] @@ -3422,298 +3439,6 @@ int explicit_maps_struct_fields(int a){ #pragma omp target map(s.a) { s.a++; } -// Region 02 -// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}) -// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] -// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] - -// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 -// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], 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: [[SEC0]] = getelementptr {{.*}}[[SB]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 - -// CK24: call void [[CALL02:@.+]]([[SC]]* {{[^,]+}}) -#pragma omp target map(s.s.s) - { s.a++; } - -// Region 03 -// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE03]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}) -// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] -// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] - -// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 -// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], 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: [[SEC0]] = getelementptr {{.*}}[[SA]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0 -// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SB]]* [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[SEC000]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 - -// CK24: call void [[CALL03:@.+]]([[SC]]* {{[^,]+}}) -#pragma omp target map(s.s.s.a) - { s.a++; } - -// Region 04 -// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE04]]{{.+}}) -// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] -// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] - -// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 -// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], 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: [[SEC0]] = getelementptr {{.*}}[10 x i32]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0 -// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 3 - -// CK24: call void [[CALL04:@.+]]([[SC]]* {{[^,]+}}) -#pragma omp target map(s.b[:5]) - { s.a++; } - -// Region 05 -// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE05]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE05]]{{.+}}) -// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] -// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] - -// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 -// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 -// CK24-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SC]]** -// CK24-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SB]]*** -// CK24-DAG: store [[SC]]* [[VAR0:%.+]], [[SC]]** [[CBP0]] -// CK24-DAG: store [[SB]]** [[SEC0:%.+]], [[SB]]*** [[CP0]] -// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2 - -// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SB]]*** -// CK24-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to [[SB]]** -// CK24-DAG: store [[SB]]** [[SEC0]], [[SB]]*** [[CBP1]] -// CK24-DAG: store [[SB]]* [[SEC1:%.+]], [[SB]]** [[CP1]] -// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SB]]* [[SEC11:%[^,]+]], i{{.+}} 0 -// CK24-DAG: [[SEC11]] = load [[SB]]*, [[SB]]** [[SEC111:%[^,]+]], -// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2 - -// CK24: call void [[CALL05:@.+]]([[SC]]* {{[^,]+}}) -#pragma omp target map(s.p[:5]) - { s.a++; } - -// Region 06 -// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE06]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}}) -// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] -// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] - -// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 -// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], 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: [[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]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 - -// CK24: call void [[CALL06:@.+]]([[SC]]* {{[^,]+}}) -#pragma omp target map(s.s.sa[3].a) - { s.a++; } - -// Region 07 -// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE07]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE07]]{{.+}}) -// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] -// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] - -// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 -// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], 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: [[SEC0]] = getelementptr {{.*}}[10 x [[SA]]*]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 3 -// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SB]]* [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 3 -// CK24-DAG: [[SEC000]] = getelementptr {{.*}}[[SC]]* [[VAR0]], 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: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SA]]*** -// CK24-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** -// CK24-DAG: store [[SA]]** [[SEC0]], [[SA]]*** [[CBP1]] -// CK24-DAG: store i32* [[SEC1:%.+]], i32** [[CP1]] -// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SA]]* [[SEC11:%[^,]+]], i{{.+}} 0, i{{.+}} 0 -// CK24-DAG: [[SEC11]] = load [[SA]]*, [[SA]]** [[SEC111:%[^,]+]], -// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[10 x [[SA]]*]* [[SEC1111:%[^,]+]], i{{.+}} 0, i{{.+}} 3 -// CK24-DAG: [[SEC1111]] = getelementptr {{.*}}[[SB]]* [[SEC11111:%[^,]+]], i{{.+}} 0, i{{.+}} 3 -// CK24-DAG: [[SEC11111]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 - -// CK24: call void [[CALL07:@.+]]([[SC]]* {{[^,]+}}) -#pragma omp target map(s.s.sp[3]->a) - { s.a++; } - -// Region 08 -// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE08]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE08]]{{.+}}) -// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] -// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] - -// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 -// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 -// CK24-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SC]]** -// CK24-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SB]]*** -// CK24-DAG: store [[SC]]* [[VAR0:%.+]], [[SC]]** [[CBP0]] -// CK24-DAG: store [[SB]]** [[SEC0:%.+]], [[SB]]*** [[CP0]] -// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2 - -// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SB]]*** -// CK24-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** -// CK24-DAG: store [[SB]]** [[SEC0]], [[SB]]*** [[CBP1]] -// CK24-DAG: store i32* [[SEC1:%.+]], i32** [[CP1]] -// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SB]]* [[SEC11:%[^,]+]], i{{.+}} 0 -// CK24-DAG: [[SEC11]] = load [[SB]]*, [[SB]]** [[SEC111:%[^,]+]], -// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2 - -// CK24: call void [[CALL08:@.+]]([[SC]]* {{[^,]+}}) -#pragma omp target map(s.p->a) - { s.a++; } - -// Region 09 -// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE09]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE09]]{{.+}}) -// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] -// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] - -// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 -// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], 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: [[SEC0]] = getelementptr {{.*}}[[SB]]* [[SEC00:[^,]+]], i{{.+}} 0, i{{.+}} 4 -// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR0]], 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: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SA]]*** -// CK24-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** -// CK24-DAG: store [[SA]]** [[SEC0]], [[SA]]*** [[CBP1]] -// CK24-DAG: store i32* [[SEC1:%.+]], i32** [[CP1]] -// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SA]]* [[SEC11:%[^,]+]], i{{.+}} 0 -// CK24-DAG: [[SEC11]] = load [[SA]]*, [[SA]]** [[SEC111:%[^,]+]], -// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[[SB]]* [[SEC1111:[^,]+]], i{{.+}} 0, i{{.+}} 4 -// CK24-DAG: [[SEC1111]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 - -// CK24: call void [[CALL09:@.+]]([[SC]]* {{[^,]+}}) -#pragma omp target map(s.s.p->a) - { s.a++; } - -// Region 10 -// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE10]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE10]]{{.+}}) -// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] -// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] - -// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 -// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], 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: [[SEC0]] = getelementptr {{.*}}[10 x i32]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 0 -// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SA]]* [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 2 -// CK24-DAG: [[SEC000]] = getelementptr {{.*}}[[SB]]* [[SEC0000:%[^,]+]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[SEC0000]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 - -// CK24: call void [[CALL10:@.+]]([[SC]]* {{[^,]+}}) -#pragma omp target map(s.s.s.b[:2]) - { s.a++; } - -// Region 11 -// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE11]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE11]]{{.+}}) -// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] -// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] - -// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 -// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], 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: [[SEC0]] = getelementptr {{.*}}[[SB]]* [[SEC00:%[^,]+]], i{{.+}} 0, i{{.+}} 4 -// CK24-DAG: [[SEC00]] = getelementptr {{.*}}[[SC]]* [[VAR0]], 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: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SA]]*** -// CK24-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** -// CK24-DAG: store [[SA]]** [[SEC0]], [[SA]]*** [[CBP1]] -// CK24-DAG: store i32* [[SEC1:%.+]], i32** [[CP1]] -// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[10 x i32]* [[SEC11:%[^,]+]], i{{.+}} 0, i{{.+}} 0 -// CK24-DAG: [[SEC11]] = getelementptr {{.*}}[[SA]]* [[SEC111:%[^,]+]], i{{.+}} 0, i{{.+}} 2 -// CK24-DAG: [[SEC111]] = load [[SA]]*, [[SA]]** [[SEC1111:%[^,]+]], -// CK24-DAG: [[SEC1111]] = getelementptr {{.*}}[[SB]]* [[SEC11111:%[^,]+]], i{{.+}} 0, i{{.+}} 4 -// CK24-DAG: [[SEC11111]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 - -// CK24: call void [[CALL11:@.+]]([[SC]]* {{[^,]+}}) -#pragma omp target map(s.s.p->b[:2]) - { s.a++; } - -// Region 12 -// CK24-DAG: call i32 @__tgt_target(i32 {{[^,]+}}, i8* {{[^,]+}}, i32 4, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[4 x i{{.+}}]* [[SIZE12]], {{.+}}getelementptr {{.+}}[4 x i{{.+}}]* [[MTYPE12]]{{.+}}) -// CK24-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] -// CK24-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] - -// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 -// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 -// CK24-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[SC]]** -// CK24-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [[SB]]*** -// CK24-DAG: store [[SC]]* [[VAR0:%.+]], [[SC]]** [[CBP0]] -// CK24-DAG: store [[SB]]** [[SEC0:%.+]], [[SB]]*** [[CP0]] -// CK24-DAG: [[SEC0]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2 - -// CK24-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[SB]]*** -// CK24-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to [[SA]]*** -// CK24-DAG: store [[SB]]** [[SEC0:%.+]], [[SB]]*** [[CBP1]] -// CK24-DAG: store [[SA]]** [[SEC1:%.+]], [[SA]]*** [[CP1]] -// CK24-DAG: [[SEC1]] = getelementptr {{.*}}[[SB]]* [[SEC11:%[^,]+]], i{{.+}} 0, i{{.+}} 4 -// CK24-DAG: [[SEC11]] = load [[SB]]*, [[SB]]** [[SEC111:%[^,]+]], -// CK24-DAG: [[SEC111]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2 - -// CK24-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 -// CK24-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 -// CK24-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [[SA]]*** -// CK24-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to [[SA]]*** -// CK24-DAG: store [[SA]]** [[SEC1:%.+]], [[SA]]*** [[CBP2]] -// CK24-DAG: store [[SA]]** [[SEC2:%.+]], [[SA]]*** [[CP2]] -// CK24-DAG: [[SEC2]] = getelementptr {{.*}}[[SA]]* [[SEC22:%[^,]+]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[SEC22]] = load [[SA]]*, [[SA]]** [[SEC222:%[^,]+]], -// CK24-DAG: [[SEC222]] = getelementptr {{.*}}[[SB]]* [[SEC2222:%[^,]+]], i{{.+}} 0, i{{.+}} 4 -// CK24-DAG: [[SEC2222]] = load [[SB]]*, [[SB]]** [[SEC22222:%[^,]+]], -// CK24-DAG: [[SEC22222]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2 - -// CK24-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3 -// CK24-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3 -// CK24-DAG: [[CBP3:%.+]] = bitcast i8** [[BP3]] to [[SA]]*** -// CK24-DAG: [[CP3:%.+]] = bitcast i8** [[P3]] to i32** -// CK24-DAG: store [[SA]]** [[SEC2]], [[SA]]*** [[CBP3]] -// CK24-DAG: store i32* [[SEC3:%.+]], i32** [[CP3]] -// CK24-DAG: [[SEC3]] = getelementptr {{.*}}[[SA]]* [[SEC33:%[^,]+]], i{{.+}} 0, i{{.+}} 0 -// CK24-DAG: [[SEC33]] = load [[SA]]*, [[SA]]** [[SEC333:%[^,]+]], -// CK24-DAG: [[SEC333]] = getelementptr {{.*}}[[SA]]* [[SEC3333:%[^,]+]], i{{.+}} 0, i{{.+}} 1 -// CK24-DAG: [[SEC3333]] = load [[SA]]*, [[SA]]** [[SEC33333:%[^,]+]], -// CK24-DAG: [[SEC33333]] = getelementptr {{.*}}[[SB]]* [[SEC333333:%[^,]+]], i{{.+}} 0, i{{.+}} 4 -// CK24-DAG: [[SEC333333]] = load [[SB]]*, [[SB]]** [[SEC3333333:%[^,]+]], -// CK24-DAG: [[SEC3333333]] = getelementptr {{.*}}[[SC]]* [[VAR0]], i{{.+}} 0, i{{.+}} 2 - -// CK24: call void [[CALL12:@.+]]([[SC]]* {{[^,]+}}) -#pragma omp target map(s.p->p->p->a) - { s.a++; } - // // Same thing but starting from a pointer. // @@ -4074,17 +3799,6 @@ int explicit_maps_struct_fields(int a){ } // CK24: define {{.+}}[[CALL01]] -// CK24: define {{.+}}[[CALL02]] -// CK24: define {{.+}}[[CALL03]] -// CK24: define {{.+}}[[CALL04]] -// CK24: define {{.+}}[[CALL05]] -// CK24: define {{.+}}[[CALL06]] -// CK24: define {{.+}}[[CALL07]] -// CK24: define {{.+}}[[CALL08]] -// CK24: define {{.+}}[[CALL09]] -// CK24: define {{.+}}[[CALL10]] -// CK24: define {{.+}}[[CALL11]] -// CK24: define {{.+}}[[CALL12]] // CK24: define {{.+}}[[CALL13]] // CK24: define {{.+}}[[CALL14]] // CK24: define {{.+}}[[CALL15]] diff --git a/clang/test/OpenMP/target_map_messages.cpp b/clang/test/OpenMP/target_map_messages.cpp index 635d2c3..b7c1b07 100644 --- a/clang/test/OpenMP/target_map_messages.cpp +++ b/clang/test/OpenMP/target_map_messages.cpp @@ -454,6 +454,25 @@ T tmain(T argc) { return 0; } +struct SA1{ + int a; + struct SA1 *p; + int b[10]; +}; +struct SB1{ + int a; + struct SA1 s; + struct SA1 sa[10]; + struct SA1 *sp[10]; + struct SA1 *p; +}; +struct SC1{ + int a; + struct SB1 s; + struct SB1 *p; + int b[10]; +}; + int main(int argc, char **argv) { const int d = 5; const int da[5] = { 0 }; @@ -467,6 +486,8 @@ int main(int argc, char **argv) { int y; int to, tofrom, always; const int (&l)[5] = da; + SC1 s; + SC1 *p; #pragma omp target data map // expected-error {{expected '(' after 'map'}} expected-error {{expected at least one 'map' or 'use_device_ptr' clause for '#pragma omp target data'}} #pragma omp target data map( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}} #pragma omp target data map() // expected-error {{expected expression}} @@ -527,6 +548,51 @@ int main(int argc, char **argv) { {} #pragma omp target map(m) {} +// expected-note@+1 {{used here}} +#pragma omp target map(s.s.s) +// expected-error@+1 {{variable already marked as mapped in current construct}} + { s.a++; } +// expected-note@+1 {{used here}} +#pragma omp target map(s.s.s.a) +// expected-error@+1 {{variable already marked as mapped in current construct}} + { s.a++; } +// expected-note@+1 {{used here}} +#pragma omp target map(s.b[:5]) +// expected-error@+1 {{variable already marked as mapped in current construct}} + { s.a++; } +// expected-note@+1 {{used here}} +#pragma omp target map(s.p[:5]) +// expected-error@+1 {{variable already marked as mapped in current construct}} + { s.a++; } +// expected-note@+1 {{used here}} +#pragma omp target map(s.s.sa[3].a) +// expected-error@+1 {{variable already marked as mapped in current construct}} + { s.a++; } +// expected-note@+1 {{used here}} +#pragma omp target map(s.s.sp[3]->a) +// expected-error@+1 {{variable already marked as mapped in current construct}} + { s.a++; } +// expected-note@+1 {{used here}} +#pragma omp target map(s.p->a) +// expected-error@+1 {{variable already marked as mapped in current construct}} + { s.a++; } +// expected-note@+1 {{used here}} +#pragma omp target map(s.s.p->a) +// expected-error@+1 {{variable already marked as mapped in current construct}} + { s.a++; } +// expected-note@+1 {{used here}} +#pragma omp target map(s.s.s.b[:2]) +// expected-error@+1 {{variable already marked as mapped in current construct}} + { s.a++; } +// expected-note@+1 {{used here}} +#pragma omp target map(s.s.p->b[:2]) +// expected-error@+1 {{variable already marked as mapped in current construct}} + { s.a++; } +// expected-note@+1 {{used here}} +#pragma omp target map(s.p->p->p->a) +// expected-error@+1 {{variable already marked as mapped in current construct}} + { s.a++; } + return tmain(argc)+tmain(argc); // expected-note {{in instantiation of function template specialization 'tmain' requested here}} expected-note {{in instantiation of function template specialization 'tmain' requested here}} } #endif diff --git a/clang/test/OpenMP/target_parallel_codegen.cpp b/clang/test/OpenMP/target_parallel_codegen.cpp index 3063ab1..c94924e 100644 --- a/clang/test/OpenMP/target_parallel_codegen.cpp +++ b/clang/test/OpenMP/target_parallel_codegen.cpp @@ -42,12 +42,12 @@ // CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [1 x i32] [i32 288] // CHECK-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [2 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2] // CHECK-DAG: [[MAPT3:@.+]] = private unnamed_addr constant [2 x i32] [i32 288, i32 288] -// CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [9 x i32] [i32 288, i32 35, i32 288, i32 35, i32 35, i32 288, i32 288, i32 35, i32 35] +// CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [9 x i32] [i32 288, i32 547, i32 288, i32 547, i32 547, i32 288, i32 288, i32 547, i32 547] // CHECK-DAG: [[SIZET5:@.+]] = private unnamed_addr constant [3 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2, i[[SZ]] 40] -// CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i32] [i32 288, i32 288, i32 35] +// CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i32] [i32 288, i32 288, i32 547] // CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [4 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2, i[[SZ]] 1, i[[SZ]] 40] -// CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i32] [i32 288, i32 288, i32 288, i32 35] -// CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [5 x i32] [i32 35, i32 288, i32 288, i32 288, i32 35] +// CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i32] [i32 288, i32 288, i32 288, i32 547] +// CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [5 x i32] [i32 547, i32 288, i32 288, i32 288, i32 547] // CHECK-DAG: @{{.*}} = private constant i8 0 // CHECK-DAG: @{{.*}} = private constant i8 0 // CHECK-DAG: @{{.*}} = private constant i8 0 diff --git a/clang/test/OpenMP/target_teams_codegen.cpp b/clang/test/OpenMP/target_teams_codegen.cpp index 87304aa..9def589 100644 --- a/clang/test/OpenMP/target_teams_codegen.cpp +++ b/clang/test/OpenMP/target_teams_codegen.cpp @@ -42,12 +42,12 @@ // CHECK-DAG: [[MAPT2:@.+]] = private unnamed_addr constant [1 x i32] [i32 288] // CHECK-DAG: [[SIZET3:@.+]] = private unnamed_addr constant [2 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2] // CHECK-DAG: [[MAPT3:@.+]] = private unnamed_addr constant [2 x i32] [i32 288, i32 288] -// CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [9 x i32] [i32 288, i32 35, i32 288, i32 35, i32 35, i32 288, i32 288, i32 35, i32 35] +// CHECK-DAG: [[MAPT4:@.+]] = private unnamed_addr constant [9 x i32] [i32 288, i32 547, i32 288, i32 547, i32 547, i32 288, i32 288, i32 547, i32 547] // CHECK-DAG: [[SIZET5:@.+]] = private unnamed_addr constant [3 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2, i[[SZ]] 40] -// CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i32] [i32 288, i32 288, i32 35] +// CHECK-DAG: [[MAPT5:@.+]] = private unnamed_addr constant [3 x i32] [i32 288, i32 288, i32 547] // CHECK-DAG: [[SIZET6:@.+]] = private unnamed_addr constant [4 x i[[SZ]]] [i[[SZ]] 4, i[[SZ]] 2, i[[SZ]] 1, i[[SZ]] 40] -// CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i32] [i32 288, i32 288, i32 288, i32 35] -// CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [5 x i32] [i32 35, i32 288, i32 288, i32 288, i32 35] +// CHECK-DAG: [[MAPT6:@.+]] = private unnamed_addr constant [4 x i32] [i32 288, i32 288, i32 288, i32 547] +// CHECK-DAG: [[MAPT7:@.+]] = private unnamed_addr constant [5 x i32] [i32 547, i32 288, i32 288, i32 288, i32 547] // CHECK-DAG: @{{.*}} = private constant i8 0 // CHECK-DAG: @{{.*}} = private constant i8 0 // CHECK-DAG: @{{.*}} = private constant i8 0 diff --git a/clang/test/OpenMP/teams_distribute_parallel_for_shared_messages.cpp b/clang/test/OpenMP/teams_distribute_parallel_for_shared_messages.cpp index 8edcd36..1a1af09 100644 --- a/clang/test/OpenMP/teams_distribute_parallel_for_shared_messages.cpp +++ b/clang/test/OpenMP/teams_distribute_parallel_for_shared_messages.cpp @@ -84,7 +84,7 @@ int main(int argc, char **argv) { #pragma omp teams distribute parallel for shared (S1) // expected-error {{'S1' does not refer to a value}} for (int j=0; j<100; j++) foo(); #pragma omp target - #pragma omp teams distribute parallel for shared (a, b, c, d, f) + #pragma omp teams distribute parallel for shared (a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} for (int j=0; j<100; j++) foo(); #pragma omp target #pragma omp teams distribute parallel for shared (argv[1]) // expected-error {{expected variable name}} diff --git a/clang/test/OpenMP/teams_distribute_parallel_for_simd_shared_messages.cpp b/clang/test/OpenMP/teams_distribute_parallel_for_simd_shared_messages.cpp index f131466..1089555 100644 --- a/clang/test/OpenMP/teams_distribute_parallel_for_simd_shared_messages.cpp +++ b/clang/test/OpenMP/teams_distribute_parallel_for_simd_shared_messages.cpp @@ -84,7 +84,7 @@ int main(int argc, char **argv) { #pragma omp teams distribute parallel for simd shared (S1) // expected-error {{'S1' does not refer to a value}} for (int j=0; j<100; j++) foo(); #pragma omp target - #pragma omp teams distribute parallel for simd shared (a, b, c, d, f) + #pragma omp teams distribute parallel for simd shared (a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} for (int j=0; j<100; j++) foo(); #pragma omp target #pragma omp teams distribute parallel for simd shared (argv[1]) // expected-error {{expected variable name}} diff --git a/clang/test/OpenMP/teams_distribute_shared_messages.cpp b/clang/test/OpenMP/teams_distribute_shared_messages.cpp index b9e096c..ecf6f2e 100644 --- a/clang/test/OpenMP/teams_distribute_shared_messages.cpp +++ b/clang/test/OpenMP/teams_distribute_shared_messages.cpp @@ -84,7 +84,7 @@ int main(int argc, char **argv) { #pragma omp teams distribute shared (S1) // expected-error {{'S1' does not refer to a value}} for (int j=0; j<100; j++) foo(); #pragma omp target - #pragma omp teams distribute shared (a, b, c, d, f) + #pragma omp teams distribute shared (a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} for (int j=0; j<100; j++) foo(); #pragma omp target #pragma omp teams distribute shared (argv[1]) // expected-error {{expected variable name}} diff --git a/clang/test/OpenMP/teams_distribute_simd_shared_messages.cpp b/clang/test/OpenMP/teams_distribute_simd_shared_messages.cpp index cd963e8..6b06f5b 100644 --- a/clang/test/OpenMP/teams_distribute_simd_shared_messages.cpp +++ b/clang/test/OpenMP/teams_distribute_simd_shared_messages.cpp @@ -84,7 +84,7 @@ int main(int argc, char **argv) { #pragma omp teams distribute simd shared (S1) // expected-error {{'S1' does not refer to a value}} for (int j=0; j<100; j++) foo(); #pragma omp target - #pragma omp teams distribute simd shared (a, b, c, d, f) + #pragma omp teams distribute simd shared (a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} for (int j=0; j<100; j++) foo(); #pragma omp target #pragma omp teams distribute simd shared (argv[1]) // expected-error {{expected variable name}} diff --git a/clang/test/OpenMP/teams_shared_messages.cpp b/clang/test/OpenMP/teams_shared_messages.cpp index e14ba1e..6b013d0 100644 --- a/clang/test/OpenMP/teams_shared_messages.cpp +++ b/clang/test/OpenMP/teams_shared_messages.cpp @@ -84,7 +84,7 @@ int main(int argc, char **argv) { #pragma omp teams shared (S1) // expected-error {{'S1' does not refer to a value}} foo(); #pragma omp target - #pragma omp teams shared (a, b, c, d, f) + #pragma omp teams shared (a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} foo(); #pragma omp target #pragma omp teams shared (argv[1]) // expected-error {{expected variable name}}