Index: include/clang-c/Index.h =================================================================== --- include/clang-c/Index.h +++ include/clang-c/Index.h @@ -2227,17 +2227,21 @@ /** \brief OpenMP taskgroup directive. */ - CXCursor_OMPTaskgroupDirective = 254, + CXCursor_OMPTaskgroupDirective = 254, /** \brief OpenMP cancellation point directive. */ - CXCursor_OMPCancellationPointDirective = 255, + CXCursor_OMPCancellationPointDirective = 255, /** \brief OpenMP cancel directive. */ - CXCursor_OMPCancelDirective = 256, + CXCursor_OMPCancelDirective = 256, - CXCursor_LastStmt = CXCursor_OMPCancelDirective, + /** \brief OpenMP target data directive. + */ + CXCursor_OMPTargetDataDirective = 257, + + CXCursor_LastStmt = CXCursor_OMPTargetDataDirective, /** * \brief Cursor that represents the translation unit itself. Index: include/clang/AST/DataRecursiveASTVisitor.h =================================================================== --- include/clang/AST/DataRecursiveASTVisitor.h +++ include/clang/AST/DataRecursiveASTVisitor.h @@ -2405,6 +2405,9 @@ DEF_TRAVERSE_STMT(OMPTargetDirective, { TRY_TO(TraverseOMPExecutableDirective(S)); }) +DEF_TRAVERSE_STMT(OMPTargetDataDirective, + { TRY_TO(TraverseOMPExecutableDirective(S)); }) + DEF_TRAVERSE_STMT(OMPTeamsDirective, { TRY_TO(TraverseOMPExecutableDirective(S)); }) Index: include/clang/AST/RecursiveASTVisitor.h =================================================================== --- include/clang/AST/RecursiveASTVisitor.h +++ include/clang/AST/RecursiveASTVisitor.h @@ -2437,6 +2437,9 @@ DEF_TRAVERSE_STMT(OMPTargetDirective, { TRY_TO(TraverseOMPExecutableDirective(S)); }) +DEF_TRAVERSE_STMT(OMPTargetDataDirective, + { 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 @@ -1798,6 +1798,64 @@ } }; +/// \brief This represents '#pragma omp target data' directive. +/// +/// \code +/// #pragma omp target data device(0) if(a) map(b[:]) +/// \endcode +/// In this example directive '#pragma omp target data' has clauses 'device' +/// with the value '0', 'if' with condition 'a' and 'map' with array +/// section 'b[:]'. +/// +class OMPTargetDataDirective : 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 N The number of clauses. + /// + OMPTargetDataDirective(SourceLocation StartLoc, SourceLocation EndLoc, + unsigned NumClauses) + : OMPExecutableDirective(this, OMPTargetDataDirectiveClass, + OMPD_target_data, StartLoc, EndLoc, NumClauses, + 1) {} + + /// \brief Build an empty directive. + /// + /// \param N Number of clauses. + /// + explicit OMPTargetDataDirective(unsigned NumClauses) + : OMPExecutableDirective(this, OMPTargetDataDirectiveClass, + OMPD_target_data, SourceLocation(), + SourceLocation(), NumClauses, 1) {} + +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 OMPTargetDataDirective * + Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc, + ArrayRef Clauses, Stmt *AssociatedStmt); + + /// \brief Creates an empty directive with the place for \a N clauses. + /// + /// \param C AST context. + /// \param N The number of clauses. + /// + static OMPTargetDataDirective *CreateEmpty(const ASTContext &C, unsigned N, + EmptyShell); + + static bool classof(const Stmt *T) { + return T->getStmtClass() == OMPTargetDataDirectiveClass; + } +}; + /// \brief This represents '#pragma omp teams' directive. /// /// \code Index: include/clang/Basic/OpenMPKinds.def =================================================================== --- include/clang/Basic/OpenMPKinds.def +++ include/clang/Basic/OpenMPKinds.def @@ -57,6 +57,9 @@ #ifndef OPENMP_TARGET_CLAUSE # define OPENMP_TARGET_CLAUSE(Name) #endif +#ifndef OPENMP_TARGET_DATA_CLAUSE +# define OPENMP_TARGET_DATA_CLAUSE(Name) +#endif #ifndef OPENMP_TEAMS_CLAUSE # define OPENMP_TEAMS_CLAUSE(Name) #endif @@ -94,6 +97,7 @@ OPENMP_DIRECTIVE(target) OPENMP_DIRECTIVE(teams) OPENMP_DIRECTIVE(cancel) +OPENMP_DIRECTIVE_EXT(target_data, "target data") OPENMP_DIRECTIVE_EXT(parallel_for, "parallel for") OPENMP_DIRECTIVE_EXT(parallel_for_simd, "parallel for simd") OPENMP_DIRECTIVE_EXT(parallel_sections, "parallel sections") @@ -272,6 +276,10 @@ // TODO More clauses for 'target' directive. OPENMP_TARGET_CLAUSE(if) +// Clauses allowed for OpenMP directive 'target data'. +// TODO More clauses for 'target data' directive. +OPENMP_TARGET_DATA_CLAUSE(if) + // Clauses allowed for OpenMP directive 'teams'. // TODO More clauses for 'teams' directive. OPENMP_TEAMS_CLAUSE(default) @@ -296,6 +304,7 @@ #undef OPENMP_TASK_CLAUSE #undef OPENMP_ATOMIC_CLAUSE #undef OPENMP_TARGET_CLAUSE +#undef OPENMP_TARGET_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 @@ -204,6 +204,7 @@ def OMPOrderedDirective : DStmt; def OMPAtomicDirective : DStmt; def OMPTargetDirective : DStmt; +def OMPTargetDataDirective : 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 @@ -7855,6 +7855,11 @@ StmtResult ActOnOpenMPTargetDirective(ArrayRef Clauses, Stmt *AStmt, SourceLocation StartLoc, SourceLocation EndLoc); + /// \brief Called on well-formed '\#pragma omp target data' after parsing of + /// the associated statement. + StmtResult ActOnOpenMPTargetDataDirective(ArrayRef Clauses, + Stmt *AStmt, 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 @@ -1397,6 +1397,7 @@ STMT_OMP_ORDERED_DIRECTIVE, STMT_OMP_ATOMIC_DIRECTIVE, STMT_OMP_TARGET_DIRECTIVE, + STMT_OMP_TARGET_DATA_DIRECTIVE, STMT_OMP_TEAMS_DIRECTIVE, STMT_OMP_TASKGROUP_DIRECTIVE, STMT_OMP_CANCELLATION_POINT_DIRECTIVE, Index: lib/AST/Stmt.cpp =================================================================== --- lib/AST/Stmt.cpp +++ lib/AST/Stmt.cpp @@ -2225,6 +2225,30 @@ return new (Mem) OMPTargetDirective(NumClauses); } +OMPTargetDataDirective *OMPTargetDataDirective::Create( + const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc, + ArrayRef Clauses, Stmt *AssociatedStmt) { + void *Mem = + C.Allocate(llvm::RoundUpToAlignment(sizeof(OMPTargetDataDirective), + llvm::alignOf()) + + sizeof(OMPClause *) * Clauses.size() + sizeof(Stmt *)); + OMPTargetDataDirective *Dir = + new (Mem) OMPTargetDataDirective(StartLoc, EndLoc, Clauses.size()); + Dir->setClauses(Clauses); + Dir->setAssociatedStmt(AssociatedStmt); + return Dir; +} + +OMPTargetDataDirective *OMPTargetDataDirective::CreateEmpty(const ASTContext &C, + unsigned N, + EmptyShell) { + void *Mem = + C.Allocate(llvm::RoundUpToAlignment(sizeof(OMPTargetDataDirective), + llvm::alignOf()) + + sizeof(OMPClause *) * N + sizeof(Stmt *)); + return new (Mem) OMPTargetDataDirective(N); +} + OMPTeamsDirective *OMPTeamsDirective::Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc, Index: lib/AST/StmtPrinter.cpp =================================================================== --- lib/AST/StmtPrinter.cpp +++ lib/AST/StmtPrinter.cpp @@ -946,6 +946,11 @@ PrintOMPExecutableDirective(Node); } +void StmtPrinter::VisitOMPTargetDataDirective(OMPTargetDataDirective *Node) { + Indent() << "#pragma omp target 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 @@ -533,6 +533,10 @@ VisitOMPExecutableDirective(S); } +void StmtProfiler::VisitOMPTargetDataDirective(const OMPTargetDataDirective *S) { + VisitOMPExecutableDirective(S); +} + void StmtProfiler::VisitOMPTeamsDirective(const OMPTeamsDirective *S) { VisitOMPExecutableDirective(S); } Index: lib/Basic/OpenMPKinds.cpp =================================================================== --- lib/Basic/OpenMPKinds.cpp +++ lib/Basic/OpenMPKinds.cpp @@ -328,6 +328,16 @@ break; } break; + case OMPD_target_data: + switch (CKind) { +#define OPENMP_TARGET_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 @@ -246,6 +246,9 @@ case Stmt::OMPCancelDirectiveClass: EmitOMPCancelDirective(cast(*S)); break; + case Stmt::OMPTargetDataDirectiveClass: + EmitOMPTargetDataDirective(cast(*S)); + break; } } Index: lib/CodeGen/CGStmtOpenMP.cpp =================================================================== --- lib/CodeGen/CGStmtOpenMP.cpp +++ lib/CodeGen/CGStmtOpenMP.cpp @@ -2121,3 +2121,12 @@ : BreakContinueStack.back().BreakBlock; return JumpDest(); } + +// Generate the instructions for '#pragma omp target data' directive. +void CodeGenFunction::EmitOMPTargetDataDirective( + const OMPTargetDataDirective &S) { + + // emit the code inside the construct for now + auto CS = cast(S.getAssociatedStmt()); + EmitStmt(CS->getCapturedStmt()); +} Index: lib/CodeGen/CodeGenFunction.h =================================================================== --- lib/CodeGen/CodeGenFunction.h +++ lib/CodeGen/CodeGenFunction.h @@ -2213,6 +2213,7 @@ void EmitOMPOrderedDirective(const OMPOrderedDirective &S); void EmitOMPAtomicDirective(const OMPAtomicDirective &S); void EmitOMPTargetDirective(const OMPTargetDirective &S); + void EmitOMPTargetDataDirective(const OMPTargetDataDirective &S); void EmitOMPTeamsDirective(const OMPTeamsDirective &S); void EmitOMPCancellationPointDirective(const OMPCancellationPointDirective &S); Index: lib/Parse/ParseOpenMP.cpp =================================================================== --- lib/Parse/ParseOpenMP.cpp +++ lib/Parse/ParseOpenMP.cpp @@ -19,6 +19,7 @@ #include "clang/Parse/Parser.h" #include "clang/Sema/Scope.h" #include "llvm/ADT/PointerIntPair.h" + using namespace clang; //===----------------------------------------------------------------------===// @@ -32,6 +33,7 @@ const OpenMPDirectiveKind F[][3] = { {OMPD_unknown /*cancellation*/, OMPD_unknown /*point*/, OMPD_cancellation_point}, + {OMPD_target, OMPD_unknown /*data*/, OMPD_target_data}, {OMPD_for, OMPD_simd, OMPD_for_simd}, {OMPD_parallel, OMPD_for, OMPD_parallel_for}, {OMPD_parallel_for, OMPD_simd, OMPD_parallel_for_simd}, @@ -41,6 +43,7 @@ Tok.isAnnotation() ? OMPD_unknown : getOpenMPDirectiveKind(P.getPreprocessor().getSpelling(Tok)); + bool TokenMatched = false; for (unsigned i = 0; i < llvm::array_lengthof(F); ++i) { if (!Tok.isAnnotation() && DKind == OMPD_unknown) { @@ -50,18 +53,24 @@ } else { TokenMatched = DKind == F[i][0] && DKind != OMPD_unknown; } + if (TokenMatched) { Tok = P.getPreprocessor().LookAhead(0); + auto TokenIsAnnotation = Tok.isAnnotation(); auto SDKind = - Tok.isAnnotation() + TokenIsAnnotation ? OMPD_unknown : getOpenMPDirectiveKind(P.getPreprocessor().getSpelling(Tok)); - if (!Tok.isAnnotation() && DKind == OMPD_unknown) { + + if (!TokenIsAnnotation && SDKind == OMPD_unknown) { TokenMatched = - (i == 0) && !P.getPreprocessor().getSpelling(Tok).compare("point"); + (i == 0) && + !P.getPreprocessor().getSpelling(Tok).compare("point") || + (i == 1) && !P.getPreprocessor().getSpelling(Tok).compare("data"); } else { TokenMatched = SDKind == F[i][1] && SDKind != OMPD_unknown; } + if (TokenMatched) { P.ConsumeToken(); DKind = F[i][2]; @@ -146,8 +155,8 @@ /// 'section' | 'single' | 'master' | 'critical' [ '(' ')' ] | /// 'parallel for' | 'parallel sections' | 'task' | 'taskyield' | /// 'barrier' | 'taskwait' | 'flush' | 'ordered' | 'atomic' | -/// 'for simd' | 'parallel for simd' | 'target' | 'teams' | 'taskgroup' -/// {clause} +/// 'for simd' | 'parallel for simd' | 'target' | 'target data' | +/// 'taskgroup' | 'teams' {clause} /// annot_pragma_openmp_end /// StmtResult @@ -221,7 +230,8 @@ case OMPD_atomic: case OMPD_target: case OMPD_teams: - case OMPD_taskgroup: { + case OMPD_taskgroup: + case OMPD_target_data: { ConsumeToken(); // Parse directive name of the 'critical' directive if any. if (DKind == OMPD_critical) { @@ -385,7 +395,7 @@ /// lastprivate-clause | reduction-clause | proc_bind-clause | /// schedule-clause | copyin-clause | copyprivate-clause | untied-clause | /// mergeable-clause | flush-clause | read-clause | write-clause | -/// update-clause | capture-clause | seq_cst-clause +/// update-clause | capture-clause | seq_cst-clause /// OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind, OpenMPClauseKind CKind, bool FirstClause) { Index: lib/Sema/SemaOpenMP.cpp =================================================================== --- lib/Sema/SemaOpenMP.cpp +++ lib/Sema/SemaOpenMP.cpp @@ -1300,6 +1300,7 @@ Params); break; } + case OMPD_target_data: case OMPD_target: { Sema::CapturedParamNameType Params[] = { std::make_pair(StringRef(), QualType()) // __context with shared vars @@ -2124,6 +2125,10 @@ "No associated statement allowed for 'omp cancel' directive"); Res = ActOnOpenMPCancelDirective(StartLoc, EndLoc, CancelRegion); break; + case OMPD_target_data: + Res = ActOnOpenMPTargetDataDirective(ClausesWithImplicit, AStmt, StartLoc, + EndLoc); + break; case OMPD_threadprivate: llvm_unreachable("OpenMP Directive is not allowed"); case OMPD_unknown: @@ -4320,6 +4325,16 @@ return OMPTargetDirective::Create(Context, StartLoc, EndLoc, Clauses, AStmt); } +StmtResult Sema::ActOnOpenMPTargetDataDirective(ArrayRef Clauses, + Stmt *AStmt, + SourceLocation StartLoc, + SourceLocation EndLoc) { + getCurFunction()->setHasBranchProtectedScope(); + + return OMPTargetDataDirective::Create(Context, StartLoc, EndLoc, Clauses, + AStmt); +} + StmtResult Sema::ActOnOpenMPTeamsDirective(ArrayRef Clauses, Stmt *AStmt, SourceLocation StartLoc, SourceLocation EndLoc) { @@ -6593,4 +6608,3 @@ return OMPDependClause::Create(Context, StartLoc, LParenLoc, EndLoc, DepKind, DepLoc, ColonLoc, Vars); } - Index: lib/Sema/TreeTransform.h =================================================================== --- lib/Sema/TreeTransform.h +++ lib/Sema/TreeTransform.h @@ -7115,6 +7115,17 @@ } template +StmtResult TreeTransform::TransformOMPTargetDataDirective( + OMPTargetDataDirective *D) { + DeclarationNameInfo DirName; + getDerived().getSema().StartOpenMPDSABlock(OMPD_target_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 @@ -2242,6 +2242,12 @@ VisitOMPExecutableDirective(D); } +void ASTStmtReader::VisitOMPTargetDataDirective(OMPTargetDataDirective *D) { + VisitStmt(D); + ++Idx; + VisitOMPExecutableDirective(D); +} + void ASTStmtReader::VisitOMPTeamsDirective(OMPTeamsDirective *D) { VisitStmt(D); // The NumClauses field was read in ReadStmtFromStream. @@ -2860,6 +2866,11 @@ Context, Record[ASTStmtReader::NumStmtFields], Empty); break; + case STMT_OMP_TARGET_DATA_DIRECTIVE: + S = OMPTargetDataDirective::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 @@ -2063,6 +2063,13 @@ Code = serialization::STMT_OMP_TARGET_DIRECTIVE; } +void ASTStmtWriter::VisitOMPTargetDataDirective(OMPTargetDataDirective *D) { + VisitStmt(D); + Record.push_back(D->getNumClauses()); + VisitOMPExecutableDirective(D); + Code = serialization::STMT_OMP_TARGET_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 @@ -821,6 +821,7 @@ case Stmt::OMPOrderedDirectiveClass: case Stmt::OMPAtomicDirectiveClass: case Stmt::OMPTargetDirectiveClass: + case Stmt::OMPTargetDataDirectiveClass: case Stmt::OMPTeamsDirectiveClass: case Stmt::OMPCancellationPointDirectiveClass: case Stmt::OMPCancelDirectiveClass: Index: test/OpenMP/target_data_ast_print.cpp =================================================================== --- /dev/null +++ test/OpenMP/target_data_ast_print.cpp @@ -0,0 +1,33 @@ +// RUN: %clang_cc1 -verify -fopenmp -ast-print %s | FileCheck %s +// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -std=c++11 -include-pch %t -fsyntax-only -verify %s -ast-print | FileCheck %s +// expected-no-diagnostics + +#ifndef HEADER +#define HEADER + +void foo() {} + +int main (int argc, char **argv) { + int b = argc, c, d, e, f, g; + static int a; +// CHECK: static int a; + +#pragma omp target data +// CHECK: #pragma omp target data + a=2; +// CHECK-NEXT: a = 2; +#pragma omp target data if (b) +// CHECK: #pragma omp target data if(b) + foo(); +// CHECK-NEXT: foo(); + +#pragma omp target data if (b > g) +// CHECK: #pragma omp target data if(b > g) + foo(); +// CHECK-NEXT: foo(); + + return (0); +} + +#endif Index: test/OpenMP/target_data_if_messages.cpp =================================================================== --- /dev/null +++ test/OpenMP/target_data_if_messages.cpp @@ -0,0 +1,26 @@ +// 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) { + #pragma omp target data if // expected-error {{expected '(' after 'if'}} + #pragma omp target data if ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target data if () // expected-error {{expected expression}} + #pragma omp target data if (argc // expected-error {{expected ')'}} expected-note {{to match this '('}} + #pragma omp target data if (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target data' are ignored}} + #pragma omp target data if (argc > 0 ? argv[1] : argv[2]) + #pragma omp target data if (argc + argc) + #pragma omp target data if (foobool(argc)), if (true) // expected-error {{directive '#pragma omp target data' cannot contain more than one 'if' clause}} + #pragma omp target data if (S1) // expected-error {{'S1' does not refer to a value}} + #pragma omp target data if (argv[1]=2) // expected-error {{expected ')'}} expected-note {{to match this '('}} + foo(); + + return 0; +} Index: test/OpenMP/target_data_messages.c =================================================================== --- /dev/null +++ test/OpenMP/target_data_messages.c @@ -0,0 +1,27 @@ +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -ferror-limit 100 -o - %s + +void foo() { } + +int main(int argc, char **argv) { + L1: + foo(); + #pragma omp target data + { + foo(); + goto L1; // expected-error {{use of undeclared label 'L1'}} + } + goto L2; // expected-error {{use of undeclared label 'L2'}} + #pragma omp target data + L2: + foo(); + + #pragma omp target data(i) // expected-warning {{extra tokens at the end of '#pragma omp target data' are ignored}} + { + foo(); + } + #pragma omp target unknown // expected-warning {{extra tokens at the end of '#pragma omp target' are ignored}} + { + foo(); + } + return 0; +} Index: tools/libclang/CIndex.cpp =================================================================== --- tools/libclang/CIndex.cpp +++ tools/libclang/CIndex.cpp @@ -1930,6 +1930,7 @@ void VisitOMPOrderedDirective(const OMPOrderedDirective *D); void VisitOMPAtomicDirective(const OMPAtomicDirective *D); void VisitOMPTargetDirective(const OMPTargetDirective *D); + void VisitOMPTargetDataDirective(const OMPTargetDataDirective *D); void VisitOMPTeamsDirective(const OMPTeamsDirective *D); private: @@ -2549,6 +2550,11 @@ VisitOMPExecutableDirective(D); } +void EnqueueVisitor::VisitOMPTargetDataDirective(const + OMPTargetDataDirective *D) { + VisitOMPExecutableDirective(D); +} + void EnqueueVisitor::VisitOMPTeamsDirective(const OMPTeamsDirective *D) { VisitOMPExecutableDirective(D); } @@ -4349,6 +4355,8 @@ return cxstring::createRef("OMPAtomicDirective"); case CXCursor_OMPTargetDirective: return cxstring::createRef("OMPTargetDirective"); + case CXCursor_OMPTargetDataDirective: + return cxstring::createRef("OMPTargetDataDirective"); case CXCursor_OMPTeamsDirective: return cxstring::createRef("OMPTeamsDirective"); case CXCursor_OMPCancellationPointDirective: Index: tools/libclang/CXCursor.cpp =================================================================== --- tools/libclang/CXCursor.cpp +++ tools/libclang/CXCursor.cpp @@ -585,6 +585,9 @@ case Stmt::OMPTargetDirectiveClass: K = CXCursor_OMPTargetDirective; break; + case Stmt::OMPTargetDataDirectiveClass: + K = CXCursor_OMPTargetDataDirective; + break; case Stmt::OMPTeamsDirectiveClass: K = CXCursor_OMPTeamsDirective; break;