Index: include/clang-c/Index.h =================================================================== --- include/clang-c/Index.h +++ include/clang-c/Index.h @@ -2274,7 +2274,11 @@ */ CXCursor_OMPDistributeDirective = 260, - CXCursor_LastStmt = CXCursor_OMPDistributeDirective, + /** \brief OpenMP target enter data directive. + */ + CXCursor_OMPTargetEnterDataDirective = 261, + + CXCursor_LastStmt = CXCursor_OMPTargetEnterDataDirective, /** * \brief Cursor that represents the translation unit itself. Index: include/clang/AST/RecursiveASTVisitor.h =================================================================== --- include/clang/AST/RecursiveASTVisitor.h +++ include/clang/AST/RecursiveASTVisitor.h @@ -2433,6 +2433,9 @@ DEF_TRAVERSE_STMT(OMPTargetDataDirective, { TRY_TO(TraverseOMPExecutableDirective(S)); }) +DEF_TRAVERSE_STMT(OMPTargetEnterDataDirective, + { TRY_TO(TraverseOMPExecutableDirective(S)); }) + DEF_TRAVERSE_STMT(OMPTeamsDirective, { TRY_TO(TraverseOMPExecutableDirective(S)); }) Index: include/clang/AST/StmtOpenMP.h =================================================================== --- include/clang/AST/StmtOpenMP.h +++ include/clang/AST/StmtOpenMP.h @@ -2039,6 +2039,65 @@ } }; +/// \brief This represents '#pragma omp target enter data' directive. +/// +/// \code +/// #pragma omp target enter data device(0) if(a) map(b[:]) +/// \endcode +/// In this example directive '#pragma omp target enter data' has clauses +/// 'device' with the value '0', 'if' with condition 'a' and 'map' with array +/// section 'b[:]'. +/// +class OMPTargetEnterDataDirective : 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. + /// + OMPTargetEnterDataDirective(SourceLocation StartLoc, SourceLocation EndLoc, + unsigned NumClauses) + : OMPExecutableDirective(this, OMPTargetEnterDataDirectiveClass, + OMPD_target_enter_data, StartLoc, EndLoc, + NumClauses, /*NumChildren=*/0) {} + + /// \brief Build an empty directive. + /// + /// \param NumClauses Number of clauses. + /// + explicit OMPTargetEnterDataDirective(unsigned NumClauses) + : OMPExecutableDirective(this, OMPTargetEnterDataDirectiveClass, + OMPD_target_enter_data, SourceLocation(), + SourceLocation(), NumClauses, + /*NumChildren=*/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. + /// \param AssociatedStmt Statement, associated with the directive. + /// + static OMPTargetEnterDataDirective * + Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc, + ArrayRef Clauses); + + /// \brief Creates an empty directive with the place for \a N clauses. + /// + /// \param C AST context. + /// \param N The number of clauses. + /// + static OMPTargetEnterDataDirective *CreateEmpty(const ASTContext &C, + unsigned N, EmptyShell); + + static bool classof(const Stmt *T) { + return T->getStmtClass() == OMPTargetEnterDataDirectiveClass; + } +}; + /// \brief This represents '#pragma omp teams' directive. /// /// \code Index: include/clang/Basic/DiagnosticSemaKinds.td =================================================================== --- include/clang/Basic/DiagnosticSemaKinds.td +++ include/clang/Basic/DiagnosticSemaKinds.td @@ -7930,6 +7930,10 @@ "variable already marked as mapped in current construct">; def err_omp_not_mappable_type : Error< "type %0 is not mappable to target">; +def err_omp_invalid_map_type_for_directive : Error< + "%select{map type '%1' is not allowed|map type must be specified}0 for '#pragma omp %2'">; +def err_omp_no_map_for_directive : Error< + "expected at least one map clause for '#pragma omp %0'">; def note_omp_polymorphic_in_target : Note< "mappable type cannot be polymorphic">; def note_omp_static_member_in_target : Note< Index: include/clang/Basic/OpenMPKinds.def =================================================================== --- include/clang/Basic/OpenMPKinds.def +++ include/clang/Basic/OpenMPKinds.def @@ -60,6 +60,9 @@ #ifndef OPENMP_TARGET_DATA_CLAUSE # define OPENMP_TARGET_DATA_CLAUSE(Name) #endif +#ifndef OPENMP_TARGET_ENTER_DATA_CLAUSE +#define OPENMP_TARGET_ENTER_DATA_CLAUSE(Name) +#endif #ifndef OPENMP_TEAMS_CLAUSE # define OPENMP_TEAMS_CLAUSE(Name) #endif @@ -125,6 +128,7 @@ OPENMP_DIRECTIVE(teams) OPENMP_DIRECTIVE(cancel) OPENMP_DIRECTIVE_EXT(target_data, "target data") +OPENMP_DIRECTIVE_EXT(target_enter_data, "target enter data") OPENMP_DIRECTIVE_EXT(parallel_for, "parallel for") OPENMP_DIRECTIVE_EXT(parallel_for_simd, "parallel for simd") OPENMP_DIRECTIVE_EXT(parallel_sections, "parallel sections") @@ -349,6 +353,12 @@ OPENMP_TARGET_DATA_CLAUSE(device) OPENMP_TARGET_DATA_CLAUSE(map) +// Clauses allowed for OpenMP directive 'target enter data'. +// TODO More clauses for 'target enter data' directive. +OPENMP_TARGET_ENTER_DATA_CLAUSE(if) +OPENMP_TARGET_ENTER_DATA_CLAUSE(device) +OPENMP_TARGET_ENTER_DATA_CLAUSE(map) + // Clauses allowed for OpenMP directive 'teams'. // TODO More clauses for 'teams' directive. OPENMP_TEAMS_CLAUSE(default) @@ -443,6 +453,7 @@ #undef OPENMP_ATOMIC_CLAUSE #undef OPENMP_TARGET_CLAUSE #undef OPENMP_TARGET_DATA_CLAUSE +#undef OPENMP_TARGET_ENTER_DATA_CLAUSE #undef OPENMP_TEAMS_CLAUSE #undef OPENMP_SIMD_CLAUSE #undef OPENMP_FOR_CLAUSE Index: include/clang/Basic/StmtNodes.td =================================================================== --- include/clang/Basic/StmtNodes.td +++ include/clang/Basic/StmtNodes.td @@ -216,6 +216,7 @@ def OMPAtomicDirective : DStmt; def OMPTargetDirective : DStmt; def OMPTargetDataDirective : DStmt; +def OMPTargetEnterDataDirective : 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 @@ -7941,6 +7941,11 @@ StmtResult ActOnOpenMPTargetDataDirective(ArrayRef Clauses, Stmt *AStmt, SourceLocation StartLoc, SourceLocation EndLoc); + /// \brief Called on well-formed '\#pragma omp target enter data' after + /// parsing of the associated statement. + StmtResult ActOnOpenMPTargetEnterDataDirective(ArrayRef Clauses, + SourceLocation StartLoc, + SourceLocation EndLoc); /// \brief Called on well-formed '\#pragma omp teams' after parsing of the /// associated statement. StmtResult ActOnOpenMPTeamsDirective(ArrayRef Clauses, Index: include/clang/Serialization/ASTBitCodes.h =================================================================== --- include/clang/Serialization/ASTBitCodes.h +++ include/clang/Serialization/ASTBitCodes.h @@ -1445,6 +1445,7 @@ STMT_OMP_ATOMIC_DIRECTIVE, STMT_OMP_TARGET_DIRECTIVE, STMT_OMP_TARGET_DATA_DIRECTIVE, + STMT_OMP_TARGET_ENTER_DATA_DIRECTIVE, STMT_OMP_TEAMS_DIRECTIVE, STMT_OMP_TASKGROUP_DIRECTIVE, STMT_OMP_CANCELLATION_POINT_DIRECTIVE, Index: lib/AST/StmtOpenMP.cpp =================================================================== --- lib/AST/StmtOpenMP.cpp +++ lib/AST/StmtOpenMP.cpp @@ -718,6 +718,28 @@ return new (Mem) OMPTargetDataDirective(N); } +OMPTargetEnterDataDirective *OMPTargetEnterDataDirective::Create( + const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc, + ArrayRef Clauses) { + void *Mem = + C.Allocate(llvm::RoundUpToAlignment(sizeof(OMPTargetEnterDataDirective), + llvm::alignOf()) + + sizeof(OMPClause *) * Clauses.size()); + OMPTargetEnterDataDirective *Dir = + new (Mem) OMPTargetEnterDataDirective(StartLoc, EndLoc, Clauses.size()); + Dir->setClauses(Clauses); + return Dir; +} + +OMPTargetEnterDataDirective *OMPTargetEnterDataDirective::CreateEmpty( + const ASTContext &C, unsigned N, EmptyShell) { + void *Mem = + C.Allocate(llvm::RoundUpToAlignment(sizeof(OMPTargetEnterDataDirective), + llvm::alignOf()) + + sizeof(OMPClause *) * N); + return new (Mem) OMPTargetEnterDataDirective(N); +} + OMPTeamsDirective *OMPTeamsDirective::Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc, Index: lib/AST/StmtPrinter.cpp =================================================================== --- lib/AST/StmtPrinter.cpp +++ lib/AST/StmtPrinter.cpp @@ -1051,6 +1051,12 @@ PrintOMPExecutableDirective(Node); } +void StmtPrinter::VisitOMPTargetEnterDataDirective( + OMPTargetEnterDataDirective *Node) { + Indent() << "#pragma omp target enter data "; + PrintOMPExecutableDirective(Node); +} + void StmtPrinter::VisitOMPTeamsDirective(OMPTeamsDirective *Node) { Indent() << "#pragma omp teams "; PrintOMPExecutableDirective(Node); Index: lib/AST/StmtProfile.cpp =================================================================== --- lib/AST/StmtProfile.cpp +++ lib/AST/StmtProfile.cpp @@ -584,6 +584,11 @@ VisitOMPExecutableDirective(S); } +void StmtProfiler::VisitOMPTargetEnterDataDirective( + const OMPTargetEnterDataDirective *S) { + VisitOMPExecutableDirective(S); +} + void StmtProfiler::VisitOMPTeamsDirective(const OMPTeamsDirective *S) { VisitOMPExecutableDirective(S); } Index: lib/Basic/OpenMPKinds.cpp =================================================================== --- lib/Basic/OpenMPKinds.cpp +++ lib/Basic/OpenMPKinds.cpp @@ -398,6 +398,16 @@ break; } break; + case OMPD_target_enter_data: + switch (CKind) { +#define OPENMP_TARGET_ENTER_DATA_CLAUSE(Name) \ + case OMPC_##Name: \ + return true; +#include "clang/Basic/OpenMPKinds.def" + default: + break; + } + break; case OMPD_teams: switch (CKind) { #define OPENMP_TEAMS_CLAUSE(Name) \ Index: lib/CodeGen/CGStmt.cpp =================================================================== --- lib/CodeGen/CGStmt.cpp +++ lib/CodeGen/CGStmt.cpp @@ -256,6 +256,9 @@ case Stmt::OMPTargetDataDirectiveClass: EmitOMPTargetDataDirective(cast(*S)); break; + case Stmt::OMPTargetEnterDataDirectiveClass: + EmitOMPTargetEnterDataDirective(cast(*S)); + break; case Stmt::OMPTaskLoopDirectiveClass: EmitOMPTaskLoopDirective(cast(*S)); break; Index: lib/CodeGen/CGStmtOpenMP.cpp =================================================================== --- lib/CodeGen/CGStmtOpenMP.cpp +++ lib/CodeGen/CGStmtOpenMP.cpp @@ -2660,6 +2660,11 @@ [&CS](CodeGenFunction &CGF) { CGF.EmitStmt(CS->getCapturedStmt()); }); } +void CodeGenFunction::EmitOMPTargetEnterDataDirective( + const OMPTargetEnterDataDirective &S) { + // TODO: codegen for target enter data. +} + void CodeGenFunction::EmitOMPTaskLoopDirective(const OMPTaskLoopDirective &S) { // emit the code inside the construct for now auto CS = cast(S.getAssociatedStmt()); Index: lib/CodeGen/CodeGenFunction.h =================================================================== --- lib/CodeGen/CodeGenFunction.h +++ lib/CodeGen/CodeGenFunction.h @@ -2335,6 +2335,7 @@ void EmitOMPAtomicDirective(const OMPAtomicDirective &S); void EmitOMPTargetDirective(const OMPTargetDirective &S); void EmitOMPTargetDataDirective(const OMPTargetDataDirective &S); + void EmitOMPTargetEnterDataDirective(const OMPTargetEnterDataDirective &S); void EmitOMPTeamsDirective(const OMPTeamsDirective &S); void EmitOMPCancellationPointDirective(const OMPCancellationPointDirective &S); Index: lib/Parse/ParseOpenMP.cpp =================================================================== --- lib/Parse/ParseOpenMP.cpp +++ lib/Parse/ParseOpenMP.cpp @@ -34,6 +34,10 @@ {OMPD_unknown /*cancellation*/, OMPD_unknown /*point*/, OMPD_cancellation_point}, {OMPD_target, OMPD_unknown /*data*/, OMPD_target_data}, + {OMPD_target, OMPD_unknown /*enter/exit*/, + OMPD_unknown /*target enter/exit*/}, + {OMPD_unknown /*target enter*/, OMPD_unknown /*data*/, + OMPD_target_enter_data}, {OMPD_for, OMPD_simd, OMPD_for_simd}, {OMPD_parallel, OMPD_for, OMPD_parallel_for}, {OMPD_parallel_for, OMPD_simd, OMPD_parallel_for_simd}, @@ -49,8 +53,10 @@ for (unsigned i = 0; i < llvm::array_lengthof(F); ++i) { if (!Tok.isAnnotation() && DKind == OMPD_unknown) { TokenMatched = - (i == 0) && - !P.getPreprocessor().getSpelling(Tok).compare("cancellation"); + ((i == 0) && + !P.getPreprocessor().getSpelling(Tok).compare("cancellation")) || + ((i == 3) && + !P.getPreprocessor().getSpelling(Tok).compare("enter")); } else { TokenMatched = DKind == F[i][0] && DKind != OMPD_unknown; } @@ -67,7 +73,10 @@ TokenMatched = ((i == 0) && !P.getPreprocessor().getSpelling(Tok).compare("point")) || - ((i == 1) && !P.getPreprocessor().getSpelling(Tok).compare("data")); + ((i == 1 || i == 3) && + !P.getPreprocessor().getSpelling(Tok).compare("data")) || + ((i == 2) && + !P.getPreprocessor().getSpelling(Tok).compare("enter")); } else { TokenMatched = SDKind == F[i][1] && SDKind != OMPD_unknown; } @@ -138,6 +147,7 @@ case OMPD_cancellation_point: case OMPD_cancel: case OMPD_target_data: + case OMPD_target_enter_data: case OMPD_taskloop: case OMPD_taskloop_simd: case OMPD_distribute: @@ -162,8 +172,7 @@ /// 'barrier' | 'taskwait' | 'flush' | 'ordered' | 'atomic' | /// 'for simd' | 'parallel for simd' | 'target' | 'target data' | /// 'taskgroup' | 'teams' | 'taskloop' | 'taskloop simd' {clause} | -/// 'distribute' -/// annot_pragma_openmp_end +/// 'distribute' | 'target enter data' | annot_pragma_openmp_end /// StmtResult Parser::ParseOpenMPDeclarativeOrExecutableDirective(bool StandAloneAllowed) { @@ -213,6 +222,7 @@ case OMPD_taskwait: case OMPD_cancellation_point: case OMPD_cancel: + case OMPD_target_enter_data: if (!StandAloneAllowed) { Diag(Tok, diag::err_omp_immediate_directive) << getOpenMPDirectiveName(DKind) << 0; Index: lib/Sema/SemaOpenMP.cpp =================================================================== --- lib/Sema/SemaOpenMP.cpp +++ lib/Sema/SemaOpenMP.cpp @@ -1610,6 +1610,7 @@ case OMPD_cancellation_point: case OMPD_cancel: case OMPD_flush: + case OMPD_target_enter_data: llvm_unreachable("OpenMP Directive is not allowed"); case OMPD_unknown: llvm_unreachable("Unknown OpenMP directive"); @@ -1724,6 +1725,8 @@ // | parallel | ordered | + | // | parallel | atomic | * | // | parallel | target | * | + // | parallel | target enter | * | + // | | data | | // | parallel | teams | + | // | parallel | cancellation | | // | | point | ! | @@ -1753,6 +1756,8 @@ // | for | ordered | * (if construct is ordered) | // | for | atomic | * | // | for | target | * | + // | for | target enter | * | + // | | data | | // | for | teams | + | // | for | cancellation | | // | | point | ! | @@ -1782,6 +1787,8 @@ // | master | ordered | + | // | master | atomic | * | // | master | target | * | + // | master | target enter | * | + // | | data | | // | master | teams | + | // | master | cancellation | | // | | point | | @@ -1810,6 +1817,8 @@ // | critical | ordered | + | // | critical | atomic | * | // | critical | target | * | + // | critical | target enter | * | + // | | data | | // | critical | teams | + | // | critical | cancellation | | // | | point | | @@ -1839,6 +1848,8 @@ // | simd | ordered | + (with simd clause) | // | simd | atomic | | // | simd | target | | + // | simd | target enter | | + // | | data | | // | simd | teams | | // | simd | cancellation | | // | | point | | @@ -1868,6 +1879,8 @@ // | for simd | ordered | + (with simd clause) | // | for simd | atomic | | // | for simd | target | | + // | for simd | target enter | | + // | | data | | // | for simd | teams | | // | for simd | cancellation | | // | | point | | @@ -1897,6 +1910,8 @@ // | parallel for simd| ordered | + (with simd clause) | // | parallel for simd| atomic | | // | parallel for simd| target | | + // | parallel for simd| target enter | | + // | | data | | // | parallel for simd| teams | | // | parallel for simd| cancellation | | // | | point | | @@ -1926,6 +1941,8 @@ // | sections | ordered | + | // | sections | atomic | * | // | sections | target | * | + // | sections | target enter | * | + // | | data | | // | sections | teams | + | // | sections | cancellation | | // | | point | ! | @@ -1955,6 +1972,8 @@ // | section | ordered | + | // | section | atomic | * | // | section | target | * | + // | section | target enter | * | + // | | data | | // | section | teams | + | // | section | cancellation | | // | | point | ! | @@ -1984,6 +2003,8 @@ // | single | ordered | + | // | single | atomic | * | // | single | target | * | + // | single | target enter | * | + // | | data | | // | single | teams | + | // | single | cancellation | | // | | point | | @@ -2013,6 +2034,8 @@ // | parallel for | ordered | * (if construct is ordered) | // | parallel for | atomic | * | // | parallel for | target | * | + // | parallel for | target enter | * | + // | | data | | // | parallel for | teams | + | // | parallel for | cancellation | | // | | point | ! | @@ -2042,6 +2065,8 @@ // | parallel sections| ordered | + | // | parallel sections| atomic | * | // | parallel sections| target | * | + // | parallel sections| target enter | * | + // | | data | | // | parallel sections| teams | + | // | parallel sections| cancellation | | // | | point | ! | @@ -2071,6 +2096,8 @@ // | task | ordered | + | // | task | atomic | * | // | task | target | * | + // | task | target enter | * | + // | | data | | // | task | teams | + | // | task | cancellation | | // | | point | ! | @@ -2100,6 +2127,8 @@ // | ordered | ordered | + | // | ordered | atomic | * | // | ordered | target | * | + // | ordered | target enter | * | + // | | data | | // | ordered | teams | + | // | ordered | cancellation | | // | | point | | @@ -2129,6 +2158,8 @@ // | atomic | ordered | | // | atomic | atomic | | // | atomic | target | | + // | atomic | target enter | | + // | | data | | // | atomic | teams | | // | atomic | cancellation | | // | | point | | @@ -2158,6 +2189,8 @@ // | target | ordered | * | // | target | atomic | * | // | target | target | * | + // | target | target enter | * | + // | | data | | // | target | teams | * | // | target | cancellation | | // | | point | | @@ -2187,6 +2220,8 @@ // | teams | ordered | + | // | teams | atomic | + | // | teams | target | + | + // | teams | target enter | + | + // | | data | | // | teams | teams | + | // | teams | cancellation | | // | | point | | @@ -2216,6 +2251,8 @@ // | taskloop | ordered | + | // | taskloop | atomic | * | // | taskloop | target | * | + // | taskloop | target enter | * | + // | | data | | // | taskloop | teams | + | // | taskloop | cancellation | | // | | point | | @@ -2244,6 +2281,8 @@ // | taskloop simd | ordered | + (with simd clause) | // | taskloop simd | atomic | | // | taskloop simd | target | | + // | taskloop simd | target enter | | + // | | data | | // | taskloop simd | teams | | // | taskloop simd | cancellation | | // | | point | | @@ -2273,6 +2312,8 @@ // | distribute | ordered | + | // | distribute | atomic | * | // | distribute | target | | + // | distribute | target enter | | + // | | data | | // | distribute | teams | | // | distribute | cancellation | + | // | | point | | @@ -2697,6 +2738,11 @@ EndLoc); AllowedNameModifiers.push_back(OMPD_target_data); break; + case OMPD_target_enter_data: + Res = ActOnOpenMPTargetEnterDataDirective(ClausesWithImplicit, StartLoc, + EndLoc); + AllowedNameModifiers.push_back(OMPD_target_enter_data); + break; case OMPD_taskloop: Res = ActOnOpenMPTaskLoopDirective(ClausesWithImplicit, AStmt, StartLoc, EndLoc, VarsWithInheritedDSA); @@ -5433,6 +5479,18 @@ return OMPTargetDirective::Create(Context, StartLoc, EndLoc, Clauses, AStmt); } +/// \brief Check for existence of a map clause in the list of clauses. +static bool HasMapClause(ArrayRef Clauses) { + for (ArrayRef::iterator I = Clauses.begin(), E = Clauses.end(); + I != E; ++I) { + if(*I != nullptr && (*I)->getClauseKind() == OMPC_map) { + return true; + } + } + + return false; +} + StmtResult Sema::ActOnOpenMPTargetDataDirective(ArrayRef Clauses, Stmt *AStmt, SourceLocation StartLoc, @@ -5448,6 +5506,20 @@ AStmt); } +StmtResult Sema::ActOnOpenMPTargetEnterDataDirective( + ArrayRef Clauses, SourceLocation StartLoc, + SourceLocation EndLoc) { + // OpenMP [2.10.2, Restrictions, p. 99] + // At least one map clause must appear on the directive. + if (!HasMapClause(Clauses)) { + Diag(StartLoc, diag::err_omp_no_map_for_directive) << + getOpenMPDirectiveName(OMPD_target_enter_data); + return StmtError(); + } + + return OMPTargetEnterDataDirective::Create(Context, StartLoc, EndLoc, Clauses); +} + StmtResult Sema::ActOnOpenMPTeamsDirective(ArrayRef Clauses, Stmt *AStmt, SourceLocation StartLoc, SourceLocation EndLoc) { @@ -8440,6 +8512,22 @@ DSAStack, 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) << + // TODO: Need to determine if map type is implicitly determined + 0 << + getOpenMPSimpleClauseTypeName(OMPC_map, MapType) << + getOpenMPDirectiveName(DKind); + // Proceed to add the variable in a map clause anyway, to prevent + // further spurious messages + } + Vars.push_back(RE); MI.RefExpr = RE; DSAStack->addMapInfoForVar(VD, MI); Index: lib/Sema/TreeTransform.h =================================================================== --- lib/Sema/TreeTransform.h +++ lib/Sema/TreeTransform.h @@ -7378,6 +7378,17 @@ } template +StmtResult TreeTransform::TransformOMPTargetEnterDataDirective( + OMPTargetEnterDataDirective *D) { + DeclarationNameInfo DirName; + getDerived().getSema().StartOpenMPDSABlock(OMPD_target_enter_data, DirName, + nullptr, D->getLocStart()); + StmtResult Res = getDerived().TransformOMPExecutableDirective(D); + getDerived().getSema().EndOpenMPDSABlock(Res.get()); + return Res; +} + +template StmtResult TreeTransform::TransformOMPTeamsDirective(OMPTeamsDirective *D) { DeclarationNameInfo DirName; Index: lib/Serialization/ASTReaderStmt.cpp =================================================================== --- lib/Serialization/ASTReaderStmt.cpp +++ lib/Serialization/ASTReaderStmt.cpp @@ -2442,6 +2442,13 @@ VisitOMPExecutableDirective(D); } +void ASTStmtReader::VisitOMPTargetEnterDataDirective( + OMPTargetEnterDataDirective *D) { + VisitStmt(D); + ++Idx; + VisitOMPExecutableDirective(D); +} + void ASTStmtReader::VisitOMPTeamsDirective(OMPTeamsDirective *D) { VisitStmt(D); // The NumClauses field was read in ReadStmtFromStream. @@ -3085,6 +3092,11 @@ Context, Record[ASTStmtReader::NumStmtFields], Empty); break; + case STMT_OMP_TARGET_ENTER_DATA_DIRECTIVE: + S = OMPTargetEnterDataDirective::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 @@ -2205,6 +2205,14 @@ Code = serialization::STMT_OMP_TARGET_DATA_DIRECTIVE; } +void ASTStmtWriter::VisitOMPTargetEnterDataDirective( + OMPTargetEnterDataDirective *D) { + VisitStmt(D); + Record.push_back(D->getNumClauses()); + VisitOMPExecutableDirective(D); + Code = serialization::STMT_OMP_TARGET_ENTER_DATA_DIRECTIVE; +} + void ASTStmtWriter::VisitOMPTaskyieldDirective(OMPTaskyieldDirective *D) { VisitStmt(D); VisitOMPExecutableDirective(D); Index: lib/StaticAnalyzer/Core/ExprEngine.cpp =================================================================== --- lib/StaticAnalyzer/Core/ExprEngine.cpp +++ lib/StaticAnalyzer/Core/ExprEngine.cpp @@ -830,6 +830,7 @@ case Stmt::OMPAtomicDirectiveClass: case Stmt::OMPTargetDirectiveClass: case Stmt::OMPTargetDataDirectiveClass: + case Stmt::OMPTargetEnterDataDirectiveClass: 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 @@ -97,6 +97,11 @@ } #pragma omp parallel { +#pragma omp target enter data map(to: a) + ++a; + } +#pragma omp parallel + { #pragma omp teams // expected-error {{region cannot be closely nested inside 'parallel' region; perhaps you forget to enclose 'omp teams' directive into a target region?}} ++a; } @@ -236,6 +241,11 @@ } #pragma omp simd for (int i = 0; i < 10; ++i) { +#pragma omp target enter data map(to: a) // expected-error {{OpenMP constructs may not be nested inside a simd region}} + ++a; + } +#pragma omp simd + for (int i = 0; i < 10; ++i) { #pragma omp teams // expected-error {{OpenMP constructs may not be nested inside a simd region}} ++a; } @@ -398,6 +408,11 @@ } #pragma omp for for (int i = 0; i < 10; ++i) { +#pragma omp target enter data map(to: a) + ++a; + } +#pragma omp for + for (int i = 0; i < 10; ++i) { #pragma omp teams // expected-error {{region cannot be closely nested inside 'for' region; perhaps you forget to enclose 'omp teams' directive into a target region?}} ++a; } @@ -537,6 +552,11 @@ } #pragma omp for simd for (int i = 0; i < 10; ++i) { +#pragma omp target enter data map(to: a) // expected-error {{OpenMP constructs may not be nested inside a simd region}} + ++a; + } +#pragma omp for simd + for (int i = 0; i < 10; ++i) { #pragma omp teams // expected-error {{OpenMP constructs may not be nested inside a simd region}} ++a; } @@ -706,6 +726,10 @@ } #pragma omp sections { +#pragma omp target enter data map(to: a) + } +#pragma omp sections + { #pragma omp teams // expected-error {{region cannot be closely nested inside 'sections' region; perhaps you forget to enclose 'omp teams' directive into a target region?}} ++a; } @@ -911,6 +935,14 @@ #pragma omp sections { #pragma omp section + { +#pragma omp target enter data map(to: a) + ++a; + } + } +#pragma omp sections + { +#pragma omp section #pragma omp teams // expected-error {{region cannot be closely nested inside 'section' region; perhaps you forget to enclose 'omp teams' directive into a target region?}} ++a; } @@ -1065,6 +1097,11 @@ } #pragma omp single { +#pragma omp target enter data map(to: a) + ++a; + } +#pragma omp single + { #pragma omp teams // expected-error {{region cannot be closely nested inside 'single' region; perhaps you forget to enclose 'omp teams' directive into a target region?}} ++a; } @@ -1217,6 +1254,11 @@ } #pragma omp master { +#pragma omp target enter data map(to: a) + ++a; + } +#pragma omp master + { #pragma omp teams // expected-error {{region cannot be closely nested inside 'master' region; perhaps you forget to enclose 'omp teams' directive into a target region?}} ++a; } @@ -1383,6 +1425,11 @@ } #pragma omp critical { +#pragma omp target enter data map(to: a) + ++a; + } +#pragma omp critical + { #pragma omp teams // expected-error {{region cannot be closely nested inside 'critical' region; perhaps you forget to enclose 'omp teams' directive into a target region?}} ++a; } @@ -1550,6 +1597,11 @@ } #pragma omp parallel for for (int i = 0; i < 10; ++i) { +#pragma omp target enter data map(to: a) + ++a; + } +#pragma omp parallel for + for (int i = 0; i < 10; ++i) { #pragma omp teams // expected-error {{region cannot be closely nested inside 'parallel for' region; perhaps you forget to enclose 'omp teams' directive into a target region?}} ++a; } @@ -1717,6 +1769,11 @@ } #pragma omp parallel for simd for (int i = 0; i < 10; ++i) { +#pragma omp target enter data map(to: a) // expected-error {{OpenMP constructs may not be nested inside a simd region}} + ++a; + } +#pragma omp parallel for simd + for (int i = 0; i < 10; ++i) { #pragma omp teams // expected-error {{OpenMP constructs may not be nested inside a simd region}} ++a; } @@ -1875,6 +1932,10 @@ } #pragma omp parallel sections { +#pragma omp target enter data map(to: a) + } +#pragma omp parallel sections + { #pragma omp teams // expected-error {{region cannot be closely nested inside 'parallel sections' region; perhaps you forget to enclose 'omp teams' directive into a target region?}} ++a; } @@ -1979,6 +2040,11 @@ } #pragma omp task { +#pragma omp target enter data map(to: a) + ++a; + } +#pragma omp task + { #pragma omp teams // expected-error {{region cannot be closely nested inside 'task' region; perhaps you forget to enclose 'omp teams' directive into a target region?}} ++a; } @@ -2141,6 +2207,11 @@ } #pragma omp ordered { +#pragma omp target enter data map(to: a) + ++a; + } +#pragma omp ordered + { #pragma omp teams // expected-error {{region cannot be closely nested inside 'ordered' region; perhaps you forget to enclose 'omp teams' directive into a target region?}} ++a; } @@ -2322,6 +2393,13 @@ // 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 enter data map(to: a) // expected-error {{OpenMP constructs may not be nested inside an atomic region}} + ++a; + } +#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 teams // expected-error {{OpenMP constructs may not be nested inside an atomic region}} ++a; } @@ -2458,6 +2536,10 @@ for (int i = 0; i < 10; ++i) ; } +#pragma omp target + { +#pragma omp target enter data map(to: a) + } // TEAMS DIRECTIVE #pragma omp target @@ -2575,6 +2657,12 @@ #pragma omp target #pragma omp teams { +#pragma omp target enter data map(to: a) // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp target enter data' directive into a parallel region?}} + ++a; + } +#pragma omp target +#pragma omp teams + { #pragma omp teams // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp teams' directive into a target region?}} ++a; } @@ -2740,6 +2828,11 @@ } #pragma omp taskloop for (int i = 0; i < 10; ++i) { +#pragma omp target enter data map(to: a) + ++a; + } +#pragma omp taskloop + for (int i = 0; i < 10; ++i) { #pragma omp teams // expected-error {{region cannot be closely nested inside 'taskloop' region; perhaps you forget to enclose 'omp teams' directive into a target region?}} ++a; } @@ -2934,6 +3027,13 @@ #pragma omp teams #pragma omp distribute for (int i = 0; i < 10; ++i) { +#pragma omp target enter data map(to: a) + ++a; + } +#pragma omp target +#pragma omp teams +#pragma omp distribute + for (int i = 0; i < 10; ++i) { #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; } @@ -3033,6 +3133,11 @@ } #pragma omp parallel { +#pragma omp target enter data map(to: a) + ++a; + } +#pragma omp parallel + { #pragma omp teams // expected-error {{region cannot be closely nested inside 'parallel' region; perhaps you forget to enclose 'omp teams' directive into a target region?}} ++a; } @@ -3165,6 +3270,11 @@ } #pragma omp simd for (int i = 0; i < 10; ++i) { +#pragma omp target enter data map(to: a) // expected-error {{OpenMP constructs may not be nested inside a simd region}} + ++a; + } +#pragma omp simd + for (int i = 0; i < 10; ++i) { #pragma omp teams // expected-error {{OpenMP constructs may not be nested inside a simd region}} ++a; } @@ -3317,6 +3427,11 @@ } #pragma omp for for (int i = 0; i < 10; ++i) { +#pragma omp target enter data map(to: a) + ++a; + } +#pragma omp for + for (int i = 0; i < 10; ++i) { #pragma omp teams // expected-error {{region cannot be closely nested inside 'for' region; perhaps you forget to enclose 'omp teams' directive into a target region?}} ++a; } @@ -3449,6 +3564,11 @@ } #pragma omp for simd for (int i = 0; i < 10; ++i) { +#pragma omp target enter data map(to: a) // expected-error {{OpenMP constructs may not be nested inside a simd region}} + ++a; + } +#pragma omp for simd + for (int i = 0; i < 10; ++i) { #pragma omp teams // expected-error {{OpenMP constructs may not be nested inside a simd region}} ++a; } @@ -3593,6 +3713,10 @@ } #pragma omp sections { +#pragma omp target enter data map(to: a) + } +#pragma omp sections + { #pragma omp teams // expected-error {{region cannot be closely nested inside 'sections' region; perhaps you forget to enclose 'omp teams' directive into a target region?}} ++a; } @@ -3803,6 +3927,14 @@ { #pragma omp section { +#pragma omp target enter data map(to: a) + ++a; + } + } +#pragma omp sections + { +#pragma omp section + { #pragma omp teams // expected-error {{region cannot be closely nested inside 'section' region; perhaps you forget to enclose 'omp teams' directive into a target region?}} ++a; } @@ -3950,6 +4082,11 @@ } #pragma omp single { +#pragma omp target enter data map(to: a) + ++a; + } +#pragma omp single + { #pragma omp teams // expected-error {{region cannot be closely nested inside 'single' region; perhaps you forget to enclose 'omp teams' directive into a target region?}} ++a; } @@ -4102,6 +4239,11 @@ } #pragma omp master { +#pragma omp target enter data map(to: a) + ++a; + } +#pragma omp master + { #pragma omp teams // expected-error {{region cannot be closely nested inside 'master' region; perhaps you forget to enclose 'omp teams' directive into a target region?}} ++a; } @@ -4273,6 +4415,11 @@ } #pragma omp critical { +#pragma omp target enter data map(to: a) + ++a; + } +#pragma omp critical + { #pragma omp teams // expected-error {{region cannot be closely nested inside 'critical' region; perhaps you forget to enclose 'omp teams' directive into a target region?}} ++a; } @@ -4440,6 +4587,11 @@ } #pragma omp parallel for for (int i = 0; i < 10; ++i) { +#pragma omp target enter data map(to: a) + ++a; + } +#pragma omp parallel for + for (int i = 0; i < 10; ++i) { #pragma omp teams // expected-error {{region cannot be closely nested inside 'parallel for' region; perhaps you forget to enclose 'omp teams' directive into a target region?}} ++a; } @@ -4607,6 +4759,11 @@ } #pragma omp parallel for simd for (int i = 0; i < 10; ++i) { +#pragma omp target enter data map(to: a) // expected-error {{OpenMP constructs may not be nested inside a simd region}} + ++a; + } +#pragma omp parallel for simd + for (int i = 0; i < 10; ++i) { #pragma omp teams // expected-error {{OpenMP constructs may not be nested inside a simd region}} ++a; } @@ -4761,6 +4918,10 @@ } #pragma omp parallel sections { +#pragma omp target enter data map(to: a) + } +#pragma omp parallel sections + { #pragma omp teams // expected-error {{region cannot be closely nested inside 'parallel sections' region; perhaps you forget to enclose 'omp teams' directive into a target region?}} ++a; } @@ -4864,6 +5025,11 @@ } #pragma omp task { +#pragma omp target enter data map(to: a) + ++a; + } +#pragma omp task + { #pragma omp teams // expected-error {{region cannot be closely nested inside 'task' region; perhaps you forget to enclose 'omp teams' directive into a target region?}} ++a; } @@ -5045,6 +5211,13 @@ // 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 enter data map(to: a) // expected-error {{OpenMP constructs may not be nested inside an atomic region}} + ++a; + } +#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 teams // expected-error {{OpenMP constructs may not be nested inside an atomic region}} ++a; } @@ -5160,6 +5333,10 @@ } #pragma omp target { +#pragma omp target enter data map(to: a) + } +#pragma omp target + { #pragma omp teams ++a; } @@ -5298,6 +5475,11 @@ #pragma omp target #pragma omp teams { +#pragma omp target enter data map(to: a) // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp target enter data' directive into a parallel region?}} + } +#pragma omp target +#pragma omp teams + { #pragma omp teams // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp teams' directive into a target region?}} ++a; } @@ -5463,6 +5645,11 @@ } #pragma omp taskloop for (int i = 0; i < 10; ++i) { +#pragma omp target enter data map(to: a) + ++a; + } +#pragma omp taskloop + for (int i = 0; i < 10; ++i) { #pragma omp teams // expected-error {{region cannot be closely nested inside 'taskloop' region; perhaps you forget to enclose 'omp teams' directive into a target region?}} ++a; } @@ -5662,5 +5849,11 @@ ++a; } return foo(); +#pragma omp target +#pragma omp teams +#pragma omp distribute + for (int i = 0; i < 10; ++i) { +#pragma omp target enter data map(to: a) + ++a; + } } - Index: test/OpenMP/target_enter_data_ast_print.cpp =================================================================== --- /dev/null +++ test/OpenMP/target_enter_data_ast_print.cpp @@ -0,0 +1,100 @@ +// 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 + +template +T tmain(T argc, T *argv) { + T i, j, b, c, d, e, x[20]; + + i = argc; +#pragma omp target enter data map(to: i) + +#pragma omp target enter data map(to: i) if (target enter data: j > 0) + +#pragma omp target enter data map(to: i) if (b) + +#pragma omp target enter data map(to: c) + +#pragma omp target enter data map(to: c) if(b>e) + +#pragma omp target enter data map(alloc: x[0:10], c) + +#pragma omp target enter data map(to: c) map(alloc: d) + +#pragma omp target enter data map(always,alloc: e) + + return 0; +} + +// CHECK: template int tmain(int argc, int *argv) { +// CHECK-NEXT: int i, j, b, c, d, e, x[20]; +// CHECK-NEXT: i = argc; +// CHECK-NEXT: #pragma omp target enter data map(to: i) +// CHECK-NEXT: #pragma omp target enter data map(to: i) if(target enter data: j > 0) +// CHECK-NEXT: #pragma omp target enter data map(to: i) if(b) +// CHECK-NEXT: #pragma omp target enter data map(to: c) +// CHECK-NEXT: #pragma omp target enter data map(to: c) if(b > e) +// CHECK-NEXT: #pragma omp target enter data map(alloc: x[0:10],c) +// CHECK-NEXT: #pragma omp target enter data map(to: c) map(alloc: d) +// CHECK-NEXT: #pragma omp target enter data map(always,alloc: e) +// CHECK: template char tmain(char argc, char *argv) { +// CHECK-NEXT: char i, j, b, c, d, e, x[20]; +// CHECK-NEXT: i = argc; +// CHECK-NEXT: #pragma omp target enter data map(to: i) +// CHECK-NEXT: #pragma omp target enter data map(to: i) if(target enter data: j > 0) +// CHECK-NEXT: #pragma omp target enter data map(to: i) if(b) +// CHECK-NEXT: #pragma omp target enter data map(to: c) +// CHECK-NEXT: #pragma omp target enter data map(to: c) if(b > e) +// CHECK-NEXT: #pragma omp target enter data map(alloc: x[0:10],c) +// CHECK-NEXT: #pragma omp target enter data map(to: c) map(alloc: d) +// CHECK-NEXT: #pragma omp target enter data map(always,alloc: e) +// CHECK: template T tmain(T argc, T *argv) { +// CHECK-NEXT: T i, j, b, c, d, e, x[20]; +// CHECK-NEXT: i = argc; +// CHECK-NEXT: #pragma omp target enter data map(to: i) +// CHECK-NEXT: #pragma omp target enter data map(to: i) if(target enter data: j > 0) +// CHECK-NEXT: #pragma omp target enter data map(to: i) if(b) +// CHECK-NEXT: #pragma omp target enter data map(to: c) +// CHECK-NEXT: #pragma omp target enter data map(to: c) if(b > e) +// CHECK-NEXT: #pragma omp target enter data map(alloc: x[0:10],c) +// CHECK-NEXT: #pragma omp target enter data map(to: c) map(alloc: d) +// CHECK-NEXT: #pragma omp target enter data map(always,alloc: e) + +int main (int argc, char **argv) { + int b = argc, c, d, e, f, g, x[20]; + static int a; +// CHECK: static int a; + +#pragma omp target enter data map(to: a) +// CHECK: #pragma omp target enter data map(to: a) + a=2; +// CHECK-NEXT: a = 2; +#pragma omp target enter data map(to: a) if (target enter data: b) +// CHECK: #pragma omp target enter data map(to: a) if(target enter data: b) + +#pragma omp target enter data map(to: a) if (b > g) +// CHECK: #pragma omp target enter data map(to: a) if(b > g) + +#pragma omp target enter data map(to: c) +// CHECK-NEXT: #pragma omp target enter data map(to: c) + +#pragma omp target enter data map(alloc: c) if(b>g) +// CHECK-NEXT: #pragma omp target enter data map(alloc: c) if(b > g) + +#pragma omp target enter data map(to: x[0:10], c) +// CHECK-NEXT: #pragma omp target enter data map(to: x[0:10],c) + +#pragma omp target enter data map(to: c) map(alloc: d) +// CHECK-NEXT: #pragma omp target enter data map(to: c) map(alloc: d) + +#pragma omp target enter data map(always,alloc: e) +// CHECK-NEXT: #pragma omp target enter data map(always,alloc: e) + + return tmain(argc, &argc) + tmain(argv[0][0], argv[0]); +} + +#endif Index: test/OpenMP/target_enter_data_device_messages.cpp =================================================================== --- /dev/null +++ test/OpenMP/target_enter_data_device_messages.cpp @@ -0,0 +1,29 @@ +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - %s + +void foo() { +} + +bool foobool(int argc) { + return argc; +} + +struct S1; // expected-note {{declared here}} + +int main(int argc, char **argv) { + int i; + #pragma omp target enter data map(to: i) device // expected-error {{expected '(' after 'device'}} + #pragma omp target enter data map(to: i) device ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target enter data map(to: i) device () // expected-error {{expected expression}} + #pragma omp target enter data map(to: i) device (argc // expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target enter data map(to: i) device (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target enter data' are ignored}} +#pragma omp target enter data map(to: i) device (argc > 0 ? argv[1] : argv[2]) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}} + #pragma omp target enter data map(to: i) device (argc + argc) + #pragma omp target enter data map(to: i) device (argc), device (argc+1) // expected-error {{directive '#pragma omp target enter data' cannot contain more than one 'device' clause}} + #pragma omp target enter data map(to: i) device (S1) // expected-error {{'S1' does not refer to a value}} + #pragma omp target enter data map(to: i) device (-2) // expected-error {{argument to 'device' clause must be a non-negative integer value}} + #pragma omp target enter data map(to: i) device (-10u) + #pragma omp target enter data map(to: i) device (3.14) // expected-error {{expression must have integral or unscoped enumeration type, not 'double'}} + foo(); + + return 0; +} Index: test/OpenMP/target_enter_data_if_messages.cpp =================================================================== --- /dev/null +++ test/OpenMP/target_enter_data_if_messages.cpp @@ -0,0 +1,35 @@ +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - %s + +void foo() { +} + +bool foobool(int argc) { + return argc; +} + +struct S1; // expected-note {{declared here}} + +int main(int argc, char **argv) { + int i; + #pragma omp target enter data map(to: i) if // expected-error {{expected '(' after 'if'}} + #pragma omp target enter data map(to: i) if ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target enter data map(to: i) if () // expected-error {{expected expression}} + #pragma omp target enter data map(to: i) if (argc // expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target enter data map(to: i) if (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target enter data' are ignored}} + #pragma omp target enter data map(to: i) if (argc > 0 ? argv[1] : argv[2]) + #pragma omp target enter data map(to: i) if (argc + argc) + #pragma omp target enter data map(to: i) if (foobool(argc)), if (true) // expected-error {{directive '#pragma omp target enter data' cannot contain more than one 'if' clause}} + #pragma omp target enter data map(to: i) if (S1) // expected-error {{'S1' does not refer to a value}} + #pragma omp target enter data map(to: i) if (argv[1]=2) // expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target enter data map(to: i) if(target data : true) // expected-error {{directive name modifier 'target data' is not allowed for '#pragma omp target enter data'}} + #pragma omp target enter data map(to: i) if(target enter data : // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target enter data map(to: i) if(target enter data : // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target enter data map(to: i) if(target enter data : argc // expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target enter data map(to: i) if(target enter data : argc) + #pragma omp target enter data map(to: i) if(target enter data : argc) if (for:argc) // expected-error {{directive name modifier 'for' is not allowed for '#pragma omp target enter data'}} + #pragma omp target enter data map(to: i) if(target enter data : argc) if (target enter data:argc) // expected-error {{directive '#pragma omp target enter data' cannot contain more than one 'if' clause with 'target enter data' name modifier}} + #pragma omp target enter data map(to: i) if(target enter data : argc) if (argc) // expected-error {{no more 'if' clause is allowed}} expected-note {{previous clause with directive name modifier specified here}} + foo(); + + return 0; +} Index: test/OpenMP/target_enter_data_map_messages.c =================================================================== --- /dev/null +++ test/OpenMP/target_enter_data_map_messages.c @@ -0,0 +1,17 @@ +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - %s + +int main(int argc, char **argv) { + + int r; + #pragma omp target enter data // expected-error {{expected at least one map clause for '#pragma omp target enter data'}} + + #pragma omp target enter data map(tofrom: r) // expected-error {{map type 'tofrom' is not allowed for '#pragma omp target enter data'}} + + #pragma omp target enter data map(always, to: r) + #pragma omp target enter data map(always, alloc: r) + #pragma omp target enter data map(always, from: r) // expected-error {{map type 'from' is not allowed for '#pragma omp target enter data'}} + #pragma omp target enter data map(release: r) // expected-error {{map type 'release' is not allowed for '#pragma omp target enter data'}} + #pragma omp target enter data map(delete: r) // expected-error {{map type 'delete' is not allowed for '#pragma omp target enter data'}} + + return 0; +} Index: tools/libclang/CIndex.cpp =================================================================== --- tools/libclang/CIndex.cpp +++ tools/libclang/CIndex.cpp @@ -1953,6 +1953,7 @@ void VisitOMPAtomicDirective(const OMPAtomicDirective *D); void VisitOMPTargetDirective(const OMPTargetDirective *D); void VisitOMPTargetDataDirective(const OMPTargetDataDirective *D); + void VisitOMPTargetEnterDataDirective(const OMPTargetEnterDataDirective *D); void VisitOMPTeamsDirective(const OMPTeamsDirective *D); void VisitOMPTaskLoopDirective(const OMPTaskLoopDirective *D); void VisitOMPTaskLoopSimdDirective(const OMPTaskLoopSimdDirective *D); @@ -2628,6 +2629,11 @@ VisitOMPExecutableDirective(D); } +void EnqueueVisitor::VisitOMPTargetEnterDataDirective( + const OMPTargetEnterDataDirective *D) { + VisitOMPExecutableDirective(D); +} + void EnqueueVisitor::VisitOMPTeamsDirective(const OMPTeamsDirective *D) { VisitOMPExecutableDirective(D); } @@ -4510,6 +4516,8 @@ return cxstring::createRef("OMPTargetDirective"); case CXCursor_OMPTargetDataDirective: return cxstring::createRef("OMPTargetDataDirective"); + case CXCursor_OMPTargetEnterDataDirective: + return cxstring::createRef("OMPTargetEnterDataDirective"); case CXCursor_OMPTeamsDirective: return cxstring::createRef("OMPTeamsDirective"); case CXCursor_OMPCancellationPointDirective: Index: tools/libclang/CXCursor.cpp =================================================================== --- tools/libclang/CXCursor.cpp +++ tools/libclang/CXCursor.cpp @@ -600,6 +600,9 @@ case Stmt::OMPTargetDataDirectiveClass: K = CXCursor_OMPTargetDataDirective; break; + case Stmt::OMPTargetEnterDataDirectiveClass: + K = CXCursor_OMPTargetEnterDataDirective; + break; case Stmt::OMPTeamsDirectiveClass: K = CXCursor_OMPTeamsDirective; break;