Index: include/clang-c/Index.h =================================================================== --- include/clang-c/Index.h +++ include/clang-c/Index.h @@ -2281,7 +2281,7 @@ */ CXCursor_OMPTaskLoopSimdDirective = 259, - /** \brief OpenMP distribute directive. + /** \brief OpenMP distribute directive. */ CXCursor_OMPDistributeDirective = 260, @@ -2301,7 +2301,11 @@ */ CXCursor_OMPTargetParallelForDirective = 264, - CXCursor_LastStmt = CXCursor_OMPTargetParallelForDirective, + /** \brief OpenMP target update directive. + */ + CXCursor_OMPTargetUpdateDirective = 265, + + 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 @@ -3471,6 +3471,128 @@ return child_range(child_iterator(), child_iterator()); } }; + +/// \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 @@ -2496,6 +2496,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)); }) @@ -2874,6 +2877,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 @@ -2660,7 +2660,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, @@ -2682,6 +2682,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 @@ -7823,8 +7823,8 @@ "expected variable name%select{|, data member of current class}0, array element or array section">; def err_omp_expected_named_var_member_or_array_expression: Error< "expected expression containing only member accesses and/or array sections based on named variables">; -def err_omp_bit_fields_forbidden_in_map_clause : Error< - "bit fields cannot be used to specify storage in a map clause">; +def err_omp_bit_fields_forbidden_in_clause : Error< + "bit fields cannot be used to specify storage in a '%0' clause">; def err_array_section_does_not_specify_contiguous_storage : Error< "array section does not specify contiguous storage">; def err_omp_union_type_not_allowed : Error< @@ -7949,6 +7949,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< @@ -8082,8 +8084,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< @@ -8128,6 +8130,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 @@ -72,6 +72,9 @@ #ifndef OPENMP_TARGET_PARALLEL_FOR_CLAUSE # define OPENMP_TARGET_PARALLEL_FOR_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 @@ -150,6 +153,7 @@ OPENMP_DIRECTIVE_EXT(target_exit_data, "target exit data") OPENMP_DIRECTIVE_EXT(target_parallel, "target parallel") OPENMP_DIRECTIVE_EXT(target_parallel_for, "target parallel for") +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") @@ -203,6 +207,8 @@ OPENMP_CLAUSE(hint, OMPHintClause) OPENMP_CLAUSE(dist_schedule, OMPDistScheduleClause) OPENMP_CLAUSE(defaultmap, OMPDefaultmapClause) +OPENMP_CLAUSE(to, OMPToClause) +OPENMP_CLAUSE(from, OMPFromClause) // Clauses allowed for OpenMP directive 'parallel'. OPENMP_PARALLEL_CLAUSE(if) @@ -439,6 +445,13 @@ OPENMP_TARGET_PARALLEL_FOR_CLAUSE(ordered) OPENMP_TARGET_PARALLEL_FOR_CLAUSE(linear) +// 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) @@ -550,3 +563,4 @@ #undef OPENMP_DIST_SCHEDULE_KIND #undef OPENMP_DEFAULTMAP_KIND #undef OPENMP_DEFAULTMAP_MODIFIER +#undef OPENMP_TARGET_UPDATE_CLAUSE Index: include/clang/Basic/StmtNodes.td =================================================================== --- include/clang/Basic/StmtNodes.td +++ include/clang/Basic/StmtNodes.td @@ -220,6 +220,7 @@ def OMPTargetExitDataDirective : DStmt; def OMPTargetParallelDirective : DStmt; def OMPTargetParallelForDirective : 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 @@ -8074,6 +8074,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, @@ -8299,6 +8303,16 @@ OpenMPDefaultmapClauseModifier M, OpenMPDefaultmapClauseKind Kind, SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation MLoc, SourceLocation KindLoc, 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 @@ -1479,6 +1479,7 @@ STMT_OMP_TASKLOOP_DIRECTIVE, STMT_OMP_TASKLOOP_SIMD_DIRECTIVE, STMT_OMP_DISTRIBUTE_DIRECTIVE, + STMT_OMP_TARGET_UPDATE_DIRECTIVE, EXPR_OMP_ARRAY_SECTION, // ARC Index: lib/AST/OpenMPClause.cpp =================================================================== --- lib/AST/OpenMPClause.cpp +++ lib/AST/OpenMPClause.cpp @@ -85,6 +85,8 @@ case OMPC_num_tasks: case OMPC_hint: case OMPC_defaultmap: + case OMPC_to: + case OMPC_from: case OMPC_unknown: break; } @@ -145,6 +147,8 @@ case OMPC_num_tasks: case OMPC_hint: case OMPC_defaultmap: + case OMPC_to: + case OMPC_from: case OMPC_unknown: break; } @@ -550,3 +554,34 @@ 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(totalSizeToAlloc(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(totalSizeToAlloc(N)); + return new (Mem) OMPToClause(N); +} + +OMPFromClause * +OMPFromClause::Create(const ASTContext &C, SourceLocation StartLoc, + SourceLocation LParenLoc, SourceLocation EndLoc, + ArrayRef VL) { + void *Mem = C.Allocate(totalSizeToAlloc(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(totalSizeToAlloc(N)); + return new (Mem) OMPFromClause(N); +} Index: lib/AST/StmtOpenMP.cpp =================================================================== --- lib/AST/StmtOpenMP.cpp +++ lib/AST/StmtOpenMP.cpp @@ -995,3 +995,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::alignTo(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::alignTo(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 @@ -912,6 +912,22 @@ } } +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 << ")"; + } +} + void OMPClausePrinter::VisitOMPDistScheduleClause(OMPDistScheduleClause *Node) { OS << "dist_schedule(" << getOpenMPSimpleClauseTypeName( OMPC_dist_schedule, Node->getDistScheduleKind()); @@ -1131,6 +1147,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 @@ -491,6 +491,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 @@ -656,6 +662,11 @@ void OMPClauseProfiler::VisitOMPDefaultmapClause(const OMPDefaultmapClause *) {} +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 @@ -158,6 +158,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"); @@ -292,6 +294,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"); @@ -475,6 +479,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) \ @@ -596,7 +610,7 @@ bool clang::isOpenMPTargetDataManagementDirective(OpenMPDirectiveKind DKind) { // TODO add target update directive check. return DKind == OMPD_target_data || DKind == OMPD_target_enter_data || - DKind == OMPD_target_exit_data; + DKind == OMPD_target_exit_data || DKind == OMPD_target_update; } bool clang::isOpenMPTeamsDirective(OpenMPDirectiveKind DKind) { Index: lib/CodeGen/CGStmt.cpp =================================================================== --- lib/CodeGen/CGStmt.cpp +++ lib/CodeGen/CGStmt.cpp @@ -277,6 +277,9 @@ case Stmt::OMPDistributeDirectiveClass: EmitOMPDistributeDirective(cast(*S)); break; + case Stmt::OMPTargetUpdateDirectiveClass: + EmitOMPTargetUpdateDirective(cast(*S)); + break; } } Index: lib/CodeGen/CGStmtOpenMP.cpp =================================================================== --- lib/CodeGen/CGStmtOpenMP.cpp +++ lib/CodeGen/CGStmtOpenMP.cpp @@ -2824,6 +2824,8 @@ case OMPC_hint: case OMPC_dist_schedule: case OMPC_defaultmap: + case OMPC_to: + case OMPC_from: llvm_unreachable("Clause is not allowed in 'omp atomic'."); } } @@ -3038,3 +3040,9 @@ [&CS](CodeGenFunction &CGF) { CGF.EmitStmt(CS->getCapturedStmt()); }); } +// Generate the instructions for '#pragma omp target update' directive. +void CodeGenFunction::EmitOMPTargetUpdateDirective( + const OMPTargetUpdateDirective &S) { + // TODO: codegen for target update +} + Index: lib/CodeGen/CodeGenFunction.h =================================================================== --- lib/CodeGen/CodeGenFunction.h +++ lib/CodeGen/CodeGenFunction.h @@ -2354,6 +2354,7 @@ void EmitOMPTargetDataDirective(const OMPTargetDataDirective &S); void EmitOMPTargetEnterDataDirective(const OMPTargetEnterDataDirective &S); void EmitOMPTargetExitDataDirective(const OMPTargetExitDataDirective &S); + void EmitOMPTargetUpdateDirective(const OMPTargetUpdateDirective &S); void EmitOMPTargetParallelDirective(const OMPTargetParallelDirective &S); void EmitOMPTargetParallelForDirective(const OMPTargetParallelForDirective &S); Index: lib/Parse/ParseOpenMP.cpp =================================================================== --- lib/Parse/ParseOpenMP.cpp +++ lib/Parse/ParseOpenMP.cpp @@ -36,7 +36,8 @@ OMPD_point, OMPD_reduction, OMPD_target_enter, - OMPD_target_exit + OMPD_target_exit, + OMPD_update }; } // namespace @@ -55,6 +56,7 @@ .Case("exit", OMPD_exit) .Case("point", OMPD_point) .Case("reduction", OMPD_reduction) + .Case("update", OMPD_update) .Default(OMPD_unknown); } @@ -70,6 +72,7 @@ { OMPD_target, OMPD_exit, OMPD_target_exit }, { OMPD_target_enter, OMPD_data, OMPD_target_enter_data }, { OMPD_target_exit, OMPD_data, OMPD_target_exit_data }, + { OMPD_target, OMPD_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 }, @@ -398,6 +401,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; @@ -426,9 +430,10 @@ /// 'parallel for' | 'parallel sections' | 'task' | 'taskyield' | /// 'barrier' | 'taskwait' | 'flush' | 'ordered' | 'atomic' | /// 'for simd' | 'parallel for simd' | 'target' | 'target data' | -/// 'taskgroup' | 'teams' | 'taskloop' | 'taskloop simd' | +/// 'taskgroup' | 'teams' | 'taskloop' | 'taskloop simd' | /// 'distribute' | 'target enter data' | 'target exit data' | -/// 'target parallel' | 'target parallel for' {clause} +/// 'target parallel' | 'target parallel for' | +/// 'target update' {clause} /// annot_pragma_openmp_end /// StmtResult Parser::ParseOpenMPDeclarativeOrExecutableDirective( @@ -501,6 +506,7 @@ case OMPD_cancel: case OMPD_target_enter_data: case OMPD_target_exit_data: + case OMPD_target_update: if (Allowed == ACK_StatementsOpenMPNonStandalone) { Diag(Tok, diag::err_omp_immediate_directive) << getOpenMPDirectiveName(DKind) << 0; @@ -706,7 +712,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) { @@ -830,6 +837,8 @@ case OMPC_flush: case OMPC_depend: case OMPC_map: + case OMPC_to: + case OMPC_from: Clause = ParseOpenMPVarListClause(DKind, CKind); break; case OMPC_unknown: @@ -1147,8 +1156,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 ')' @@ -1173,6 +1181,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 @@ -1700,6 +1700,7 @@ case OMPD_target_enter_data: case OMPD_target_exit_data: case OMPD_declare_reduction: + case OMPD_target_update: llvm_unreachable("OpenMP Directive is not allowed"); case OMPD_unknown: llvm_unreachable("Unknown OpenMP directive"); @@ -3140,6 +3141,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: case OMPD_declare_reduction: llvm_unreachable("OpenMP Directive is not allowed"); @@ -6059,6 +6066,21 @@ return OMPTargetExitDataDirective::Create(Context, StartLoc, EndLoc, Clauses); } +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(); + } + return OMPTargetUpdateDirective::Create(Context, StartLoc, EndLoc, Clauses); +} + StmtResult Sema::ActOnOpenMPTeamsDirective(ArrayRef Clauses, Stmt *AStmt, SourceLocation StartLoc, SourceLocation EndLoc) { @@ -6312,6 +6334,8 @@ case OMPC_nogroup: case OMPC_dist_schedule: case OMPC_defaultmap: + case OMPC_to: + case OMPC_from: case OMPC_unknown: llvm_unreachable("Clause is not allowed."); } @@ -6597,6 +6621,8 @@ case OMPC_hint: case OMPC_dist_schedule: case OMPC_defaultmap: + case OMPC_to: + case OMPC_from: case OMPC_unknown: llvm_unreachable("Clause is not allowed."); } @@ -6747,6 +6773,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."); } @@ -6934,6 +6962,8 @@ case OMPC_hint: case OMPC_dist_schedule: case OMPC_defaultmap: + case OMPC_to: + case OMPC_from: case OMPC_unknown: llvm_unreachable("Clause is not allowed."); } @@ -7048,6 +7078,12 @@ 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: @@ -9195,7 +9231,8 @@ // Return the expression of the base of the map clause or null if it cannot // be determined and do all the necessary checks to see if the expression is // valid as a standalone map clause expression. -static Expr *CheckMapClauseExpressionBase(Sema &SemaRef, Expr *E) { +static Expr *CheckMapClauseExpressionBase(Sema &SemaRef, Expr *E, + OpenMPClauseKind CKind) { SourceLocation ELoc = E->getExprLoc(); SourceRange ERange = E->getSourceRange(); @@ -9277,8 +9314,8 @@ // A bit-field cannot appear in a map clause. // if (FD->isBitField()) { - SemaRef.Diag(ELoc, diag::err_omp_bit_fields_forbidden_in_map_clause) - << CurE->getSourceRange(); + SemaRef.Diag(ELoc, diag::err_omp_bit_fields_forbidden_in_clause) << + CurE->getSourceRange() << getOpenMPClauseName(CKind); break; } @@ -9386,7 +9423,8 @@ // Return true if expression E associated with value VD has conflicts with other // map information. static bool CheckMapConflicts(Sema &SemaRef, DSAStackTy *DSAS, ValueDecl *VD, - Expr *E, bool CurrentRegionOnly) { + Expr *E, bool CurrentRegionOnly, + OpenMPClauseKind CKind) { assert(VD && E); // Types used to organize the components of a valid map clause. @@ -9506,7 +9544,14 @@ // other, it means they are sharing storage. if (CI == CE && SI == SE) { if (CurrentRegionOnly) { - SemaRef.Diag(ELoc, diag::err_omp_map_shared_storage) << ERange; + if (CKind == OMPC_map) + SemaRef.Diag(ELoc, diag::err_omp_map_shared_storage) << ERange; + else { + assert(CKind == OMPC_to || CKind == OMPC_from); + SemaRef.Diag(ELoc, + diag::err_omp_once_referenced_in_target_update) + << ERange; + } SemaRef.Diag(RE->getExprLoc(), diag::note_used_here) << RE->getSourceRange(); return true; @@ -9560,9 +9605,16 @@ // // An expression is a subset of the other. if (CurrentRegionOnly && (CI == CE || SI == SE)) { - SemaRef.Diag(ELoc, diag::err_omp_map_shared_storage) << ERange; + if (CKind == OMPC_map) + SemaRef.Diag(ELoc, diag::err_omp_map_shared_storage) << ERange; + else { + assert(CKind == OMPC_to || CKind == OMPC_from); + SemaRef.Diag(ELoc, + diag::err_omp_once_referenced_in_target_update) + << ERange; + } SemaRef.Diag(RE->getExprLoc(), diag::note_used_here) - << RE->getSourceRange(); + << RE->getSourceRange(); return true; } @@ -9605,16 +9657,15 @@ return FoundError; } -OMPClause * -Sema::ActOnOpenMPMapClause(OpenMPMapClauseKind MapTypeModifier, - OpenMPMapClauseKind MapType, bool IsMapTypeImplicit, - SourceLocation MapLoc, SourceLocation ColonLoc, - ArrayRef VarList, SourceLocation StartLoc, - SourceLocation LParenLoc, SourceLocation EndLoc) { - SmallVector Vars; - +static void checkOpenMPToFromMapVars(Sema &SemaRef, DSAStackTy *DSAS, + ArrayRef VarList, + SmallVector &Vars, + SourceLocation StartLoc, + OpenMPClauseKind CKind, + OpenMPMapClauseKind MapType, + bool IsMapTypeImplicit) { for (auto &RE : VarList) { - assert(RE && "Null expr in omp map"); + assert(RE && "Null expr in omp to/from/map clause"); if (isa(RE)) { // It will be analyzed later. Vars.push_back(RE); @@ -9636,18 +9687,19 @@ auto *SimpleExpr = RE->IgnoreParenCasts(); if (!RE->IgnoreParenImpCasts()->isLValue()) { - Diag(ELoc, diag::err_omp_expected_named_var_member_or_array_expression) + SemaRef.Diag(ELoc, + diag::err_omp_expected_named_var_member_or_array_expression) << RE->getSourceRange(); continue; } // Obtain the array or member expression bases if required. - auto *BE = CheckMapClauseExpressionBase(*this, SimpleExpr); + auto *BE = CheckMapClauseExpressionBase(SemaRef, SimpleExpr, CKind); if (!BE) continue; // If the base is a reference to a variable, we rely on that variable for - // the following checks. If it is a 'this' expression we rely on the field. + // the following checks. if it is a 'this' expression we rely on the field. ValueDecl *D = nullptr; if (auto *DRE = dyn_cast(BE)) { D = DRE->getDecl(); @@ -9656,7 +9708,7 @@ assert(isa(ME->getBase()) && "Unexpected expression!"); D = ME->getMemberDecl(); } - assert(D && "Null decl on map clause."); + assert(D && "Null decl on to/from/map clause."); auto *VD = dyn_cast(D); auto *FD = dyn_cast(D); @@ -9665,15 +9717,17 @@ (void)FD; // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.10] - // threadprivate variables cannot appear in a map clause. - if (VD && DSAStack->isThreadPrivate(VD)) { - auto DVar = DSAStack->getTopDSA(VD, false); - Diag(ELoc, diag::err_omp_threadprivate_in_map); - ReportOriginalDSA(*this, DSAStack, VD, DVar); + // threadprivate variables cannot appear in a map clause. + // OpenMP 4.5 [2.10.5, target update Construct] + // threadprivate variables cannot appear in a from clause. + if (VD && DSAS->isThreadPrivate(VD)) { + auto DVar = DSAS->getTopDSA(VD, false); + SemaRef.Diag(ELoc, diag::err_omp_threadprivate_in_clause) + << getOpenMPClauseName(CKind); + ReportOriginalDSA(SemaRef, DSAS, VD, DVar); continue; } - // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.9] // A list item cannot appear in both a map clause and a data-sharing // attribute clause on the same construct. // @@ -9681,61 +9735,78 @@ // missing implementation of the other data sharing clauses in target // directives. - // Check conflicts with other map clause expressions. We check the conflicts - // with the current construct separately from the enclosing data - // environment, because the restrictions are different. - if (CheckMapConflicts(*this, DSAStack, D, SimpleExpr, - /*CurrentRegionOnly=*/true)) + // Check conflicts with other map, to or from clause expressions. We check + // the conflicts with the current construct separately from the enclosing + // data environment, because the restrictions are different. + if (CheckMapConflicts(SemaRef, DSAS, D, SimpleExpr, + /*CurrentRegionOnly=*/true, CKind)) break; - if (CheckMapConflicts(*this, DSAStack, D, SimpleExpr, - /*CurrentRegionOnly=*/false)) + if (CKind == OMPC_map && + CheckMapConflicts(SemaRef, DSAS, D, SimpleExpr, + /*CurrentRegionOnly=*/false, CKind)) break; + // OpenMP 4.5 [2.10.5, target update Construct] // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, C++, p.1] - // If the type of a list item is a reference to a type T then the type will - // be considered to be T for all purposes of this clause. - QualType Type = D->getType(); + // If the type of a list item is a reference to a type T then the type will + // be considered to be T for all purpose of this clause. + auto Type = D->getType(); if (Type->isReferenceType()) Type = Type->getPointeeType(); + // OpenMP 4.5 [2.10.5, target update Construct, Restrictions, p.4] + // A list item in a to or from clause must have a mappable type. // OpenMP 4.5 [2.15.5.1, map Clause, Restrictions, p.9] - // A list item must have a mappable type. - if (!CheckTypeMappable(VE->getExprLoc(), VE->getSourceRange(), *this, - DSAStack, Type)) + // A list item must have a mappable type. + if (!CheckTypeMappable(VE->getExprLoc(), VE->getSourceRange(), SemaRef, + DSAS, Type)) continue; - // target enter data - // OpenMP [2.10.2, Restrictions, p. 99] - // A map-type must be specified in all map clauses and must be either - // to or alloc. - OpenMPDirectiveKind DKind = DSAStack->getCurrentDirective(); - if (DKind == OMPD_target_enter_data && - !(MapType == OMPC_MAP_to || MapType == OMPC_MAP_alloc)) { - Diag(StartLoc, diag::err_omp_invalid_map_type_for_directive) - << (IsMapTypeImplicit ? 1 : 0) - << getOpenMPSimpleClauseTypeName(OMPC_map, MapType) - << getOpenMPDirectiveName(DKind); - continue; - } + if (CKind == OMPC_map) { + // target enter data + // OpenMP [2.10.2, Restrictions, p. 99] + // A map-type must be specified in all map clauses and must be either + // to or alloc. + OpenMPDirectiveKind DKind = DSAS->getCurrentDirective(); + if (DKind == OMPD_target_enter_data && + !(MapType == OMPC_MAP_to || MapType == OMPC_MAP_alloc)) { + SemaRef.Diag(StartLoc, diag::err_omp_invalid_map_type_for_directive) + << (IsMapTypeImplicit ? 1 : 0) + << getOpenMPSimpleClauseTypeName(OMPC_map, MapType) + << getOpenMPDirectiveName(DKind); + continue; + } - // target exit_data - // OpenMP [2.10.3, Restrictions, p. 102] - // A map-type must be specified in all map clauses and must be either - // from, release, or delete. - DKind = DSAStack->getCurrentDirective(); - if (DKind == OMPD_target_exit_data && - !(MapType == OMPC_MAP_from || MapType == OMPC_MAP_release || - MapType == OMPC_MAP_delete)) { - Diag(StartLoc, diag::err_omp_invalid_map_type_for_directive) - << (IsMapTypeImplicit ? 1 : 0) - << getOpenMPSimpleClauseTypeName(OMPC_map, MapType) - << getOpenMPDirectiveName(DKind); - continue; + // target exit_data + // OpenMP [2.10.3, Restrictions, p. 102] + // A map-type must be specified in all map clauses and must be either + // from, release, or delete. + DKind = DSAS->getCurrentDirective(); + if (DKind == OMPD_target_exit_data && + !(MapType == OMPC_MAP_from || MapType == OMPC_MAP_release || + MapType == OMPC_MAP_delete)) { + SemaRef.Diag(StartLoc, diag::err_omp_invalid_map_type_for_directive) + << (IsMapTypeImplicit ? 1 : 0) + << getOpenMPSimpleClauseTypeName(OMPC_map, MapType) + << getOpenMPDirectiveName(DKind); + continue; + } } Vars.push_back(RE); - DSAStack->addExprToVarMapInfo(D, RE); + DSAS->addExprToVarMapInfo(D, RE); } +} + +OMPClause * +Sema::ActOnOpenMPMapClause(OpenMPMapClauseKind MapTypeModifier, + OpenMPMapClauseKind MapType, bool IsMapTypeImplicit, + SourceLocation MapLoc, SourceLocation ColonLoc, + ArrayRef VarList, SourceLocation StartLoc, + SourceLocation LParenLoc, SourceLocation EndLoc) { + SmallVector Vars; + checkOpenMPToFromMapVars(*this, DSAStack, VarList, Vars, StartLoc, OMPC_map, + MapType, IsMapTypeImplicit); // We need to produce a map clause even if we don't have variables so that // other diagnostics related with non-existing map clauses are accurate. @@ -10146,3 +10217,29 @@ return new (Context) OMPDefaultmapClause(StartLoc, LParenLoc, MLoc, KindLoc, EndLoc, Kind, M); } + +OMPClause *Sema::ActOnOpenMPToClause(ArrayRef VarList, + SourceLocation StartLoc, + SourceLocation LParenLoc, + SourceLocation EndLoc) { + SmallVector Vars; + checkOpenMPToFromMapVars(*this, DSAStack, VarList, Vars, StartLoc, OMPC_to, + OMPC_MAP_unknown, /*IsMapTypeImplicit=*/false); + 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; + checkOpenMPToFromMapVars(*this, DSAStack, VarList, Vars, StartLoc, OMPC_from, + OMPC_MAP_unknown, /*IsMapTypeImplicit=*/false); + 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 @@ -1750,6 +1750,29 @@ Kind, ChunkSize, StartLoc, LParenLoc, KindLoc, CommaLoc, 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. @@ -7468,6 +7491,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; @@ -7996,6 +8030,34 @@ return C; } +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 @@ -1885,6 +1885,12 @@ case OMPC_defaultmap: C = new (Context) OMPDefaultmapClause(); 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)); @@ -2280,6 +2286,28 @@ C->setDefaultmapKindLoc(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. //===----------------------------------------------------------------------===// @@ -2545,6 +2573,12 @@ VisitOMPLoopDirective(D); } +void ASTStmtReader::VisitOMPTargetUpdateDirective(OMPTargetUpdateDirective *D) { + VisitStmt(D); + ++Idx; + VisitOMPExecutableDirective(D); +} + //===----------------------------------------------------------------------===// // ASTReader Implementation //===----------------------------------------------------------------------===// @@ -3177,6 +3211,11 @@ 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 @@ -2069,6 +2069,20 @@ Writer->Writer.AddSourceLocation(C->getDefaultmapKindLoc(), 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. //===----------------------------------------------------------------------===// @@ -2346,6 +2360,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 @@ -835,6 +835,7 @@ case Stmt::OMPTargetExitDataDirectiveClass: case Stmt::OMPTargetParallelDirectiveClass: case Stmt::OMPTargetParallelForDirectiveClass: + case Stmt::OMPTargetUpdateDirectiveClass: case Stmt::OMPTeamsDirectiveClass: case Stmt::OMPCancellationPointDirectiveClass: case Stmt::OMPCancelDirectiveClass: Index: test/OpenMP/nesting_of_regions.cpp =================================================================== --- test/OpenMP/nesting_of_regions.cpp +++ test/OpenMP/nesting_of_regions.cpp @@ -133,6 +133,10 @@ for (int i = 0; i < 10; ++i) ; } +#pragma omp parallel + { +#pragma omp target update to(a) + } // SIMD DIRECTIVE #pragma omp simd @@ -293,6 +297,10 @@ for (int j = 0; j < 10; ++j) ; } +#pragma omp simd + for (int i = 0; i < 10; ++i) { +#pragma omp target update to(a) // expected-error {{OpenMP constructs may not be nested inside a simd region}} + } // FOR DIRECTIVE #pragma omp for @@ -476,6 +484,10 @@ for (int j = 0; j < 10; ++j) ; } +#pragma omp for + for (int i = 0; i < 10; ++i) { +#pragma omp target update to(a) + } // FOR SIMD DIRECTIVE #pragma omp for simd @@ -636,6 +648,11 @@ for (int j = 0; j < 10; ++j) ; } +#pragma omp for simd + for (int i = 0; i < 10; ++i) { +#pragma omp target update to(a) // expected-error {{OpenMP constructs may not be nested inside a simd region}} + bar(); + } // SECTIONS DIRECTIVE #pragma omp sections @@ -824,6 +841,10 @@ for (int i = 0; i < 10; ++i) ; } +#pragma omp sections + { +#pragma omp target update to(a) + } // SECTION DIRECTIVE #pragma omp section // expected-error {{orphaned 'omp section' directives are prohibited, it must be closely nested to a sections region}} @@ -1062,6 +1083,14 @@ for (int i = 0; i < 10; ++i) ; } +#pragma omp sections + { +#pragma omp section + { + bar(); +#pragma omp target update to(a) + } + } // SINGLE DIRECTIVE #pragma omp single @@ -1235,6 +1264,11 @@ for (int i = 0; i < 10; ++i) ; } +#pragma omp single + { +#pragma omp target update to(a) + bar(); + } // MASTER DIRECTIVE #pragma omp master @@ -1408,6 +1442,11 @@ for (int i = 0; i < 10; ++i) ; } +#pragma omp master + { +#pragma omp target update to(a) + bar(); + } // CRITICAL DIRECTIVE #pragma omp critical @@ -1595,6 +1634,11 @@ for (int i = 0; i < 10; ++i) ; } +#pragma omp critical + { +#pragma omp target update to(a) + bar(); + } // PARALLEL FOR DIRECTIVE #pragma omp parallel for @@ -1783,6 +1827,10 @@ for (int j = 0; j < 10; ++j) ; } +#pragma omp parallel for + for (int i = 0; i < 10; ++i) { +#pragma omp target update from(a) + } // PARALLEL FOR SIMD DIRECTIVE #pragma omp parallel for simd @@ -1971,6 +2019,11 @@ for (int j = 0; j < 10; ++j) ; } +#pragma omp parallel for simd + for (int i = 0; i < 10; ++i) { +#pragma omp target update to(a) // expected-error {{OpenMP constructs may not be nested inside a simd region}} + bar(); + } // PARALLEL SECTIONS DIRECTIVE #pragma omp parallel sections @@ -2148,6 +2201,10 @@ for (int i = 0; i < 10; ++i) ; } +#pragma omp parallel sections + { +#pragma omp target update to(a) + } // TASK DIRECTIVE #pragma omp task @@ -2271,6 +2328,11 @@ for (int i = 0; i < 10; ++i) ; } +#pragma omp task + { +#pragma omp target update to(a) + bar(); + } // ORDERED DIRECTIVE #pragma omp ordered @@ -2464,6 +2526,12 @@ for (int i = 0; i < 10; ++i) ; } +#pragma omp ordered + { + bar(); +#pragma omp target update to(a) + bar(); + } // ATOMIC DIRECTIVE #pragma omp atomic @@ -2678,6 +2746,13 @@ for (int i = 0; i < 10; ++i) ; } +#pragma omp atomic + // expected-error@+2 {{the statement for 'atomic' must be an expression statement of form '++x;', '--x;', 'x++;', 'x--;', 'x binop= expr;', 'x = x binop expr' or 'x = expr binop x', where x is an l-value expression with scalar type}} + // expected-note@+1 {{expected an expression statement}} + { +#pragma omp target update to(a) // expected-error {{OpenMP constructs may not be nested inside an atomic region}} + bar(); + } // TARGET DIRECTIVE #pragma omp target @@ -2812,6 +2887,10 @@ { #pragma omp target exit data map(from: a) // expected-error {{region cannot be nested inside 'target' region}} } +#pragma omp target + { +#pragma omp target update to(a) // expected-error {{region cannot be nested inside 'target' region}} + } // TARGET PARALLEL DIRECTIVE #pragma omp target parallel @@ -2946,6 +3025,10 @@ { #pragma omp target exit data map(from: a) // expected-error {{region cannot be nested inside 'target parallel' region}} } +#pragma omp target parallel + { +#pragma omp target update from(a) // expected-error {{region cannot be nested inside 'target parallel' region}} + } // TARGET PARALLEL FOR DIRECTIVE #pragma omp target parallel for @@ -3134,6 +3217,10 @@ for (int j = 0; j < 10; ++j) ; } +#pragma omp target parallel for + for (int i = 0; i < 10; ++i) { +#pragma omp target update from(a) // expected-error {{region cannot be nested inside 'target parallel for' region}} + } // TEAMS DIRECTIVE #pragma omp target @@ -3297,6 +3384,11 @@ #pragma omp distribute for (int j = 0; j < 10; ++j) ; +#pragma omp target +#pragma omp teams + { +#pragma omp target update to(a) // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp target update' directive into a parallel region?}} + } // TASKLOOP DIRECTIVE #pragma omp taskloop @@ -3469,6 +3561,13 @@ for (int i = 0; i < 10; ++i) ++a; } +#pragma omp taskloop + for (int i = 0; i < 10; ++i) { +#pragma omp target update to(a) + bar(); + } + + // DISTRIBUTE DIRECTIVE #pragma omp target #pragma omp teams @@ -3686,6 +3785,13 @@ #pragma omp teams // expected-error {{region cannot be closely nested inside 'distribute' region; perhaps you forget to enclose 'omp teams' directive into a target region?}} ++a; } +#pragma omp target +#pragma omp teams +#pragma omp distribute + for (int i = 0; i < 10; ++i) { +#pragma omp target update to(a) // expected-error {{region cannot be nested inside 'target' region}} + ++a; + } } void foo() { @@ -3816,6 +3922,11 @@ for (int i = 0; i < 10; ++i) ; } +#pragma omp parallel + { +#pragma omp target update to(a) + a++; + } // SIMD DIRECTIVE #pragma omp simd @@ -3969,6 +4080,11 @@ for (int j = 0; j < 10; ++j) ; } +#pragma omp simd + for (int i = 0; i < 10; ++i) { +#pragma omp target update from(a) // expected-error {{OpenMP constructs may not be nested inside a simd region}} + a++; + } // FOR DIRECTIVE #pragma omp for @@ -4142,6 +4258,11 @@ for (int j = 0; j < 10; ++j) ; } +#pragma omp for + for (int i = 0; i < 10; ++i) { +#pragma omp target update from(a) + ++a; + } // FOR SIMD DIRECTIVE #pragma omp for simd @@ -4295,6 +4416,11 @@ for (int j = 0; j < 10; ++j) ; } +#pragma omp for simd + for (int i = 0; i < 10; ++i) { +#pragma omp target update from(a) // expected-error {{OpenMP constructs may not be nested inside a simd region}} + ++a; + } // SECTIONS DIRECTIVE #pragma omp sections @@ -4458,6 +4584,10 @@ for (int i = 0; i < 10; ++i) ; } +#pragma omp sections + { +#pragma omp target update from(a) + } // SECTION DIRECTIVE #pragma omp section // expected-error {{orphaned 'omp section' directives are prohibited, it must be closely nested to a sections region}} @@ -4706,6 +4836,14 @@ for (int i = 0; i < 10; ++i) ; } +#pragma omp sections + { +#pragma omp section + { +#pragma omp target update from(a) + a++; + } + } // SINGLE DIRECTIVE #pragma omp single @@ -4869,6 +5007,11 @@ for (int i = 0; i < 10; ++i) ; } +#pragma omp single + { +#pragma omp target update from(a) + a++; + } // MASTER DIRECTIVE #pragma omp master @@ -5042,6 +5185,11 @@ for (int i = 0; i < 10; ++i) ; } +#pragma omp master + { +#pragma omp target update from(a) + ++a; + } // CRITICAL DIRECTIVE #pragma omp critical @@ -5234,6 +5382,11 @@ for (int i = 0; i < 10; ++i) ; } +#pragma omp critical + { +#pragma omp target update from(a) + a++; + } // PARALLEL FOR DIRECTIVE #pragma omp parallel for @@ -5422,6 +5575,11 @@ for (int j = 0; j < 10; ++j) ; } +#pragma omp parallel for + for (int i = 0; i < 10; ++i) { +#pragma omp target update from(a) + a++; + } // PARALLEL FOR SIMD DIRECTIVE #pragma omp parallel for simd @@ -5610,6 +5768,11 @@ for (int j = 0; j < 10; ++j) ; } +#pragma omp parallel for simd + for (int i = 0; i < 10; ++i) { +#pragma omp target update from(a) // expected-error {{OpenMP constructs may not be nested inside a simd region}} + a++; + } // PARALLEL SECTIONS DIRECTIVE #pragma omp parallel sections @@ -5783,6 +5946,10 @@ for (int i = 0; i < 10; ++i) ; } +#pragma omp parallel sections + { +#pragma omp target update from(a) + } // TASK DIRECTIVE #pragma omp task @@ -5905,6 +6072,11 @@ for (int i = 0; i < 10; ++i) ; } +#pragma omp task + { +#pragma omp target update from(a) + a++; + } // ATOMIC DIRECTIVE #pragma omp atomic @@ -6119,6 +6291,12 @@ for (int i = 0; i < 10; ++i) ; } +#pragma omp atomic + // expected-error@+2 {{the statement for 'atomic' must be an expression statement of form '++x;', '--x;', 'x++;', 'x--;', 'x binop= expr;', 'x = x binop expr' or 'x = expr binop x', where x is an l-value expression with scalar type}} + // expected-note@+1 {{expected an expression statement}} + { +#pragma omp target update // expected-error {{OpenMP constructs may not be nested inside an atomic region}} + } // TARGET DIRECTIVE #pragma omp target @@ -6253,6 +6431,13 @@ for (int i = 0; i < 10; ++i) ; } +#pragma omp atomic + // expected-error@+2 {{the statement for 'atomic' must be an expression statement of form '++x;', '--x;', 'x++;', 'x--;', 'x binop= expr;', 'x = x binop expr' or 'x = expr binop x', where x is an l-value expression with scalar type}} + // expected-note@+1 {{expected an expression statement}} + { +#pragma omp target update from(a) // expected-error {{OpenMP constructs may not be nested inside an atomic region}} + a++; + } // TARGET PARALLEL DIRECTIVE #pragma omp target parallel @@ -6387,6 +6572,11 @@ { #pragma omp target exit data map(from: a) // expected-error {{region cannot be nested inside 'target parallel' region}} } +#pragma omp target parallel + { +#pragma omp target update from(a) // expected-error {{region cannot be nested inside 'target parallel' region}} + } + // TARGET PARALLEL FOR DIRECTIVE #pragma omp target parallel for @@ -6575,6 +6765,11 @@ for (int j = 0; j < 10; ++j) ; } +#pragma omp target parallel for + for (int i = 0; i < 10; ++i) { +#pragma omp target update from(a) // expected-error {{region cannot be nested inside 'target parallel for' region}} + a++; + } // TEAMS DIRECTIVE #pragma omp target @@ -6736,6 +6931,12 @@ #pragma omp distribute for (int j = 0; j < 10; ++j) ; +#pragma omp target +#pragma omp teams + { +#pragma omp target update from(a) // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp target update' directive into a parallel region?}} + ++a; + } // TASKLOOP DIRECTIVE #pragma omp taskloop @@ -6908,6 +7109,11 @@ for (int i = 0; i < 10; ++i) ++a; } +#pragma omp taskloop + for (int i = 0; i < 10; ++i) { +#pragma omp target update from(a) + ++a; + } // DISTRIBUTE DIRECTIVE #pragma omp target @@ -7127,4 +7333,11 @@ #pragma omp target exit data map(from: a) // expected-error {{region cannot be nested inside 'target' region}} ++a; } +#pragma omp target +#pragma omp teams +#pragma omp distribute + for (int i = 0; i < 10; ++i) { +#pragma omp target update from(a) // expected-error {{region cannot be nested inside 'target' region}} + ++a; + } } Index: test/OpenMP/target_map_messages.cpp =================================================================== --- test/OpenMP/target_map_messages.cpp +++ test/OpenMP/target_map_messages.cpp @@ -40,7 +40,7 @@ #pragma omp target map(arg,a,d[:2]) // expected-error {{subscripted value is not an array or pointer}} {} - #pragma omp target map(to:ss) // expected-error {{threadprivate variables are not allowed in map clause}} + #pragma omp target map(to:ss) // expected-error {{threadprivate variables are not allowed in 'map' clause}} {} #pragma omp target map(to:b,e) @@ -239,7 +239,7 @@ {} #pragma omp target map(r.C, t.C) {} - #pragma omp target map(r.A) // expected-error {{bit fields cannot be used to specify storage in a map clause}} + #pragma omp target map(r.A) // expected-error {{bit fields cannot be used to specify storage in a 'map' clause}} {} #pragma omp target map(r.Arr) {} @@ -407,7 +407,7 @@ #pragma omp target data map(S2::S2s) #pragma omp target data map(S2::S2sc) #pragma omp target data map(e, g) -#pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in map clause}} +#pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}} #pragma omp target data 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 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}} foo(); @@ -476,7 +476,7 @@ #pragma omp target data map(S2::S2s) #pragma omp target data map(S2::S2sc) #pragma omp target data map(e, g) -#pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in map clause}} +#pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}} #pragma omp target data 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 {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}} foo(); Index: test/OpenMP/target_parallel_for_map_messages.cpp =================================================================== --- test/OpenMP/target_parallel_for_map_messages.cpp +++ test/OpenMP/target_parallel_for_map_messages.cpp @@ -126,7 +126,7 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for map(e, g) for (i = 0; i < argc; ++i) foo(); -#pragma omp target parallel for map(h) // expected-error {{threadprivate variables are not allowed in map clause}} +#pragma omp target parallel for map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}} for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} for (i = 0; i < argc; ++i) foo(); @@ -230,7 +230,7 @@ for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for map(e, g) for (i = 0; i < argc; ++i) foo(); -#pragma omp target parallel for map(h) // expected-error {{threadprivate variables are not allowed in map clause}} +#pragma omp target parallel for map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}} for (i = 0; i < argc; ++i) foo(); #pragma omp target parallel for map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}} for (i = 0; i < argc; ++i) foo(); Index: test/OpenMP/target_parallel_map_messages.cpp =================================================================== --- test/OpenMP/target_parallel_map_messages.cpp +++ test/OpenMP/target_parallel_map_messages.cpp @@ -126,7 +126,7 @@ foo(); #pragma omp target parallel map(e, g) foo(); -#pragma omp target parallel map(h) // expected-error {{threadprivate variables are not allowed in map clause}} +#pragma omp target parallel map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}} foo(); #pragma omp target parallel map(k), map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}} foo(); @@ -229,7 +229,7 @@ foo(); #pragma omp target parallel map(e, g) foo(); -#pragma omp target parallel map(h) // expected-error {{threadprivate variables are not allowed in map clause}} +#pragma omp target parallel map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}} foo(); #pragma omp target parallel map(k), map(k) // 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,176 @@ +// 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) { } +}; +struct S6 { + int ii; + int aa[30]; + float xx; + double *pp; +}; +struct S7 { + int i; + int a[50]; + float x; + S6 s6[5]; + double *p; + unsigned bfa : 4; +}; + +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; + T *m; + S7 s7; + +#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 expression containing only member accesses and/or array sections based on named variables}} +#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 expression containing only member accesses and/or array sections based on named variables}} +#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(t), from(t[: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]) + +#pragma omp target update from(x, a[:2]) // expected-error {{subscripted value is not an array or pointer}} +#pragma omp target update from(x, c[:]) // expected-error {{subscripted value is not an array or pointer}} +#pragma omp target update from(x, (m+1)[2]) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target update from(s7.i, s7.a[:3]) +#pragma omp target update from(s7.s6[1].aa[0:5]) +#pragma omp target update from(x, s7.s6[:5].aa[6]) // expected-error {{OpenMP array section is not allowed here}} +#pragma omp target update from(x, s7.s6[:5].aa[:6]) // expected-error {{OpenMP array section is not allowed here}} +#pragma omp target update from(s7.p[:10]) +#pragma omp target update from(x, s7.bfa) // expected-error {{bit fields cannot be used to specify storage in a 'from' clause}} +#pragma omp target update from(x, s7.p[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} +#pragma omp target data map(to: s7.i) + { +#pragma omp target update from(s7.x) + } + + 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; + int *m; + S7 s7; + +#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 expression containing only member accesses and/or array sections based on named variables}} 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(t), from(t[: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]) + +#pragma omp target update from(x, a[:2]) // expected-error {{subscripted value is not an array or pointer}} +#pragma omp target update from(x, c[:]) // expected-error {{subscripted value is not an array or pointer}} +#pragma omp target update from(x, (m+1)[2]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target update from(s7.i, s7.a[:3]) +#pragma omp target update from(s7.s6[1].aa[0:5]) +#pragma omp target update from(x, s7.s6[:5].aa[6]) // expected-error {{OpenMP array section is not allowed here}} +#pragma omp target update from(x, s7.s6[:5].aa[:6]) // expected-error {{OpenMP array section is not allowed here}} +#pragma omp target update from(s7.p[:10]) +#pragma omp target update from(x, s7.bfa) // expected-error {{bit fields cannot be used to specify storage in a 'from' clause}} +#pragma omp target update from(x, s7.p[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} +#pragma omp target data map(to: s7.i) + { +#pragma omp target update from(s7.x) + } + + 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,58 @@ +// 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 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}} + return 0; +} + +int main(int argc, char **argv) { + int m; +#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}} + return tmain(argc, argv); +} Index: test/OpenMP/target_update_messages.cpp =================================================================== --- /dev/null +++ test/OpenMP/target_update_messages.cpp @@ -0,0 +1,31 @@ +// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 %s + +void foo() { +} + +bool foobool(int argc) { + return argc; +} + +struct S1; // Aexpected-note {{declared here}} + +template // Aexpected-note {{declared here}} +int tmain(T argc, S **argv) { + int n; + return 0; +} + +int main(int argc, char **argv) { + int m; + #pragma omp target update to(m) { // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} + #pragma omp target update to(m) ( // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} + #pragma omp target update to(m) [ // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} + #pragma omp target update to(m) ] // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} + #pragma omp target update to(m) ) // expected-warning {{extra tokens at the end of '#pragma omp target update' are ignored}} + + #pragma omp target update from(m) // OK + { + 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,175 @@ +// 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) { } +}; +struct S6 { + int ii; + int aa[30]; + float xx; + double *pp; +}; +struct S7 { + int i; + int a[50]; + float x; + S6 s6[5]; + double *p; + unsigned bfa : 4; +}; + +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 *m; + T i, t[20]; + T &j = i; + T *k = &j; + T x; + T y; + T to; + const T (&l)[5] = da; + S7 s7; + +#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 expression containing only member accesses and/or array sections based on named variables}} +#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 expression containing only member accesses and/or array sections based on named variables}} +#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(t), to(t[: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]) + +#pragma omp target update to(x, a[:2]) // expected-error {{subscripted value is not an array or pointer}} +#pragma omp target update to(x, c[:]) // expected-error {{subscripted value is not an array or pointer}} +#pragma omp target update to(x, (m+1)[2]) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target update to(s7.i, s7.a[:3]) +#pragma omp target update to(s7.s6[1].aa[0:5]) +#pragma omp target update to(x, s7.s6[:5].aa[6]) // expected-error {{OpenMP array section is not allowed here}} +#pragma omp target update to(x, s7.s6[:5].aa[:6]) // expected-error {{OpenMP array section is not allowed here}} +#pragma omp target update to(s7.p[:10]) +#pragma omp target update to(x, s7.bfa) // expected-error {{bit fields cannot be used to specify storage in a 'to' clause}} +#pragma omp target update to(x, s7.p[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} +#pragma omp target data map(to: s7.i) + { +#pragma omp target update to(s7.x) + } + 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; + S7 s7; + int *m; + +#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 expression containing only member accesses and/or array sections based on named variables}} 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(t), to(t[: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]) + +#pragma omp target update to(x, a[:2]) // expected-error {{subscripted value is not an array or pointer}} +#pragma omp target update to(x, c[:]) // expected-error {{subscripted value is not an array or pointer}} +#pragma omp target update to(x, (m+1)[2]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}} +#pragma omp target update to(s7.i, s7.a[:3]) +#pragma omp target update to(s7.s6[1].aa[0:5]) +#pragma omp target update to(x, s7.s6[:5].aa[6]) // expected-error {{OpenMP array section is not allowed here}} +#pragma omp target update to(x, s7.s6[:5].aa[:6]) // expected-error {{OpenMP array section is not allowed here}} +#pragma omp target update to(s7.p[:10]) +#pragma omp target update to(x, s7.bfa) // expected-error {{bit fields cannot be used to specify storage in a 'to' clause}} +#pragma omp target update to(x, s7.p[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} +#pragma omp target data map(to: s7.i) + { +#pragma omp target update to(s7.x) + } + + 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 @@ -2255,6 +2255,12 @@ } void OMPClauseEnqueue::VisitOMPDefaultmapClause( const OMPDefaultmapClause * /*C*/) {} +void OMPClauseEnqueue::VisitOMPFromClause(const OMPFromClause *C) { + VisitOMPClauseList(C); +} +void OMPClauseEnqueue::VisitOMPToClause(const OMPToClause *C) { + VisitOMPClauseList(C); +} } void EnqueueVisitor::EnqueueChildren(const OMPClause *S) { @@ -4825,6 +4831,8 @@ return cxstring::createRef("OMPTargetParallelDirective"); case CXCursor_OMPTargetParallelForDirective: return cxstring::createRef("OMPTargetParallelForDirective"); + 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 @@ -612,6 +612,9 @@ case Stmt::OMPTargetParallelForDirectiveClass: K = CXCursor_OMPTargetParallelForDirective; break; + case Stmt::OMPTargetUpdateDirectiveClass: + K = CXCursor_OMPTargetUpdateDirective; + break; case Stmt::OMPTeamsDirectiveClass: K = CXCursor_OMPTeamsDirective; break;