Index: include/clang/AST/DataRecursiveASTVisitor.h =================================================================== --- include/clang/AST/DataRecursiveASTVisitor.h +++ include/clang/AST/DataRecursiveASTVisitor.h @@ -2703,6 +2703,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,117 @@ } }; +/// \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; + friend class OMPClauseWriter; + friend class Sema; + + /// \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 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. + /// + 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 @@ -2765,6 +2765,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 @@ -991,6 +991,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 @@ -7780,6 +7780,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,15 @@ OMPC_LINEAR_unknown }; +/// \brief OpenMP mapping kind for 'map' clause. +enum OpenMPMapClauseKind { + OMPC_MAP_unknown = 0, +#define OPENMP_MAP_KIND(Name) \ + OMPC_MAP_##Name, +#include "clang/Basic/OpenMPKinds.def" + NUM_OPENMP_MAP_KIND +}; + 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 @@ -8009,7 +8009,9 @@ SourceLocation ColonLoc, SourceLocation EndLoc, CXXScopeSpec &ReductionIdScopeSpec, const DeclarationNameInfo &ReductionId, OpenMPDependClauseKind DepKind, - OpenMPLinearClauseKind LinKind, SourceLocation DepLinLoc); + OpenMPLinearClauseKind LinKind, OpenMPMapClauseKind MapTypeModifier, + OpenMPMapClauseKind MapType, SourceLocation DepLinLoc, + SourceLocation MapLoc); /// \brief Called on well-formed 'private' clause. OMPClause *ActOnOpenMPPrivateClause(ArrayRef VarList, SourceLocation StartLoc, @@ -8075,7 +8077,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,26 @@ OS << ")"; } } + +void OMPClausePrinter::VisitOMPMapClause(OMPMapClause *Node) { + if (!Node->varlist_empty()) { + OS << "map("; + if (Node->getMapTypeModifier() != OMPC_MAP_unknown) { + OS << getOpenMPSimpleClauseTypeName(OMPC_map, Node->getMapTypeModifier()); + OS << ','; + } + OS << getOpenMPSimpleClauseTypeName(OMPC_map, Node->getMapType()); + OS << ':'; + + for (OMPMapClause::varlist_iterator I = Node->varlist_begin(), + E = Node->varlist_end(); + I != E; ++I) { + OS << (I == Node->varlist_begin() ? ' ' : ',') + << *cast(cast(*I)->getDecl()); + } + 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,12 @@ // 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; + OpenMPMapClauseKind MapType = OMPC_MAP_unknown; + OpenMPMapClauseKind MapTypeModifier = OMPC_MAP_unknown; + bool MapTypeModifierSpecified = false; + bool UnexpectedId = false; SourceLocation DepLinLoc; + SourceLocation MapLoc; // Parse '('. BalancedDelimiterTracker T(*this, tok::l_paren, tok::annot_pragma_openmp_end); @@ -824,12 +833,75 @@ 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) : "")); + MapLoc = 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 +951,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, DepLinLoc, MapLoc); } 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 Tmp = {0}; + for (auto Cnt = Stack.size() - 1; Cnt > 0; --Cnt) { + if (Stack[Cnt].MappedDecls.count(VD)) { + Tmp = Stack[Cnt].MappedDecls[VD]; + break; + } + } + return Tmp; + } + + 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 Tmp = {0}; + if (Stack.size() > 1 && Stack.back().MappedDecls.count(VD)) { + Tmp = Stack.back().MappedDecls[VD]; + } + return Tmp; + } }; 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,9 @@ 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 DepLinLoc, + SourceLocation MapLoc) { OMPClause *Res = nullptr; switch (Kind) { case OMPC_private: @@ -5702,6 +5741,10 @@ Res = ActOnOpenMPDependClause(DepKind, DepLinLoc, ColonLoc, VarList, StartLoc, LParenLoc, EndLoc); break; + case OMPC_map: + Res = ActOnOpenMPMapClause(MapTypeModifier, MapType, MapLoc, ColonLoc, + VarList, StartLoc, LParenLoc, EndLoc); + break; case OMPC_if: case OMPC_final: case OMPC_num_threads: @@ -7422,3 +7465,185 @@ } 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 (DeclContext::decl_iterator I = DC->noload_decls_begin(), + E = DC->noload_decls_end(); + I != E; ++I) { + if (*I) { + if (CXXMethodDecl *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 (VarDecl *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 (CXXRecordDecl::base_class_iterator I = RD->bases_begin(), + E = RD->bases_end(); + I != E; ++I) { + 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. + Expr *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 @@ -1839,6 +1839,8 @@ break; case OMPC_device: C = new (Context) OMPDeviceClause(); + case OMPC_map: + C = OMPMapClause::CreateEmpty(Context, Record[Idx++]); break; } Visit(C); @@ -2147,6 +2149,26 @@ 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)); + 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); + Vars.clear(); + for (unsigned i = 0; i != NumVars; ++i) { + Vars.push_back(Reader->Reader.ReadSubExpr()); + } +} + //===----------------------------------------------------------------------===// // 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_map_messages.cpp =================================================================== --- /dev/null +++ test/OpenMP/target_map_messages.cpp @@ -0,0 +1,119 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s + +void foo() { +} + +bool foobool(int argc) { + return argc; +} + +struct S1; // expected-note {{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 2 {{mappable type cannot contain static members}} + static const float S2sc; // expected-note 2 {{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 {{defined as threadprivate or thread local}} + +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 0; +} + Index: tools/libclang/CIndex.cpp =================================================================== --- tools/libclang/CIndex.cpp +++ tools/libclang/CIndex.cpp @@ -2171,6 +2171,9 @@ void OMPClauseEnqueue::VisitOMPDependClause(const OMPDependClause *C) { VisitOMPClauseList(C); } +void OMPClauseEnqueue::VisitOMPMapClause(const OMPMapClause *C) { + VisitOMPClauseList(C); +} } void EnqueueVisitor::EnqueueChildren(const OMPClause *S) {