Index: include/clang-c/Index.h =================================================================== --- include/clang-c/Index.h +++ include/clang-c/Index.h @@ -2270,11 +2270,15 @@ */ CXCursor_OMPTaskLoopSimdDirective = 259, - /** \brief OpenMP distribute directive. + /** \brief OpenMP distribute directive. */ CXCursor_OMPDistributeDirective = 260, - CXCursor_LastStmt = CXCursor_OMPDistributeDirective, + /** \brief OpenMP target update directive. + */ + CXCursor_OMPTargetUpdateDirective = 261, + + CXCursor_LastStmt = CXCursor_OMPTargetUpdateDirective, /** * \brief Cursor that represents the translation unit itself. Index: include/clang/AST/OpenMPClause.h =================================================================== --- include/clang/AST/OpenMPClause.h +++ include/clang/AST/OpenMPClause.h @@ -3193,6 +3193,127 @@ child_range children() { return child_range(&Hint, &Hint + 1); } }; +/// \brief This represents clause 'from' in the '#pragma omp ...' +/// directives. +/// +/// \code +/// #pragma omp target update from(a,b) +/// \endcode +/// In this example directive '#pragma omp target update' has clause 'from' +/// with the variables 'a' and 'b'. +/// +class OMPFromClause final : public OMPVarListClause, + private llvm::TrailingObjects { + friend TrailingObjects; + friend OMPVarListClause; + friend class OMPClauseReader; + + /// \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 OMPFromClause(SourceLocation StartLoc, SourceLocation LParenLoc, + SourceLocation EndLoc, unsigned N) + : OMPVarListClause(OMPC_from, StartLoc, LParenLoc, EndLoc, N) + {} + /// \brief Build an empty clause. + /// + /// \param N Number of variables. + /// + explicit OMPFromClause(unsigned N) + : OMPVarListClause(OMPC_from, SourceLocation(), + SourceLocation(), SourceLocation(), N) {} + +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 OMPFromClause *Create(const ASTContext &C, SourceLocation StartLoc, + SourceLocation LParenLoc, SourceLocation EndLoc, + ArrayRef VL); + /// \brief Creates an empty clause with the place for \a N variables. + /// + /// \param C AST context. + /// \param N The number of variables. + /// + static OMPFromClause *CreateEmpty(const ASTContext &C, unsigned N); + + static bool classof(const OMPClause *T) { + return T->getClauseKind() == OMPC_from; + } + + child_range children() { + return child_range(reinterpret_cast(varlist_begin()), + reinterpret_cast(varlist_end())); + } +}; + +/// \brief This represents clause 'to' in the '#pragma omp ...' +/// directives. +/// +/// \code +/// #pragma omp target update to(a,b) +/// \endcode +/// In this example directive '#pragma omp target update' has clause 'to' +/// with the variables 'a' and 'b'. +/// +class OMPToClause final : public OMPVarListClause, + private llvm::TrailingObjects { + friend TrailingObjects; + friend OMPVarListClause; + friend class OMPClauseReader; + + /// \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 OMPToClause(SourceLocation StartLoc, SourceLocation LParenLoc, + SourceLocation EndLoc, unsigned N) + : OMPVarListClause(OMPC_to, StartLoc, LParenLoc, EndLoc, N) {} + /// \brief Build an empty clause. + /// + /// \param N Number of variables. + /// + explicit OMPToClause(unsigned N) + : OMPVarListClause(OMPC_to, SourceLocation(), SourceLocation(), + SourceLocation(), N) {} + +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 OMPToClause *Create(const ASTContext &C, SourceLocation StartLoc, + SourceLocation LParenLoc, SourceLocation EndLoc, + ArrayRef VL); + /// \brief Creates an empty clause with the place for \a N variables. + /// + /// \param C AST context. + /// \param N The number of variables. + /// + static OMPToClause *CreateEmpty(const ASTContext &C, unsigned N); + + static bool classof(const OMPClause *T) { + return T->getClauseKind() == OMPC_to; + } + + 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 @@ -2432,6 +2432,9 @@ DEF_TRAVERSE_STMT(OMPTeamsDirective, { TRY_TO(TraverseOMPExecutableDirective(S)); }) +DEF_TRAVERSE_STMT(OMPTargetUpdateDirective, + { TRY_TO(TraverseOMPExecutableDirective(S)); }) + DEF_TRAVERSE_STMT(OMPTaskLoopDirective, { TRY_TO(TraverseOMPExecutableDirective(S)); }) @@ -2777,6 +2780,18 @@ return true; } +template +bool RecursiveASTVisitor::VisitOMPFromClause(OMPFromClause *C) { + TRY_TO(VisitOMPClauseList(C)); + return true; +} + +template +bool RecursiveASTVisitor::VisitOMPToClause(OMPToClause *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/StmtOpenMP.h =================================================================== --- include/clang/AST/StmtOpenMP.h +++ include/clang/AST/StmtOpenMP.h @@ -2395,7 +2395,7 @@ /// \param Clauses List of clauses. /// \param AssociatedStmt Statement, associated with the directive. /// \param Exprs Helper expressions for CodeGen. - /// + /// static OMPDistributeDirective * Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc, unsigned CollapsedNum, ArrayRef Clauses, @@ -2417,6 +2417,64 @@ } }; +/// \brief This represents '#pragma omp target update' directive. +/// +/// \code +/// #pragma omp target update to(a) from(b) device(1) +/// \endcode +/// In this example directive '#pragma omp target update' has clause 'to' with +/// argument 'a', clause 'from' with argument 'b' and clause 'device' with +/// argument '1'. +/// +class OMPTargetUpdateDirective : public OMPExecutableDirective { + friend class ASTStmtReader; + /// \brief Build directive with the given start and end location. + /// + /// \param StartLoc Starting location of the directive kind. + /// \param EndLoc Ending Location of the directive. + /// \param NumClauses The number of clauses. + /// + OMPTargetUpdateDirective(SourceLocation StartLoc, SourceLocation EndLoc, + unsigned NumClauses) + : OMPExecutableDirective(this, OMPTargetUpdateDirectiveClass, + OMPD_target_update, StartLoc, EndLoc, NumClauses, + 0) {} + + /// \brief Build an empty directive. + /// + /// \param NumClauses Number of clauses. + /// + explicit OMPTargetUpdateDirective(unsigned NumClauses) + : OMPExecutableDirective(this, OMPTargetUpdateDirectiveClass, + OMPD_target_update, SourceLocation(), + SourceLocation(), NumClauses, 0) {} +public: + /// \brief Creates directive with a list of \a Clauses. + /// + /// \param C AST context. + /// \param StartLoc Starting location of the directive kind. + /// \param EndLoc Ending Location of the directive. + /// \param Clauses List of clauses. + /// + static OMPTargetUpdateDirective *Create(const ASTContext &C, + SourceLocation StartLoc, + SourceLocation EndLoc, + ArrayRef Clauses); + + /// \brief Creates an empty directive with the place for \a NumClauses clauses. + /// + /// \param C AST context. + /// \param NumClauses The number of clauses. + /// + static OMPTargetUpdateDirective *CreateEmpty(const ASTContext &C, + unsigned NumClauses, + EmptyShell); + + static bool classof(const Stmt *T) { + return T->getStmtClass() == OMPTargetUpdateDirectiveClass; + } +}; + } // end namespace clang #endif Index: include/clang/Basic/DiagnosticSemaKinds.td =================================================================== --- include/clang/Basic/DiagnosticSemaKinds.td +++ include/clang/Basic/DiagnosticSemaKinds.td @@ -7805,6 +7805,8 @@ "arguments of OpenMP clause 'reduction' with bitwise operators cannot be of floating type">; def err_omp_once_referenced : Error< "variable can appear only once in OpenMP '%0' clause">; +def err_omp_once_referenced_in_target_update : Error< + "variable can appear only once in OpenMP 'target update' construct">; def note_omp_referenced : Note< "previously referenced here">; def err_omp_reduction_in_task : Error< @@ -7931,8 +7933,8 @@ "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">; +def err_omp_threadprivate_in_clause : Error< + "threadprivate variables are not allowed in '%0' clause">; def err_omp_wrong_ordered_loop_count : Error< "the parameter of the 'ordered' clause must be greater than or equal to the parameter of the 'collapse' clause">; def note_collapse_loop_count : Note< @@ -7977,6 +7979,8 @@ "'schedule' clause with 'nonmonotonic' modifier cannot be specified if an 'ordered' clause is specified">; def err_omp_ordered_simd : Error< "'ordered' clause with a parameter can not be specified in '#pragma omp %0' directive">; +def err_omp_at_least_one_motion_clause_required : Error< + "expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'">; } // end of OpenMP category let CategoryName = "Related Result Type Issue" in { Index: include/clang/Basic/OpenMPKinds.def =================================================================== --- include/clang/Basic/OpenMPKinds.def +++ include/clang/Basic/OpenMPKinds.def @@ -60,6 +60,9 @@ #ifndef OPENMP_TARGET_DATA_CLAUSE # define OPENMP_TARGET_DATA_CLAUSE(Name) #endif +#ifndef OPENMP_TARGET_UPDATE_CLAUSE +# define OPENMP_TARGET_UPDATE_CLAUSE(Name) +#endif #ifndef OPENMP_TEAMS_CLAUSE # define OPENMP_TEAMS_CLAUSE(Name) #endif @@ -125,6 +128,7 @@ OPENMP_DIRECTIVE(teams) OPENMP_DIRECTIVE(cancel) OPENMP_DIRECTIVE_EXT(target_data, "target data") +OPENMP_DIRECTIVE_EXT(target_update, "target update") OPENMP_DIRECTIVE_EXT(parallel_for, "parallel for") OPENMP_DIRECTIVE_EXT(parallel_for_simd, "parallel for simd") OPENMP_DIRECTIVE_EXT(parallel_sections, "parallel sections") @@ -175,6 +179,8 @@ OPENMP_CLAUSE(nogroup, OMPNogroupClause) OPENMP_CLAUSE(num_tasks, OMPNumTasksClause) OPENMP_CLAUSE(hint, OMPHintClause) +OPENMP_CLAUSE(to, OMPToClause) +OPENMP_CLAUSE(from, OMPFromClause) // Clauses allowed for OpenMP directive 'parallel'. OPENMP_PARALLEL_CLAUSE(if) @@ -349,6 +355,13 @@ OPENMP_TARGET_DATA_CLAUSE(device) OPENMP_TARGET_DATA_CLAUSE(map) +// Clauses allowed for OpenMP directive 'target update'. +// TODO More clauses for 'target update' directive. +OPENMP_TARGET_UPDATE_CLAUSE(if) +OPENMP_TARGET_UPDATE_CLAUSE(device) +OPENMP_TARGET_UPDATE_CLAUSE(to) +OPENMP_TARGET_UPDATE_CLAUSE(from) + // Clauses allowed for OpenMP directive 'teams'. // TODO More clauses for 'teams' directive. OPENMP_TEAMS_CLAUSE(default) @@ -449,3 +462,5 @@ #undef OPENMP_FOR_SIMD_CLAUSE #undef OPENMP_MAP_KIND #undef OPENMP_DISTRIBUTE_CLAUSE +#undef OPENMP_TARGET_UPDATE_CLAUSE + Index: include/clang/Basic/StmtNodes.td =================================================================== --- include/clang/Basic/StmtNodes.td +++ include/clang/Basic/StmtNodes.td @@ -216,6 +216,7 @@ def OMPAtomicDirective : DStmt; def OMPTargetDirective : DStmt; def OMPTargetDataDirective : DStmt; +def OMPTargetUpdateDirective : DStmt; def OMPTeamsDirective : DStmt; def OMPCancellationPointDirective : DStmt; def OMPCancelDirective : DStmt; Index: include/clang/Sema/Sema.h =================================================================== --- include/clang/Sema/Sema.h +++ include/clang/Sema/Sema.h @@ -7971,6 +7971,10 @@ ArrayRef Clauses, Stmt *AStmt, SourceLocation StartLoc, SourceLocation EndLoc, llvm::DenseMap &VarsWithImplicitDSA); + /// \brief Called on well-formed '\#pragma omp target update'. + StmtResult ActOnOpenMPTargetUpdateDirective(ArrayRef Clauses, + SourceLocation StartLoc, + SourceLocation EndLoc); OMPClause *ActOnOpenMPSingleExprClause(OpenMPClauseKind Kind, Expr *Expr, @@ -8183,6 +8187,16 @@ OMPClause *ActOnOpenMPPriorityClause(Expr *Priority, SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc); + /// \brief Called on well-formed 'to' clause. + OMPClause *ActOnOpenMPToClause(ArrayRef VarList, + SourceLocation StartLoc, + SourceLocation LParenLoc, + SourceLocation EndLoc); + /// \brief Called on well-formed 'from' clause. + OMPClause *ActOnOpenMPFromClause(ArrayRef VarList, + SourceLocation StartLoc, + SourceLocation LParenLoc, + SourceLocation EndLoc); /// \brief The kind of conversion being performed. enum CheckedConversionKind { Index: include/clang/Serialization/ASTBitCodes.h =================================================================== --- include/clang/Serialization/ASTBitCodes.h +++ include/clang/Serialization/ASTBitCodes.h @@ -1451,6 +1451,7 @@ STMT_OMP_TASKLOOP_SIMD_DIRECTIVE, STMT_OMP_DISTRIBUTE_DIRECTIVE, EXPR_OMP_ARRAY_SECTION, + STMT_OMP_TARGET_UPDATE_DIRECTIVE, // ARC EXPR_OBJC_BRIDGED_CAST, // ObjCBridgedCastExpr Index: lib/AST/OpenMPClause.cpp =================================================================== --- lib/AST/OpenMPClause.cpp +++ lib/AST/OpenMPClause.cpp @@ -415,3 +415,42 @@ void *Mem = C.Allocate(totalSizeToAlloc(N)); return new (Mem) OMPMapClause(N); } + +OMPToClause *OMPToClause::Create(const ASTContext &C, SourceLocation StartLoc, + SourceLocation LParenLoc, SourceLocation EndLoc, + ArrayRef VL) { + void *Mem = C.Allocate( + llvm::RoundUpToAlignment(sizeof(OMPToClause), llvm::alignOf()) + + sizeof(Expr *) * VL.size()); + OMPToClause *Clause = new (Mem) OMPToClause(StartLoc, LParenLoc, EndLoc, + VL.size()); + Clause->setVarRefs(VL); + return Clause; +} + +OMPToClause *OMPToClause::CreateEmpty(const ASTContext &C, unsigned N) { + void *Mem = C.Allocate( + llvm::RoundUpToAlignment(sizeof(OMPToClause), llvm::alignOf()) + + sizeof(Expr *) * N); + return new (Mem) OMPToClause(N); +} + +OMPFromClause * +OMPFromClause::Create(const ASTContext &C, SourceLocation StartLoc, + SourceLocation LParenLoc, SourceLocation EndLoc, + ArrayRef VL) { + void *Mem = C.Allocate( + llvm::RoundUpToAlignment(sizeof(OMPFromClause), llvm::alignOf()) + + sizeof(Expr *) * VL.size()); + OMPFromClause *Clause = new (Mem) OMPFromClause(StartLoc, LParenLoc, EndLoc, + VL.size()); + Clause->setVarRefs(VL); + return Clause; +} + +OMPFromClause *OMPFromClause::CreateEmpty(const ASTContext &C, unsigned N) { + void *Mem = C.Allocate( + llvm::RoundUpToAlignment(sizeof(OMPFromClause), llvm::alignOf()) + + sizeof(Expr *) * N); + return new (Mem) OMPFromClause(N); +} Index: lib/AST/StmtOpenMP.cpp =================================================================== --- lib/AST/StmtOpenMP.cpp +++ lib/AST/StmtOpenMP.cpp @@ -882,3 +882,25 @@ numLoopChildren(CollapsedNum, OMPD_distribute)); return new (Mem) OMPDistributeDirective(CollapsedNum, NumClauses); } + +OMPTargetUpdateDirective * +OMPTargetUpdateDirective::Create(const ASTContext &C, SourceLocation StartLoc, + SourceLocation EndLoc, + ArrayRef Clauses) { + unsigned Size = llvm::RoundUpToAlignment(sizeof(OMPTargetUpdateDirective), + llvm::alignOf()); + void *Mem = C.Allocate(Size + sizeof(OMPClause *) * Clauses.size()); + OMPTargetUpdateDirective *Dir = + new (Mem) OMPTargetUpdateDirective(StartLoc, EndLoc, Clauses.size()); + Dir->setClauses(Clauses); + return Dir; +} + +OMPTargetUpdateDirective * +OMPTargetUpdateDirective::CreateEmpty(const ASTContext &C, unsigned NumClauses, + EmptyShell) { + unsigned Size = llvm::RoundUpToAlignment(sizeof(OMPTargetUpdateDirective), + llvm::alignOf()); + void *Mem = C.Allocate(Size + sizeof(OMPClause *) * NumClauses); + return new (Mem) OMPTargetUpdateDirective(NumClauses); +} Index: lib/AST/StmtPrinter.cpp =================================================================== --- lib/AST/StmtPrinter.cpp +++ lib/AST/StmtPrinter.cpp @@ -909,6 +909,22 @@ OS << ")"; } } + +void OMPClausePrinter::VisitOMPToClause(OMPToClause *Node) { + if (!Node->varlist_empty()) { + OS << "to"; + VisitOMPClauseList(Node, '('); + OS << ")"; + } +} + +void OMPClausePrinter::VisitOMPFromClause(OMPFromClause *Node) { + if (!Node->varlist_empty()) { + OS << "from"; + VisitOMPClauseList(Node, '('); + OS << ")"; + } +} } //===----------------------------------------------------------------------===// @@ -1085,6 +1101,11 @@ PrintOMPExecutableDirective(Node); } +void StmtPrinter::VisitOMPTargetUpdateDirective(OMPTargetUpdateDirective *Node) { + Indent() << "#pragma omp target update "; + PrintOMPExecutableDirective(Node); +} + //===----------------------------------------------------------------------===// // Expr printing methods. //===----------------------------------------------------------------------===// Index: lib/AST/StmtProfile.cpp =================================================================== --- lib/AST/StmtProfile.cpp +++ lib/AST/StmtProfile.cpp @@ -475,6 +475,12 @@ void OMPClauseProfiler::VisitOMPHintClause(const OMPHintClause *C) { Profiler->VisitStmt(C->getHint()); } +void OMPClauseProfiler::VisitOMPToClause(const OMPToClause *C) { + VisitOMPClauseList(C); +} +void OMPClauseProfiler::VisitOMPFromClause(const OMPFromClause *C) { + VisitOMPClauseList(C); +} } void @@ -611,6 +617,11 @@ VisitOMPLoopDirective(S); } +void StmtProfiler::VisitOMPTargetUpdateDirective( + const OMPTargetUpdateDirective *S) { + VisitOMPExecutableDirective(S); +} + void StmtProfiler::VisitExpr(const Expr *S) { VisitStmt(S); } Index: lib/Basic/OpenMPKinds.cpp =================================================================== --- lib/Basic/OpenMPKinds.cpp +++ lib/Basic/OpenMPKinds.cpp @@ -145,6 +145,8 @@ case OMPC_nogroup: case OMPC_num_tasks: case OMPC_hint: + case OMPC_to: + case OMPC_from: break; } llvm_unreachable("Invalid OpenMP simple clause kind"); @@ -255,6 +257,8 @@ case OMPC_nogroup: case OMPC_num_tasks: case OMPC_hint: + case OMPC_to: + case OMPC_from: break; } llvm_unreachable("Invalid OpenMP simple clause kind"); @@ -398,6 +402,16 @@ break; } break; + case OMPD_target_update: + switch (CKind) { +#define OPENMP_TARGET_UPDATE_CLAUSE(Name) \ + case OMPC_##Name: \ + return true; +#include "clang/Basic/OpenMPKinds.def" + default: + break; + } + break; case OMPD_teams: switch (CKind) { #define OPENMP_TEAMS_CLAUSE(Name) \ Index: lib/CodeGen/CGStmt.cpp =================================================================== --- lib/CodeGen/CGStmt.cpp +++ lib/CodeGen/CGStmt.cpp @@ -262,9 +262,12 @@ case Stmt::OMPTaskLoopSimdDirectiveClass: EmitOMPTaskLoopSimdDirective(cast(*S)); break; -case Stmt::OMPDistributeDirectiveClass: + case Stmt::OMPDistributeDirectiveClass: EmitOMPDistributeDirective(cast(*S)); - break; + break; + case Stmt::OMPTargetUpdateDirectiveClass: + EmitOMPTargetUpdateDirective(cast(*S)); + break; } } Index: lib/CodeGen/CGStmtOpenMP.cpp =================================================================== --- lib/CodeGen/CGStmtOpenMP.cpp +++ lib/CodeGen/CGStmtOpenMP.cpp @@ -2525,6 +2525,8 @@ case OMPC_nogroup: case OMPC_num_tasks: case OMPC_hint: + case OMPC_to: + case OMPC_from: llvm_unreachable("Clause is not allowed in 'omp atomic'."); } } @@ -2677,3 +2679,9 @@ [&CS](CodeGenFunction &CGF) { CGF.EmitStmt(CS->getCapturedStmt()); }); } +// Generate the instructions for '#pragma omp target update' directive. +void CodeGenFunction::EmitOMPTargetUpdateDirective( + const OMPTargetUpdateDirective &S) { + llvm_unreachable("CodeGen for 'omp target update' is not supported yet."); +} + Index: lib/CodeGen/CodeGenFunction.h =================================================================== --- lib/CodeGen/CodeGenFunction.h +++ lib/CodeGen/CodeGenFunction.h @@ -2342,6 +2342,7 @@ void EmitOMPTaskLoopDirective(const OMPTaskLoopDirective &S); void EmitOMPTaskLoopSimdDirective(const OMPTaskLoopSimdDirective &S); void EmitOMPDistributeDirective(const OMPDistributeDirective &S); + void EmitOMPTargetUpdateDirective(const OMPTargetUpdateDirective &S); /// \brief Emit inner loop of the worksharing/simd construct. /// Index: lib/Parse/ParseOpenMP.cpp =================================================================== --- lib/Parse/ParseOpenMP.cpp +++ lib/Parse/ParseOpenMP.cpp @@ -34,6 +34,7 @@ {OMPD_unknown /*cancellation*/, OMPD_unknown /*point*/, OMPD_cancellation_point}, {OMPD_target, OMPD_unknown /*data*/, OMPD_target_data}, + {OMPD_target, OMPD_unknown /*update*/, OMPD_target_update}, {OMPD_for, OMPD_simd, OMPD_for_simd}, {OMPD_parallel, OMPD_for, OMPD_parallel_for}, {OMPD_parallel_for, OMPD_simd, OMPD_parallel_for_simd}, @@ -67,7 +68,10 @@ TokenMatched = ((i == 0) && !P.getPreprocessor().getSpelling(Tok).compare("point")) || - ((i == 1) && !P.getPreprocessor().getSpelling(Tok).compare("data")); + ((i == 1) && + !P.getPreprocessor().getSpelling(Tok).compare("data")) || + ((i == 2) && + !P.getPreprocessor().getSpelling(Tok).compare("update")); } else { TokenMatched = SDKind == F[i][1] && SDKind != OMPD_unknown; } @@ -141,6 +145,7 @@ case OMPD_taskloop: case OMPD_taskloop_simd: case OMPD_distribute: + case OMPD_target_update: Diag(Tok, diag::err_omp_unexpected_directive) << getOpenMPDirectiveName(DKind); break; @@ -161,8 +166,8 @@ /// 'parallel for' | 'parallel sections' | 'task' | 'taskyield' | /// 'barrier' | 'taskwait' | 'flush' | 'ordered' | 'atomic' | /// 'for simd' | 'parallel for simd' | 'target' | 'target data' | -/// 'taskgroup' | 'teams' | 'taskloop' | 'taskloop simd' {clause} | -/// 'distribute' +/// 'taskgroup' | 'teams' | 'taskloop' | 'taskloop simd' | +/// 'distribute' | 'target update' {clause} /// annot_pragma_openmp_end /// StmtResult @@ -213,6 +218,7 @@ case OMPD_taskwait: case OMPD_cancellation_point: case OMPD_cancel: + case OMPD_target_update: if (!StandAloneAllowed) { Diag(Tok, diag::err_omp_immediate_directive) << getOpenMPDirectiveName(DKind) << 0; @@ -416,7 +422,8 @@ /// update-clause | capture-clause | seq_cst-clause | device-clause | /// simdlen-clause | threads-clause | simd-clause | num_teams-clause | /// thread_limit-clause | priority-clause | grainsize-clause | -/// nogroup-clause | num_tasks-clause | hint-clause +/// nogroup-clause | num_tasks-clause | hint-clause | to-clause | +/// from-clause /// OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind, OpenMPClauseKind CKind, bool FirstClause) { @@ -536,6 +543,8 @@ case OMPC_flush: case OMPC_depend: case OMPC_map: + case OMPC_to: + case OMPC_from: Clause = ParseOpenMPVarListClause(DKind, CKind); break; case OMPC_unknown: @@ -820,8 +829,7 @@ TemplateKWLoc, ReductionId); } -/// \brief Parsing of OpenMP clause 'private', 'firstprivate', 'lastprivate', -/// 'shared', 'copyin', 'copyprivate', 'flush' or 'reduction'. +/// \brief Parsing of OpenMP clauses /// /// private-clause: /// 'private' '(' list ')' @@ -846,6 +854,10 @@ /// map-clause: /// 'map' '(' [ [ always , ] /// to | from | tofrom | alloc | release | delete ':' ] list ')'; +/// to-clause: +/// 'to' '(' list ')' +/// from-clause: +/// 'from' '(' list ')' /// /// For 'linear' clause linear-list may have the following forms: /// list Index: lib/Sema/SemaOpenMP.cpp =================================================================== --- lib/Sema/SemaOpenMP.cpp +++ lib/Sema/SemaOpenMP.cpp @@ -1610,6 +1610,7 @@ case OMPD_cancellation_point: case OMPD_cancel: case OMPD_flush: + case OMPD_target_update: llvm_unreachable("OpenMP Directive is not allowed"); case OMPD_unknown: llvm_unreachable("Unknown OpenMP directive"); @@ -2711,6 +2712,12 @@ Res = ActOnOpenMPDistributeDirective(ClausesWithImplicit, AStmt, StartLoc, EndLoc, VarsWithInheritedDSA); break; + case OMPD_target_update: + assert(!AStmt && "Statement is not allowed for target update"); + Res = + ActOnOpenMPTargetUpdateDirective(ClausesWithImplicit, StartLoc, EndLoc); + AllowedNameModifiers.push_back(OMPD_target_update); + break; case OMPD_threadprivate: llvm_unreachable("OpenMP Directive is not allowed"); case OMPD_unknown: @@ -5448,6 +5455,23 @@ AStmt); } +StmtResult Sema::ActOnOpenMPTargetUpdateDirective(ArrayRef Clauses, + SourceLocation StartLoc, + SourceLocation EndLoc) { + bool seenMotionClause = false; + for (auto *C : Clauses) { + if (C->getClauseKind() == OMPC_to || C->getClauseKind() == OMPC_from) + seenMotionClause = true; + } + if (!seenMotionClause) { + Diag(StartLoc, diag::err_omp_at_least_one_motion_clause_required); + return StmtError(); + } + + getCurFunction()->setHasBranchProtectedScope(); + return OMPTargetUpdateDirective::Create(Context, StartLoc, EndLoc, Clauses); +} + StmtResult Sema::ActOnOpenMPTeamsDirective(ArrayRef Clauses, Stmt *AStmt, SourceLocation StartLoc, SourceLocation EndLoc) { @@ -5699,6 +5723,8 @@ case OMPC_simd: case OMPC_map: case OMPC_nogroup: + case OMPC_to: + case OMPC_from: case OMPC_unknown: llvm_unreachable("Clause is not allowed."); } @@ -5982,6 +6008,8 @@ case OMPC_nogroup: case OMPC_num_tasks: case OMPC_hint: + case OMPC_to: + case OMPC_from: case OMPC_unknown: llvm_unreachable("Clause is not allowed."); } @@ -6119,6 +6147,8 @@ case OMPC_nogroup: case OMPC_num_tasks: case OMPC_hint: + case OMPC_to: + case OMPC_from: case OMPC_unknown: llvm_unreachable("Clause is not allowed."); } @@ -6303,6 +6333,8 @@ case OMPC_grainsize: case OMPC_num_tasks: case OMPC_hint: + case OMPC_to: + case OMPC_from: case OMPC_unknown: llvm_unreachable("Clause is not allowed."); } @@ -6415,6 +6447,12 @@ Res = ActOnOpenMPMapClause(MapTypeModifier, MapType, DepLinMapLoc, ColonLoc, VarList, StartLoc, LParenLoc, EndLoc); break; + case OMPC_to: + Res = ActOnOpenMPToClause(VarList, StartLoc, LParenLoc, EndLoc); + break; + case OMPC_from: + Res = ActOnOpenMPFromClause(VarList, StartLoc, LParenLoc, EndLoc); + break; case OMPC_if: case OMPC_final: case OMPC_num_threads: @@ -8370,7 +8408,7 @@ (ASE && !ASE->getBase()->getType()->isAnyPointerType() && !ASE->getBase()->getType()->isArrayType())) { Diag(ELoc, diag::err_omp_expected_var_name_or_array_item) - << RE->getSourceRange(); + << RE->getSourceRange(); continue; } @@ -8391,7 +8429,7 @@ // 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); + Diag(ELoc, diag::err_omp_threadprivate_in_clause) << "map"; ReportOriginalDSA(*this, DSAStack, VD, DVar); continue; } @@ -8542,3 +8580,188 @@ OMPHintClause(HintExpr.get(), StartLoc, LParenLoc, EndLoc); } +OMPClause *Sema::ActOnOpenMPToClause(ArrayRef VarList, + SourceLocation StartLoc, + SourceLocation LParenLoc, + SourceLocation EndLoc) { + SmallVector Vars; + for (auto &RE : VarList) { + assert(RE && "Null expr in omp to clause"); + if (isa(RE)) { + // It will be analyzed later. + Vars.push_back(RE); + continue; + } + SourceLocation ELoc = RE->getExprLoc(); + + // OpenMP [2.9.3, 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 + // as a list item in a clause of a target update construct. + 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 to clause."); + auto *VD = cast(D); + + // threadprivate variables cannot appear in a to clause. + if (DSAStack->isThreadPrivate(VD)) { + auto DVar = DSAStack->getTopDSA(VD, false); + Diag(ELoc, diag::err_omp_threadprivate_in_clause) << "to"; + ReportOriginalDSA(*this, DSAStack, VD, DVar); + continue; + } + + // OpenMP [2.9.3, Restrictions, p.6] + // A list item in a to or from clause must have a mappable type. + QualType Type = VD->getType(); + if (!CheckTypeMappable(VE->getExprLoc(), VE->getSourceRange(), *this, + DSAStack, Type)) { + continue; + } + + // OpenMP [2.9.3, Restrictions, p.6] + // A list item can only appear in a to or from clause, but not both. + DSAStackTy::MapInfo MI = DSAStack->IsMappedInCurrentRegion(VD); + if (MI.RefExpr) { + Diag(ELoc, diag::err_omp_once_referenced_in_target_update) + << VD->getSourceRange(); + Diag(MI.RefExpr->getExprLoc(), diag::note_used_here) + << MI.RefExpr->getSourceRange(); + continue; + } + + Vars.push_back(RE); + MI.RefExpr = RE; + DSAStack->addMapInfoForVar(VD, MI); + } + + if (Vars.empty()) + return nullptr; + + return OMPToClause::Create(Context, StartLoc, LParenLoc, EndLoc, Vars); +} + +OMPClause *Sema::ActOnOpenMPFromClause(ArrayRef VarList, + SourceLocation StartLoc, + SourceLocation LParenLoc, + SourceLocation EndLoc) { + SmallVector Vars; + for (auto &RE : VarList) { + assert(RE && "Null expr in omp from clause"); + if (isa(RE)) { + // It will be analyzed later. + Vars.push_back(RE); + continue; + } + SourceLocation ELoc = RE->getExprLoc(); + + // OpenMP [2.9.3, 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 + // as a list item in a clause of a target update construct. + 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 from clause."); + auto *VD = cast(D); + + // threadprivate variables cannot appear in a from clause. + if (DSAStack->isThreadPrivate(VD)) { + auto DVar = DSAStack->getTopDSA(VD, false); + Diag(ELoc, diag::err_omp_threadprivate_in_clause) << "from"; + ReportOriginalDSA(*this, DSAStack, VD, DVar); + continue; + } + + // OpenMP [2.9.3, Restrictions, p.6] + // A list item in a to or from clause must have a mappable type. + QualType Type = VD->getType(); + if (!CheckTypeMappable(VE->getExprLoc(), VE->getSourceRange(), *this, + DSAStack, Type)) { + continue; + } + + // OpenMP [2.9.3, Restrictions, p.6] + // A list item can only appear in a to or from clause, but not both. + DSAStackTy::MapInfo MI = DSAStack->IsMappedInCurrentRegion(VD); + if (MI.RefExpr) { + Diag(ELoc, diag::err_omp_once_referenced_in_target_update) + << VD->getSourceRange(); + Diag(MI.RefExpr->getExprLoc(), diag::note_used_here) + << MI.RefExpr->getSourceRange(); + continue; + } + + Vars.push_back(RE); + MI.RefExpr = RE; + DSAStack->addMapInfoForVar(VD, MI); + } + + if (Vars.empty()) + return nullptr; + + return OMPFromClause::Create(Context, StartLoc, LParenLoc, EndLoc, Vars); +} Index: lib/Sema/TreeTransform.h =================================================================== --- lib/Sema/TreeTransform.h +++ lib/Sema/TreeTransform.h @@ -1731,6 +1731,29 @@ return getSema().ActOnOpenMPHintClause(Hint, StartLoc, LParenLoc, EndLoc); } + /// \brief Build a new OpenMP 'to' clause. + /// + /// By default, performs semantic analysis to build the new statement. + /// Subclasses may override this routine to provide different behavior. + OMPClause *RebuildOMPToClause(ArrayRef VarList, + SourceLocation StartLoc, + SourceLocation LParenLoc, + SourceLocation EndLoc) { + return getSema().ActOnOpenMPToClause(VarList, StartLoc, LParenLoc, EndLoc); + } + + /// \brief Build a new OpenMP 'from' clause. + /// + /// By default, performs semantic analysis to build the new statement. + /// Subclasses may override this routine to provide different behavior. + OMPClause *RebuildOMPFromClause(ArrayRef VarList, + SourceLocation StartLoc, + SourceLocation LParenLoc, + SourceLocation EndLoc) { + return getSema().ActOnOpenMPFromClause(VarList, StartLoc, LParenLoc, + EndLoc); + } + /// \brief Rebuild the operand to an Objective-C \@synchronized statement. /// /// By default, performs semantic analysis to build the new statement. @@ -7353,6 +7376,17 @@ } template +StmtResult TreeTransform::TransformOMPTargetUpdateDirective( + OMPTargetUpdateDirective *D) { + DeclarationNameInfo DirName; + getDerived().getSema().StartOpenMPDSABlock(OMPD_target_update, DirName, + nullptr, D->getLocStart()); + StmtResult Res = getDerived().TransformOMPExecutableDirective(D); + getDerived().getSema().EndOpenMPDSABlock(Res.get()); + return Res; +} + +template StmtResult TreeTransform::TransformOMPTeamsDirective(OMPTeamsDirective *D) { DeclarationNameInfo DirName; @@ -7864,6 +7898,34 @@ C->getLParenLoc(), C->getLocEnd()); } +template +OMPClause *TreeTransform::TransformOMPToClause(OMPToClause *C) { + llvm::SmallVector Vars; + Vars.reserve(C->varlist_size()); + for (auto *VE : C->varlists()) { + ExprResult EVar = getDerived().TransformExpr(cast(VE)); + if (EVar.isInvalid()) + return 0; + Vars.push_back(EVar.get()); + } + return getDerived().RebuildOMPToClause(Vars, C->getLocStart(), + C->getLParenLoc(), C->getLocEnd()); +} + +template +OMPClause *TreeTransform::TransformOMPFromClause(OMPFromClause *C) { + llvm::SmallVector Vars; + Vars.reserve(C->varlist_size()); + for (auto *VE : C->varlists()) { + ExprResult EVar = getDerived().TransformExpr(cast(VE)); + if (EVar.isInvalid()) + return 0; + Vars.push_back(EVar.get()); + } + return getDerived().RebuildOMPFromClause(Vars, C->getLocStart(), + C->getLParenLoc(), C->getLocEnd()); +} + //===----------------------------------------------------------------------===// // Expression transformation //===----------------------------------------------------------------------===// Index: lib/Serialization/ASTReaderStmt.cpp =================================================================== --- lib/Serialization/ASTReaderStmt.cpp +++ lib/Serialization/ASTReaderStmt.cpp @@ -1881,6 +1881,12 @@ case OMPC_hint: C = new (Context) OMPHintClause(); break; + case OMPC_to: + C = OMPToClause::CreateEmpty(Context, Record[Idx++]); + break; + case OMPC_from: + C = OMPFromClause::CreateEmpty(Context, Record[Idx++]); + break; } Visit(C); C->setLocStart(Reader->ReadSourceLocation(Record, Idx)); @@ -2243,6 +2249,28 @@ C->setLParenLoc(Reader->ReadSourceLocation(Record, Idx)); } +void OMPClauseReader::VisitOMPToClause(OMPToClause *C) { + C->setLParenLoc(Reader->ReadSourceLocation(Record, Idx)); + unsigned 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); +} + +void OMPClauseReader::VisitOMPFromClause(OMPFromClause *C) { + C->setLParenLoc(Reader->ReadSourceLocation(Record, Idx)); + unsigned 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. //===----------------------------------------------------------------------===// @@ -2479,6 +2507,12 @@ VisitOMPLoopDirective(D); } +void ASTStmtReader::VisitOMPTargetUpdateDirective(OMPTargetUpdateDirective *D) { + VisitStmt(D); + ++Idx; + VisitOMPExecutableDirective(D); +} + //===----------------------------------------------------------------------===// // ASTReader Implementation //===----------------------------------------------------------------------===// @@ -3088,6 +3122,11 @@ Context, Record[ASTStmtReader::NumStmtFields], Empty); break; + case STMT_OMP_TARGET_UPDATE_DIRECTIVE: + S = OMPTargetUpdateDirective::CreateEmpty( + Context, Record[ASTStmtReader::NumStmtFields], Empty); + break; + case STMT_OMP_TEAMS_DIRECTIVE: S = OMPTeamsDirective::CreateEmpty( Context, Record[ASTStmtReader::NumStmtFields], Empty); Index: lib/Serialization/ASTWriterStmt.cpp =================================================================== --- lib/Serialization/ASTWriterStmt.cpp +++ lib/Serialization/ASTWriterStmt.cpp @@ -2044,6 +2044,20 @@ Writer->Writer.AddSourceLocation(C->getLParenLoc(), Record); } +void OMPClauseWriter::VisitOMPToClause(OMPToClause *C) { + Record.push_back(C->varlist_size()); + Writer->Writer.AddSourceLocation(C->getLParenLoc(), Record); + for (auto *VE : C->varlists()) + Writer->Writer.AddStmt(VE); +} + +void OMPClauseWriter::VisitOMPFromClause(OMPFromClause *C) { + Record.push_back(C->varlist_size()); + Writer->Writer.AddSourceLocation(C->getLParenLoc(), Record); + for (auto *VE : C->varlists()) + Writer->Writer.AddStmt(VE); +} + //===----------------------------------------------------------------------===// // OpenMP Directives. //===----------------------------------------------------------------------===// @@ -2288,6 +2302,13 @@ Code = serialization::STMT_OMP_DISTRIBUTE_DIRECTIVE; } +void ASTStmtWriter::VisitOMPTargetUpdateDirective(OMPTargetUpdateDirective *D) { + VisitStmt(D); + Record.push_back(D->getNumClauses()); + VisitOMPExecutableDirective(D); + Code = serialization::STMT_OMP_TARGET_UPDATE_DIRECTIVE; +} + //===----------------------------------------------------------------------===// // ASTWriter Implementation //===----------------------------------------------------------------------===// Index: lib/StaticAnalyzer/Core/ExprEngine.cpp =================================================================== --- lib/StaticAnalyzer/Core/ExprEngine.cpp +++ lib/StaticAnalyzer/Core/ExprEngine.cpp @@ -830,6 +830,7 @@ case Stmt::OMPAtomicDirectiveClass: case Stmt::OMPTargetDirectiveClass: case Stmt::OMPTargetDataDirectiveClass: + case Stmt::OMPTargetUpdateDirectiveClass: case Stmt::OMPTeamsDirectiveClass: case Stmt::OMPCancellationPointDirectiveClass: case Stmt::OMPCancelDirectiveClass: Index: test/OpenMP/target_map_messages.cpp =================================================================== --- test/OpenMP/target_map_messages.cpp +++ test/OpenMP/target_map_messages.cpp @@ -107,7 +107,7 @@ #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(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(); @@ -177,7 +177,7 @@ #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(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(); Index: test/OpenMP/target_update_ast_print.cpp =================================================================== --- /dev/null +++ test/OpenMP/target_update_ast_print.cpp @@ -0,0 +1,52 @@ +// RUN: %clang_cc1 -verify -fopenmp -ast-print %s | FileCheck %s +// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s +// expected-no-diagnostics + +#ifndef HEADER +#define HEADER + +void foo() {} + +template +T foo(T targ, U uarg) { + static T a; + U b; + int l; +#pragma omp target update to(a) if(l>5) device(l) + +#pragma omp target update from(b) if(l<5) device(l-1) + return a + targ + (T)b; +} +// CHECK: static int a; +// CHECK-NEXT: float b; +// CHECK-NEXT: int l; +// CHECK-NEXT: #pragma omp target update to(a) if(l > 5) device(l) +// CHECK-NEXT: #pragma omp target update from(b) if(l < 5) device(l - 1) +// CHECK: static char a; +// CHECK-NEXT: float b; +// CHECK-NEXT: int l; +// CHECK-NEXT: #pragma omp target update to(a) if(l > 5) device(l) +// CHECK-NEXT: #pragma omp target update from(b) if(l < 5) device(l - 1) +// CHECK: static T a; +// CHECK-NEXT: U b; +// CHECK-NEXT: int l; +// CHECK-NEXT: #pragma omp target update to(a) if(l > 5) device(l) +// CHECK-NEXT: #pragma omp target update from(b) if(l < 5) device(l - 1) + +int main(int argc, char **argv) { + static int a; + int n; + float f; + +// CHECK: static int a; +// CHECK-NEXT: int n; +// CHECK-NEXT: float f; +#pragma omp target update to(a) if(f>0.0) device(n) + // CHECK-NEXT: #pragma omp target update to(a) if(f > 0.) device(n) +#pragma omp target update from(f) if(f<0.0) device(n+1) + // CHECK-NEXT: #pragma omp target update from(f) if(f < 0.) device(n + 1) + return foo(argc, f) + foo(argv[0][0], f) + a; +} + +#endif Index: test/OpenMP/target_update_device_messages.cpp =================================================================== --- /dev/null +++ test/OpenMP/target_update_device_messages.cpp @@ -0,0 +1,43 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s + +void foo() { +} + +bool foobool(int argc) { + return argc; +} + +struct S1; // expected-note 2 {{declared here}} + +template +int tmain(T argc, S **argv) { + int i; + #pragma omp target update to(i) device // expected-error {{expected '(' after 'device'}} + #pragma omp target update to(i) device ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target update to(i) device () // expected-error {{expected expression}} + #pragma omp target update to(i) device (argc // expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target update to(i) device (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} + #pragma omp target update from(i) device (argc > 0 ? argv[1] : argv[2]) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}} + #pragma omp target update from(i) device (argc + argc) + #pragma omp target update from(i) device (argc), device (argc+1) // expected-error {{directive '#pragma omp target update' cannot contain more than one 'device' clause}} + #pragma omp target update from(i) device (S1) // expected-error {{'S1' does not refer to a value}} + #pragma omp target update from(i) device (3.14) // expected-error 2 {{expression must have integral or unscoped enumeration type, not 'double'}} + #pragma omp target update from(i) device (-2) // expected-error {{argument to 'device' clause must be a non-negative integer value}} +} + +int main(int argc, char **argv) { + int j; + #pragma omp target update to(j) device // expected-error {{expected '(' after 'device'}} + #pragma omp target update from(j) device ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target update to(j) device () // expected-error {{expected expression}} + #pragma omp target update from(j) device (argc // expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target update to(j) device (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} + #pragma omp target update from(j) device (argc > 0 ? argv[1] : argv[2]) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}} + #pragma omp target update to(j) device (argc + argc) + #pragma omp target update from(j) device (argc), device (argc+1) // expected-error {{directive '#pragma omp target update' cannot contain more than one 'device' clause}} + #pragma omp target update to(j) device (S1) // expected-error {{'S1' does not refer to a value}} + #pragma omp target update from(j) device (-2) // expected-error {{argument to 'device' clause must be a non-negative integer value}} + #pragma omp target update to(j) device (3.14) // expected-error {{expression must have integral or unscoped enumeration type, not 'double'}} + + return tmain(argc, argv); // expected-note {{in instantiation of function template specialization 'tmain' requested here}} +} Index: test/OpenMP/target_update_from_messages.cpp =================================================================== --- /dev/null +++ test/OpenMP/target_update_from_messages.cpp @@ -0,0 +1,127 @@ +// 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 to; + +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 from; + const T (&l)[5] = da; + +#pragma omp target update from // expected-error {{expected '(' after 'from'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from() // expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update() // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(alloc) // expected-error {{use of undeclared identifier 'alloc'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(x) +#pragma omp target update from(t[:I]) +#pragma omp target update from(T) // expected-error {{'T' does not refer to a value}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(I) // expected-error 2 {{expected variable name, array element or array section}} +#pragma omp target update from(S2::S2s) +#pragma omp target update from(S2::S2sc) +#pragma omp target update from(from) +#pragma omp target update from(y x) // expected-error {{expected ',' or ')' in 'from' clause}} +#pragma omp target update from(argc > 0 ? x : y) // expected-error 2 {{expected variable name, array element or array section}} +#pragma omp target update from(S1) // expected-error {{'S1' does not refer to a value}}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(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 update from(ba) // expected-error 2 {{type 'S2' is not mappable to target}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(h) // expected-error {{threadprivate variables are not allowed in 'from' clause}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(k), to(k) // expected-error 2 {{variable can appear only once in OpenMP 'target update' construct}} expected-note 2 {{used here}} +#pragma omp target update from(k), from(k[:5]) // expected-error 2 {{variable can appear only once in OpenMP 'target update' construct}} expected-note 2 {{used here}} +#pragma omp target update from(da) +#pragma omp target update from(da[:4]) + 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, t[20]; + int &j = i; + int *k = &j; + int x; + int y; + int from; + const int (&l)[5] = da; + +#pragma omp target update from // expected-error {{expected '(' after 'from'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from() // expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update() // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(alloc) // expected-error {{use of undeclared identifier 'alloc'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(x) +#pragma omp target update from(t[:i]) +#pragma omp target update from(S2::S2s) +#pragma omp target update from(S2::S2sc) +#pragma omp target update from(from) +#pragma omp target update from(y x) // expected-error {{expected ',' or ')' in 'from' clause}} +#pragma omp target update from(argc > 0 ? x : y) // expected-error {{expected variable name, array element or array section}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(S1) // expected-error {{'S1' does not refer to a value}}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(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 update from(ba) // expected-error 2 {{type 'S2' is not mappable to target}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(h) // expected-error {{threadprivate variables are not allowed in 'from' clause}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update from(k), to(k) // expected-error {{variable can appear only once in OpenMP 'target update' construct}} expected-note {{used here}} +#pragma omp target update from(k), from(k[:5]) // expected-error {{variable can appear only once in OpenMP 'target update' construct}} expected-note {{used here}} +#pragma omp target update from(da) +#pragma omp target update from(da[:4]) + + 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: test/OpenMP/target_update_if_messages.cpp =================================================================== --- /dev/null +++ test/OpenMP/target_update_if_messages.cpp @@ -0,0 +1,68 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s + +void foo() { +} + +bool foobool(int argc) { + return argc; +} + +struct S1; // expected-note {{declared here}} + +template // expected-note {{declared here}} +int tmain(T argc, S **argv) { + int n; + #pragma omp target + { + #pragma omp target update to(n) if // expected-error {{expected '(' after 'if'}} + #pragma omp target update from(n) if ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target update to(n) if () // expected-error {{expected expression}} + #pragma omp target update from(n) if (argc // expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target update to(n) if (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} + #pragma omp target update from(n) if (argc > 0 ? argv[1] : argv[2]) + #pragma omp target update to(n) if (foobool(argc)), if (true) // expected-error {{directive '#pragma omp target update' cannot contain more than one 'if' clause}} + #pragma omp target update from(n) if (S) // expected-error {{'S' does not refer to a value}} + #pragma omp target update to(n) if (argv[1]=2) // expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target update from(n) if (argc argc) // expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target update to(n) if(argc) + #pragma omp target update from(n) if(target update // expected-warning {{missing ':' after directive name modifier - ignoring}} expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target update to(n) if(target update : // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target update from(n) if(target update : argc // expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target update to(n) if(target update : argc) + #pragma omp target update from(n) if(target update : argc) if (for:argc) // expected-error {{directive name modifier 'for' is not allowed for '#pragma omp target update'}} + #pragma omp target update to(n) if(target update : argc) if (target update:argc) // expected-error {{directive '#pragma omp target update' cannot contain more than one 'if' clause with 'target update' name modifier}} + #pragma omp target update from(n) if(target update : argc) if (argc) // expected-error {{no more 'if' clause is allowed}} expected-note {{previous clause with directive name modifier specified here}} + foo(); + } + + return 0; +} + +int main(int argc, char **argv) { + int m; + #pragma omp parallel + { + #pragma omp target update to(m) if // expected-error {{expected '(' after 'if'}} + #pragma omp target update from(m) if ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target update to(m) if () // expected-error {{expected expression}} + #pragma omp target update from(m) if (argc // expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target update to(m) if (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} + #pragma omp target update from(m) if (argc > 0 ? argv[1] : argv[2]) + #pragma omp target update to(m) if (foobool(argc)), if (true) // expected-error {{directive '#pragma omp target update' cannot contain more than one 'if' clause}} + #pragma omp target update from(m) if (S1) // expected-error {{'S1' does not refer to a value}} + #pragma omp target update to(m) if (argv[1]=2) // expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target update from(m) if (argc argc) // expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target update to(m) if (1 0) // expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target update from(m) if(if(tmain(argc, argv) // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target update to(m) if(target update // expected-warning {{missing ':' after directive name modifier - ignoring}} expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target update from(m) if(target update : // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target update to(m) if(target update : argc // expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target update from(m) if(target update : argc) + #pragma omp target update to(m) if(target update : argc) if (for:argc) // expected-error {{directive name modifier 'for' is not allowed for '#pragma omp target update'}} + #pragma omp target update from(m) if(target update : argc) if (target update:argc) // expected-error {{directive '#pragma omp target update' cannot contain more than one 'if' clause with 'target update' name modifier}} + #pragma omp target update to(m) if(target update : argc) if (argc) // expected-error {{no more 'if' clause is allowed}} expected-note {{previous clause with directive name modifier specified here}} + foo(); + } + + return tmain(argc, argv); +} Index: test/OpenMP/target_update_to_messages.cpp =================================================================== --- /dev/null +++ test/OpenMP/target_update_to_messages.cpp @@ -0,0 +1,127 @@ +// 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; + const T (&l)[5] = da; + +#pragma omp target update to // expected-error {{expected '(' after 'to'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to() // expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update() // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(alloc) // expected-error {{use of undeclared identifier 'alloc'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(x) +#pragma omp target update to(t[:I]) +#pragma omp target update to(T) // expected-error {{'T' does not refer to a value}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(I) // expected-error 2 {{expected variable name, array element or array section}} +#pragma omp target update to(S2::S2s) +#pragma omp target update to(S2::S2sc) +#pragma omp target update to(to) +#pragma omp target update to(y x) // expected-error {{expected ',' or ')' in 'to' clause}} +#pragma omp target update to(argc > 0 ? x : y) // expected-error 2 {{expected variable name, array element or array section}} +#pragma omp target update to(S1) // expected-error {{'S1' does not refer to a value}}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(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 update to(ba) // expected-error 2 {{type 'S2' is not mappable to target}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(h) // expected-error {{threadprivate variables are not allowed in 'to' clause}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(k), from(k) // expected-error 2 {{variable can appear only once in OpenMP 'target update' construct}} expected-note 2 {{used here}} +#pragma omp target update to(k), to(k[:5]) // expected-error 2 {{variable can appear only once in OpenMP 'target update' construct}} expected-note 2 {{used here}} +#pragma omp target update to(da) +#pragma omp target update to(da[:4]) + 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, t[20]; + int &j = i; + int *k = &j; + int x; + int y; + int to; + const int (&l)[5] = da; + +#pragma omp target update to // expected-error {{expected '(' after 'to'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to() // expected-error {{expected expression}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update() // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(alloc) // expected-error {{use of undeclared identifier 'alloc'}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(x) +#pragma omp target update to(t[:i]) +#pragma omp target update to(S2::S2s) +#pragma omp target update to(S2::S2sc) +#pragma omp target update to(to) +#pragma omp target update to(y x) // expected-error {{expected ',' or ')' in 'to' clause}} +#pragma omp target update to(argc > 0 ? x : y) // expected-error {{expected variable name, array element or array section}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(S1) // expected-error {{'S1' does not refer to a value}}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(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 update to(ba) // expected-error 2 {{type 'S2' is not mappable to target}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(h) // expected-error {{threadprivate variables are not allowed in 'to' clause}} expected-error {{expected at least one 'to' clause or 'from' clause specified to '#pragma omp target update'}} +#pragma omp target update to(k), from(k) // expected-error {{variable can appear only once in OpenMP 'target update' construct}} expected-note {{used here}} +#pragma omp target update to(k), to(k[:5]) // expected-error {{variable can appear only once in OpenMP 'target update' construct}} expected-note {{used here}} +#pragma omp target update to(da) +#pragma omp target update to(da[:4]) + + 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 @@ -2218,6 +2218,12 @@ void OMPClauseEnqueue::VisitOMPMapClause(const OMPMapClause *C) { VisitOMPClauseList(C); } +void OMPClauseEnqueue::VisitOMPFromClause(const OMPFromClause *C) { + VisitOMPClauseList(C); +} +void OMPClauseEnqueue::VisitOMPToClause(const OMPToClause *C) { + VisitOMPClauseList(C); +} } void EnqueueVisitor::EnqueueChildren(const OMPClause *S) { @@ -4506,6 +4512,8 @@ return cxstring::createRef("OMPTargetDirective"); case CXCursor_OMPTargetDataDirective: return cxstring::createRef("OMPTargetDataDirective"); + case CXCursor_OMPTargetUpdateDirective: + return cxstring::createRef("OMPTargetUpdateDirective"); case CXCursor_OMPTeamsDirective: return cxstring::createRef("OMPTeamsDirective"); case CXCursor_OMPCancellationPointDirective: Index: tools/libclang/CXCursor.cpp =================================================================== --- tools/libclang/CXCursor.cpp +++ tools/libclang/CXCursor.cpp @@ -600,6 +600,9 @@ case Stmt::OMPTargetDataDirectiveClass: K = CXCursor_OMPTargetDataDirective; break; + case Stmt::OMPTargetUpdateDirectiveClass: + K = CXCursor_OMPTargetUpdateDirective; + break; case Stmt::OMPTeamsDirectiveClass: K = CXCursor_OMPTeamsDirective; break;