Index: clang/include/clang/AST/OpenMPClause.h =================================================================== --- clang/include/clang/AST/OpenMPClause.h +++ clang/include/clang/AST/OpenMPClause.h @@ -6329,8 +6329,20 @@ friend OMPVarListClause; friend TrailingObjects; + /// Motion-modifiers for the 'to' clause. + OpenMPMotionModifierKind MotionModifiers[NumberOfOMPMotionModifiers] = { + OMPC_MOTION_MODIFIER_unknown, OMPC_MOTION_MODIFIER_unknown}; + + /// Location of motion-modifiers for the 'to' clause. + SourceLocation MotionModifiersLoc[NumberOfOMPMotionModifiers]; + + /// Colon location. + SourceLocation ColonLoc; + /// Build clause with number of variables \a NumVars. /// + /// \param TheMotionModifiers Motion-modifiers. + /// \param TheMotionModifiersLoc Locations of motion-modifiers. /// \param MapperQualifierLoc C++ nested name specifier for the associated /// user-defined mapper. /// \param MapperIdInfo The identifier of associated user-defined mapper. @@ -6342,13 +6354,24 @@ /// NumUniqueDeclarations: number of unique base declarations in this clause; /// 3) NumComponentLists: number of component lists in this clause; and 4) /// NumComponents: total number of expression components in the clause. - explicit OMPToClause(NestedNameSpecifierLoc MapperQualifierLoc, + explicit OMPToClause(ArrayRef TheMotionModifiers, + ArrayRef TheMotionModifiersLoc, + NestedNameSpecifierLoc MapperQualifierLoc, DeclarationNameInfo MapperIdInfo, const OMPVarListLocTy &Locs, const OMPMappableExprListSizeTy &Sizes) : OMPMappableExprListClause(llvm::omp::OMPC_to, Locs, Sizes, /*SupportsMapper=*/true, &MapperQualifierLoc, - &MapperIdInfo) {} + &MapperIdInfo) { + assert(llvm::array_lengthof(MotionModifiers) == TheMotionModifiers.size() && + "Unexpected number of motion modifiers."); + llvm::copy(TheMotionModifiers, std::begin(MotionModifiers)); + + assert(llvm::array_lengthof(MotionModifiersLoc) == + TheMotionModifiersLoc.size() && + "Unexpected number of motion modifier locations."); + llvm::copy(TheMotionModifiersLoc, std::begin(MotionModifiersLoc)); + } /// Build an empty clause. /// @@ -6361,6 +6384,29 @@ : OMPMappableExprListClause(llvm::omp::OMPC_to, OMPVarListLocTy(), Sizes, /*SupportsMapper=*/true) {} + /// Set motion-modifier for the clause. + /// + /// \param I index for motion-modifier. + /// \param T motion-modifier for the clause. + void setMotionModifier(unsigned I, OpenMPMotionModifierKind T) { + assert(I < NumberOfOMPMotionModifiers && + "Unexpected index to store motion modifier, exceeds array size."); + MotionModifiers[I] = T; + } + + /// Set location for the motion-modifier. + /// + /// \param I index for motion-modifier location. + /// \param TLoc motion-modifier location. + void setMotionModifierLoc(unsigned I, SourceLocation TLoc) { + assert(I < NumberOfOMPMotionModifiers && + "Index to store motion modifier location exceeds array size."); + MotionModifiersLoc[I] = TLoc; + } + + /// Set colon location. + void setColonLoc(SourceLocation Loc) { ColonLoc = Loc; } + /// Define the sizes of each trailing object array except the last one. This /// is required for TrailingObjects to work properly. size_t numTrailingObjects(OverloadToken) const { @@ -6385,6 +6431,8 @@ /// \param Vars The original expression used in the clause. /// \param Declarations Declarations used in the clause. /// \param ComponentLists Component lists used in the clause. + /// \param MotionModifiers Motion-modifiers. + /// \param MotionModifiersLoc Location of motion-modifiers. /// \param UDMapperRefs References to user-defined mappers associated with /// expressions used in the clause. /// \param UDMQualifierLoc C++ nested name specifier for the associated @@ -6395,6 +6443,8 @@ ArrayRef Declarations, MappableExprComponentListsRef ComponentLists, ArrayRef UDMapperRefs, + ArrayRef MotionModifiers, + ArrayRef MotionModifiersLoc, NestedNameSpecifierLoc UDMQualifierLoc, DeclarationNameInfo MapperId); @@ -6409,6 +6459,38 @@ static OMPToClause *CreateEmpty(const ASTContext &C, const OMPMappableExprListSizeTy &Sizes); + /// Fetches the motion-modifier at 'Cnt' index of array of modifiers. + /// + /// \param Cnt index for motion-modifier. + OpenMPMotionModifierKind getMotionModifier(unsigned Cnt) const LLVM_READONLY { + assert(Cnt < NumberOfOMPMotionModifiers && + "Requested modifier exceeds the total number of modifiers."); + return MotionModifiers[Cnt]; + } + + /// Fetches the motion-modifier location at 'Cnt' index of array of modifiers' + /// locations. + /// + /// \param Cnt index for motion-modifier location. + SourceLocation getMotionModifierLoc(unsigned Cnt) const LLVM_READONLY { + assert(Cnt < NumberOfOMPMotionModifiers && + "Requested modifier location exceeds total number of modifiers."); + return MotionModifiersLoc[Cnt]; + } + + /// Fetches ArrayRef of motion-modifiers. + ArrayRef getMotionModifiers() const LLVM_READONLY { + return llvm::makeArrayRef(MotionModifiers); + } + + /// Fetches ArrayRef of location of motion-modifiers. + ArrayRef getMotionModifiersLoc() const LLVM_READONLY { + return llvm::makeArrayRef(MotionModifiersLoc); + } + + /// Get colon location. + SourceLocation getColonLoc() const { return ColonLoc; } + child_range children() { return child_range(reinterpret_cast(varlist_begin()), reinterpret_cast(varlist_end())); @@ -6449,8 +6531,20 @@ friend OMPVarListClause; friend TrailingObjects; + /// Motion-modifiers for the 'from' clause. + OpenMPMotionModifierKind MotionModifiers[NumberOfOMPMotionModifiers] = { + OMPC_MOTION_MODIFIER_unknown, OMPC_MOTION_MODIFIER_unknown}; + + /// Location of motion-modifiers for the 'from' clause. + SourceLocation MotionModifiersLoc[NumberOfOMPMotionModifiers]; + + /// Colon location. + SourceLocation ColonLoc; + /// Build clause with number of variables \a NumVars. /// + /// \param TheMotionModifiers Motion-modifiers. + /// \param TheMotionModifiersLoc Locations of motion-modifiers. /// \param MapperQualifierLoc C++ nested name specifier for the associated /// user-defined mapper. /// \param MapperIdInfo The identifier of associated user-defined mapper. @@ -6462,13 +6556,24 @@ /// NumUniqueDeclarations: number of unique base declarations in this clause; /// 3) NumComponentLists: number of component lists in this clause; and 4) /// NumComponents: total number of expression components in the clause. - explicit OMPFromClause(NestedNameSpecifierLoc MapperQualifierLoc, + explicit OMPFromClause(ArrayRef TheMotionModifiers, + ArrayRef TheMotionModifiersLoc, + NestedNameSpecifierLoc MapperQualifierLoc, DeclarationNameInfo MapperIdInfo, const OMPVarListLocTy &Locs, const OMPMappableExprListSizeTy &Sizes) : OMPMappableExprListClause(llvm::omp::OMPC_from, Locs, Sizes, /*SupportsMapper=*/true, &MapperQualifierLoc, - &MapperIdInfo) {} + &MapperIdInfo) { + assert(llvm::array_lengthof(MotionModifiers) == TheMotionModifiers.size() && + "Unexpected number of motion modifiers."); + llvm::copy(TheMotionModifiers, std::begin(MotionModifiers)); + + assert(llvm::array_lengthof(MotionModifiersLoc) == + TheMotionModifiersLoc.size() && + "Unexpected number of motion modifier locations."); + llvm::copy(TheMotionModifiersLoc, std::begin(MotionModifiersLoc)); + } /// Build an empty clause. /// @@ -6481,6 +6586,29 @@ : OMPMappableExprListClause(llvm::omp::OMPC_from, OMPVarListLocTy(), Sizes, /*SupportsMapper=*/true) {} + /// Set motion-modifier for the clause. + /// + /// \param I index for motion-modifier. + /// \param T motion-modifier for the clause. + void setMotionModifier(unsigned I, OpenMPMotionModifierKind T) { + assert(I < NumberOfOMPMotionModifiers && + "Unexpected index to store motion modifier, exceeds array size."); + MotionModifiers[I] = T; + } + + /// Set location for the motion-modifier. + /// + /// \param I index for motion-modifier location. + /// \param TLoc motion-modifier location. + void setMotionModifierLoc(unsigned I, SourceLocation TLoc) { + assert(I < NumberOfOMPMotionModifiers && + "Index to store motion modifier location exceeds array size."); + MotionModifiersLoc[I] = TLoc; + } + + /// Set colon location. + void setColonLoc(SourceLocation Loc) { ColonLoc = Loc; } + /// Define the sizes of each trailing object array except the last one. This /// is required for TrailingObjects to work properly. size_t numTrailingObjects(OverloadToken) const { @@ -6505,18 +6633,21 @@ /// \param Vars The original expression used in the clause. /// \param Declarations Declarations used in the clause. /// \param ComponentLists Component lists used in the clause. + /// \param MotionModifiers Motion-modifiers. + /// \param MotionModifiersLoc Location of motion-modifiers. /// \param UDMapperRefs References to user-defined mappers associated with /// expressions used in the clause. /// \param UDMQualifierLoc C++ nested name specifier for the associated /// user-defined mapper. /// \param MapperId The identifier of associated user-defined mapper. - static OMPFromClause *Create(const ASTContext &C, const OMPVarListLocTy &Locs, - ArrayRef Vars, - ArrayRef Declarations, - MappableExprComponentListsRef ComponentLists, - ArrayRef UDMapperRefs, - NestedNameSpecifierLoc UDMQualifierLoc, - DeclarationNameInfo MapperId); + static OMPFromClause * + Create(const ASTContext &C, const OMPVarListLocTy &Locs, + ArrayRef Vars, ArrayRef Declarations, + MappableExprComponentListsRef ComponentLists, + ArrayRef UDMapperRefs, + ArrayRef MotionModifiers, + ArrayRef MotionModifiersLoc, + NestedNameSpecifierLoc UDMQualifierLoc, DeclarationNameInfo MapperId); /// Creates an empty clause with the place for \a NumVars variables. /// @@ -6529,6 +6660,38 @@ static OMPFromClause *CreateEmpty(const ASTContext &C, const OMPMappableExprListSizeTy &Sizes); + /// Fetches the motion-modifier at 'Cnt' index of array of modifiers. + /// + /// \param Cnt index for motion-modifier. + OpenMPMotionModifierKind getMotionModifier(unsigned Cnt) const LLVM_READONLY { + assert(Cnt < NumberOfOMPMotionModifiers && + "Requested modifier exceeds the total number of modifiers."); + return MotionModifiers[Cnt]; + } + + /// Fetches the motion-modifier location at 'Cnt' index of array of modifiers' + /// locations. + /// + /// \param Cnt index for motion-modifier location. + SourceLocation getMotionModifierLoc(unsigned Cnt) const LLVM_READONLY { + assert(Cnt < NumberOfOMPMotionModifiers && + "Requested modifier location exceeds total number of modifiers."); + return MotionModifiersLoc[Cnt]; + } + + /// Fetches ArrayRef of motion-modifiers. + ArrayRef getMotionModifiers() const LLVM_READONLY { + return llvm::makeArrayRef(MotionModifiers); + } + + /// Fetches ArrayRef of location of motion-modifiers. + ArrayRef getMotionModifiersLoc() const LLVM_READONLY { + return llvm::makeArrayRef(MotionModifiersLoc); + } + + /// Get colon location. + SourceLocation getColonLoc() const { return ColonLoc; } + child_range children() { return child_range(reinterpret_cast(varlist_begin()), reinterpret_cast(varlist_end())); @@ -7623,6 +7786,8 @@ /// Process clauses with list of variables. template void VisitOMPClauseList(T *Node, char StartSym); + /// Process motion clauses. + template void VisitOMPMotionClause(T *Node); public: OMPClausePrinter(raw_ostream &OS, const PrintingPolicy &Policy) Index: clang/include/clang/Basic/DiagnosticSemaKinds.td =================================================================== --- clang/include/clang/Basic/DiagnosticSemaKinds.td +++ clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -9807,6 +9807,8 @@ "multiple array elements associated with the same variable are not allowed in map clauses of the same construct">; def err_omp_duplicate_map_type_modifier : Error< "same map type modifier has been specified more than once">; +def err_omp_duplicate_motion_modifier : Error< + "same motion modifier has been specified more than once">; def err_omp_pointer_mapped_along_with_derived_section : Error< "pointer cannot be mapped along with a section derived from itself">; def err_omp_original_storage_is_shared_and_does_not_contain : Error< Index: clang/include/clang/Basic/OpenMPKinds.h =================================================================== --- clang/include/clang/Basic/OpenMPKinds.h +++ clang/include/clang/Basic/OpenMPKinds.h @@ -94,6 +94,10 @@ OMPC_MOTION_MODIFIER_unknown }; +/// Number of allowed motion-modifiers. +static constexpr unsigned NumberOfOMPMotionModifiers = + OMPC_MOTION_MODIFIER_unknown; + /// OpenMP attributes for 'dist_schedule' clause. enum OpenMPDistScheduleClauseKind { #define OPENMP_DIST_SCHEDULE_KIND(Name) OMPC_DIST_SCHEDULE_##Name, Index: clang/include/clang/Basic/OpenMPKinds.def =================================================================== --- clang/include/clang/Basic/OpenMPKinds.def +++ clang/include/clang/Basic/OpenMPKinds.def @@ -125,6 +125,7 @@ // Modifiers for 'to' or 'from' clause. OPENMP_MOTION_MODIFIER_KIND(mapper) +OPENMP_MOTION_MODIFIER_KIND(present) // Static attributes for 'dist_schedule' clause. OPENMP_DIST_SCHEDULE_KIND(static) Index: clang/include/clang/Parse/Parser.h =================================================================== --- clang/include/clang/Parse/Parser.h +++ clang/include/clang/Parse/Parser.h @@ -3221,6 +3221,9 @@ MapTypeModifiers; SmallVector MapTypeModifiersLoc; + SmallVector + MotionModifiers; + SmallVector MotionModifiersLoc; bool IsMapTypeImplicit = false; SourceLocation ExtraModifierLoc; }; Index: clang/include/clang/Sema/Sema.h =================================================================== --- clang/include/clang/Sema/Sema.h +++ clang/include/clang/Sema/Sema.h @@ -10720,7 +10720,9 @@ DeclarationNameInfo &ReductionOrMapperId, int ExtraModifier, ArrayRef MapTypeModifiers, ArrayRef MapTypeModifiersLoc, bool IsMapTypeImplicit, - SourceLocation ExtraModifierLoc); + SourceLocation ExtraModifierLoc, + ArrayRef MotionModifiers, + ArrayRef MotionModifiersLoc); /// Called on well-formed 'inclusive' clause. OMPClause *ActOnOpenMPInclusiveClause(ArrayRef VarList, SourceLocation StartLoc, @@ -10857,15 +10859,20 @@ SourceLocation KindLoc, SourceLocation EndLoc); /// Called on well-formed 'to' clause. OMPClause * - ActOnOpenMPToClause(ArrayRef VarList, CXXScopeSpec &MapperIdScopeSpec, - DeclarationNameInfo &MapperId, - const OMPVarListLocTy &Locs, + ActOnOpenMPToClause(ArrayRef MotionModifiers, + ArrayRef MotionModifiersLoc, + CXXScopeSpec &MapperIdScopeSpec, + DeclarationNameInfo &MapperId, SourceLocation ColonLoc, + ArrayRef VarList, const OMPVarListLocTy &Locs, ArrayRef UnresolvedMappers = llvm::None); /// Called on well-formed 'from' clause. - OMPClause *ActOnOpenMPFromClause( - ArrayRef VarList, CXXScopeSpec &MapperIdScopeSpec, - DeclarationNameInfo &MapperId, const OMPVarListLocTy &Locs, - ArrayRef UnresolvedMappers = llvm::None); + OMPClause * + ActOnOpenMPFromClause(ArrayRef MotionModifiers, + ArrayRef MotionModifiersLoc, + CXXScopeSpec &MapperIdScopeSpec, + DeclarationNameInfo &MapperId, SourceLocation ColonLoc, + ArrayRef VarList, const OMPVarListLocTy &Locs, + ArrayRef UnresolvedMappers = llvm::None); /// Called on well-formed 'use_device_ptr' clause. OMPClause *ActOnOpenMPUseDevicePtrClause(ArrayRef VarList, const OMPVarListLocTy &Locs); Index: clang/lib/AST/OpenMPClause.cpp =================================================================== --- clang/lib/AST/OpenMPClause.cpp +++ clang/lib/AST/OpenMPClause.cpp @@ -1097,6 +1097,8 @@ const ASTContext &C, const OMPVarListLocTy &Locs, ArrayRef Vars, ArrayRef Declarations, MappableExprComponentListsRef ComponentLists, ArrayRef UDMapperRefs, + ArrayRef MotionModifiers, + ArrayRef MotionModifiersLoc, NestedNameSpecifierLoc UDMQualifierLoc, DeclarationNameInfo MapperId) { OMPMappableExprListSizeTy Sizes; Sizes.NumVars = Vars.size(); @@ -1121,7 +1123,8 @@ Sizes.NumUniqueDeclarations + Sizes.NumComponentLists, Sizes.NumComponents)); - auto *Clause = new (Mem) OMPToClause(UDMQualifierLoc, MapperId, Locs, Sizes); + auto *Clause = new (Mem) OMPToClause(MotionModifiers, MotionModifiersLoc, + UDMQualifierLoc, MapperId, Locs, Sizes); Clause->setVarRefs(Vars); Clause->setUDMapperRefs(UDMapperRefs); @@ -1144,6 +1147,8 @@ const ASTContext &C, const OMPVarListLocTy &Locs, ArrayRef Vars, ArrayRef Declarations, MappableExprComponentListsRef ComponentLists, ArrayRef UDMapperRefs, + ArrayRef MotionModifiers, + ArrayRef MotionModifiersLoc, NestedNameSpecifierLoc UDMQualifierLoc, DeclarationNameInfo MapperId) { OMPMappableExprListSizeTy Sizes; Sizes.NumVars = Vars.size(); @@ -1169,7 +1174,8 @@ Sizes.NumComponents)); auto *Clause = - new (Mem) OMPFromClause(UDMQualifierLoc, MapperId, Locs, Sizes); + new (Mem) OMPFromClause(MotionModifiers, MotionModifiersLoc, + UDMQualifierLoc, MapperId, Locs, Sizes); Clause->setVarRefs(Vars); Clause->setUDMapperRefs(UDMapperRefs); @@ -1936,6 +1942,17 @@ OS << ")"; } +template +static void PrintMapper(raw_ostream &OS, T *Node, + const PrintingPolicy &Policy) { + OS << '('; + NestedNameSpecifier *MapperNNS = + Node->getMapperQualifierLoc().getNestedNameSpecifier(); + if (MapperNNS) + MapperNNS->print(OS, Policy); + OS << Node->getMapperIdInfo() << ')'; +} + void OMPClausePrinter::VisitOMPMapClause(OMPMapClause *Node) { if (!Node->varlist_empty()) { OS << "map("; @@ -1944,14 +1961,8 @@ if (Node->getMapTypeModifier(I) != OMPC_MAP_MODIFIER_unknown) { OS << getOpenMPSimpleClauseTypeName(OMPC_map, Node->getMapTypeModifier(I)); - if (Node->getMapTypeModifier(I) == OMPC_MAP_MODIFIER_mapper) { - OS << '('; - NestedNameSpecifier *MapperNNS = - Node->getMapperQualifierLoc().getNestedNameSpecifier(); - if (MapperNNS) - MapperNNS->print(OS, Policy); - OS << Node->getMapperIdInfo() << ')'; - } + if (Node->getMapTypeModifier(I) == OMPC_MAP_MODIFIER_mapper) + PrintMapper(OS, Node, Policy); OS << ','; } } @@ -1963,44 +1974,41 @@ } } -void OMPClausePrinter::VisitOMPToClause(OMPToClause *Node) { - if (!Node->varlist_empty()) { - OS << "to"; - DeclarationNameInfo MapperId = Node->getMapperIdInfo(); - if (MapperId.getName() && !MapperId.getName().isEmpty()) { - OS << '('; - OS << "mapper("; - NestedNameSpecifier *MapperNNS = - Node->getMapperQualifierLoc().getNestedNameSpecifier(); - if (MapperNNS) - MapperNNS->print(OS, Policy); - OS << MapperId << "):"; - VisitOMPClauseList(Node, ' '); - } else { - VisitOMPClauseList(Node, '('); +template void OMPClausePrinter::VisitOMPMotionClause(T *Node) { + if (Node->varlist_empty()) + return; + OS << getOpenMPClauseName(Node->getClauseKind()); + unsigned ModifierCount = 0; + for (unsigned I = 0; I < NumberOfOMPMotionModifiers; ++I) { + if (Node->getMotionModifier(I) != OMPC_MOTION_MODIFIER_unknown) + ++ModifierCount; + } + if (ModifierCount) { + OS << '('; + for (unsigned I = 0; I < NumberOfOMPMotionModifiers; ++I) { + if (Node->getMotionModifier(I) != OMPC_MOTION_MODIFIER_unknown) { + OS << getOpenMPSimpleClauseTypeName(Node->getClauseKind(), + Node->getMotionModifier(I)); + if (Node->getMotionModifier(I) == OMPC_MOTION_MODIFIER_mapper) + PrintMapper(OS, Node, Policy); + if (I < ModifierCount - 1) + OS << ", "; + } } - OS << ")"; + OS << ':'; + VisitOMPClauseList(Node, ' '); + } else { + VisitOMPClauseList(Node, '('); } + OS << ")"; +} + +void OMPClausePrinter::VisitOMPToClause(OMPToClause *Node) { + VisitOMPMotionClause(Node); } void OMPClausePrinter::VisitOMPFromClause(OMPFromClause *Node) { - if (!Node->varlist_empty()) { - OS << "from"; - DeclarationNameInfo MapperId = Node->getMapperIdInfo(); - if (MapperId.getName() && !MapperId.getName().isEmpty()) { - OS << '('; - OS << "mapper("; - NestedNameSpecifier *MapperNNS = - Node->getMapperQualifierLoc().getNestedNameSpecifier(); - if (MapperNNS) - MapperNNS->print(OS, Policy); - OS << MapperId << "):"; - VisitOMPClauseList(Node, ' '); - } else { - VisitOMPClauseList(Node, '('); - } - OS << ")"; - } + VisitOMPMotionClause(Node); } void OMPClausePrinter::VisitOMPDistScheduleClause(OMPDistScheduleClause *Node) { Index: clang/lib/Basic/OpenMPKinds.cpp =================================================================== --- clang/lib/Basic/OpenMPKinds.cpp +++ clang/lib/Basic/OpenMPKinds.cpp @@ -64,12 +64,16 @@ return Type; } case OMPC_to: - case OMPC_from: - return llvm::StringSwitch(Str) + case OMPC_from: { + unsigned Type = llvm::StringSwitch(Str) #define OPENMP_MOTION_MODIFIER_KIND(Name) \ .Case(#Name, static_cast(OMPC_MOTION_MODIFIER_##Name)) #include "clang/Basic/OpenMPKinds.def" .Default(OMPC_MOTION_MODIFIER_unknown); + if (OpenMPVersion < 51 && Type != OMPC_MOTION_MODIFIER_mapper) + return OMPC_MOTION_MODIFIER_unknown; + return Type; + } case OMPC_dist_schedule: return llvm::StringSwitch(Str) #define OPENMP_DIST_SCHEDULE_KIND(Name) .Case(#Name, OMPC_DIST_SCHEDULE_##Name) Index: clang/lib/CodeGen/CGOpenMPRuntime.cpp =================================================================== --- clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -7920,6 +7920,20 @@ Layout.push_back(Data.get()); } } + static void translateMotionModifiers( + ArrayRef MotionModifiers, + SmallVectorImpl &MapModifiers) { + for (OpenMPMotionModifierKind MotionMod : MotionModifiers) { + switch (MotionMod) { + case OMPC_MOTION_MODIFIER_present: + MapModifiers.push_back(OMPC_MAP_MODIFIER_present); + break; + case OMPC_MOTION_MODIFIER_mapper: + case OMPC_MOTION_MODIFIER_unknown: + break; + } + } + } public: MappableExprsHandler(const OMPExecutableDirective &Dir, CodeGenFunction &CGF) @@ -8038,12 +8052,16 @@ } for (const auto *C : CurExecDir->getClausesOfKind()) for (const auto L : C->component_lists()) { - InfoGen(std::get<0>(L), std::get<1>(L), OMPC_MAP_to, llvm::None, + SmallVector MapModifiers; + translateMotionModifiers(C->getMotionModifiers(), MapModifiers); + InfoGen(std::get<0>(L), std::get<1>(L), OMPC_MAP_to, MapModifiers, /*ReturnDevicePointer=*/false, C->isImplicit(), std::get<2>(L)); } for (const auto *C : CurExecDir->getClausesOfKind()) for (const auto L : C->component_lists()) { - InfoGen(std::get<0>(L), std::get<1>(L), OMPC_MAP_from, llvm::None, + SmallVector MapModifiers; + translateMotionModifiers(C->getMotionModifiers(), MapModifiers); + InfoGen(std::get<0>(L), std::get<1>(L), OMPC_MAP_from, MapModifiers, /*ReturnDevicePointer=*/false, C->isImplicit(), std::get<2>(L)); } Index: clang/lib/Parse/ParseOpenMP.cpp =================================================================== --- clang/lib/Parse/ParseOpenMP.cpp +++ clang/lib/Parse/ParseOpenMP.cpp @@ -3441,25 +3441,43 @@ if (Tok.is(tok::colon)) Data.ColonLoc = ConsumeToken(); } else if (Kind == OMPC_to || Kind == OMPC_from) { - if (Tok.is(tok::identifier)) { + while (Tok.is(tok::identifier)) { auto Modifier = static_cast(getOpenMPSimpleClauseType( Kind, PP.getSpelling(Tok), getLangOpts().OpenMP)); + if (Modifier == OMPC_MOTION_MODIFIER_unknown) + break; + Data.MotionModifiers.push_back(Modifier); + Data.MotionModifiersLoc.push_back(Tok.getLocation()); + ConsumeToken(); if (Modifier == OMPC_MOTION_MODIFIER_mapper) { - // Parse the mapper modifier. - ConsumeToken(); IsInvalidMapperModifier = parseMapperModifier(Data); - if (Tok.isNot(tok::colon)) { - if (!IsInvalidMapperModifier) - Diag(Tok, diag::warn_pragma_expected_colon) << ")"; - SkipUntil(tok::colon, tok::r_paren, tok::annot_pragma_openmp_end, - StopBeforeMatch); - } - // Consume ':'. - if (Tok.is(tok::colon)) - ConsumeToken(); + if (IsInvalidMapperModifier) + break; } + // OpenMP < 5.1 doesn't permit a ',' or additional modifiers. + if (getLangOpts().OpenMP < 51) + break; + // OpenMP 5.1 accepts an optional ',' even if the next character is ':'. + // TODO: Is that intentional? + if (Tok.is(tok::comma)) + ConsumeToken(); + } + if (!Data.MotionModifiers.empty() && Tok.isNot(tok::colon)) { + if (!IsInvalidMapperModifier) { + if (getLangOpts().OpenMP < 51) + Diag(Tok, diag::warn_pragma_expected_colon) << ")"; + else + Diag(Tok, diag::warn_pragma_expected_colon) << "motion modifier"; + } + SkipUntil(tok::colon, tok::r_paren, tok::annot_pragma_openmp_end, + StopBeforeMatch); } + // OpenMP 5.1 permits a ':' even without a preceding modifier. TODO: Is + // that intentional? + if ((!Data.MotionModifiers.empty() || getLangOpts().OpenMP >= 51) && + Tok.is(tok::colon)) + Data.ColonLoc = ConsumeToken(); } else if (Kind == OMPC_allocate || (Kind == OMPC_affinity && Tok.is(tok::identifier) && PP.getSpelling(Tok) == "iterator")) { @@ -3635,6 +3653,7 @@ Kind, Vars, Data.DepModOrTailExpr, Locs, Data.ColonLoc, Data.ReductionOrMapperIdScopeSpec, Data.ReductionOrMapperId, Data.ExtraModifier, Data.MapTypeModifiers, Data.MapTypeModifiersLoc, - Data.IsMapTypeImplicit, Data.ExtraModifierLoc); + Data.IsMapTypeImplicit, Data.ExtraModifierLoc, Data.MotionModifiers, + Data.MotionModifiersLoc); } Index: clang/lib/Sema/SemaOpenMP.cpp =================================================================== --- clang/lib/Sema/SemaOpenMP.cpp +++ clang/lib/Sema/SemaOpenMP.cpp @@ -13536,7 +13536,9 @@ DeclarationNameInfo &ReductionOrMapperId, int ExtraModifier, ArrayRef MapTypeModifiers, ArrayRef MapTypeModifiersLoc, bool IsMapTypeImplicit, - SourceLocation ExtraModifierLoc) { + SourceLocation ExtraModifierLoc, + ArrayRef MotionModifiers, + ArrayRef MotionModifiersLoc) { SourceLocation StartLoc = Locs.StartLoc; SourceLocation LParenLoc = Locs.LParenLoc; SourceLocation EndLoc = Locs.EndLoc; @@ -13613,12 +13615,14 @@ IsMapTypeImplicit, ExtraModifierLoc, ColonLoc, VarList, Locs); break; case OMPC_to: - Res = ActOnOpenMPToClause(VarList, ReductionOrMapperIdScopeSpec, - ReductionOrMapperId, Locs); + Res = ActOnOpenMPToClause(MotionModifiers, MotionModifiersLoc, + ReductionOrMapperIdScopeSpec, ReductionOrMapperId, + ColonLoc, VarList, Locs); break; case OMPC_from: - Res = ActOnOpenMPFromClause(VarList, ReductionOrMapperIdScopeSpec, - ReductionOrMapperId, Locs); + Res = ActOnOpenMPFromClause(MotionModifiers, MotionModifiersLoc, + ReductionOrMapperIdScopeSpec, + ReductionOrMapperId, ColonLoc, VarList, Locs); break; case OMPC_use_device_ptr: Res = ActOnOpenMPUseDevicePtrClause(VarList, Locs); @@ -18576,11 +18580,31 @@ checkDeclInTargetContext(E->getExprLoc(), E->getSourceRange(), *this, D); } -OMPClause *Sema::ActOnOpenMPToClause(ArrayRef VarList, - CXXScopeSpec &MapperIdScopeSpec, - DeclarationNameInfo &MapperId, - const OMPVarListLocTy &Locs, - ArrayRef UnresolvedMappers) { +OMPClause *Sema::ActOnOpenMPToClause( + ArrayRef MotionModifiers, + ArrayRef MotionModifiersLoc, + CXXScopeSpec &MapperIdScopeSpec, DeclarationNameInfo &MapperId, + SourceLocation ColonLoc, ArrayRef VarList, + const OMPVarListLocTy &Locs, ArrayRef UnresolvedMappers) { + OpenMPMotionModifierKind Modifiers[] = {OMPC_MOTION_MODIFIER_unknown, + OMPC_MOTION_MODIFIER_unknown}; + SourceLocation ModifiersLoc[NumberOfOMPMotionModifiers]; + + // Process motion-modifiers, flag errors for duplicate modifiers. + unsigned Count = 0; + for (unsigned I = 0, E = MotionModifiers.size(); I < E; ++I) { + if (MotionModifiers[I] != OMPC_MOTION_MODIFIER_unknown && + llvm::find(Modifiers, MotionModifiers[I]) != std::end(Modifiers)) { + Diag(MotionModifiersLoc[I], diag::err_omp_duplicate_motion_modifier); + continue; + } + assert(Count < NumberOfOMPMotionModifiers && + "Modifiers exceed the allowed number of motion modifiers"); + Modifiers[Count] = MotionModifiers[I]; + ModifiersLoc[Count] = MotionModifiersLoc[I]; + ++Count; + } + MappableVarListInfo MVLI(VarList); checkMappableExpressionList(*this, DSAStack, OMPC_to, MVLI, Locs.StartLoc, MapperIdScopeSpec, MapperId, UnresolvedMappers); @@ -18589,15 +18613,35 @@ return OMPToClause::Create( Context, Locs, MVLI.ProcessedVarList, MVLI.VarBaseDeclarations, - MVLI.VarComponents, MVLI.UDMapperList, + MVLI.VarComponents, MVLI.UDMapperList, Modifiers, ModifiersLoc, MapperIdScopeSpec.getWithLocInContext(Context), MapperId); } -OMPClause *Sema::ActOnOpenMPFromClause(ArrayRef VarList, - CXXScopeSpec &MapperIdScopeSpec, - DeclarationNameInfo &MapperId, - const OMPVarListLocTy &Locs, - ArrayRef UnresolvedMappers) { +OMPClause *Sema::ActOnOpenMPFromClause( + ArrayRef MotionModifiers, + ArrayRef MotionModifiersLoc, + CXXScopeSpec &MapperIdScopeSpec, DeclarationNameInfo &MapperId, + SourceLocation ColonLoc, ArrayRef VarList, + const OMPVarListLocTy &Locs, ArrayRef UnresolvedMappers) { + OpenMPMotionModifierKind Modifiers[] = {OMPC_MOTION_MODIFIER_unknown, + OMPC_MOTION_MODIFIER_unknown}; + SourceLocation ModifiersLoc[NumberOfOMPMotionModifiers]; + + // Process motion-modifiers, flag errors for duplicate modifiers. + unsigned Count = 0; + for (unsigned I = 0, E = MotionModifiers.size(); I < E; ++I) { + if (MotionModifiers[I] != OMPC_MOTION_MODIFIER_unknown && + llvm::find(Modifiers, MotionModifiers[I]) != std::end(Modifiers)) { + Diag(MotionModifiersLoc[I], diag::err_omp_duplicate_motion_modifier); + continue; + } + assert(Count < NumberOfOMPMotionModifiers && + "Modifiers exceed the allowed number of motion modifiers"); + Modifiers[Count] = MotionModifiers[I]; + ModifiersLoc[Count] = MotionModifiersLoc[I]; + ++Count; + } + MappableVarListInfo MVLI(VarList); checkMappableExpressionList(*this, DSAStack, OMPC_from, MVLI, Locs.StartLoc, MapperIdScopeSpec, MapperId, UnresolvedMappers); @@ -18606,7 +18650,7 @@ return OMPFromClause::Create( Context, Locs, MVLI.ProcessedVarList, MVLI.VarBaseDeclarations, - MVLI.VarComponents, MVLI.UDMapperList, + MVLI.VarComponents, MVLI.UDMapperList, Modifiers, ModifiersLoc, MapperIdScopeSpec.getWithLocInContext(Context), MapperId); } Index: clang/lib/Sema/TreeTransform.h =================================================================== --- clang/lib/Sema/TreeTransform.h +++ clang/lib/Sema/TreeTransform.h @@ -2006,26 +2006,32 @@ /// /// By default, performs semantic analysis to build the new statement. /// Subclasses may override this routine to provide different behavior. - OMPClause *RebuildOMPToClause(ArrayRef VarList, - CXXScopeSpec &MapperIdScopeSpec, - DeclarationNameInfo &MapperId, - const OMPVarListLocTy &Locs, - ArrayRef UnresolvedMappers) { - return getSema().ActOnOpenMPToClause(VarList, MapperIdScopeSpec, MapperId, - Locs, UnresolvedMappers); + OMPClause * + RebuildOMPToClause(ArrayRef MotionModifiers, + ArrayRef MotionModifiersLoc, + CXXScopeSpec &MapperIdScopeSpec, + DeclarationNameInfo &MapperId, SourceLocation ColonLoc, + ArrayRef VarList, const OMPVarListLocTy &Locs, + ArrayRef UnresolvedMappers) { + return getSema().ActOnOpenMPToClause(MotionModifiers, MotionModifiersLoc, + MapperIdScopeSpec, MapperId, ColonLoc, + VarList, Locs, UnresolvedMappers); } /// Build a new OpenMP 'from' clause. /// /// By default, performs semantic analysis to build the new statement. /// Subclasses may override this routine to provide different behavior. - OMPClause *RebuildOMPFromClause(ArrayRef VarList, - CXXScopeSpec &MapperIdScopeSpec, - DeclarationNameInfo &MapperId, - const OMPVarListLocTy &Locs, - ArrayRef UnresolvedMappers) { - return getSema().ActOnOpenMPFromClause(VarList, MapperIdScopeSpec, MapperId, - Locs, UnresolvedMappers); + OMPClause * + RebuildOMPFromClause(ArrayRef MotionModifiers, + ArrayRef MotionModifiersLoc, + CXXScopeSpec &MapperIdScopeSpec, + DeclarationNameInfo &MapperId, SourceLocation ColonLoc, + ArrayRef VarList, const OMPVarListLocTy &Locs, + ArrayRef UnresolvedMappers) { + return getSema().ActOnOpenMPFromClause( + MotionModifiers, MotionModifiersLoc, MapperIdScopeSpec, MapperId, + ColonLoc, VarList, Locs, UnresolvedMappers); } /// Build a new OpenMP 'use_device_ptr' clause. @@ -9738,8 +9744,9 @@ if (transformOMPMappableExprListClause( *this, C, Vars, MapperIdScopeSpec, MapperIdInfo, UnresolvedMappers)) return nullptr; - return getDerived().RebuildOMPToClause(Vars, MapperIdScopeSpec, MapperIdInfo, - Locs, UnresolvedMappers); + return getDerived().RebuildOMPToClause( + C->getMotionModifiers(), C->getMotionModifiersLoc(), MapperIdScopeSpec, + MapperIdInfo, C->getColonLoc(), Vars, Locs, UnresolvedMappers); } template @@ -9753,7 +9760,8 @@ *this, C, Vars, MapperIdScopeSpec, MapperIdInfo, UnresolvedMappers)) return nullptr; return getDerived().RebuildOMPFromClause( - Vars, MapperIdScopeSpec, MapperIdInfo, Locs, UnresolvedMappers); + C->getMotionModifiers(), C->getMotionModifiersLoc(), MapperIdScopeSpec, + MapperIdInfo, C->getColonLoc(), Vars, Locs, UnresolvedMappers); } template Index: clang/lib/Serialization/ASTReader.cpp =================================================================== --- clang/lib/Serialization/ASTReader.cpp +++ clang/lib/Serialization/ASTReader.cpp @@ -12601,8 +12601,14 @@ void OMPClauseReader::VisitOMPToClause(OMPToClause *C) { C->setLParenLoc(Record.readSourceLocation()); + for (unsigned I = 0; I < NumberOfOMPMotionModifiers; ++I) { + C->setMotionModifier( + I, static_cast(Record.readInt())); + C->setMotionModifierLoc(I, Record.readSourceLocation()); + } C->setMapperQualifierLoc(Record.readNestedNameSpecifierLoc()); C->setMapperIdInfo(Record.readDeclarationNameInfo()); + C->setColonLoc(Record.readSourceLocation()); auto NumVars = C->varlist_size(); auto UniqueDecls = C->getUniqueDeclarationsNum(); auto TotalLists = C->getTotalComponentListNum(); @@ -12651,8 +12657,14 @@ void OMPClauseReader::VisitOMPFromClause(OMPFromClause *C) { C->setLParenLoc(Record.readSourceLocation()); + for (unsigned I = 0; I < NumberOfOMPMotionModifiers; ++I) { + C->setMotionModifier( + I, static_cast(Record.readInt())); + C->setMotionModifierLoc(I, Record.readSourceLocation()); + } C->setMapperQualifierLoc(Record.readNestedNameSpecifierLoc()); C->setMapperIdInfo(Record.readDeclarationNameInfo()); + C->setColonLoc(Record.readSourceLocation()); auto NumVars = C->varlist_size(); auto UniqueDecls = C->getUniqueDeclarationsNum(); auto TotalLists = C->getTotalComponentListNum(); Index: clang/lib/Serialization/ASTWriter.cpp =================================================================== --- clang/lib/Serialization/ASTWriter.cpp +++ clang/lib/Serialization/ASTWriter.cpp @@ -6581,8 +6581,13 @@ Record.push_back(C->getTotalComponentListNum()); Record.push_back(C->getTotalComponentsNum()); Record.AddSourceLocation(C->getLParenLoc()); + for (unsigned I = 0; I < NumberOfOMPMotionModifiers; ++I) { + Record.push_back(C->getMotionModifier(I)); + Record.AddSourceLocation(C->getMotionModifierLoc(I)); + } Record.AddNestedNameSpecifierLoc(C->getMapperQualifierLoc()); Record.AddDeclarationNameInfo(C->getMapperIdInfo()); + Record.AddSourceLocation(C->getColonLoc()); for (auto *E : C->varlists()) Record.AddStmt(E); for (auto *E : C->mapperlists()) @@ -6605,8 +6610,13 @@ Record.push_back(C->getTotalComponentListNum()); Record.push_back(C->getTotalComponentsNum()); Record.AddSourceLocation(C->getLParenLoc()); + for (unsigned I = 0; I < NumberOfOMPMotionModifiers; ++I) { + Record.push_back(C->getMotionModifier(I)); + Record.AddSourceLocation(C->getMotionModifierLoc(I)); + } Record.AddNestedNameSpecifierLoc(C->getMapperQualifierLoc()); Record.AddDeclarationNameInfo(C->getMapperIdInfo()); + Record.AddSourceLocation(C->getColonLoc()); for (auto *E : C->varlists()) Record.AddStmt(E); for (auto *E : C->mapperlists()) Index: clang/test/OpenMP/declare_mapper_ast_print.c =================================================================== --- clang/test/OpenMP/declare_mapper_ast_print.c +++ clang/test/OpenMP/declare_mapper_ast_print.c @@ -5,6 +5,14 @@ // RUN: %clang_cc1 -verify -fopenmp-simd -ast-print %s | FileCheck %s // RUN: %clang_cc1 -fopenmp-simd -emit-pch -o %t %s // RUN: %clang_cc1 -fopenmp-simd -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s + +// RUN: %clang_cc1 -DOMP51 -verify -fopenmp -fopenmp-version=51 -ast-print %s | FileCheck -check-prefixes=CHECK,OMP51 %s +// RUN: %clang_cc1 -DOMP51 -fopenmp -fopenmp-version=51 -emit-pch -o %t %s +// RUN: %clang_cc1 -DOMP51 -fopenmp -fopenmp-version=51 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck -check-prefixes=CHECK,OMP51 %s + +// RUN: %clang_cc1 -DOMP51 -verify -fopenmp-simd -fopenmp-version=51 -ast-print %s | FileCheck -check-prefixes=CHECK,OMP51 %s +// RUN: %clang_cc1 -DOMP51 -fopenmp-simd -fopenmp-version=51 -emit-pch -o %t %s +// RUN: %clang_cc1 -DOMP51 -fopenmp-simd -fopenmp-version=51 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck -check-prefixes=CHECK,OMP51 %s // expected-no-diagnostics #ifndef HEADER @@ -50,6 +58,12 @@ { dd[0].i++; } #pragma omp target update to(mapper(id): vv) from(mapper(default): dd[0:10]) // CHECK: #pragma omp target update to(mapper(id): vv) from(mapper(default): dd[0:10]) +#ifdef OMP51 +#pragma omp target update to(mapper(id) present: vv) from(mapper(default), present: dd[0:10]) +// OMP51: #pragma omp target update to(mapper(id), present: vv) from(mapper(default), present: dd[0:10]) +#pragma omp target update to(present mapper(id): vv) from(present, mapper(default): dd[0:10]) +// OMP51: #pragma omp target update to(present, mapper(id): vv) from(present, mapper(default): dd[0:10]) +#endif } return 0; } Index: clang/test/OpenMP/declare_mapper_codegen.cpp =================================================================== --- clang/test/OpenMP/declare_mapper_codegen.cpp +++ clang/test/OpenMP/declare_mapper_codegen.cpp @@ -839,4 +839,242 @@ #endif // CK3 +///==========================================================================/// +// RUN: %clang_cc1 -DCK4 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK4 --check-prefix CK4-64 %s +// RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s +// RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK4 --check-prefix CK4-64 %s +// RUN: %clang_cc1 -DCK4 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix CK4 --check-prefix CK4-32 %s +// RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s +// RUN: %clang_cc1 -DCK4 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix CK4 --check-prefix CK4-32 %s + +// RUN: %clang_cc1 -DCK4 -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s +// RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s +// RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s +// RUN: %clang_cc1 -DCK4 -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm -femit-all-decls -disable-llvm-passes %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s +// RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -femit-all-decls -disable-llvm-passes -o %t %s +// RUN: %clang_cc1 -DCK4 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -femit-all-decls -disable-llvm-passes -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s + +#ifdef CK4 +// Mapper function code generation and runtime interface. + +// CK4-64: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16] +// CK4-32: [[TSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8] +// PRESENT=0x1000 | TARGET_PARAM=0x20 | TO=0x1 = 0x1021 +// CK4: [[TTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1021]]] + +// CK4-64: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16] +// CK4-32: [[FSIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8] +// PRESENT=0x1000 | TARGET_PARAM=0x20 | FROM=0x2 = 0x1022 +// CK4: [[FTYPES:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1022]]] + +class C { +public: + int a; + double *b; +}; + +#pragma omp declare mapper(id: C s) map(s.a, s.b[0:2]) + +// CK4: define {{.*}}void [[MPRFUNC:@[.]omp_mapper[.].*C[.]id]](i8*{{.*}}, i8*{{.*}}, i8*{{.*}}, i64{{.*}}, i64{{.*}}) +// CK4: store i8* %{{[^,]+}}, i8** [[HANDLEADDR:%[^,]+]] +// CK4: store i8* %{{[^,]+}}, i8** [[BPTRADDR:%[^,]+]] +// CK4: store i8* %{{[^,]+}}, i8** [[VPTRADDR:%[^,]+]] +// CK4: store i64 %{{[^,]+}}, i{{64|32}}* [[SIZEADDR:%[^,]+]] +// CK4: store i64 %{{[^,]+}}, i64* [[TYPEADDR:%[^,]+]] +// CK4-DAG: [[BYTESIZE:%.+]] = load i64, i64* [[SIZEADDR]] +// CK4-64-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 16 +// CK4-32-DAG: [[SIZE:%.+]] = udiv exact i64 [[BYTESIZE]], 8 +// CK4-DAG: [[TYPE:%.+]] = load i64, i64* [[TYPEADDR]] +// CK4-DAG: [[HANDLE:%.+]] = load i8*, i8** [[HANDLEADDR]] +// CK4-DAG: [[PTRBEGIN:%.+]] = bitcast i8** [[VPTRADDR]] to %class.C** +// CK4-DAG: [[PTREND:%.+]] = getelementptr %class.C*, %class.C** [[PTRBEGIN]], i64 [[SIZE]] +// CK4-DAG: [[BPTR:%.+]] = load i8*, i8** [[BPTRADDR]] +// CK4-DAG: [[BEGIN:%.+]] = load i8*, i8** [[VPTRADDR]] +// CK4: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1 +// CK4: br i1 [[ISARRAY]], label %[[INITEVALDEL:[^,]+]], label %[[LHEAD:[^,]+]] + +// CK4: [[INITEVALDEL]] +// CK4: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 +// CK4: [[ISNOTDEL:%.+]] = icmp eq i64 [[TYPEDEL]], 0 +// CK4: br i1 [[ISNOTDEL]], label %[[INIT:[^,]+]], label %[[LHEAD:[^,]+]] +// CK4: [[INIT]] +// CK4-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16 +// CK4-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8 +// CK4-DAG: [[ITYPE:%.+]] = and i64 [[TYPE]], -4 +// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[ITYPE]]) +// CK4: br label %[[LHEAD:[^,]+]] + +// CK4: [[LHEAD]] +// CK4: [[ISEMPTY:%.+]] = icmp eq %class.C** [[PTRBEGIN]], [[PTREND]] +// CK4: br i1 [[ISEMPTY]], label %[[DONE:[^,]+]], label %[[LBODY:[^,]+]] +// CK4: [[LBODY]] +// CK4: [[PTR:%.+]] = phi %class.C** [ [[PTRBEGIN]], %[[LHEAD]] ], [ [[PTRNEXT:%.+]], %[[LCORRECT:[^,]+]] ] +// CK4: [[OBJ:%.+]] = load %class.C*, %class.C** [[PTR]] +// CK4-DAG: [[ABEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 0 +// CK4-DAG: [[BBEGIN:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1 +// CK4-DAG: [[BBEGIN2:%.+]] = getelementptr inbounds %class.C, %class.C* [[OBJ]], i32 0, i32 1 +// CK4-DAG: [[BARRBEGIN:%.+]] = load double*, double** [[BBEGIN2]] +// CK4-DAG: [[BARRBEGINGEP:%.+]] = getelementptr inbounds double, double* [[BARRBEGIN]], i[[sz:64|32]] 0 +// CK4-DAG: [[BEND:%.+]] = getelementptr double*, double** [[BBEGIN]], i32 1 +// CK4-DAG: [[ABEGINV:%.+]] = bitcast i32* [[ABEGIN]] to i8* +// CK4-DAG: [[BENDV:%.+]] = bitcast double** [[BEND]] to i8* +// CK4-DAG: [[ABEGINI:%.+]] = ptrtoint i8* [[ABEGINV]] to i64 +// CK4-DAG: [[BENDI:%.+]] = ptrtoint i8* [[BENDV]] to i64 +// CK4-DAG: [[CSIZE:%.+]] = sub i64 [[BENDI]], [[ABEGINI]] +// CK4-DAG: [[CUSIZE:%.+]] = sdiv exact i64 [[CSIZE]], ptrtoint (i8* getelementptr (i8, i8* null, i32 1) to i64) +// CK4-DAG: [[BPTRADDR0BC:%.+]] = bitcast %class.C* [[OBJ]] to i8* +// CK4-DAG: [[PTRADDR0BC:%.+]] = bitcast i32* [[ABEGIN]] to i8* +// CK4-DAG: [[PRESIZE:%.+]] = call i64 @__tgt_mapper_num_components(i8* [[HANDLE]]) +// CK4-DAG: [[SHIPRESIZE:%.+]] = shl i64 [[PRESIZE]], 48 +// CK4-DAG: br label %[[MEMBER:[^,]+]] +// CK4-DAG: [[MEMBER]] +// CK4-DAG: br i1 true, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]] +// CK4-DAG: [[MEMBERCOM]] +// CK4-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 32, [[SHIPRESIZE]] +// CK4-DAG: br label %[[LTYPE]] +// CK4-DAG: [[LTYPE]] +// CK4-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 32, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ] +// CK4-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3 +// CK4-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0 +// CK4-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]] +// CK4-DAG: [[ALLOC]] +// CK4-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4 +// CK4-DAG: br label %[[TYEND:[^,]+]] +// CK4-DAG: [[ALLOCELSE]] +// CK4-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1 +// CK4-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]] +// CK4-DAG: [[TO]] +// CK4-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3 +// CK4-DAG: br label %[[TYEND]] +// CK4-DAG: [[TOELSE]] +// CK4-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2 +// CK4-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]] +// CK4-DAG: [[FROM]] +// CK4-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2 +// CK4-DAG: br label %[[TYEND]] +// CK4-DAG: [[TYEND]] +// CK4-DAG: [[PHITYPE0:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ] +// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR0BC]], i8* [[PTRADDR0BC]], i64 [[CUSIZE]], i64 [[PHITYPE0]]) +// CK4-DAG: [[BPTRADDR1BC:%.+]] = bitcast %class.C* [[OBJ]] to i8* +// CK4-DAG: [[PTRADDR1BC:%.+]] = bitcast i32* [[ABEGIN]] to i8* +// CK4-DAG: br label %[[MEMBER:[^,]+]] +// CK4-DAG: [[MEMBER]] +// CK4-DAG: br i1 false, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]] +// CK4-DAG: [[MEMBERCOM]] +// 281474976710659 == 0x1,000,000,003 +// CK4-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 281474976710659, [[SHIPRESIZE]] +// CK4-DAG: br label %[[LTYPE]] +// CK4-DAG: [[LTYPE]] +// CK4-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 281474976710659, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ] +// CK4-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3 +// CK4-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0 +// CK4-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]] +// CK4-DAG: [[ALLOC]] +// CK4-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4 +// CK4-DAG: br label %[[TYEND:[^,]+]] +// CK4-DAG: [[ALLOCELSE]] +// CK4-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1 +// CK4-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]] +// CK4-DAG: [[TO]] +// CK4-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3 +// CK4-DAG: br label %[[TYEND]] +// CK4-DAG: [[TOELSE]] +// CK4-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2 +// CK4-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]] +// CK4-DAG: [[FROM]] +// CK4-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2 +// CK4-DAG: br label %[[TYEND]] +// CK4-DAG: [[TYEND]] +// CK4-DAG: [[TYPE1:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ] +// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR1BC]], i8* [[PTRADDR1BC]], i64 4, i64 [[TYPE1]]) +// CK4-DAG: [[BPTRADDR2BC:%.+]] = bitcast double** [[BBEGIN]] to i8* +// CK4-DAG: [[PTRADDR2BC:%.+]] = bitcast double* [[BARRBEGINGEP]] to i8* +// CK4-DAG: br label %[[MEMBER:[^,]+]] +// CK4-DAG: [[MEMBER]] +// CK4-DAG: br i1 false, label %[[LTYPE:[^,]+]], label %[[MEMBERCOM:[^,]+]] +// CK4-DAG: [[MEMBERCOM]] +// 281474976710675 == 0x1,000,000,013 +// CK4-DAG: [[MEMBERCOMTYPE:%.+]] = add nuw i64 281474976710675, [[SHIPRESIZE]] +// CK4-DAG: br label %[[LTYPE]] +// CK4-DAG: [[LTYPE]] +// CK4-DAG: [[MEMBERTYPE:%.+]] = phi i64 [ 281474976710675, %[[MEMBER]] ], [ [[MEMBERCOMTYPE]], %[[MEMBERCOM]] ] +// CK4-DAG: [[TYPETF:%.+]] = and i64 [[TYPE]], 3 +// CK4-DAG: [[ISALLOC:%.+]] = icmp eq i64 [[TYPETF]], 0 +// CK4-DAG: br i1 [[ISALLOC]], label %[[ALLOC:[^,]+]], label %[[ALLOCELSE:[^,]+]] +// CK4-DAG: [[ALLOC]] +// CK4-DAG: [[ALLOCTYPE:%.+]] = and i64 [[MEMBERTYPE]], -4 +// CK4-DAG: br label %[[TYEND:[^,]+]] +// CK4-DAG: [[ALLOCELSE]] +// CK4-DAG: [[ISTO:%.+]] = icmp eq i64 [[TYPETF]], 1 +// CK4-DAG: br i1 [[ISTO]], label %[[TO:[^,]+]], label %[[TOELSE:[^,]+]] +// CK4-DAG: [[TO]] +// CK4-DAG: [[TOTYPE:%.+]] = and i64 [[MEMBERTYPE]], -3 +// CK4-DAG: br label %[[TYEND]] +// CK4-DAG: [[TOELSE]] +// CK4-DAG: [[ISFROM:%.+]] = icmp eq i64 [[TYPETF]], 2 +// CK4-DAG: br i1 [[ISFROM]], label %[[FROM:[^,]+]], label %[[TYEND]] +// CK4-DAG: [[FROM]] +// CK4-DAG: [[FROMTYPE:%.+]] = and i64 [[MEMBERTYPE]], -2 +// CK4-DAG: br label %[[TYEND]] +// CK4-DAG: [[TYEND]] +// CK4-DAG: [[TYPE2:%.+]] = phi i64 [ [[ALLOCTYPE]], %[[ALLOC]] ], [ [[TOTYPE]], %[[TO]] ], [ [[FROMTYPE]], %[[FROM]] ], [ [[MEMBERTYPE]], %[[TOELSE]] ] +// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTRADDR2BC]], i8* [[PTRADDR2BC]], i64 16, i64 [[TYPE2]]) +// CK4: [[PTRNEXT]] = getelementptr %class.C*, %class.C** [[PTR]], i32 1 +// CK4: [[ISDONE:%.+]] = icmp eq %class.C** [[PTRNEXT]], [[PTREND]] +// CK4: br i1 [[ISDONE]], label %[[LEXIT:[^,]+]], label %[[LBODY]] + +// CK4: [[LEXIT]] +// CK4: [[ISARRAY:%.+]] = icmp sge i64 [[SIZE]], 1 +// CK4: br i1 [[ISARRAY]], label %[[EVALDEL:[^,]+]], label %[[DONE]] +// CK4: [[EVALDEL]] +// CK4: [[TYPEDEL:%.+]] = and i64 [[TYPE]], 8 +// CK4: [[ISDEL:%.+]] = icmp ne i64 [[TYPEDEL]], 0 +// CK4: br i1 [[ISDEL]], label %[[DEL:[^,]+]], label %[[DONE]] +// CK4: [[DEL]] +// CK4-64-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 16 +// CK4-32-DAG: [[ARRSIZE:%.+]] = mul nuw i64 [[SIZE]], 8 +// CK4-DAG: [[DTYPE:%.+]] = and i64 [[TYPE]], -4 +// CK4: call void @__tgt_push_mapper_component(i8* [[HANDLE]], i8* [[BPTR]], i8* [[BEGIN]], i64 [[ARRSIZE]], i64 [[DTYPE]]) +// CK4: br label %[[DONE]] +// CK4: [[DONE]] +// CK4: ret void + + +// CK4-LABEL: define {{.*}}void @{{.*}}foo{{.*}} +void foo(int a){ + int i = a; + C c; + c.a = a; + + // CK4-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[TGEPBP:%.+]], i8** [[TGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[TSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[TTYPES]]{{.+}}, i8** [[TMPRGEP:%.+]]) + // CK4-DAG: [[TGEPBP]] = getelementptr inbounds {{.+}}[[TBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0 + // CK4-DAG: [[TGEPP]] = getelementptr inbounds {{.+}}[[TP:%[^,]+]], i{{.+}} 0, i{{.+}} 0 + // CK4-DAG: [[TMPRGEP]] = bitcast [1 x i8*]* [[TMPR:%[^,]+]] to i8** + // CK4-DAG: [[TBP0:%.+]] = getelementptr inbounds {{.+}}[[TBP]], i{{.+}} 0, i{{.+}} 0 + // CK4-DAG: [[TP0:%.+]] = getelementptr inbounds {{.+}}[[TP]], i{{.+}} 0, i{{.+}} 0 + // CK4-DAG: [[TMPR1:%.+]] = getelementptr inbounds {{.+}}[[TMPR]], i[[sz]] 0, i[[sz]] 0 + // CK4-DAG: [[TCBP0:%.+]] = bitcast i8** [[TBP0]] to %class.C** + // CK4-DAG: [[TCP0:%.+]] = bitcast i8** [[TP0]] to %class.C** + // CK4-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[TCBP0]] + // CK4-DAG: store %class.C* [[VAL]], %class.C** [[TCP0]] + // CK4-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[TMPR1]] + #pragma omp target update to(present, mapper(id): c) + + // CK4-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[FGEPBP:%.+]], i8** [[FGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[FSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[FTYPES]]{{.+}}, i8** [[FMPRGEP:%.+]]) + // CK4-DAG: [[FGEPBP]] = getelementptr inbounds {{.+}}[[FBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0 + // CK4-DAG: [[FGEPP]] = getelementptr inbounds {{.+}}[[FP:%[^,]+]], i{{.+}} 0, i{{.+}} 0 + // CK4-DAG: [[FMPRGEP]] = bitcast [1 x i8*]* [[FMPR:%[^,]+]] to i8** + // CK4-DAG: [[FBP0:%.+]] = getelementptr inbounds {{.+}}[[FBP]], i{{.+}} 0, i{{.+}} 0 + // CK4-DAG: [[FP0:%.+]] = getelementptr inbounds {{.+}}[[FP]], i{{.+}} 0, i{{.+}} 0 + // CK4-DAG: [[FMPR1:%.+]] = getelementptr inbounds {{.+}}[[FMPR]], i[[sz]] 0, i[[sz]] 0 + // CK4-DAG: [[FCBP0:%.+]] = bitcast i8** [[FBP0]] to %class.C** + // CK4-DAG: [[FCP0:%.+]] = bitcast i8** [[FP0]] to %class.C** + // CK4-DAG: store %class.C* [[VAL]], %class.C** [[FCBP0]] + // CK4-DAG: store %class.C* [[VAL]], %class.C** [[FCP0]] + // CK4-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[FMPR1]] + #pragma omp target update from(mapper(id), present: c) +} + +#endif // CK4 + #endif // HEADER Index: clang/test/OpenMP/target_update_ast_print.cpp =================================================================== --- clang/test/OpenMP/target_update_ast_print.cpp +++ clang/test/OpenMP/target_update_ast_print.cpp @@ -5,6 +5,14 @@ // RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=50 -ast-print %s | FileCheck %s // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -x c++ -std=c++11 -emit-pch -o %t %s // RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s + +// RUN: %clang_cc1 -DOMP51 -verify -fopenmp -fopenmp-version=51 -ast-print %s | FileCheck -check-prefixes=CHECK,OMP51 %s +// RUN: %clang_cc1 -DOMP51 -fopenmp -fopenmp-version=51 -x c++ -std=c++11 -emit-pch -o %t %s +// RUN: %clang_cc1 -DOMP51 -fopenmp -fopenmp-version=51 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck -check-prefixes=CHECK,OMP51 %s + +// RUN: %clang_cc1 -DOMP51 -verify -fopenmp-simd -fopenmp-version=51 -ast-print %s | FileCheck -check-prefixes=CHECK,OMP51 %s +// RUN: %clang_cc1 -DOMP51 -fopenmp-simd -fopenmp-version=51 -x c++ -std=c++11 -emit-pch -o %t %s +// RUN: %clang_cc1 -DOMP51 -fopenmp-simd -fopenmp-version=51 -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck -check-prefixes=CHECK,OMP51 %s // expected-no-diagnostics #ifndef HEADER @@ -25,6 +33,13 @@ #pragma omp target update to(arr[2][0:1:2]) #pragma omp target update from(arr[2][0:1:2]) + +#ifdef OMP51 +#pragma omp target update to(present: arr[2][0:1:2]) + +#pragma omp target update from(present: arr[2][0:1:2], a) +#endif + return a + targ + (T)b; } // CHECK: static T a, *p; @@ -45,6 +60,8 @@ // CHECK: int arr[100][100]; // CHECK-NEXT: #pragma omp target update to(arr[2][0:1:2]) // CHECK-NEXT: #pragma omp target update from(arr[2][0:1:2]) +// OMP5-NEXT: #pragma omp target update to(present: arr[2][0:1:2]) +// OMP5-NEXT: #pragma omp target update from(present: arr[2][0:1:2], a) int main(int argc, char **argv) { static int a; @@ -62,6 +79,13 @@ // CHECK-NEXT: #pragma omp target update to(argv[2][0:1:2]) #pragma omp target update from(argv[2][0:1:2]) // CHECK-NEXT: #pragma omp target update from(argv[2][0:1:2]) +#ifdef OMP51 +#pragma omp target update to(present: argv[2][0:1:2]) +// OMP5-NEXT: #pragma omp target update to(present: arr[2][0:1:2]) +#pragma omp target update from(argv[2][0:1:2], a) +// OMP5-NEXT: #pragma omp target update from(present: arr[2][0:1:2], a) +#endif + return foo(argc, f) + foo(argv[0][0], f) + a; } Index: clang/test/OpenMP/target_update_codegen.cpp =================================================================== --- clang/test/OpenMP/target_update_codegen.cpp +++ clang/test/OpenMP/target_update_codegen.cpp @@ -1084,5 +1084,68 @@ #pragma omp target update from(([sa][5])f) } +#endif +///==========================================================================/// +// RUN: %clang_cc1 -DCK19 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK19 --check-prefix CK19-64 +// RUN: %clang_cc1 -DCK19 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK19 --check-prefix CK19-64 +// RUN: %clang_cc1 -DCK19 -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK19 --check-prefix CK19-32 +// RUN: %clang_cc1 -DCK19 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK19 --check-prefix CK19-32 + +// RUN: %clang_cc1 -DCK19 -verify -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s +// RUN: %clang_cc1 -DCK19 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s +// RUN: %clang_cc1 -DCK19 -verify -fopenmp-version=51 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s +// RUN: %clang_cc1 -DCK19 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s +// SIMD-ONLY0-NOT: {{__kmpc|__tgt}} +#ifdef CK19 + +// PRESENT=0x1000 | TARGET_PARAM=0x20 | TO=0x1 = 0x1021 +// CK19: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1021]]] + +// PRESENT=0x1000 | TARGET_PARAM=0x20 | FROM=0x2 = 0x1022 +// CK19: [[MTYPE01:@.+]] = {{.+}}constant [1 x i64] [i64 [[#0x1022]]] + +// CK19-LABEL: _Z13check_presenti +void check_present(int arg) { + int la; + float lb[arg]; + + // Region 00 + // CK19-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] + + // 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 float** + // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float** + // CK19-DAG: store float* [[VAL0:%[^,]+]], float** [[CBP0]] + // CK19-DAG: store float* [[VAL0]], float** [[CP0]] + // CK19-DAG: store i64 [[CSVAL0:%[^,]+]], i64* [[S0]] + #pragma omp target update to(present: lb) + ; + + // Region 01 + // CK19-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}}, i8** null) + // CK19-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] + // CK19-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] + // CK19-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] + + // 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 float** + // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float** + // CK19-DAG: store float* [[VAL0:%[^,]+]], float** [[CBP0]] + // CK19-DAG: store float* [[VAL0]], float** [[CP0]] + // CK19-DAG: store i64 [[CSVAL0:%[^,]+]], i64* [[S0]] + #pragma omp target update from(present: lb) + ; +} #endif #endif Index: clang/test/OpenMP/target_update_messages.cpp =================================================================== --- clang/test/OpenMP/target_update_messages.cpp +++ clang/test/OpenMP/target_update_messages.cpp @@ -1,8 +1,10 @@ -// RUN: %clang_cc1 -verify=expected,omp45 -fopenmp -fopenmp-version=45 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized -// RUN: %clang_cc1 -verify=expected,omp50 -fopenmp -fopenmp-version=50 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp -fopenmp-version=45 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp -fopenmp-version=50 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp -fopenmp-version=51 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized -// RUN: %clang_cc1 -verify=expected,omp45 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized -// RUN: %clang_cc1 -verify=expected,omp50 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,lt50,lt51 -fopenmp-simd -fopenmp-version=45 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,ge50,lt51 -fopenmp-simd -fopenmp-version=50 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized +// RUN: %clang_cc1 -verify=expected,ge50,ge51 -fopenmp-simd -fopenmp-version=51 -ferror-limit 100 -o - -std=c++11 %s -Wuninitialized void xxx(int argc) { int x; // expected-note {{initialize the variable 'x' to silence this warning}} @@ -25,6 +27,10 @@ return 0; } +struct S { + int i; +}; + int main(int argc, char **argv) { int m; #pragma omp target update // expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} @@ -34,15 +40,114 @@ #pragma omp target update to(m) ] // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} #pragma omp target update to(m) ) // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} + #pragma omp declare mapper(id: S s) map(s.i) + S s; + + // Check parsing with no modifiers. + // lt51-error@+2 {{expected expression}} + // lt51-error@+1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + #pragma omp target update to(: s) + // expected-error@+2 {{expected expression}} + // expected-error@+1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + #pragma omp target update to(:) + // expected-error@+2 2 {{expected expression}} + // expected-error@+1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + #pragma omp target update to(,:) + + // Check parsing with one modifier. + // expected-error@+2 {{use of undeclared identifier 'foobar'}} + // expected-error@+1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + #pragma omp target update to(foobar: s) + // expected-error@+3 {{expected ',' or ')' in 'to' clause}} + // expected-error@+2 {{expected ')'}} + // expected-note@+1 {{to match this '('}} + #pragma omp target update to(m: s) + #pragma omp target update to(mapper(id): s) + // lt51-error@+2 {{use of undeclared identifier 'present'}} + // lt51-error@+1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + #pragma omp target update to(present: s) + // ge51-warning@+4 {{missing ':' after motion modifier - ignoring}} + // lt51-warning@+3 {{missing ':' after ) - ignoring}} + // expected-error@+2 {{expected expression}} + // expected-error@+1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + #pragma omp target update to(mapper(id) s) + // ge51-warning@+4 {{missing ':' after motion modifier - ignoring}} + // ge51-error@+3 {{expected expression}} + // lt51-error@+2 {{use of undeclared identifier 'present'}} + // expected-error@+1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + #pragma omp target update to(present s) + // ge51-warning@+4 {{missing ':' after motion modifier - ignoring}} + // lt51-warning@+3 {{missing ':' after ) - ignoring}} + // expected-error@+2 {{expected expression}} + // expected-error@+1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + #pragma omp target update to(mapper(id)) + // ge51-warning@+4 {{missing ':' after motion modifier - ignoring}} + // ge51-error@+3 {{expected expression}} + // lt51-error@+2 {{use of undeclared identifier 'present'}} + // expected-error@+1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + #pragma omp target update to(present) + // expected-error@+2 {{expected expression}} + // expected-error@+1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + #pragma omp target update to(mapper(id):) + // ge51-error@+3 {{expected expression}} + // lt51-error@+2 {{use of undeclared identifier 'present'}} + // expected-error@+1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + #pragma omp target update to(present:) + + // Check parsing with two modifiers. + // lt51-warning@+1 {{missing ':' after ) - ignoring}} + #pragma omp target update to(mapper(id), present: s) + // lt51-error@+3 {{use of undeclared identifier 'present'}} + // lt51-error@+2 {{use of undeclared identifier 'id'}} + // lt51-error@+1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + #pragma omp target update to(present, mapper(id): s) + // lt51-warning@+1 {{missing ':' after ) - ignoring}} + #pragma omp target update to(mapper(id) present: s) + // lt51-error@+2 {{use of undeclared identifier 'present'}} + // lt51-error@+1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + #pragma omp target update to(present mapper(id): s) + + // Check parsing with unnecessary commas. + // lt51-warning@+1 {{missing ':' after ) - ignoring}} + #pragma omp target update to(mapper(id),: s) + // lt51-error@+3 {{use of undeclared identifier 'present'}} + // lt51-error@+2 {{expected expression}} + // lt51-error@+1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + #pragma omp target update to(present , : s) + // ge51-warning@+2 {{missing ':' after motion modifier - ignoring}} + // lt51-warning@+1 {{missing ':' after ) - ignoring}} + #pragma omp target update to(mapper(id),,: s) + // ge51-warning@+5 {{missing ':' after motion modifier - ignoring}} + // lt51-error@+4 {{use of undeclared identifier 'present'}} + // lt51-error@+3 {{expected expression}} + // lt51-error@+2 {{expected expression}} + // lt51-error@+1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + #pragma omp target update to(present,,: s) + // lt51-warning@+1 {{missing ':' after ) - ignoring}} + #pragma omp target update to(mapper(id), present,: s) + // lt51-error@+4 {{use of undeclared identifier 'present'}} + // lt51-error@+3 {{use of undeclared identifier 'id'}} + // lt51-error@+2 {{expected expression}} + // lt51-error@+1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} + #pragma omp target update to(present, mapper(id),: s) + #pragma omp target update from(m) allocate(m) // expected-error {{unexpected OpenMP clause 'allocate' in directive '#pragma omp target update'}} { foo(); } int iarr[5][5]; -#pragma omp target update to(iarr[0:][1:2:-1]) // omp50-error {{section stride is evaluated to a non-positive value -1}} omp45-error {{expected ']'}} omp45-note {{to match this '['}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +// ge50-error@+4 {{section stride is evaluated to a non-positive value -1}} +// lt50-error@+3 {{expected ']'}} +// lt50-note@+2 {{to match this '['}} +// expected-error@+1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(iarr[0:][1:2:-1]) {} -#pragma omp target update from(iarr[0:][1:2:-1]) // omp50-error {{section stride is evaluated to a non-positive value -1}} omp45-error {{expected ']'}} omp45-note {{to match this '['}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +// ge50-error@+4 {{section stride is evaluated to a non-positive value -1}} +// lt50-error@+3 {{expected ']'}} +// lt50-note@+2 {{to match this '['}} +// expected-error@+1 {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(iarr[0:][1:2:-1]) return tmain(argc, argv); }