diff --git a/clang/include/clang-c/Index.h b/clang/include/clang-c/Index.h index 69a98d7978c9..c33ff952151f 100644 --- a/clang/include/clang-c/Index.h +++ b/clang/include/clang-c/Index.h @@ -2274,7 +2274,11 @@ enum CXCursorKind { */ 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. diff --git a/clang/include/clang/AST/RecursiveASTVisitor.h b/clang/include/clang/AST/RecursiveASTVisitor.h index f1ae76b59c23..1928d4612d9c 100644 --- a/clang/include/clang/AST/RecursiveASTVisitor.h +++ b/clang/include/clang/AST/RecursiveASTVisitor.h @@ -2433,6 +2433,9 @@ DEF_TRAVERSE_STMT(OMPTargetDirective, DEF_TRAVERSE_STMT(OMPTargetDataDirective, { TRY_TO(TraverseOMPExecutableDirective(S)); }) +DEF_TRAVERSE_STMT(OMPTargetEnterDataDirective, + { TRY_TO(TraverseOMPExecutableDirective(S)); }) + DEF_TRAVERSE_STMT(OMPTeamsDirective, { TRY_TO(TraverseOMPExecutableDirective(S)); }) diff --git a/clang/include/clang/AST/StmtOpenMP.h b/clang/include/clang/AST/StmtOpenMP.h index c82aeda2aeec..d216204b2057 100644 --- a/clang/include/clang/AST/StmtOpenMP.h +++ b/clang/include/clang/AST/StmtOpenMP.h @@ -2038,6 +2038,66 @@ public: } }; +/// \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 diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 8cbb71b021c8..ade42225ebc4 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -7937,6 +7937,10 @@ def err_omp_map_shared_storage : Error< "variable already marked as mapped in current construct">; def err_omp_not_mappable_type : Error< "type %0 is not mappable to target">; +def 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< diff --git a/clang/include/clang/Basic/OpenMPKinds.def b/clang/include/clang/Basic/OpenMPKinds.def index c388be3991a0..ffa1872e769a 100644 --- a/clang/include/clang/Basic/OpenMPKinds.def +++ b/clang/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 @@ -128,6 +131,7 @@ OPENMP_DIRECTIVE(target) 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") @@ -354,6 +358,12 @@ OPENMP_TARGET_DATA_CLAUSE(if) 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) @@ -452,6 +462,7 @@ OPENMP_DIST_SCHEDULE_KIND(static) #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 diff --git a/clang/include/clang/Basic/StmtNodes.td b/clang/include/clang/Basic/StmtNodes.td index 36519ea29c98..05fc806ce147 100644 --- a/clang/include/clang/Basic/StmtNodes.td +++ b/clang/include/clang/Basic/StmtNodes.td @@ -216,6 +216,7 @@ def OMPOrderedDirective : DStmt; def OMPAtomicDirective : DStmt; def OMPTargetDirective : DStmt; def OMPTargetDataDirective : DStmt; +def OMPTargetEnterDataDirective : DStmt; def OMPTeamsDirective : DStmt; def OMPCancellationPointDirective : DStmt; def OMPCancelDirective : DStmt; diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index c60d364f15d4..e3136b45b3d6 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -7950,6 +7950,11 @@ public: 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, diff --git a/clang/include/clang/Serialization/ASTBitCodes.h b/clang/include/clang/Serialization/ASTBitCodes.h index 0dfb8cf37146..03bc505895b2 100644 --- a/clang/include/clang/Serialization/ASTBitCodes.h +++ b/clang/include/clang/Serialization/ASTBitCodes.h @@ -1445,6 +1445,7 @@ namespace clang { 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, diff --git a/clang/lib/AST/StmtOpenMP.cpp b/clang/lib/AST/StmtOpenMP.cpp index 72e62b7001dd..b3dfd2516925 100644 --- a/clang/lib/AST/StmtOpenMP.cpp +++ b/clang/lib/AST/StmtOpenMP.cpp @@ -716,6 +716,27 @@ OMPTargetDataDirective *OMPTargetDataDirective::CreateEmpty(const ASTContext &C, return new (Mem) OMPTargetDataDirective(N); } +OMPTargetEnterDataDirective *OMPTargetEnterDataDirective::Create( + const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc, + ArrayRef Clauses) { + void *Mem = C.Allocate(llvm::alignTo(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::alignTo(sizeof(OMPTargetEnterDataDirective), + llvm::alignOf()) + + sizeof(OMPClause *) * N); + return new (Mem) OMPTargetEnterDataDirective(N); +} + OMPTeamsDirective *OMPTeamsDirective::Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc, diff --git a/clang/lib/AST/StmtPrinter.cpp b/clang/lib/AST/StmtPrinter.cpp index 0d7063e7a206..2e8da360c943 100644 --- a/clang/lib/AST/StmtPrinter.cpp +++ b/clang/lib/AST/StmtPrinter.cpp @@ -1061,6 +1061,12 @@ void StmtPrinter::VisitOMPTargetDataDirective(OMPTargetDataDirective *Node) { 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); diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp index f8aa4db2c065..9ab9c8e12172 100644 --- a/clang/lib/AST/StmtProfile.cpp +++ b/clang/lib/AST/StmtProfile.cpp @@ -584,6 +584,11 @@ void StmtProfiler::VisitOMPTargetDataDirective(const OMPTargetDataDirective *S) VisitOMPExecutableDirective(S); } +void StmtProfiler::VisitOMPTargetEnterDataDirective( + const OMPTargetEnterDataDirective *S) { + VisitOMPExecutableDirective(S); +} + void StmtProfiler::VisitOMPTeamsDirective(const OMPTeamsDirective *S) { VisitOMPExecutableDirective(S); } diff --git a/clang/lib/Basic/OpenMPKinds.cpp b/clang/lib/Basic/OpenMPKinds.cpp index 687bf4c5f529..167a700d27f0 100644 --- a/clang/lib/Basic/OpenMPKinds.cpp +++ b/clang/lib/Basic/OpenMPKinds.cpp @@ -408,6 +408,16 @@ bool clang::isAllowedClauseForDirective(OpenMPDirectiveKind DKind, #define OPENMP_TARGET_DATA_CLAUSE(Name) \ case OMPC_##Name: \ return true; +#include "clang/Basic/OpenMPKinds.def" + default: + 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; diff --git a/clang/lib/CodeGen/CGStmt.cpp b/clang/lib/CodeGen/CGStmt.cpp index cc4fa2ec5972..df6766a768e1 100644 --- a/clang/lib/CodeGen/CGStmt.cpp +++ b/clang/lib/CodeGen/CGStmt.cpp @@ -256,6 +256,9 @@ void CodeGenFunction::EmitStmt(const Stmt *S) { case Stmt::OMPTargetDataDirectiveClass: EmitOMPTargetDataDirective(cast(*S)); break; + case Stmt::OMPTargetEnterDataDirectiveClass: + EmitOMPTargetEnterDataDirective(cast(*S)); + break; case Stmt::OMPTaskLoopDirectiveClass: EmitOMPTaskLoopDirective(cast(*S)); break; diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp index bf00b04563c6..f6b8e5907d47 100644 --- a/clang/lib/CodeGen/CGStmtOpenMP.cpp +++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp @@ -2661,6 +2661,11 @@ void CodeGenFunction::EmitOMPTargetDataDirective( [&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()); diff --git a/clang/lib/CodeGen/CodeGenFunction.h b/clang/lib/CodeGen/CodeGenFunction.h index 3dc0f35b0db2..d1cf0c9206e8 100644 --- a/clang/lib/CodeGen/CodeGenFunction.h +++ b/clang/lib/CodeGen/CodeGenFunction.h @@ -2336,6 +2336,7 @@ public: 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); diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp index 0531847b47cc..4baf5d266169 100644 --- a/clang/lib/Parse/ParseOpenMP.cpp +++ b/clang/lib/Parse/ParseOpenMP.cpp @@ -34,6 +34,10 @@ static OpenMPDirectiveKind ParseOpenMPDirectiveKind(Parser &P) { {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,9 @@ static OpenMPDirectiveKind ParseOpenMPDirectiveKind(Parser &P) { 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 +72,10 @@ static OpenMPDirectiveKind ParseOpenMPDirectiveKind(Parser &P) { 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 +146,7 @@ Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirective() { 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 +171,7 @@ Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirective() { /// '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( AllowedContsructsKind Allowed) { @@ -217,6 +225,7 @@ StmtResult Parser::ParseOpenMPDeclarativeOrExecutableDirective( case OMPD_taskwait: case OMPD_cancellation_point: case OMPD_cancel: + case OMPD_target_enter_data: if (Allowed == ACK_StatementsOpenMPNonStandalone) { Diag(Tok, diag::err_omp_immediate_directive) << getOpenMPDirectiveName(DKind) << 0; diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp index acc54c1f9c16..a3e33bc65802 100644 --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -1610,6 +1610,7 @@ void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) { 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,13 +1725,15 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack, // | parallel | ordered | + | // | parallel | atomic | * | // | parallel | target | * | + // | parallel | target enter | * | + // | | data | | // | parallel | teams | + | // | parallel | cancellation | | // | | point | ! | // | parallel | cancel | ! | // | parallel | taskloop | * | // | parallel | taskloop simd | * | - // | parallel | distribute | | + // | parallel | distribute | | // +------------------+-----------------+------------------------------------+ // | for | parallel | * | // | for | for | + | @@ -1753,6 +1756,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack, // | for | ordered | * (if construct is ordered) | // | for | atomic | * | // | for | target | * | + // | for | target enter | * | + // | | data | | // | for | teams | + | // | for | cancellation | | // | | point | ! | @@ -1782,6 +1787,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack, // | master | ordered | + | // | master | atomic | * | // | master | target | * | + // | master | target enter | * | + // | | data | | // | master | teams | + | // | master | cancellation | | // | | point | | @@ -1810,6 +1817,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack, // | critical | ordered | + | // | critical | atomic | * | // | critical | target | * | + // | critical | target enter | * | + // | | data | | // | critical | teams | + | // | critical | cancellation | | // | | point | | @@ -1839,6 +1848,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack, // | simd | ordered | + (with simd clause) | // | simd | atomic | | // | simd | target | | + // | simd | target enter | | + // | | data | | // | simd | teams | | // | simd | cancellation | | // | | point | | @@ -1868,6 +1879,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack, // | 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 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack, // | 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 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack, // | sections | ordered | + | // | sections | atomic | * | // | sections | target | * | + // | sections | target enter | * | + // | | data | | // | sections | teams | + | // | sections | cancellation | | // | | point | ! | @@ -1955,6 +1972,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack, // | section | ordered | + | // | section | atomic | * | // | section | target | * | + // | section | target enter | * | + // | | data | | // | section | teams | + | // | section | cancellation | | // | | point | ! | @@ -1984,6 +2003,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack, // | single | ordered | + | // | single | atomic | * | // | single | target | * | + // | single | target enter | * | + // | | data | | // | single | teams | + | // | single | cancellation | | // | | point | | @@ -2013,6 +2034,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack, // | 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,13 +2065,15 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack, // | parallel sections| ordered | + | // | parallel sections| atomic | * | // | parallel sections| target | * | + // | parallel sections| target enter | * | + // | | data | | // | parallel sections| teams | + | // | parallel sections| cancellation | | // | | point | ! | // | parallel sections| cancel | ! | // | parallel sections| taskloop | * | // | parallel sections| taskloop simd | * | - // | parallel sections| distribute | | + // | parallel sections| distribute | | // +------------------+-----------------+------------------------------------+ // | task | parallel | * | // | task | for | + | @@ -2071,6 +2096,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack, // | task | ordered | + | // | task | atomic | * | // | task | target | * | + // | task | target enter | * | + // | | data | | // | task | teams | + | // | task | cancellation | | // | | point | ! | @@ -2100,6 +2127,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack, // | ordered | ordered | + | // | ordered | atomic | * | // | ordered | target | * | + // | ordered | target enter | * | + // | | data | | // | ordered | teams | + | // | ordered | cancellation | | // | | point | | @@ -2129,13 +2158,15 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack, // | atomic | ordered | | // | atomic | atomic | | // | atomic | target | | + // | atomic | target enter | | + // | | data | | // | atomic | teams | | // | atomic | cancellation | | // | | point | | // | atomic | cancel | | // | atomic | taskloop | | // | atomic | taskloop simd | | - // | atomic | distribute | | + // | atomic | distribute | | // +------------------+-----------------+------------------------------------+ // | target | parallel | * | // | target | for | * | @@ -2158,6 +2189,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack, // | target | ordered | * | // | target | atomic | * | // | target | target | * | + // | target | target enter | * | + // | | data | | // | target | teams | * | // | target | cancellation | | // | | point | | @@ -2187,6 +2220,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack, // | teams | ordered | + | // | teams | atomic | + | // | teams | target | + | + // | teams | target enter | + | + // | | data | | // | teams | teams | + | // | teams | cancellation | | // | | point | | @@ -2216,6 +2251,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack, // | taskloop | ordered | + | // | taskloop | atomic | * | // | taskloop | target | * | + // | taskloop | target enter | * | + // | | data | | // | taskloop | teams | + | // | taskloop | cancellation | | // | | point | | @@ -2244,6 +2281,8 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack, // | 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 @@ static bool CheckNestingOfRegions(Sema &SemaRef, DSAStackTy *Stack, // | distribute | ordered | + | // | distribute | atomic | * | // | distribute | target | | + // | distribute | target enter | | + // | | data | | // | distribute | teams | | // | distribute | cancellation | + | // | | point | | @@ -2697,6 +2738,11 @@ StmtResult Sema::ActOnOpenMPExecutableDirective( 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 @@ StmtResult Sema::ActOnOpenMPTargetDirective(ArrayRef Clauses, 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,22 @@ StmtResult Sema::ActOnOpenMPTargetDataDirective(ArrayRef Clauses, 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) { @@ -8449,6 +8523,22 @@ OMPClause *Sema::ActOnOpenMPMapClause( 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); diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h index 60c661c45477..4b614d888c1a 100644 --- a/clang/lib/Sema/TreeTransform.h +++ b/clang/lib/Sema/TreeTransform.h @@ -7390,6 +7390,17 @@ StmtResult TreeTransform::TransformOMPTargetDataDirective( return Res; } +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) { diff --git a/clang/lib/Serialization/ASTReaderStmt.cpp b/clang/lib/Serialization/ASTReaderStmt.cpp index 3c2a98b09f7d..4b312ce75795 100644 --- a/clang/lib/Serialization/ASTReaderStmt.cpp +++ b/clang/lib/Serialization/ASTReaderStmt.cpp @@ -2455,6 +2455,13 @@ void ASTStmtReader::VisitOMPTargetDataDirective(OMPTargetDataDirective *D) { 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. @@ -3098,6 +3105,11 @@ Stmt *ASTReader::ReadStmtFromStream(ModuleFile &F) { 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); diff --git a/clang/lib/Serialization/ASTWriterStmt.cpp b/clang/lib/Serialization/ASTWriterStmt.cpp index 7b5440b2f57c..ed6044d45f32 100644 --- a/clang/lib/Serialization/ASTWriterStmt.cpp +++ b/clang/lib/Serialization/ASTWriterStmt.cpp @@ -2214,6 +2214,14 @@ void ASTStmtWriter::VisitOMPTargetDataDirective(OMPTargetDataDirective *D) { 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); diff --git a/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp b/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp index 662b0a2dd798..e4e61ab2d646 100644 --- a/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp +++ b/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp @@ -830,6 +830,7 @@ void ExprEngine::Visit(const Stmt *S, ExplodedNode *Pred, case Stmt::OMPAtomicDirectiveClass: case Stmt::OMPTargetDirectiveClass: case Stmt::OMPTargetDataDirectiveClass: + case Stmt::OMPTargetEnterDataDirectiveClass: case Stmt::OMPTeamsDirectiveClass: case Stmt::OMPCancellationPointDirectiveClass: case Stmt::OMPCancelDirectiveClass: diff --git a/clang/test/OpenMP/nesting_of_regions.cpp b/clang/test/OpenMP/nesting_of_regions.cpp index b2b87db6a158..0f83a5ee8c4d 100644 --- a/clang/test/OpenMP/nesting_of_regions.cpp +++ b/clang/test/OpenMP/nesting_of_regions.cpp @@ -97,6 +97,11 @@ void foo() { } #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 @@ void foo() { } #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 @@ void foo() { } #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 @@ void foo() { } #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 @@ void foo() { } #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; } @@ -910,6 +934,14 @@ void foo() { } #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 @@ void foo() { } #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 @@ void foo() { } #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 @@ void foo() { } #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 @@ void foo() { } #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 @@ void foo() { } #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 @@ void foo() { } #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 @@ void foo() { } #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 @@ void foo() { } #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 @@ void foo() { // 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 @@ void foo() { for (int i = 0; i < 10; ++i) ; } +#pragma omp target + { +#pragma omp target enter data map(to: a) + } // TEAMS DIRECTIVE #pragma omp target @@ -2573,6 +2655,12 @@ void foo() { ++a; } #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?}} @@ -2740,6 +2828,11 @@ void foo() { } #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; } @@ -2932,6 +3025,13 @@ void 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; + } +#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?}} @@ -3033,6 +3133,11 @@ void foo() { } #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 @@ void foo() { } #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 @@ void foo() { } #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 @@ void foo() { } #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 @@ void foo() { } #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 @@ void foo() { { #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 @@ void foo() { } #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 @@ void foo() { } #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 @@ void foo() { } #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 @@ void foo() { } #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 @@ void foo() { } #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 @@ void foo() { } #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 @@ void foo() { } #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 @@ void foo() { // 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 @@ void foo() { } #pragma omp target { +#pragma omp target enter data map(to: a) + } +#pragma omp target + { #pragma omp teams ++a; } @@ -5296,6 +5473,11 @@ void foo() { ++a; } #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?}} @@ -5463,6 +5645,11 @@ void foo() { } #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 @@ void foo() { ++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; + } } - diff --git a/clang/test/OpenMP/target_enter_data_ast_print.cpp b/clang/test/OpenMP/target_enter_data_ast_print.cpp new file mode 100644 index 000000000000..d0ec621f3dd3 --- /dev/null +++ b/clang/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 diff --git a/clang/test/OpenMP/target_enter_data_device_messages.cpp b/clang/test/OpenMP/target_enter_data_device_messages.cpp new file mode 100644 index 000000000000..d954eca319dd --- /dev/null +++ b/clang/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; +} diff --git a/clang/test/OpenMP/target_enter_data_if_messages.cpp b/clang/test/OpenMP/target_enter_data_if_messages.cpp new file mode 100644 index 000000000000..0d18af187b19 --- /dev/null +++ b/clang/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; +} diff --git a/clang/test/OpenMP/target_enter_data_map_messages.c b/clang/test/OpenMP/target_enter_data_map_messages.c new file mode 100644 index 000000000000..fcb27a0b796f --- /dev/null +++ b/clang/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; +} diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp index 18851a0dd8b7..84ba6b0ab178 100644 --- a/clang/tools/libclang/CIndex.cpp +++ b/clang/tools/libclang/CIndex.cpp @@ -1953,6 +1953,7 @@ public: 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); @@ -2633,6 +2634,11 @@ void EnqueueVisitor::VisitOMPTargetDataDirective(const VisitOMPExecutableDirective(D); } +void EnqueueVisitor::VisitOMPTargetEnterDataDirective( + const OMPTargetEnterDataDirective *D) { + VisitOMPExecutableDirective(D); +} + void EnqueueVisitor::VisitOMPTeamsDirective(const OMPTeamsDirective *D) { VisitOMPExecutableDirective(D); } @@ -4838,6 +4844,8 @@ CXString clang_getCursorKindSpelling(enum CXCursorKind Kind) { 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: diff --git a/clang/tools/libclang/CXCursor.cpp b/clang/tools/libclang/CXCursor.cpp index c766d2d69fa9..cf37c7c32e2c 100644 --- a/clang/tools/libclang/CXCursor.cpp +++ b/clang/tools/libclang/CXCursor.cpp @@ -600,6 +600,9 @@ CXCursor cxcursor::MakeCXCursor(const Stmt *S, const Decl *Parent, case Stmt::OMPTargetDataDirectiveClass: K = CXCursor_OMPTargetDataDirective; break; + case Stmt::OMPTargetEnterDataDirectiveClass: + K = CXCursor_OMPTargetEnterDataDirective; + break; case Stmt::OMPTeamsDirectiveClass: K = CXCursor_OMPTeamsDirective; break;