Index: include/clang/AST/DataRecursiveASTVisitor.h =================================================================== --- include/clang/AST/DataRecursiveASTVisitor.h +++ include/clang/AST/DataRecursiveASTVisitor.h @@ -2707,6 +2707,12 @@ return true; } +template +bool RecursiveASTVisitor::VisitOMPMapClause(OMPMapClause *C) { + TRY_TO(VisitOMPClauseList(C)); + return true; +} + // FIXME: look at the following tricky-seeming exprs to see if we // need to recurse on anything. These are ones that have methods // returning decls or qualtypes or nestednamespecifier -- though I'm Index: include/clang/AST/OpenMPClause.h =================================================================== --- include/clang/AST/OpenMPClause.h +++ include/clang/AST/OpenMPClause.h @@ -2613,6 +2613,121 @@ } }; +/// \brief This represents clause 'map' in the '#pragma omp ...' +/// directives. +/// +/// \code +/// #pragma omp target map(a,b) +/// \endcode +/// In this example directive '#pragma omp target' has clause 'map' +/// with the variables 'a' and 'b'. +/// +class OMPMapClause : public OMPVarListClause { + friend class OMPClauseReader; + + /// \brief Map type modifier for the 'map' clause. + OpenMPMapClauseKind MapTypeModifier; + /// \brief Map type for the 'map' clause. + OpenMPMapClauseKind MapType; + /// \brief Location of the map type. + SourceLocation MapLoc; + /// \brief Colon location. + SourceLocation ColonLoc; + + /// \brief Set type modifier for the clause. + /// + /// \param T Type Modifier for the clause. + /// + void setMapTypeModifier(OpenMPMapClauseKind T) { MapTypeModifier = T; } + + /// \brief Set type for the clause. + /// + /// \param T Type for the clause. + /// + void setMapType(OpenMPMapClauseKind T) { MapType = T; } + + /// \brief Set type location. + /// + /// \param TLoc Type location. + /// + void setMapLoc(SourceLocation TLoc) { MapLoc = TLoc; } + + /// \brief Set colon location. + void setColonLoc(SourceLocation Loc) { ColonLoc = Loc; } + + /// \brief Build clause with number of variables \a N. + /// + /// \param MayTypeModifier Map type modifier. + /// \param MapType Map type. + /// \param MapLoc Location of the map type. + /// \param StartLoc Starting location of the clause. + /// \param EndLoc Ending location of the clause. + /// \param N Number of the variables in the clause. + /// + explicit OMPMapClause(OpenMPMapClauseKind MapTypeModifier, + OpenMPMapClauseKind MapType, SourceLocation MapLoc, + SourceLocation StartLoc, SourceLocation LParenLoc, + SourceLocation EndLoc, unsigned N) + : OMPVarListClause(OMPC_map, StartLoc, LParenLoc, EndLoc, N), + MapTypeModifier(MapTypeModifier), MapType(MapType), MapLoc(MapLoc) {} + + /// \brief Build an empty clause. + /// + /// \param N Number of variables. + /// + explicit OMPMapClause(unsigned N) + : OMPVarListClause(OMPC_map, SourceLocation(), + SourceLocation(), SourceLocation(), N), + MapTypeModifier(OMPC_MAP_unknown), MapType(OMPC_MAP_unknown), MapLoc() {} + +public: + /// \brief Creates clause with a list of variables \a VL. + /// + /// \param C AST context. + /// \brief StartLoc Starting location of the clause. + /// \brief EndLoc Ending location of the clause. + /// \param VL List of references to the variables. + /// \param TypeModifier Map type modifier. + /// \param Type Map type. + /// \param TypeLoc Location of the map type. + /// + static OMPMapClause *Create(const ASTContext &C, SourceLocation StartLoc, + SourceLocation LParenLoc, + SourceLocation EndLoc, ArrayRef VL, + OpenMPMapClauseKind TypeModifier, + OpenMPMapClauseKind Type, SourceLocation TypeLoc); + /// \brief Creates an empty clause with the place for \a N variables. + /// + /// \param C AST context. + /// \param N The number of variables. + /// + static OMPMapClause *CreateEmpty(const ASTContext &C, unsigned N); + + /// \brief Fetches mapping kind for the clause. + OpenMPMapClauseKind getMapType() const LLVM_READONLY { return MapType; } + + /// \brief Fetches the map type modifier for the clause. + OpenMPMapClauseKind getMapTypeModifier() const LLVM_READONLY { + return MapTypeModifier; + } + + /// \brief Fetches location of clause mapping kind. + SourceLocation getMapLoc() const LLVM_READONLY { return MapLoc; } + + /// \brief Get colon location. + SourceLocation getColonLoc() const { return ColonLoc; } + + static bool classof(const OMPClause *T) { + return T->getClauseKind() == OMPC_map; + } + + child_range children() { + return child_range( + reinterpret_cast(varlist_begin()), + reinterpret_cast(varlist_end())); + } +}; + } // end namespace clang #endif // LLVM_CLANG_AST_OPENMPCLAUSE_H Index: include/clang/AST/RecursiveASTVisitor.h =================================================================== --- include/clang/AST/RecursiveASTVisitor.h +++ include/clang/AST/RecursiveASTVisitor.h @@ -2769,6 +2769,12 @@ return true; } +template +bool RecursiveASTVisitor::VisitOMPMapClause(OMPMapClause *C) { + TRY_TO(VisitOMPClauseList(C)); + return true; +} + // FIXME: look at the following tricky-seeming exprs to see if we // need to recurse on anything. These are ones that have methods // returning decls or qualtypes or nestednamespecifier -- though I'm Index: include/clang/Basic/DiagnosticParseKinds.td =================================================================== --- include/clang/Basic/DiagnosticParseKinds.td +++ include/clang/Basic/DiagnosticParseKinds.td @@ -928,6 +928,12 @@ "'#pragma omp %0' cannot be an immediate substatement">; def err_omp_expected_identifier_for_critical : Error< "expected identifier specifying the name of the 'omp critical' directive">; +def err_omp_unknown_map_type : Error< + "incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'">; +def err_omp_unknown_map_type_modifier : Error< + "incorrect map type modifier, expected 'always'">; +def err_omp_map_type_missing : Error< + "missing map type">; // Pragma loop support. def err_pragma_loop_missing_argument : Error< Index: include/clang/Basic/DiagnosticSemaKinds.td =================================================================== --- include/clang/Basic/DiagnosticSemaKinds.td +++ include/clang/Basic/DiagnosticSemaKinds.td @@ -7879,6 +7879,16 @@ "'ordered' clause with specified parameter">; def err_omp_expected_base_var_name : Error< "expected variable name as a base of the array %select{subscript|section}0">; +def err_omp_map_shared_storage : Error< + "variable already marked as mapped in current construct">; +def err_omp_not_mappable_type : Error< + "type %0 is not mappable to target">; +def note_omp_polymorphic_in_target : Note< + "mappable type cannot be polymorphic">; +def note_omp_static_member_in_target : Note< + "mappable type cannot contain static members">; +def err_omp_threadprivate_in_map : Error< + "threadprivate variables are not allowed in map clause">; } // end of OpenMP category let CategoryName = "Related Result Type Issue" in { Index: include/clang/Basic/OpenMPKinds.h =================================================================== --- include/clang/Basic/OpenMPKinds.h +++ include/clang/Basic/OpenMPKinds.h @@ -78,6 +78,14 @@ OMPC_LINEAR_unknown }; +/// \brief OpenMP mapping kind for 'map' clause. +enum OpenMPMapClauseKind { +#define OPENMP_MAP_KIND(Name) \ + OMPC_MAP_##Name, +#include "clang/Basic/OpenMPKinds.def" + OMPC_MAP_unknown +}; + OpenMPDirectiveKind getOpenMPDirectiveKind(llvm::StringRef Str); const char *getOpenMPDirectiveName(OpenMPDirectiveKind Kind); Index: include/clang/Basic/OpenMPKinds.def =================================================================== --- include/clang/Basic/OpenMPKinds.def +++ include/clang/Basic/OpenMPKinds.def @@ -84,6 +84,9 @@ #ifndef OPENMP_LINEAR_KIND #define OPENMP_LINEAR_KIND(Name) #endif +#ifndef OPENMP_MAP_KIND +#define OPENMP_MAP_KIND(Name) +#endif // OpenMP directives. OPENMP_DIRECTIVE(threadprivate) @@ -146,6 +149,7 @@ OPENMP_CLAUSE(device, OMPDeviceClause) OPENMP_CLAUSE(threads, OMPThreadsClause) OPENMP_CLAUSE(simd, OMPSIMDClause) +OPENMP_CLAUSE(map, OMPMapClause) // Clauses allowed for OpenMP directive 'parallel'. OPENMP_PARALLEL_CLAUSE(if) @@ -302,11 +306,13 @@ // TODO More clauses for 'target' directive. OPENMP_TARGET_CLAUSE(if) OPENMP_TARGET_CLAUSE(device) +OPENMP_TARGET_CLAUSE(map) // Clauses allowed for OpenMP directive 'target data'. // TODO More clauses for 'target data' directive. OPENMP_TARGET_DATA_CLAUSE(if) OPENMP_TARGET_DATA_CLAUSE(device) +OPENMP_TARGET_DATA_CLAUSE(map) // Clauses allowed for OpenMP directive 'teams'. // TODO More clauses for 'teams' directive. @@ -321,6 +327,15 @@ OPENMP_ORDERED_CLAUSE(threads) OPENMP_ORDERED_CLAUSE(simd) +// Map types and map type modifier for 'map' clause. +OPENMP_MAP_KIND(alloc) +OPENMP_MAP_KIND(to) +OPENMP_MAP_KIND(from) +OPENMP_MAP_KIND(tofrom) +OPENMP_MAP_KIND(delete) +OPENMP_MAP_KIND(release) +OPENMP_MAP_KIND(always) + #undef OPENMP_LINEAR_KIND #undef OPENMP_DEPEND_KIND #undef OPENMP_SCHEDULE_KIND @@ -345,4 +360,4 @@ #undef OPENMP_SIMD_CLAUSE #undef OPENMP_FOR_CLAUSE #undef OPENMP_FOR_SIMD_CLAUSE - +#undef OPENMP_MAP_KIND Index: include/clang/Sema/Sema.h =================================================================== --- include/clang/Sema/Sema.h +++ include/clang/Sema/Sema.h @@ -8031,7 +8031,8 @@ SourceLocation ColonLoc, SourceLocation EndLoc, CXXScopeSpec &ReductionIdScopeSpec, const DeclarationNameInfo &ReductionId, OpenMPDependClauseKind DepKind, - OpenMPLinearClauseKind LinKind, SourceLocation DepLinLoc); + OpenMPLinearClauseKind LinKind, OpenMPMapClauseKind MapTypeModifier, + OpenMPMapClauseKind MapType, SourceLocation DepLinMapLoc); /// \brief Called on well-formed 'private' clause. OMPClause *ActOnOpenMPPrivateClause(ArrayRef VarList, SourceLocation StartLoc, @@ -8097,7 +8098,12 @@ OMPClause *ActOnOpenMPDeviceClause(Expr *Device, SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc); - + /// \brief Called on well-formed 'map' clause. + OMPClause *ActOnOpenMPMapClause( + OpenMPMapClauseKind MapTypeModifier, OpenMPMapClauseKind MapType, + SourceLocation MapLoc, SourceLocation ColonLoc, ArrayRef VarList, + SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc); + /// \brief The kind of conversion being performed. enum CheckedConversionKind { /// \brief An implicit conversion. Index: lib/AST/OpenMPClause.cpp =================================================================== --- lib/AST/OpenMPClause.cpp +++ lib/AST/OpenMPClause.cpp @@ -438,3 +438,28 @@ sizeof(Expr *) * N); return new (Mem) OMPDependClause(N); } + +OMPMapClause *OMPMapClause::Create(const ASTContext &C, SourceLocation StartLoc, + SourceLocation LParenLoc, + SourceLocation EndLoc, ArrayRef VL, + OpenMPMapClauseKind TypeModifier, + OpenMPMapClauseKind Type, + SourceLocation TypeLoc) { + void *Mem = C.Allocate( + llvm::RoundUpToAlignment(sizeof(OMPMapClause), llvm::alignOf()) + + sizeof(Expr *) * VL.size()); + OMPMapClause *Clause = new (Mem) OMPMapClause( + TypeModifier, Type, TypeLoc, StartLoc, LParenLoc, EndLoc, VL.size()); + Clause->setVarRefs(VL); + Clause->setMapTypeModifier(TypeModifier); + Clause->setMapType(Type); + Clause->setMapLoc(TypeLoc); + return Clause; +} + +OMPMapClause *OMPMapClause::CreateEmpty(const ASTContext &C, unsigned N) { + void *Mem = C.Allocate( + llvm::RoundUpToAlignment(sizeof(OMPMapClause), llvm::alignOf()) + + sizeof(Expr *) * N); + return new (Mem) OMPMapClause(N); +} Index: lib/AST/StmtPrinter.cpp =================================================================== --- lib/AST/StmtPrinter.cpp +++ lib/AST/StmtPrinter.cpp @@ -842,6 +842,23 @@ OS << ")"; } } + +void OMPClausePrinter::VisitOMPMapClause(OMPMapClause *Node) { + if (!Node->varlist_empty()) { + OS << "map("; + if (Node->getMapType() != OMPC_MAP_unknown) { + if (Node->getMapTypeModifier() != OMPC_MAP_unknown) { + OS << getOpenMPSimpleClauseTypeName(OMPC_map, + Node->getMapTypeModifier()); + OS << ','; + } + OS << getOpenMPSimpleClauseTypeName(OMPC_map, Node->getMapType()); + OS << ':'; + } + VisitOMPClauseList(Node, ' '); + OS << ")"; + } +} } //===----------------------------------------------------------------------===// Index: lib/AST/StmtProfile.cpp =================================================================== --- lib/AST/StmtProfile.cpp +++ lib/AST/StmtProfile.cpp @@ -450,6 +450,9 @@ void OMPClauseProfiler::VisitOMPDeviceClause(const OMPDeviceClause *C) { Profiler->VisitStmt(C->getDevice()); } +void OMPClauseProfiler::VisitOMPMapClause(const OMPMapClause *C) { + VisitOMPClauseList(C); +} } void Index: lib/Basic/OpenMPKinds.cpp =================================================================== --- lib/Basic/OpenMPKinds.cpp +++ lib/Basic/OpenMPKinds.cpp @@ -101,6 +101,11 @@ #define OPENMP_LINEAR_KIND(Name) .Case(#Name, OMPC_LINEAR_##Name) #include "clang/Basic/OpenMPKinds.def" .Default(OMPC_LINEAR_unknown); + case OMPC_map: + return llvm::StringSwitch(Str) +#define OPENMP_MAP_KIND(Name) .Case(#Name, OMPC_MAP_##Name) +#include "clang/Basic/OpenMPKinds.def" + .Default(OMPC_MAP_unknown); case OMPC_unknown: case OMPC_threadprivate: case OMPC_if: @@ -187,6 +192,18 @@ #include "clang/Basic/OpenMPKinds.def" } llvm_unreachable("Invalid OpenMP 'linear' clause type"); + case OMPC_map: + switch (Type) { + case OMPC_MAP_unknown: + return "unknown"; +#define OPENMP_MAP_KIND(Name) \ + case OMPC_MAP_##Name: \ + return #Name; +#include "clang/Basic/OpenMPKinds.def" + default: + break; + } + llvm_unreachable("Invalid OpenMP 'map' clause type"); case OMPC_unknown: case OMPC_threadprivate: case OMPC_if: Index: lib/CodeGen/CGStmtOpenMP.cpp =================================================================== --- lib/CodeGen/CGStmtOpenMP.cpp +++ lib/CodeGen/CGStmtOpenMP.cpp @@ -2433,6 +2433,7 @@ case OMPC_device: case OMPC_threads: case OMPC_simd: + case OMPC_map: llvm_unreachable("Clause is not allowed in 'omp atomic'."); } } Index: lib/Parse/ParseOpenMP.cpp =================================================================== --- lib/Parse/ParseOpenMP.cpp +++ lib/Parse/ParseOpenMP.cpp @@ -497,6 +497,7 @@ case OMPC_copyprivate: case OMPC_flush: case OMPC_depend: + case OMPC_map: Clause = ParseOpenMPVarListClause(CKind); break; case OMPC_unknown: @@ -752,6 +753,9 @@ /// 'flush' '(' list ')' /// depend-clause: /// 'depend' '(' in | out | inout : list ')' +/// map-clause: +/// 'map' '(' [ [ always , ] +/// to | from | tofrom | alloc | release | delete ':' ] list ')'; /// /// For 'linear' clause linear-list may have the following forms: /// list @@ -769,7 +773,11 @@ // OpenMP 4.1 [2.15.3.7, linear Clause] // If no modifier is specified it is assumed to be val. OpenMPLinearClauseKind LinearModifier = OMPC_LINEAR_val; - SourceLocation DepLinLoc; + OpenMPMapClauseKind MapType = OMPC_MAP_unknown; + OpenMPMapClauseKind MapTypeModifier = OMPC_MAP_unknown; + bool MapTypeModifierSpecified = false; + bool UnexpectedId = false; + SourceLocation DepLinMapLoc; // Parse '('. BalancedDelimiterTracker T(*this, tok::l_paren, tok::annot_pragma_openmp_end); @@ -802,7 +810,7 @@ ColonProtectionRAIIObject ColonRAII(*this); DepKind = static_cast(getOpenMPSimpleClauseType( Kind, Tok.is(tok::identifier) ? PP.getSpelling(Tok) : "")); - DepLinLoc = Tok.getLocation(); + DepLinMapLoc = Tok.getLocation(); if (DepKind == OMPC_DEPEND_unknown) { SkipUntil(tok::colon, tok::r_paren, tok::annot_pragma_openmp_end, @@ -820,16 +828,79 @@ if (Tok.is(tok::identifier) && PP.LookAhead(0).is(tok::l_paren)) { LinearModifier = static_cast( getOpenMPSimpleClauseType(Kind, PP.getSpelling(Tok))); - DepLinLoc = ConsumeToken(); + DepLinMapLoc = ConsumeToken(); LinearT.consumeOpen(); NeedRParenForLinear = true; } + } else if (Kind == OMPC_map) { + // Handle map type for map clause. + ColonProtectionRAIIObject ColonRAII(*this); + + // the first identifier may be a list item, a map-type or + // a map-type-modifier + MapType = static_cast(getOpenMPSimpleClauseType( + Kind, Tok.is(tok::identifier) ? PP.getSpelling(Tok) : "")); + DepLinMapLoc = Tok.getLocation(); + bool ColonExpected = false; + + if (Tok.is(tok::identifier)) { + if (PP.LookAhead(0).is(tok::colon)) { + MapType = static_cast(getOpenMPSimpleClauseType( + Kind, Tok.is(tok::identifier) ? PP.getSpelling(Tok) : "")); + if (MapType == OMPC_MAP_unknown) { + Diag(Tok, diag::err_omp_unknown_map_type); + } else if (MapType == OMPC_MAP_always) { + Diag(Tok, diag::err_omp_map_type_missing); + } + ConsumeToken(); + } else if (PP.LookAhead(0).is(tok::comma)) { + if (PP.LookAhead(1).is(tok::identifier) && + PP.LookAhead(2).is(tok::colon)) { + MapTypeModifier = + static_cast(getOpenMPSimpleClauseType( + Kind, Tok.is(tok::identifier) ? PP.getSpelling(Tok) : "")); + if (MapTypeModifier != OMPC_MAP_always) { + Diag(Tok, diag::err_omp_unknown_map_type_modifier); + MapTypeModifier = OMPC_MAP_unknown; + } else { + MapTypeModifierSpecified = true; + } + + ConsumeToken(); + ConsumeToken(); + + MapType = static_cast(getOpenMPSimpleClauseType( + Kind, Tok.is(tok::identifier) ? PP.getSpelling(Tok) : "")); + if (MapType == OMPC_MAP_unknown || MapType == OMPC_MAP_always) { + Diag(Tok, diag::err_omp_unknown_map_type); + } + ConsumeToken(); + } else { + MapType = OMPC_MAP_tofrom; + } + } else { + MapType = OMPC_MAP_tofrom; + } + } else { + UnexpectedId = true; + } + + if (Tok.is(tok::colon)) { + ColonLoc = ConsumeToken(); + } else if (ColonExpected) { + Diag(Tok, diag::warn_pragma_expected_colon) << "map type"; + } } SmallVector Vars; - bool IsComma = ((Kind != OMPC_reduction) && (Kind != OMPC_depend)) || - ((Kind == OMPC_reduction) && !InvalidReductionId) || - ((Kind == OMPC_depend) && DepKind != OMPC_DEPEND_unknown); + bool IsComma = + ((Kind != OMPC_reduction) && (Kind != OMPC_depend) && + (Kind != OMPC_map)) || + ((Kind == OMPC_reduction) && !InvalidReductionId) || + ((Kind == OMPC_map) && (UnexpectedId || MapType != OMPC_MAP_unknown) && + (!MapTypeModifierSpecified || + (MapTypeModifierSpecified && MapTypeModifier == OMPC_MAP_always))) || + ((Kind == OMPC_depend) && DepKind != OMPC_DEPEND_unknown); const bool MayHaveTail = (Kind == OMPC_linear || Kind == OMPC_aligned); while (IsComma || (Tok.isNot(tok::r_paren) && Tok.isNot(tok::colon) && Tok.isNot(tok::annot_pragma_openmp_end))) { @@ -879,14 +950,16 @@ T.consumeClose(); if ((Kind == OMPC_depend && DepKind != OMPC_DEPEND_unknown && Vars.empty()) || (Kind != OMPC_depend && Vars.empty()) || (MustHaveTail && !TailExpr) || - InvalidReductionId) + (Kind == OMPC_map && MapType == OMPC_MAP_unknown) || + InvalidReductionId) { return nullptr; + } return Actions.ActOnOpenMPVarListClause( Kind, Vars, TailExpr, Loc, LOpen, ColonLoc, Tok.getLocation(), ReductionIdScopeSpec, ReductionId.isValid() ? Actions.GetNameFromUnqualifiedId(ReductionId) : DeclarationNameInfo(), - DepKind, LinearModifier, DepLinLoc); + DepKind, LinearModifier, MapTypeModifier, MapType, DepLinMapLoc); } Index: lib/Sema/SemaOpenMP.cpp =================================================================== --- lib/Sema/SemaOpenMP.cpp +++ lib/Sema/SemaOpenMP.cpp @@ -77,6 +77,11 @@ ImplicitDSALoc() {} }; +public: + struct MapInfo { + Expr *RefExpr; + }; + private: struct DSAInfo { OpenMPClauseKind Attributes; @@ -85,10 +90,12 @@ typedef llvm::SmallDenseMap DeclSAMapTy; typedef llvm::SmallDenseMap AlignedMapTy; typedef llvm::DenseSet LoopControlVariablesSetTy; + typedef llvm::SmallDenseMap MappedDeclsTy; struct SharingMapTy { DeclSAMapTy SharingMap; AlignedMapTy AlignedMap; + MappedDeclsTy MappedDecls; LoopControlVariablesSetTy LCVSet; DefaultDataSharingAttributes DefaultAttr; SourceLocation DefaultAttrLoc; @@ -307,6 +314,32 @@ Scope *getCurScope() const { return Stack.back().CurScope; } Scope *getCurScope() { return Stack.back().CurScope; } SourceLocation getConstructLoc() { return Stack.back().ConstructLoc; } + + MapInfo getMapInfoForVar(VarDecl *VD) { + MapInfo VarMI = {0}; + for (auto Cnt = Stack.size() - 1; Cnt > 0; --Cnt) { + if (Stack[Cnt].MappedDecls.count(VD)) { + VarMI = Stack[Cnt].MappedDecls[VD]; + break; + } + } + return VarMI; + } + + void addMapInfoForVar(VarDecl *VD, MapInfo MI) { + if (Stack.size() > 1) { + Stack.back().MappedDecls[VD] = MI; + } + } + + MapInfo IsMappedInCurrentRegion(VarDecl *VD) { + assert(Stack.size() > 1 && "Target level is 0"); + MapInfo VarMI = {0}; + if (Stack.size() > 1 && Stack.back().MappedDecls.count(VD)) { + VarMI = Stack.back().MappedDecls[VD]; + } + return VarMI; + } }; bool isParallelOrTaskRegion(OpenMPDirectiveKind DKind) { return isOpenMPParallelDirective(DKind) || DKind == OMPD_task || @@ -5078,6 +5111,7 @@ case OMPC_depend: case OMPC_threads: case OMPC_simd: + case OMPC_map: case OMPC_unknown: llvm_unreachable("Clause is not allowed."); } @@ -5340,6 +5374,7 @@ case OMPC_device: case OMPC_threads: case OMPC_simd: + case OMPC_map: case OMPC_unknown: llvm_unreachable("Clause is not allowed."); } @@ -5469,6 +5504,7 @@ case OMPC_device: case OMPC_threads: case OMPC_simd: + case OMPC_map: case OMPC_unknown: llvm_unreachable("Clause is not allowed."); } @@ -5600,6 +5636,7 @@ case OMPC_flush: case OMPC_depend: case OMPC_device: + case OMPC_map: case OMPC_unknown: llvm_unreachable("Clause is not allowed."); } @@ -5662,7 +5699,8 @@ SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation ColonLoc, SourceLocation EndLoc, CXXScopeSpec &ReductionIdScopeSpec, const DeclarationNameInfo &ReductionId, OpenMPDependClauseKind DepKind, - OpenMPLinearClauseKind LinKind, SourceLocation DepLinLoc) { + OpenMPLinearClauseKind LinKind, OpenMPMapClauseKind MapTypeModifier, + OpenMPMapClauseKind MapType, SourceLocation DepLinMapLoc) { OMPClause *Res = nullptr; switch (Kind) { case OMPC_private: @@ -5683,7 +5721,7 @@ break; case OMPC_linear: Res = ActOnOpenMPLinearClause(VarList, TailExpr, StartLoc, LParenLoc, - LinKind, DepLinLoc, ColonLoc, EndLoc); + LinKind, DepLinMapLoc, ColonLoc, EndLoc); break; case OMPC_aligned: Res = ActOnOpenMPAlignedClause(VarList, TailExpr, StartLoc, LParenLoc, @@ -5699,8 +5737,12 @@ Res = ActOnOpenMPFlushClause(VarList, StartLoc, LParenLoc, EndLoc); break; case OMPC_depend: - Res = ActOnOpenMPDependClause(DepKind, DepLinLoc, ColonLoc, VarList, StartLoc, - LParenLoc, EndLoc); + Res = ActOnOpenMPDependClause(DepKind, DepLinMapLoc, ColonLoc, VarList, + StartLoc, LParenLoc, EndLoc); + break; + case OMPC_map: + Res = ActOnOpenMPMapClause(MapTypeModifier, MapType, DepLinMapLoc, ColonLoc, + VarList, StartLoc, LParenLoc, EndLoc); break; case OMPC_if: case OMPC_final: @@ -7422,3 +7464,179 @@ } return new (Context) OMPDeviceClause(ValExpr, StartLoc, LParenLoc, EndLoc); } + +static bool IsCXXRecordForMappable(Sema &SemaRef, SourceLocation Loc, + DSAStackTy *Stack, CXXRecordDecl *RD) { + if (!RD || RD->isInvalidDecl()) + return true; + + auto QTy = SemaRef.Context.getRecordType(RD); + if (RD->isDynamicClass()) { + SemaRef.Diag(Loc, diag::err_omp_not_mappable_type) << QTy; + SemaRef.Diag(RD->getLocation(), diag::note_omp_polymorphic_in_target); + return false; + } + auto *DC = RD; + bool IsCorrect = true; + for (auto *I : DC->decls()) { + if (I) { + if (auto *MD = dyn_cast(I)) { + if (MD->isStatic()) { + SemaRef.Diag(Loc, diag::err_omp_not_mappable_type) << QTy; + SemaRef.Diag(MD->getLocation(), + diag::note_omp_static_member_in_target); + IsCorrect = false; + } + } else if (auto *VD = dyn_cast(I)) { + if (VD->isStaticDataMember()) { + SemaRef.Diag(Loc, diag::err_omp_not_mappable_type) << QTy; + SemaRef.Diag(VD->getLocation(), + diag::note_omp_static_member_in_target); + IsCorrect = false; + } + } + } + } + + for (auto &I : RD->bases()) { + if (!IsCXXRecordForMappable(SemaRef, I.getLocStart(), Stack, + I.getType()->getAsCXXRecordDecl())) + IsCorrect = false; + } + return IsCorrect; +} + +static bool CheckTypeMappable(SourceLocation SL, SourceRange SR, Sema &SemaRef, + DSAStackTy *Stack, QualType QTy) { + NamedDecl *ND; + if (QTy->isIncompleteType(&ND)) { + SemaRef.Diag(SL, diag::err_incomplete_type) << QTy << SR; + return false; + } else if (CXXRecordDecl *RD = dyn_cast_or_null(ND)) { + if (!RD->isInvalidDecl() && + !IsCXXRecordForMappable(SemaRef, SL, Stack, RD)) + return false; + } + return true; +} + +OMPClause *Sema::ActOnOpenMPMapClause( + OpenMPMapClauseKind MapTypeModifier, OpenMPMapClauseKind MapType, + SourceLocation MapLoc, SourceLocation ColonLoc, ArrayRef VarList, + SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc) { + SmallVector Vars; + + for (auto &RE : VarList) { + assert(RE && "Null expr in omp map"); + if (isa(RE)) { + // It will be analyzed later. + Vars.push_back(RE); + continue; + } + SourceLocation ELoc = RE->getExprLoc(); + + // OpenMP [2.14.5, Restrictions] + // A variable that is part of another variable (such as field of a + // structure) but is not an array element or an array section cannot appear + // in a map clause. + auto *VE = RE->IgnoreParenLValueCasts(); + + if (VE->isValueDependent() || VE->isTypeDependent() || + VE->isInstantiationDependent() || + VE->containsUnexpandedParameterPack()) { + // It will be analyzed later. + Vars.push_back(RE); + continue; + } + + auto *SimpleExpr = RE->IgnoreParenCasts(); + auto *DE = dyn_cast(SimpleExpr); + auto *ASE = dyn_cast(SimpleExpr); + auto *OASE = dyn_cast(SimpleExpr); + + if (!RE->IgnoreParenImpCasts()->isLValue() || + (!OASE && !ASE && !DE) || + (DE && !isa(DE->getDecl())) || + (ASE && !ASE->getBase()->getType()->isAnyPointerType() && + !ASE->getBase()->getType()->isArrayType())) { + Diag(ELoc, diag::err_omp_expected_var_name_or_array_item) + << RE->getSourceRange(); + continue; + } + + Decl *D = nullptr; + if (DE) { + D = DE->getDecl(); + } else if (ASE) { + auto *B = ASE->getBase()->IgnoreParenCasts(); + D = dyn_cast(B)->getDecl(); + } else if (OASE) { + auto *B = OASE->getBase(); + D = dyn_cast(B)->getDecl(); + } + assert(D && "Null decl on map clause."); + auto *VD = cast(D); + + // OpenMP [2.14.5, Restrictions, p.8] + // threadprivate variables cannot appear in a map clause. + if (DSAStack->isThreadPrivate(VD)) { + auto DVar = DSAStack->getTopDSA(VD, false); + Diag(ELoc, diag::err_omp_threadprivate_in_map); + ReportOriginalDSA(*this, DSAStack, VD, DVar); + continue; + } + + // OpenMP [2.14.5, Restrictions, p.2] + // At most one list item can be an array item derived from a given variable + // in map clauses of the same construct. + // OpenMP [2.14.5, Restrictions, p.3] + // List items of map clauses in the same construct must not share original + // storage. + // OpenMP [2.14.5, Restrictions, C/C++, p.2] + // A variable for which the type is pointer, reference to array, or + // reference to pointer and an array section derived from that variable + // must not appear as list items of map clauses of the same construct. + DSAStackTy::MapInfo MI = DSAStack->IsMappedInCurrentRegion(VD); + if (MI.RefExpr) { + Diag(ELoc, diag::err_omp_map_shared_storage) << ELoc; + Diag(MI.RefExpr->getExprLoc(), diag::note_used_here) + << MI.RefExpr->getSourceRange(); + continue; + } + + // OpenMP [2.14.5, Restrictions, C/C++, p.3,4] + // A variable for which the type is pointer, reference to array, or + // reference to pointer must not appear as a list item if the enclosing + // device data environment already contains an array section derived from + // that variable. + // An array section derived from a variable for which the type is pointer, + // reference to array, or reference to pointer must not appear as a list + // item if the enclosing device data environment already contains that + // variable. + QualType Type = VD->getType(); + MI = DSAStack->getMapInfoForVar(VD); + if (MI.RefExpr && (isa(MI.RefExpr->IgnoreParenLValueCasts()) != + isa(VE)) && + (Type->isPointerType() || Type->isReferenceType())) { + Diag(ELoc, diag::err_omp_map_shared_storage) << ELoc; + Diag(MI.RefExpr->getExprLoc(), diag::note_used_here) + << MI.RefExpr->getSourceRange(); + continue; + } + + // OpenMP [2.14.5, Restrictions, C/C++, p.7] + // A list item must have a mappable type. + if (!CheckTypeMappable(VE->getExprLoc(), VE->getSourceRange(), *this, + DSAStack, Type)) + continue; + + Vars.push_back(RE); + MI.RefExpr = RE; + DSAStack->addMapInfoForVar(VD, MI); + } + if (Vars.empty()) + return nullptr; + + return OMPMapClause::Create(Context, StartLoc, LParenLoc, EndLoc, Vars, + MapTypeModifier, MapType, MapLoc); +} Index: lib/Sema/TreeTransform.h =================================================================== --- lib/Sema/TreeTransform.h +++ lib/Sema/TreeTransform.h @@ -1652,6 +1652,20 @@ EndLoc); } + /// \brief Build a new OpenMP 'map' clause. + /// + /// By default, performs semantic analysis to build the new OpenMP clause. + /// Subclasses may override this routine to provide different behavior. + OMPClause *RebuildOMPMapClause( + OpenMPMapClauseKind MapTypeModifier, OpenMPMapClauseKind MapType, + SourceLocation MapLoc, SourceLocation ColonLoc, ArrayRef VarList, + SourceLocation StartLoc, SourceLocation LParenLoc, + SourceLocation EndLoc) { + return getSema().ActOnOpenMPMapClause(MapTypeModifier, MapType, MapLoc, + ColonLoc, VarList,StartLoc, + LParenLoc, EndLoc); + } + /// \brief Rebuild the operand to an Objective-C \@synchronized statement. /// /// By default, performs semantic analysis to build the new statement. @@ -7648,6 +7662,22 @@ E.get(), C->getLocStart(), C->getLParenLoc(), C->getLocEnd()); } +template +OMPClause *TreeTransform::TransformOMPMapClause(OMPMapClause *C) { + llvm::SmallVector Vars; + Vars.reserve(C->varlist_size()); + for (auto *VE : C->varlists()) { + ExprResult EVar = getDerived().TransformExpr(cast(VE)); + if (EVar.isInvalid()) + return nullptr; + Vars.push_back(EVar.get()); + } + return getDerived().RebuildOMPMapClause( + C->getMapTypeModifier(), C->getMapType(), C->getMapLoc(), + C->getColonLoc(), Vars, C->getLocStart(), C->getLParenLoc(), + C->getLocEnd()); +} + //===----------------------------------------------------------------------===// // Expression transformation //===----------------------------------------------------------------------===// Index: lib/Serialization/ASTReaderStmt.cpp =================================================================== --- lib/Serialization/ASTReaderStmt.cpp +++ lib/Serialization/ASTReaderStmt.cpp @@ -1840,6 +1840,9 @@ case OMPC_device: C = new (Context) OMPDeviceClause(); break; + case OMPC_map: + C = OMPMapClause::CreateEmpty(Context, Record[Idx++]); + break; } Visit(C); C->setLocStart(Reader->ReadSourceLocation(Record, Idx)); @@ -2147,6 +2150,23 @@ C->setLParenLoc(Reader->ReadSourceLocation(Record, Idx)); } +void OMPClauseReader::VisitOMPMapClause(OMPMapClause *C) { + C->setLParenLoc(Reader->ReadSourceLocation(Record, Idx)); + C->setMapTypeModifier( + static_cast(Record[Idx++])); + C->setMapType( + static_cast(Record[Idx++])); + C->setMapLoc(Reader->ReadSourceLocation(Record, Idx)); + C->setColonLoc(Reader->ReadSourceLocation(Record, Idx)); + auto NumVars = C->varlist_size(); + SmallVector Vars; + Vars.reserve(NumVars); + for (unsigned i = 0; i != NumVars; ++i) { + Vars.push_back(Reader->Reader.ReadSubExpr()); + } + C->setVarRefs(Vars); +} + //===----------------------------------------------------------------------===// // OpenMP Directives. //===----------------------------------------------------------------------===// Index: lib/Serialization/ASTWriterStmt.cpp =================================================================== --- lib/Serialization/ASTWriterStmt.cpp +++ lib/Serialization/ASTWriterStmt.cpp @@ -1982,6 +1982,17 @@ Writer->Writer.AddSourceLocation(C->getLParenLoc(), Record); } +void OMPClauseWriter::VisitOMPMapClause(OMPMapClause *C) { + Record.push_back(C->varlist_size()); + Writer->Writer.AddSourceLocation(C->getLParenLoc(), Record); + Record.push_back(C->getMapTypeModifier()); + Record.push_back(C->getMapType()); + Writer->Writer.AddSourceLocation(C->getMapLoc(), Record); + Writer->Writer.AddSourceLocation(C->getColonLoc(), Record); + for (auto *VE : C->varlists()) + Writer->Writer.AddStmt(VE); +} + //===----------------------------------------------------------------------===// // OpenMP Directives. //===----------------------------------------------------------------------===// Index: test/OpenMP/target_ast_print.cpp =================================================================== --- test/OpenMP/target_ast_print.cpp +++ test/OpenMP/target_ast_print.cpp @@ -10,39 +10,77 @@ template T tmain(T argc, T *argv) { + T i, j, a[20]; #pragma omp target foo(); #pragma omp target if (target:argc > 0) foo(); #pragma omp target if (C) foo(); +#pragma omp target map(i) + foo(); +#pragma omp target map(a[0:10], i) + foo(); +#pragma omp target map(to: i) map(from: j) + foo(); +#pragma omp target map(always,alloc: i) + foo(); return 0; } // CHECK: template int tmain(int argc, int *argv) { +// CHECK-NEXT: int i, j, a[20] // CHECK-NEXT: #pragma omp target // CHECK-NEXT: foo(); // CHECK-NEXT: #pragma omp target if(target: argc > 0) // CHECK-NEXT: foo() // CHECK-NEXT: #pragma omp target if(5) // CHECK-NEXT: foo() +// CHECK-NEXT: #pragma omp target map(tofrom: i) +// CHECK-NEXT: foo() +// CHECK-NEXT: #pragma omp target map(tofrom: a[0:10],i) +// CHECK-NEXT: foo() +// CHECK-NEXT: #pragma omp target map(to: i) map(from: j) +// CHECK-NEXT: foo() +// CHECK-NEXT: #pragma omp target map(always,alloc: i) +// CHECK-NEXT: foo() // CHECK: template char tmain(char argc, char *argv) { +// CHECK-NEXT: char i, j, a[20] // CHECK-NEXT: #pragma omp target // CHECK-NEXT: foo(); // CHECK-NEXT: #pragma omp target if(target: argc > 0) // CHECK-NEXT: foo() // CHECK-NEXT: #pragma omp target if(1) // CHECK-NEXT: foo() +// CHECK-NEXT: #pragma omp target map(tofrom: i) +// CHECK-NEXT: foo() +// CHECK-NEXT: #pragma omp target map(tofrom: a[0:10],i) +// CHECK-NEXT: foo() +// CHECK-NEXT: #pragma omp target map(to: i) map(from: j) +// CHECK-NEXT: foo() +// CHECK-NEXT: #pragma omp target map(always,alloc: i) +// CHECK-NEXT: foo() // CHECK: template T tmain(T argc, T *argv) { +// CHECK-NEXT: T i, j, a[20] // CHECK-NEXT: #pragma omp target // CHECK-NEXT: foo(); // CHECK-NEXT: #pragma omp target if(target: argc > 0) // CHECK-NEXT: foo() // CHECK-NEXT: #pragma omp target if(C) // CHECK-NEXT: foo() +// CHECK-NEXT: #pragma omp target map(tofrom: i) +// CHECK-NEXT: foo() +// CHECK-NEXT: #pragma omp target map(tofrom: a[0:10],i) +// CHECK-NEXT: foo() +// CHECK-NEXT: #pragma omp target map(to: i) map(from: j) +// CHECK-NEXT: foo() +// CHECK-NEXT: #pragma omp target map(always,alloc: i) +// CHECK-NEXT: foo() // CHECK-LABEL: int main(int argc, char **argv) { int main (int argc, char **argv) { + int i, j, a[20]; +// CHECK-NEXT: int i, j, a[20] #pragma omp target // CHECK-NEXT: #pragma omp target foo(); @@ -51,6 +89,32 @@ // CHECK-NEXT: #pragma omp target if(argc > 0) foo(); // CHECK-NEXT: foo(); + +#pragma omp target map(i) if(argc>0) +// CHECK-NEXT: #pragma omp target map(tofrom: i) if(argc > 0) + foo(); +// CHECK-NEXT: foo(); + +#pragma omp target map(i) +// CHECK-NEXT: #pragma omp target map(tofrom: i) + foo(); +// CHECK-NEXT: foo(); + +#pragma omp target map(a[0:10], i) +// CHECK-NEXT: #pragma omp target map(tofrom: a[0:10],i) + foo(); +// CHECK-NEXT: foo(); + +#pragma omp target map(to: i) map(from: j) +// CHECK-NEXT: #pragma omp target map(to: i) map(from: j) + foo(); +// CHECK-NEXT: foo(); + +#pragma omp target map(always,alloc: i) +// CHECK-NEXT: #pragma omp target map(always,alloc: i) + foo(); +// CHECK-NEXT: foo(); + return tmain(argc, &argc) + tmain(argv[0][0], argv[0]); } Index: test/OpenMP/target_data_ast_print.cpp =================================================================== --- test/OpenMP/target_data_ast_print.cpp +++ test/OpenMP/target_data_ast_print.cpp @@ -8,8 +8,113 @@ void foo() {} +template +T tmain(T argc, T *argv) { + T i, j, b, c, d, e, x[20]; + +#pragma omp target data + i = argc; + +#pragma omp target data if (target data: j > 0) + foo(); + +#pragma omp target data if (b) + foo(); + +#pragma omp target data map(c) + foo(); + +#pragma omp target data map(c) if(b>e) + foo(); + +#pragma omp target data map(x[0:10], c) + foo(); + +#pragma omp target data map(to: c) map(from: d) + foo(); + +#pragma omp target data map(always,alloc: e) + foo(); + +// nesting a target region +#pragma omp target data map(e) +{ + #pragma omp target map(always, alloc: e) + foo(); +} + + return 0; +} + +// CHECK: template int tmain(int argc, int *argv) { +// CHECK-NEXT: int i, j, b, c, d, e, x[20]; +// CHECK-NEXT: #pragma omp target data +// CHECK-NEXT: i = argc; +// CHECK-NEXT: #pragma omp target data if(target data: j > 0) +// CHECK-NEXT: foo(); +// CHECK-NEXT: #pragma omp target data if(b) +// CHECK-NEXT: foo(); +// CHECK-NEXT: #pragma omp target data map(tofrom: c) +// CHECK-NEXT: foo(); +// CHECK-NEXT: #pragma omp target data map(tofrom: c) if(b > e) +// CHECK-NEXT: foo(); +// CHECK-NEXT: #pragma omp target data map(tofrom: x[0:10],c) +// CHECK-NEXT: foo(); +// CHECK-NEXT: #pragma omp target data map(to: c) map(from: d) +// CHECK-NEXT: foo(); +// CHECK-NEXT: #pragma omp target data map(always,alloc: e) +// CHECK-NEXT: foo(); +// CHECK-NEXT: #pragma omp target data map(tofrom: e) +// CHECK-NEXT: { +// CHECK-NEXT: #pragma omp target map(always,alloc: e) +// CHECK-NEXT: foo(); +// CHECK: template char tmain(char argc, char *argv) { +// CHECK-NEXT: char i, j, b, c, d, e, x[20]; +// CHECK-NEXT: #pragma omp target data +// CHECK-NEXT: i = argc; +// CHECK-NEXT: #pragma omp target data if(target data: j > 0) +// CHECK-NEXT: foo(); +// CHECK-NEXT: #pragma omp target data if(b) +// CHECK-NEXT: foo(); +// CHECK-NEXT: #pragma omp target data map(tofrom: c) +// CHECK-NEXT: foo(); +// CHECK-NEXT: #pragma omp target data map(tofrom: c) if(b > e) +// CHECK-NEXT: foo(); +// CHECK-NEXT: #pragma omp target data map(tofrom: x[0:10],c) +// CHECK-NEXT: foo(); +// CHECK-NEXT: #pragma omp target data map(to: c) map(from: d) +// CHECK-NEXT: foo(); +// CHECK-NEXT: #pragma omp target data map(always,alloc: e) +// CHECK-NEXT: foo(); +// CHECK-NEXT: #pragma omp target data map(tofrom: e) +// CHECK-NEXT: { +// CHECK-NEXT: #pragma omp target map(always,alloc: e) +// CHECK-NEXT: foo(); +// CHECK: template T tmain(T argc, T *argv) { +// CHECK-NEXT: T i, j, b, c, d, e, x[20]; +// CHECK-NEXT: #pragma omp target data +// CHECK-NEXT: i = argc; +// CHECK-NEXT: #pragma omp target data if(target data: j > 0) +// CHECK-NEXT: foo(); +// CHECK-NEXT: #pragma omp target data if(b) +// CHECK-NEXT: foo(); +// CHECK-NEXT: #pragma omp target data map(tofrom: c) +// CHECK-NEXT: foo(); +// CHECK-NEXT: #pragma omp target data map(tofrom: c) if(b > e) +// CHECK-NEXT: foo(); +// CHECK-NEXT: #pragma omp target data map(tofrom: x[0:10],c) +// CHECK-NEXT: foo(); +// CHECK-NEXT: #pragma omp target data map(to: c) map(from: d) +// CHECK-NEXT: foo(); +// CHECK-NEXT: #pragma omp target data map(always,alloc: e) +// CHECK-NEXT: foo(); +// CHECK-NEXT: #pragma omp target data map(tofrom: e) +// CHECK-NEXT: { +// CHECK-NEXT: #pragma omp target map(always,alloc: e) +// CHECK-NEXT: foo(); + int main (int argc, char **argv) { - int b = argc, c, d, e, f, g; + int b = argc, c, d, e, f, g, x[20]; static int a; // CHECK: static int a; @@ -27,7 +132,42 @@ foo(); // CHECK-NEXT: foo(); - return (0); +#pragma omp target data map(c) +// CHECK-NEXT: #pragma omp target data map(tofrom: c) + foo(); +// CHECK-NEXT: foo(); + +#pragma omp target data map(c) if(b>g) +// CHECK-NEXT: #pragma omp target data map(tofrom: c) if(b > g) + foo(); +// CHECK-NEXT: foo(); + +#pragma omp target data map(x[0:10], c) +// CHECK-NEXT: #pragma omp target data map(tofrom: x[0:10],c) + foo(); +// CHECK-NEXT: foo(); + +#pragma omp target data map(to: c) map(from: d) +// CHECK-NEXT: #pragma omp target data map(to: c) map(from: d) + foo(); +// CHECK-NEXT: foo(); + +#pragma omp target data map(always,alloc: e) +// CHECK-NEXT: #pragma omp target data map(always,alloc: e) + foo(); +// CHECK-NEXT: foo(); + +// nesting a target region +#pragma omp target data map(e) +// CHECK-NEXT: #pragma omp target data map(tofrom: e) +{ +// CHECK-NEXT: { + #pragma omp target map(always, alloc: e) +// CHECK-NEXT: #pragma omp target map(always,alloc: e) + foo(); +// CHECK-NEXT: foo(); +} + return tmain(argc, &argc) + tmain(argv[0][0], argv[0]); } #endif Index: test/OpenMP/target_map_messages.cpp =================================================================== --- /dev/null +++ test/OpenMP/target_map_messages.cpp @@ -0,0 +1,207 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s + +void foo() { +} + +bool foobool(int argc) { + return argc; +} + +struct S1; // expected-note 2 {{declared here}} +extern S1 a; +class S2 { + mutable int a; +public: + S2():a(0) { } + S2(S2 &s2):a(s2.a) { } + static float S2s; // expected-note 4 {{mappable type cannot contain static members}} + static const float S2sc; // expected-note 4 {{mappable type cannot contain static members}} +}; +const float S2::S2sc = 0; +const S2 b; +const S2 ba[5]; +class S3 { + int a; +public: + S3():a(0) { } + S3(S3 &s3):a(s3.a) { } +}; +const S3 c; +const S3 ca[5]; +extern const int f; +class S4 { + int a; + S4(); + S4(const S4 &s4); +public: + S4(int v):a(v) { } +}; +class S5 { + int a; + S5():a(0) {} + S5(const S5 &s5):a(s5.a) { } +public: + S5(int v):a(v) { } +}; + +S3 h; +#pragma omp threadprivate(h) // expected-note 2 {{defined as threadprivate or thread local}} + +typedef int from; + +template // expected-note {{declared here}} +T tmain(T argc) { + const T d = 5; + const T da[5] = { 0 }; + S4 e(4); + S5 g(5); + T i, t[20]; + T &j = i; + T *k = &j; + T x; + T y; + T to, tofrom, always; + const T (&l)[5] = da; + + +#pragma omp target map // expected-error {{expected '(' after 'map'}} +#pragma omp target map( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}} +#pragma omp target map() // expected-error {{expected expression}} +#pragma omp target map(alloc) // expected-error {{use of undeclared identifier 'alloc'}} +#pragma omp target map(to argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected ',' or ')' in 'map' clause}} +#pragma omp target map(to:) // expected-error {{expected expression}} +#pragma omp target map(from: argc, // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} +#pragma omp target map(x: y) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} +#pragma omp target map(x) + foo(); +#pragma omp target map(tofrom: t[:I]) + foo(); +#pragma omp target map(T: a) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} + foo(); +#pragma omp target map(T) // expected-error {{'T' does not refer to a value}} + foo(); +#pragma omp target map(I) // expected-error 2 {{expected variable name, array element or array section}} + foo(); +#pragma omp target map(S2::S2s) + foo(); +#pragma omp target map(S2::S2sc) + foo(); +#pragma omp target map(x) + foo(); +#pragma omp target map(to: x) + foo(); +#pragma omp target map(to: to) + foo(); +#pragma omp target map(to) + foo(); +#pragma omp target map(to, x) + foo(); +#pragma omp target map(to x) // expected-error {{expected ',' or ')' in 'map' clause}} +#pragma omp target map(tofrom: argc > 0 ? x : y) // expected-error 2 {{expected variable name, array element or array section}} +#pragma omp target map(argc) +#pragma omp target map(S1) // expected-error {{'S1' does not refer to a value}} +#pragma omp target map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} expected-error 2 {{type 'S2' is not mappable to target}} +#pragma omp target map(ba) // expected-error 2 {{type 'S2' is not mappable to target}} +#pragma omp target map(ca) +#pragma omp target map(da) +#pragma omp target map(S2::S2s) +#pragma omp target map(S2::S2sc) +#pragma omp target map(e, g) +#pragma omp target map(h) // expected-error {{threadprivate variables are not allowed in map clause}} +#pragma omp target map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} +#pragma omp target map(k), map(k[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} + foo(); +#pragma omp target map(da) +#pragma omp target map(da[:4]) + foo(); +#pragma omp target map(k, j, l) // expected-note 4 {{used here}} +#pragma omp target map(k[:4]) // expected-error 2 {{variable already marked as mapped in current construct}} +#pragma omp target map(j) +#pragma omp target map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} + foo(); +#pragma omp target map(k[:4], j, l[:5]) // expected-note 4 {{used here}} +#pragma omp target map(k) // expected-error 2 {{variable already marked as mapped in current construct}} +#pragma omp target map(j) +#pragma omp target map(l) // expected-error 2 {{variable already marked as mapped in current construct}} + foo(); + +#pragma omp target map(always, tofrom: x) +#pragma omp target map(always: x) // expected-error {{missing map type}} +#pragma omp target map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always'}} expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} +#pragma omp target map(always, tofrom: always, tofrom, x) +#pragma omp target map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}} + foo(); + + return 0; +} + +int main(int argc, char **argv) { + const int d = 5; + const int da[5] = { 0 }; + S4 e(4); + S5 g(5); + int i; + int &j = i; + int *k = &j; + int x; + int y; + int to, tofrom, always; + const int (&l)[5] = da; +#pragma omp target map // expected-error {{expected '(' after 'map'}} +#pragma omp target map( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}} +#pragma omp target map() // expected-error {{expected expression}} +#pragma omp target map(alloc) // expected-error {{use of undeclared identifier 'alloc'}} +#pragma omp target map(to argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected ',' or ')' in 'map' clause}} +#pragma omp target map(to:) // expected-error {{expected expression}} +#pragma omp target map(from: argc, // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} +#pragma omp target map(x: y) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} +#pragma omp target map(x) + foo(); +#pragma omp target map(to: x) + foo(); +#pragma omp target map(to: to) + foo(); +#pragma omp target map(to) + foo(); +#pragma omp target map(to, x) + foo(); +#pragma omp target map(to x) // expected-error {{expected ',' or ')' in 'map' clause}} +#pragma omp target map(tofrom: argc > 0 ? argv[1] : argv[2]) // expected-error {{expected variable name, array element or array section}} +#pragma omp target map(argc) +#pragma omp target map(S1) // expected-error {{'S1' does not refer to a value}} +#pragma omp target map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} expected-error 2 {{type 'S2' is not mappable to target}} +#pragma omp target map(argv[1]) +#pragma omp target map(ba) // expected-error 2 {{type 'S2' is not mappable to target}} +#pragma omp target map(ca) +#pragma omp target map(da) +#pragma omp target map(S2::S2s) +#pragma omp target map(S2::S2sc) +#pragma omp target map(e, g) +#pragma omp target map(h) // expected-error {{threadprivate variables are not allowed in map clause}} +#pragma omp target map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} +#pragma omp target map(k), map(k[:5]) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} + foo(); +#pragma omp target map(da) +#pragma omp target map(da[:4]) + foo(); +#pragma omp target map(k, j, l) // expected-note 2 {{used here}} +#pragma omp target map(k[:4]) // expected-error {{variable already marked as mapped in current construct}} +#pragma omp target map(j) +#pragma omp target map(l[:5]) // expected-error {{variable already marked as mapped in current construct}} + foo(); +#pragma omp target map(k[:4], j, l[:5]) // expected-note 2 {{used here}} +#pragma omp target map(k) // expected-error {{variable already marked as mapped in current construct}} +#pragma omp target map(j) +#pragma omp target map(l) // expected-error {{variable already marked as mapped in current construct}} + foo(); + +#pragma omp target map(always, tofrom: x) +#pragma omp target map(always: x) // expected-error {{missing map type}} +#pragma omp target map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always'}} expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} +#pragma omp target map(always, tofrom: always, tofrom, x) +#pragma omp target map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}} + foo(); + + 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}} +} + Index: tools/libclang/CIndex.cpp =================================================================== --- tools/libclang/CIndex.cpp +++ tools/libclang/CIndex.cpp @@ -2178,6 +2178,9 @@ void OMPClauseEnqueue::VisitOMPDependClause(const OMPDependClause *C) { VisitOMPClauseList(C); } +void OMPClauseEnqueue::VisitOMPMapClause(const OMPMapClause *C) { + VisitOMPClauseList(C); +} } void EnqueueVisitor::EnqueueChildren(const OMPClause *S) {