Author: kli Date: Sun Nov 22 23:32:03 2015 New Revision: 253849 URL: http://llvm.org/viewvc/llvm-project?rev=253849&view=rev Log: [OpenMP] Parsing and sema support for map clause
http://reviews.llvm.org/D14134 Modified: cfe/trunk/include/clang/AST/DataRecursiveASTVisitor.h cfe/trunk/include/clang/AST/OpenMPClause.h cfe/trunk/include/clang/AST/RecursiveASTVisitor.h cfe/trunk/include/clang/Basic/DiagnosticParseKinds.td cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td cfe/trunk/include/clang/Basic/OpenMPKinds.def cfe/trunk/include/clang/Basic/OpenMPKinds.h cfe/trunk/include/clang/Sema/Sema.h cfe/trunk/lib/AST/OpenMPClause.cpp cfe/trunk/lib/AST/StmtPrinter.cpp cfe/trunk/lib/AST/StmtProfile.cpp cfe/trunk/lib/Basic/OpenMPKinds.cpp cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp cfe/trunk/lib/Parse/ParseOpenMP.cpp cfe/trunk/lib/Sema/SemaOpenMP.cpp cfe/trunk/lib/Sema/TreeTransform.h cfe/trunk/lib/Serialization/ASTReaderStmt.cpp cfe/trunk/lib/Serialization/ASTWriterStmt.cpp cfe/trunk/test/OpenMP/target_ast_print.cpp cfe/trunk/test/OpenMP/target_data_ast_print.cpp cfe/trunk/tools/libclang/CIndex.cpp Modified: cfe/trunk/include/clang/AST/DataRecursiveASTVisitor.h URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/AST/DataRecursiveASTVisitor.h?rev=253849&r1=253848&r2=253849&view=diff ============================================================================== --- cfe/trunk/include/clang/AST/DataRecursiveASTVisitor.h (original) +++ cfe/trunk/include/clang/AST/DataRecursiveASTVisitor.h Sun Nov 22 23:32:03 2015 @@ -2707,6 +2707,12 @@ bool RecursiveASTVisitor<Derived>::Visit return true; } +template <typename Derived> +bool RecursiveASTVisitor<Derived>::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 Modified: cfe/trunk/include/clang/AST/OpenMPClause.h URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/AST/OpenMPClause.h?rev=253849&r1=253848&r2=253849&view=diff ============================================================================== --- cfe/trunk/include/clang/AST/OpenMPClause.h (original) +++ cfe/trunk/include/clang/AST/OpenMPClause.h Sun Nov 22 23:32:03 2015 @@ -2613,6 +2613,121 @@ public: } }; +/// \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<OMPMapClause> { + 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<OMPMapClause>(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<OMPMapClause>(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<Expr *> 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<Stmt **>(varlist_begin()), + reinterpret_cast<Stmt **>(varlist_end())); + } +}; + } // end namespace clang #endif // LLVM_CLANG_AST_OPENMPCLAUSE_H Modified: cfe/trunk/include/clang/AST/RecursiveASTVisitor.h URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/AST/RecursiveASTVisitor.h?rev=253849&r1=253848&r2=253849&view=diff ============================================================================== --- cfe/trunk/include/clang/AST/RecursiveASTVisitor.h (original) +++ cfe/trunk/include/clang/AST/RecursiveASTVisitor.h Sun Nov 22 23:32:03 2015 @@ -2769,6 +2769,12 @@ bool RecursiveASTVisitor<Derived>::Visit return true; } +template <typename Derived> +bool RecursiveASTVisitor<Derived>::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 Modified: cfe/trunk/include/clang/Basic/DiagnosticParseKinds.td URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticParseKinds.td?rev=253849&r1=253848&r2=253849&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/DiagnosticParseKinds.td (original) +++ cfe/trunk/include/clang/Basic/DiagnosticParseKinds.td Sun Nov 22 23:32:03 2015 @@ -928,6 +928,12 @@ def err_omp_immediate_directive : Error< "'#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< Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=253849&r1=253848&r2=253849&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original) +++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Sun Nov 22 23:32:03 2015 @@ -7885,6 +7885,16 @@ def note_omp_ordered_param : Note< "'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 { Modified: cfe/trunk/include/clang/Basic/OpenMPKinds.def URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/OpenMPKinds.def?rev=253849&r1=253848&r2=253849&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/OpenMPKinds.def (original) +++ cfe/trunk/include/clang/Basic/OpenMPKinds.def Sun Nov 22 23:32:03 2015 @@ -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(depend, OMPDependClause) 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 @@ OPENMP_ATOMIC_CLAUSE(seq_cst) // 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_TEAMS_CLAUSE(reduction) 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 @@ OPENMP_ORDERED_CLAUSE(simd) #undef OPENMP_SIMD_CLAUSE #undef OPENMP_FOR_CLAUSE #undef OPENMP_FOR_SIMD_CLAUSE - +#undef OPENMP_MAP_KIND Modified: cfe/trunk/include/clang/Basic/OpenMPKinds.h URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/OpenMPKinds.h?rev=253849&r1=253848&r2=253849&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/OpenMPKinds.h (original) +++ cfe/trunk/include/clang/Basic/OpenMPKinds.h Sun Nov 22 23:32:03 2015 @@ -78,6 +78,14 @@ enum OpenMPLinearClauseKind { 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); Modified: cfe/trunk/include/clang/Sema/Sema.h URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=253849&r1=253848&r2=253849&view=diff ============================================================================== --- cfe/trunk/include/clang/Sema/Sema.h (original) +++ cfe/trunk/include/clang/Sema/Sema.h Sun Nov 22 23:32:03 2015 @@ -8049,7 +8049,8 @@ public: 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<Expr *> VarList, SourceLocation StartLoc, @@ -8115,7 +8116,12 @@ public: 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<Expr *> VarList, + SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc); + /// \brief The kind of conversion being performed. enum CheckedConversionKind { /// \brief An implicit conversion. Modified: cfe/trunk/lib/AST/OpenMPClause.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/OpenMPClause.cpp?rev=253849&r1=253848&r2=253849&view=diff ============================================================================== --- cfe/trunk/lib/AST/OpenMPClause.cpp (original) +++ cfe/trunk/lib/AST/OpenMPClause.cpp Sun Nov 22 23:32:03 2015 @@ -438,3 +438,28 @@ OMPDependClause *OMPDependClause::Create sizeof(Expr *) * N); return new (Mem) OMPDependClause(N); } + +OMPMapClause *OMPMapClause::Create(const ASTContext &C, SourceLocation StartLoc, + SourceLocation LParenLoc, + SourceLocation EndLoc, ArrayRef<Expr *> VL, + OpenMPMapClauseKind TypeModifier, + OpenMPMapClauseKind Type, + SourceLocation TypeLoc) { + void *Mem = C.Allocate( + llvm::RoundUpToAlignment(sizeof(OMPMapClause), llvm::alignOf<Expr *>()) + + 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<Expr *>()) + + sizeof(Expr *) * N); + return new (Mem) OMPMapClause(N); +} Modified: cfe/trunk/lib/AST/StmtPrinter.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/StmtPrinter.cpp?rev=253849&r1=253848&r2=253849&view=diff ============================================================================== --- cfe/trunk/lib/AST/StmtPrinter.cpp (original) +++ cfe/trunk/lib/AST/StmtPrinter.cpp Sun Nov 22 23:32:03 2015 @@ -842,6 +842,23 @@ void OMPClausePrinter::VisitOMPDependCla 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 << ")"; + } +} } //===----------------------------------------------------------------------===// Modified: cfe/trunk/lib/AST/StmtProfile.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/StmtProfile.cpp?rev=253849&r1=253848&r2=253849&view=diff ============================================================================== --- cfe/trunk/lib/AST/StmtProfile.cpp (original) +++ cfe/trunk/lib/AST/StmtProfile.cpp Sun Nov 22 23:32:03 2015 @@ -450,6 +450,9 @@ void OMPClauseProfiler::VisitOMPDependCl void OMPClauseProfiler::VisitOMPDeviceClause(const OMPDeviceClause *C) { Profiler->VisitStmt(C->getDevice()); } +void OMPClauseProfiler::VisitOMPMapClause(const OMPMapClause *C) { + VisitOMPClauseList(C); +} } void Modified: cfe/trunk/lib/Basic/OpenMPKinds.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/OpenMPKinds.cpp?rev=253849&r1=253848&r2=253849&view=diff ============================================================================== --- cfe/trunk/lib/Basic/OpenMPKinds.cpp (original) +++ cfe/trunk/lib/Basic/OpenMPKinds.cpp Sun Nov 22 23:32:03 2015 @@ -101,6 +101,11 @@ unsigned clang::getOpenMPSimpleClauseTyp #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<OpenMPMapClauseKind>(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 @@ const char *clang::getOpenMPSimpleClause #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: Modified: cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp?rev=253849&r1=253848&r2=253849&view=diff ============================================================================== --- cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp (original) +++ cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp Sun Nov 22 23:32:03 2015 @@ -2433,6 +2433,7 @@ static void EmitOMPAtomicExpr(CodeGenFun case OMPC_device: case OMPC_threads: case OMPC_simd: + case OMPC_map: llvm_unreachable("Clause is not allowed in 'omp atomic'."); } } Modified: cfe/trunk/lib/Parse/ParseOpenMP.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Parse/ParseOpenMP.cpp?rev=253849&r1=253848&r2=253849&view=diff ============================================================================== --- cfe/trunk/lib/Parse/ParseOpenMP.cpp (original) +++ cfe/trunk/lib/Parse/ParseOpenMP.cpp Sun Nov 22 23:32:03 2015 @@ -497,6 +497,7 @@ OMPClause *Parser::ParseOpenMPClause(Ope case OMPC_copyprivate: case OMPC_flush: case OMPC_depend: + case OMPC_map: Clause = ParseOpenMPVarListClause(CKind); break; case OMPC_unknown: @@ -752,6 +753,9 @@ static bool ParseReductionId(Parser &P, /// '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 @@ OMPClause *Parser::ParseOpenMPVarListCla // 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 @@ OMPClause *Parser::ParseOpenMPVarListCla ColonProtectionRAIIObject ColonRAII(*this); DepKind = static_cast<OpenMPDependClauseKind>(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 @@ OMPClause *Parser::ParseOpenMPVarListCla if (Tok.is(tok::identifier) && PP.LookAhead(0).is(tok::l_paren)) { LinearModifier = static_cast<OpenMPLinearClauseKind>( 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<OpenMPMapClauseKind>(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<OpenMPMapClauseKind>(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<OpenMPMapClauseKind>(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<OpenMPMapClauseKind>(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<Expr *, 5> 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 @@ OMPClause *Parser::ParseOpenMPVarListCla 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); } Modified: cfe/trunk/lib/Sema/SemaOpenMP.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOpenMP.cpp?rev=253849&r1=253848&r2=253849&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaOpenMP.cpp (original) +++ cfe/trunk/lib/Sema/SemaOpenMP.cpp Sun Nov 22 23:32:03 2015 @@ -77,6 +77,11 @@ public: ImplicitDSALoc() {} }; +public: + struct MapInfo { + Expr *RefExpr; + }; + private: struct DSAInfo { OpenMPClauseKind Attributes; @@ -85,10 +90,12 @@ private: typedef llvm::SmallDenseMap<VarDecl *, DSAInfo, 64> DeclSAMapTy; typedef llvm::SmallDenseMap<VarDecl *, DeclRefExpr *, 64> AlignedMapTy; typedef llvm::DenseSet<VarDecl *> LoopControlVariablesSetTy; + typedef llvm::SmallDenseMap<VarDecl *, MapInfo, 64> MappedDeclsTy; struct SharingMapTy { DeclSAMapTy SharingMap; AlignedMapTy AlignedMap; + MappedDeclsTy MappedDecls; LoopControlVariablesSetTy LCVSet; DefaultDataSharingAttributes DefaultAttr; SourceLocation DefaultAttrLoc; @@ -307,6 +314,32 @@ public: 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 @@ OMPClause *Sema::ActOnOpenMPSingleExprCl 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 @@ OMPClause *Sema::ActOnOpenMPSimpleClause 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 @@ OMPClause *Sema::ActOnOpenMPSingleExprWi 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 @@ OMPClause *Sema::ActOnOpenMPClause(OpenM 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 @@ OMPClause *Sema::ActOnOpenMPVarListClaus 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 @@ OMPClause *Sema::ActOnOpenMPVarListClaus 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 @@ OMPClause *Sema::ActOnOpenMPVarListClaus 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 @@ OMPClause *Sema::ActOnOpenMPDeviceClause } 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<CXXMethodDecl>(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<VarDecl>(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<CXXRecordDecl>(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<Expr *> VarList, + SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc) { + SmallVector<Expr *, 4> Vars; + + for (auto &RE : VarList) { + assert(RE && "Null expr in omp map"); + if (isa<DependentScopeDeclRefExpr>(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<DeclRefExpr>(SimpleExpr); + auto *ASE = dyn_cast<ArraySubscriptExpr>(SimpleExpr); + auto *OASE = dyn_cast<OMPArraySectionExpr>(SimpleExpr); + + if (!RE->IgnoreParenImpCasts()->isLValue() || + (!OASE && !ASE && !DE) || + (DE && !isa<VarDecl>(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<DeclRefExpr>(B)->getDecl(); + } else if (OASE) { + auto *B = OASE->getBase(); + D = dyn_cast<DeclRefExpr>(B)->getDecl(); + } + assert(D && "Null decl on map clause."); + auto *VD = cast<VarDecl>(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<DeclRefExpr>(MI.RefExpr->IgnoreParenLValueCasts()) != + isa<DeclRefExpr>(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); +} Modified: cfe/trunk/lib/Sema/TreeTransform.h URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/TreeTransform.h?rev=253849&r1=253848&r2=253849&view=diff ============================================================================== --- cfe/trunk/lib/Sema/TreeTransform.h (original) +++ cfe/trunk/lib/Sema/TreeTransform.h Sun Nov 22 23:32:03 2015 @@ -1652,6 +1652,20 @@ public: 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<Expr *> 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 @@ TreeTransform<Derived>::TransformOMPDevi E.get(), C->getLocStart(), C->getLParenLoc(), C->getLocEnd()); } +template <typename Derived> +OMPClause *TreeTransform<Derived>::TransformOMPMapClause(OMPMapClause *C) { + llvm::SmallVector<Expr *, 16> Vars; + Vars.reserve(C->varlist_size()); + for (auto *VE : C->varlists()) { + ExprResult EVar = getDerived().TransformExpr(cast<Expr>(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 //===----------------------------------------------------------------------===// Modified: cfe/trunk/lib/Serialization/ASTReaderStmt.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Serialization/ASTReaderStmt.cpp?rev=253849&r1=253848&r2=253849&view=diff ============================================================================== --- cfe/trunk/lib/Serialization/ASTReaderStmt.cpp (original) +++ cfe/trunk/lib/Serialization/ASTReaderStmt.cpp Sun Nov 22 23:32:03 2015 @@ -1840,6 +1840,9 @@ OMPClause *OMPClauseReader::readClause() 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 @@ void OMPClauseReader::VisitOMPDeviceClau C->setLParenLoc(Reader->ReadSourceLocation(Record, Idx)); } +void OMPClauseReader::VisitOMPMapClause(OMPMapClause *C) { + C->setLParenLoc(Reader->ReadSourceLocation(Record, Idx)); + C->setMapTypeModifier( + static_cast<OpenMPMapClauseKind>(Record[Idx++])); + C->setMapType( + static_cast<OpenMPMapClauseKind>(Record[Idx++])); + C->setMapLoc(Reader->ReadSourceLocation(Record, Idx)); + C->setColonLoc(Reader->ReadSourceLocation(Record, Idx)); + auto NumVars = C->varlist_size(); + SmallVector<Expr *, 16> Vars; + Vars.reserve(NumVars); + for (unsigned i = 0; i != NumVars; ++i) { + Vars.push_back(Reader->Reader.ReadSubExpr()); + } + C->setVarRefs(Vars); +} + //===----------------------------------------------------------------------===// // OpenMP Directives. //===----------------------------------------------------------------------===// Modified: cfe/trunk/lib/Serialization/ASTWriterStmt.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Serialization/ASTWriterStmt.cpp?rev=253849&r1=253848&r2=253849&view=diff ============================================================================== --- cfe/trunk/lib/Serialization/ASTWriterStmt.cpp (original) +++ cfe/trunk/lib/Serialization/ASTWriterStmt.cpp Sun Nov 22 23:32:03 2015 @@ -1982,6 +1982,17 @@ void OMPClauseWriter::VisitOMPDeviceClau 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. //===----------------------------------------------------------------------===// Modified: cfe/trunk/test/OpenMP/target_ast_print.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_ast_print.cpp?rev=253849&r1=253848&r2=253849&view=diff ============================================================================== --- cfe/trunk/test/OpenMP/target_ast_print.cpp (original) +++ cfe/trunk/test/OpenMP/target_ast_print.cpp Sun Nov 22 23:32:03 2015 @@ -10,39 +10,77 @@ void foo() {} template <typename T, int C> 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 <typename T = int, int C = 5> 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 <typename T = char, int C = 1> 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 <typename T, int C> 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 @@ int main (int argc, char **argv) { // 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<int, 5>(argc, &argc) + tmain<char, 1>(argv[0][0], argv[0]); } Modified: cfe/trunk/test/OpenMP/target_data_ast_print.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_data_ast_print.cpp?rev=253849&r1=253848&r2=253849&view=diff ============================================================================== --- cfe/trunk/test/OpenMP/target_data_ast_print.cpp (original) +++ cfe/trunk/test/OpenMP/target_data_ast_print.cpp Sun Nov 22 23:32:03 2015 @@ -8,8 +8,113 @@ void foo() {} +template <typename T, int C> +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 <typename T = int, int C = 5> 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 <typename T = char, int C = 1> 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 <typename T, int C> 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 @@ int main (int argc, char **argv) { 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<int, 5>(argc, &argc) + tmain<char, 1>(argv[0][0], argv[0]); } #endif Modified: cfe/trunk/tools/libclang/CIndex.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/tools/libclang/CIndex.cpp?rev=253849&r1=253848&r2=253849&view=diff ============================================================================== --- cfe/trunk/tools/libclang/CIndex.cpp (original) +++ cfe/trunk/tools/libclang/CIndex.cpp Sun Nov 22 23:32:03 2015 @@ -2178,6 +2178,9 @@ void OMPClauseEnqueue::VisitOMPFlushClau void OMPClauseEnqueue::VisitOMPDependClause(const OMPDependClause *C) { VisitOMPClauseList(C); } +void OMPClauseEnqueue::VisitOMPMapClause(const OMPMapClause *C) { + VisitOMPClauseList(C); +} } void EnqueueVisitor::EnqueueChildren(const OMPClause *S) { _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
