[OPENMP] Parsing and sema support for #pragma omp target data" directive.

[kkwli0]

Summary

Parsing and sema support for #pragma omp target data

[ABataev]

include/clang/AST/OpenMPClause.h

21 ↗

(On Diff #28530)

Header file is not required, remove

131 ↗

(On Diff #28530)

Method is not used, so it must be removed

include/clang/AST/StmtOpenMP.h

1821

Some of your changes are not formatted, use clang-format tool.

include/clang/Sema/Sema.h

7860–7862

Format

lib/CodeGen/CGStmtOpenMP.cpp

2103

I think all target related directives should just emit original code at first, so the user could use them.

lib/Parse/ParseOpenMP.cpp

46–58

This code is copied from clang-omp, we're using another solution in official clang. See the patch for 'declare simd' construct (http://reviews.llvm.org/D10599)

[kkwli0]

lib/Parse/ParseOpenMP.cpp

46–58

Perhaps I misunderstand it. The issue of the suggested approach is that it treats 'data' as a directive in order to have { OMPD_target, OMPD_data, OMPD_target_data }. Creating a fake OMPD_data may be confusing. That is why I took this approach.

[ABataev]

Kelvin, you don't need fake OMPD_data, instead mark it as OMPD_unknown
and check explicitly that 'data' identifier is expected

Best regards,

Alexey Bataev

Software Engineer
Intel Compiler Team

01.07.2015 18:03, Kelvin Li пишет:

Comment at: lib/Parse/ParseOpenMP.cpp:43-55
@@ -42,2 +42,15 @@

: getOpenMPDirectiveKind(P.getPreprocessor().getSpelling(Tok));

+
+ switch (DKind) {
+ case OMPD_target: {
+ auto SavedToken = P.getPreprocessor().LookAhead(0);
+ if (!SavedToken.isAnnotation()) {
+ if (P.getPreprocessor().getSpelling(SavedToken) == "data") {
+ DKind = OMPD_target_data;
+ P.ConsumeToken();
+ }
+ }
+ break;
+ }
+ default:

for (unsigned i = 0; i < llvm::array_lengthof(F); ++i) {

---

ABataev wrote:

This code is copied from clang-omp, we're using another solution in official clang. See the patch for 'declare simd' construct (http://reviews.llvm.org/D10599)

Perhaps I misunderstand it. The issue of the suggested approach is that it treats 'data' as a directive in order to have { OMPD_target, OMPD_data, OMPD_target_data }. Creating a fake OMPD_data may be confusing. That is why I took this approach.

http://reviews.llvm.org/D10765

EMAIL PREFERENCES

http://reviews.llvm.org/settings/panel/emailpreferences/

[kkwli0]

Revised based on comments.

• Fixed format
• Added code not to assert in EmitOMPTargetDataDirective
• Incorporated 'target data' parsing in the current framework

[ABataev]

Seems to me, the updated patch is not correct. It is missing a lot of required things. Besides, new tests were removed from updated patch

include/clang-c/Index.h

2230

Add a cursor for TargetData directive

include/clang/AST/DataRecursiveASTVisitor.h

2407–2410

Add recursive visitor for the new directive

include/clang/AST/RecursiveASTVisitor.h

2439–2442

Add recursive visitor for the new directive

include/clang/Basic/OpenMPKinds.def

95

Add definition of the target data directive

274

There is support for if clause, but macro OPENMP_TARGET_DATA_CLAUSE() is not defined and if clause is not associated with target data directive

include/clang/Basic/StmtNodes.td

206–207

Add a node for target data directive

include/clang/Serialization/ASTBitCodes.h

1399–1400

Add a serialization code for target data directive

lib/CodeGen/CGExpr.cpp

1961 ↗

(On Diff #29206)

Bad solution, restore original code. Instead you should capture variables used in the target data directive

lib/CodeGen/CGStmtOpenMP.cpp

2103–2105

Remove braces

2104–2121

In general, this a bad solution. Look at static void emitOMPIfClause() in CGOpenMPRuntime.cpp for better codegen. Your solution won't work with cleanups and exceptions.

lib/CodeGen/CodeGenFunction.h

2218

This one may be just a static function in CGStmtOpenMP.cpp

lib/Parse/ParseOpenMP.cpp

34

Move it to the beginning of the array.

76

Update your patch to the latest revision and sync it with it

[kkwli0]

• this is the full diff
• sync up the code in ParseOpenMPDirectiveKind
• use the existing infrastructure for generating if clause on target data

Thanks.

[ABataev]

lib/CodeGen/CGExpr.cpp

1963 ↗

(On Diff #29959)

Restore original file

lib/CodeGen/CGOpenMPRuntime.cpp

1099 ↗

(On Diff #29959)

1. If EmitOMPIfClause is now a part of CodeGenFunction, it must be moved to CGStmtOpenmp.cpp.
2. The first argument CGF is not required, so remove it and use '*this' instead.
3. This changes must be committed in a separate patch (for codegen).

lib/CodeGen/CGStmtOpenMP.cpp

2124–2152

As the first commit we can ignore all clauses, just call EmitStmt(cast<CapturedStmt>(S.getAssociateStmt())->getCapturedStmt());

lib/Parse/ParseOpenMP.cpp

65

SDKind == OMPD_unknown

68–69

TokenMatched =
    (i == 0) && !P.getPreprocessor().getSpelling(Tok).compare("point") ||
    (i == 1) && !P.getPreprocessor().getSpelling(Tok).compare("data");

[kkwli0]

lib/CodeGen/CGOpenMPRuntime.cpp

1099 ↗

(On Diff #29959)

Ok. I will leave the changes for the codegen patch.

[kkwli0]

• excluded the codegen code in EmitOMPTargetDataDirective, it currently just emits the code inside the construct
• updated ParseOpenMPDirectiveKind for handling "target data"

[• fraggamuffin]

This has been committed for Kelvin Li.

Command: Commit
Modified: D:\llvmtrunk\tools\clang\include\clang\AST\DataRecursiveASTVisitor.h
Modified: D:\llvmtrunk\tools\clang\include\clang\AST\RecursiveASTVisitor.h
Modified: D:\llvmtrunk\tools\clang\include\clang\AST\StmtOpenMP.h
Modified: D:\llvmtrunk\tools\clang\include\clang\Basic\OpenMPKinds.def
Modified: D:\llvmtrunk\tools\clang\include\clang\Basic\StmtNodes.td
Modified: D:\llvmtrunk\tools\clang\include\clang\Sema\Sema.h
Modified: D:\llvmtrunk\tools\clang\include\clang\Serialization\ASTBitCodes.h
Modified: D:\llvmtrunk\tools\clang\include\clang-c\Index.h
Modified: D:\llvmtrunk\tools\clang\lib\AST\Stmt.cpp
Modified: D:\llvmtrunk\tools\clang\lib\AST\StmtPrinter.cpp
Modified: D:\llvmtrunk\tools\clang\lib\AST\StmtProfile.cpp
Modified: D:\llvmtrunk\tools\clang\lib\Basic\OpenMPKinds.cpp
Modified: D:\llvmtrunk\tools\clang\lib\CodeGen\CGStmt.cpp
Modified: D:\llvmtrunk\tools\clang\lib\CodeGen\CGStmtOpenMP.cpp
Modified: D:\llvmtrunk\tools\clang\lib\CodeGen\CodeGenFunction.h
Modified: D:\llvmtrunk\tools\clang\lib\Parse\ParseOpenMP.cpp
Modified: D:\llvmtrunk\tools\clang\lib\Sema\SemaOpenMP.cpp
Modified: D:\llvmtrunk\tools\clang\lib\Sema\TreeTransform.h
Modified: D:\llvmtrunk\tools\clang\lib\Serialization\ASTReaderStmt.cpp
Modified: D:\llvmtrunk\tools\clang\lib\Serialization\ASTWriterStmt.cpp
Modified: D:\llvmtrunk\tools\clang\lib\StaticAnalyzer\Core\ExprEngine.cpp
Adding: D:\llvmtrunk\tools\clang\test\OpenMP\target_data_ast_print.cpp
Adding: D:\llvmtrunk\tools\clang\test\OpenMP\target_data_if_messages.cpp
Adding: D:\llvmtrunk\tools\clang\test\OpenMP\target_data_messages.c
Modified: D:\llvmtrunk\tools\clang\tools\libclang\CIndex.cpp
Modified: D:\llvmtrunk\tools\clang\tools\libclang\CXCursor.cpp
Sending content: D:\llvmtrunk\tools\clang\lib\AST\StmtPrinter.cpp
Sending content: D:\llvmtrunk\tools\clang\lib\CodeGen\CodeGenFunction.h
Sending content: D:\llvmtrunk\tools\clang\lib\CodeGen\CGStmtOpenMP.cpp
Sending content: D:\llvmtrunk\tools\clang\include\clang\Basic\OpenMPKinds.def
Sending content: D:\llvmtrunk\tools\clang\lib\Sema\TreeTransform.h
Sending content: D:\llvmtrunk\tools\clang\include\clang\AST\DataRecursiveASTVisitor.h
Sending content: D:\llvmtrunk\tools\clang\include\clang\Serialization\ASTBitCodes.h
Sending content: D:\llvmtrunk\tools\clang\include\clang\Basic\StmtNodes.td
Sending content: D:\llvmtrunk\tools\clang\lib\Sema\SemaOpenMP.cpp
Sending content: D:\llvmtrunk\tools\clang\lib\StaticAnalyzer\Core\ExprEngine.cpp
Sending content: D:\llvmtrunk\tools\clang\test\OpenMP\target_data_if_messages.cpp
Sending content: D:\llvmtrunk\tools\clang\include\clang\AST\RecursiveASTVisitor.h
Sending content: D:\llvmtrunk\tools\clang\include\clang-c\Index.h
Sending content: D:\llvmtrunk\tools\clang\lib\AST\StmtProfile.cpp
Sending content: D:\llvmtrunk\tools\clang\lib\Serialization\ASTReaderStmt.cpp
Sending content: D:\llvmtrunk\tools\clang\lib\Parse\ParseOpenMP.cpp
Sending content: D:\llvmtrunk\tools\clang\include\clang\AST\StmtOpenMP.h
Sending content: D:\llvmtrunk\tools\clang\lib\CodeGen\CGStmt.cpp
Sending content: D:\llvmtrunk\tools\clang\lib\Basic\OpenMPKinds.cpp
Sending content: D:\llvmtrunk\tools\clang\lib\Serialization\ASTWriterStmt.cpp
Sending content: D:\llvmtrunk\tools\clang\test\OpenMP\target_data_ast_print.cpp
Sending content: D:\llvmtrunk\tools\clang\test\OpenMP\target_data_messages.c
Sending content: D:\llvmtrunk\tools\clang\tools\libclang\CXCursor.cpp
Sending content: D:\llvmtrunk\tools\clang\lib\AST\Stmt.cpp
Sending content: D:\llvmtrunk\tools\clang\tools\libclang\CIndex.cpp
Sending content: D:\llvmtrunk\tools\clang\include\clang\Sema\Sema.h
Completed: At revision: 242785

include/clang-c/Index.h

enum CXCursorKind {
   CXCursor_OMPTargetDirective            = 252,
 
   /** \brief OpenMP teams directive.
    */
   CXCursor_OMPTeamsDirective             = 253,
 
   /** \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.
    *
    * The translation unit cursor exists primarily to act as the root
    * cursor for traversing the contents of a translation unit.
    */
   CXCursor_TranslationUnit               = 300,

include/clang/AST/DataRecursiveASTVisitor.h

 DEF_TRAVERSE_STMT(OMPOrderedDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
 DEF_TRAVERSE_STMT(OMPAtomicDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
 DEF_TRAVERSE_STMT(OMPTargetDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
+DEF_TRAVERSE_STMT(OMPTargetDataDirective,
+                  { TRY_TO(TraverseOMPExecutableDirective(S)); })
+
 DEF_TRAVERSE_STMT(OMPTeamsDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
 // OpenMP clauses.
 template <typename Derived>
 bool RecursiveASTVisitor<Derived>::TraverseOMPClause(OMPClause *C) {
   if (!C)
     return true;

include/clang/AST/RecursiveASTVisitor.h

 DEF_TRAVERSE_STMT(OMPOrderedDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
 DEF_TRAVERSE_STMT(OMPAtomicDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
 DEF_TRAVERSE_STMT(OMPTargetDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
+DEF_TRAVERSE_STMT(OMPTargetDataDirective,
+                  { TRY_TO(TraverseOMPExecutableDirective(S)); })
+
 DEF_TRAVERSE_STMT(OMPTeamsDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
 // OpenMP clauses.
 template <typename Derived>
 bool RecursiveASTVisitor<Derived>::TraverseOMPClause(OMPClause *C) {
   if (!C)
     return true;

include/clang/AST/StmtOpenMP.h

public:
   static OMPTargetDirective *CreateEmpty(const ASTContext &C,
                                          unsigned NumClauses, EmptyShell);
 
   static bool classof(const Stmt *T) {
     return T->getStmtClass() == OMPTargetDirectiveClass;
   }
 };
 
+/// \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<OMPClause *> 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
 /// #pragma omp teams if(a)
 /// \endcode
 /// In this example directive '#pragma omp teams' has clause 'if' with
 /// condition 'a'.
 ///

include/clang/Basic/OpenMPKinds.def

 #  define OPENMP_TASK_CLAUSE(Name)
 #endif
 #ifndef OPENMP_ATOMIC_CLAUSE
 #  define OPENMP_ATOMIC_CLAUSE(Name)
 #endif
 #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
 #ifndef OPENMP_DEFAULT_KIND
 #  define OPENMP_DEFAULT_KIND(Name)
 #endif
 #ifndef OPENMP_PROC_BIND_KIND
 #  define OPENMP_PROC_BIND_KIND(Name)

 OPENMP_DIRECTIVE(single)
 OPENMP_DIRECTIVE(master)
 OPENMP_DIRECTIVE(critical)
 OPENMP_DIRECTIVE(taskyield)
 OPENMP_DIRECTIVE(barrier)
 OPENMP_DIRECTIVE(taskwait)
 OPENMP_DIRECTIVE(taskgroup)
 OPENMP_DIRECTIVE(flush)
 OPENMP_DIRECTIVE(ordered)
 OPENMP_DIRECTIVE(atomic)
 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")
 OPENMP_DIRECTIVE_EXT(for_simd, "for simd")
 OPENMP_DIRECTIVE_EXT(cancellation_point, "cancellation point")
 
 // OpenMP clauses.
 OPENMP_CLAUSE(if, OMPIfClause)

 OPENMP_TASK_CLAUSE(depend)
 
 // Clauses allowed for OpenMP directive 'atomic'.
 OPENMP_ATOMIC_CLAUSE(read)
 OPENMP_ATOMIC_CLAUSE(write)
 OPENMP_ATOMIC_CLAUSE(update)
 OPENMP_ATOMIC_CLAUSE(capture)
 OPENMP_ATOMIC_CLAUSE(seq_cst)
 
 // Clauses allowed for OpenMP directive 'target'.
 // 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)
 OPENMP_TEAMS_CLAUSE(private)
 OPENMP_TEAMS_CLAUSE(firstprivate)
 OPENMP_TEAMS_CLAUSE(shared)
 OPENMP_TEAMS_CLAUSE(reduction)
 
 #undef OPENMP_DEPEND_KIND
 #undef OPENMP_SCHEDULE_KIND
 #undef OPENMP_PROC_BIND_KIND
 #undef OPENMP_DEFAULT_KIND
 #undef OPENMP_DIRECTIVE
 #undef OPENMP_DIRECTIVE_EXT
 #undef OPENMP_CLAUSE
 #undef OPENMP_SINGLE_CLAUSE
 #undef OPENMP_SECTIONS_CLAUSE
 #undef OPENMP_PARALLEL_CLAUSE
 #undef OPENMP_PARALLEL_FOR_CLAUSE
 #undef OPENMP_PARALLEL_FOR_SIMD_CLAUSE
 #undef OPENMP_PARALLEL_SECTIONS_CLAUSE
 #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
 #undef OPENMP_FOR_SIMD_CLAUSE
 

include/clang/Basic/StmtNodes.td

 def OMPTaskDirective : DStmt<OMPExecutableDirective>;
 def OMPTaskyieldDirective : DStmt<OMPExecutableDirective>;
 def OMPBarrierDirective : DStmt<OMPExecutableDirective>;
 def OMPTaskwaitDirective : DStmt<OMPExecutableDirective>;
 def OMPTaskgroupDirective : DStmt<OMPExecutableDirective>;
 def OMPFlushDirective : DStmt<OMPExecutableDirective>;
 def OMPOrderedDirective : DStmt<OMPExecutableDirective>;
 def OMPAtomicDirective : DStmt<OMPExecutableDirective>;
 def OMPTargetDirective : DStmt<OMPExecutableDirective>;
+def OMPTargetDataDirective : DStmt<OMPExecutableDirective>;
 def OMPTeamsDirective : DStmt<OMPExecutableDirective>;
 def OMPCancellationPointDirective : DStmt<OMPExecutableDirective>;
 def OMPCancelDirective : DStmt<OMPExecutableDirective>;

include/clang/Sema/Sema.h

public:
   StmtResult ActOnOpenMPAtomicDirective(ArrayRef<OMPClause *> Clauses,
                                         Stmt *AStmt, SourceLocation StartLoc,
                                         SourceLocation EndLoc);
   /// \brief Called on well-formed '\#pragma omp target' after parsing of the
   /// associated statement.
   StmtResult ActOnOpenMPTargetDirective(ArrayRef<OMPClause *> 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<OMPClause *> Clauses,
+                                            Stmt *AStmt, SourceLocation StartLoc,
+                                            SourceLocation EndLoc);
   /// \brief Called on well-formed '\#pragma omp teams' after parsing of the
   /// associated statement.
   StmtResult ActOnOpenMPTeamsDirective(ArrayRef<OMPClause *> Clauses,
                                        Stmt *AStmt, SourceLocation StartLoc,
                                        SourceLocation EndLoc);
   /// \brief Called on well-formed '\#pragma omp cancellation point'.
   StmtResult
   ActOnOpenMPCancellationPointDirective(SourceLocation StartLoc,

include/clang/Serialization/ASTBitCodes.h

enum StmtCode {
       STMT_OMP_PARALLEL_SECTIONS_DIRECTIVE,
       STMT_OMP_TASK_DIRECTIVE,
       STMT_OMP_TASKYIELD_DIRECTIVE,
       STMT_OMP_BARRIER_DIRECTIVE,
       STMT_OMP_TASKWAIT_DIRECTIVE,
       STMT_OMP_FLUSH_DIRECTIVE,
       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,
       STMT_OMP_CANCEL_DIRECTIVE,
 
       // ARC
       EXPR_OBJC_BRIDGED_CAST,     // ObjCBridgedCastExpr
       

lib/AST/Stmt.cpp

OMPTargetDirective *OMPTargetDirective::CreateEmpty(const ASTContext &C,
                                                     EmptyShell) {
   unsigned Size = llvm::RoundUpToAlignment(sizeof(OMPTargetDirective),
                                            llvm::alignOf<OMPClause *>());
   void *Mem =
       C.Allocate(Size + sizeof(OMPClause *) * NumClauses + sizeof(Stmt *));
   return new (Mem) OMPTargetDirective(NumClauses);
 }
 
+OMPTargetDataDirective *OMPTargetDataDirective::Create(
+    const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
+    ArrayRef<OMPClause *> Clauses, Stmt *AssociatedStmt) {
+  void *Mem =
+      C.Allocate(llvm::RoundUpToAlignment(sizeof(OMPTargetDataDirective),
+                                          llvm::alignOf<OMPClause *>()) +
+                 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<OMPClause *>()) +
+                 sizeof(OMPClause *) * N + sizeof(Stmt *));
+  return new (Mem) OMPTargetDataDirective(N);
+}
+
 OMPTeamsDirective *OMPTeamsDirective::Create(const ASTContext &C,
                                              SourceLocation StartLoc,
                                              SourceLocation EndLoc,
                                              ArrayRef<OMPClause *> Clauses,
                                              Stmt *AssociatedStmt) {
   unsigned Size = llvm::RoundUpToAlignment(sizeof(OMPTeamsDirective),
                                            llvm::alignOf<OMPClause *>());
   void *Mem =

lib/AST/StmtPrinter.cpp

void StmtPrinter::VisitOMPAtomicDirective(OMPAtomicDirective *Node) {
   PrintOMPExecutableDirective(Node);
 }
 
 void StmtPrinter::VisitOMPTargetDirective(OMPTargetDirective *Node) {
   Indent() << "#pragma omp target ";
   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);
 }
 
 void StmtPrinter::VisitOMPCancellationPointDirective(
     OMPCancellationPointDirective *Node) {
   Indent() << "#pragma omp cancellation point "

lib/AST/StmtProfile.cpp

 void StmtProfiler::VisitOMPAtomicDirective(const OMPAtomicDirective *S) {
   VisitOMPExecutableDirective(S);
 }
 
 void StmtProfiler::VisitOMPTargetDirective(const OMPTargetDirective *S) {
   VisitOMPExecutableDirective(S);
 }
 
+void StmtProfiler::VisitOMPTargetDataDirective(const OMPTargetDataDirective *S) {
+  VisitOMPExecutableDirective(S);
+}
+
 void StmtProfiler::VisitOMPTeamsDirective(const OMPTeamsDirective *S) {
   VisitOMPExecutableDirective(S);
 }
 
 void StmtProfiler::VisitOMPCancellationPointDirective(
     const OMPCancellationPointDirective *S) {
   VisitOMPExecutableDirective(S);
 }

lib/Basic/OpenMPKinds.cpp

 #define OPENMP_TARGET_CLAUSE(Name)                                             \
   case OMPC_##Name:                                                            \
     return true;
 #include "clang/Basic/OpenMPKinds.def"
     default:
       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)                                              \
   case OMPC_##Name:                                                            \
     return true;
 #include "clang/Basic/OpenMPKinds.def"
     default:
       break;

lib/CodeGen/CGStmt.cpp

case Stmt::OMPTeamsDirectiveClass:
     EmitOMPTeamsDirective(cast<OMPTeamsDirective>(*S));
     break;
   case Stmt::OMPCancellationPointDirectiveClass:
     EmitOMPCancellationPointDirective(cast<OMPCancellationPointDirective>(*S));
     break;
   case Stmt::OMPCancelDirectiveClass:
     EmitOMPCancelDirective(cast<OMPCancelDirective>(*S));
     break;
+  case Stmt::OMPTargetDataDirectiveClass:
+    EmitOMPTargetDataDirective(cast<OMPTargetDataDirective>(*S));
+    break;
   }
 }
 
 bool CodeGenFunction::EmitSimpleStmt(const Stmt *S) {
   switch (S->getStmtClass()) {
   default: return false;
   case Stmt::NullStmtClass: break;
   case Stmt::CompoundStmtClass: EmitCompoundStmt(cast<CompoundStmt>(*S)); break;

lib/CodeGen/CGStmtOpenMP.cpp

 
 void CodeGenFunction::EmitOMPTargetDirective(const OMPTargetDirective &) {
   llvm_unreachable("CodeGen for 'omp target' is not supported yet.");
 }
 
 void CodeGenFunction::EmitOMPTeamsDirective(const OMPTeamsDirective &) {
   llvm_unreachable("CodeGen for 'omp teams' is not supported yet.");
 }
 
 void CodeGenFunction::EmitOMPCancellationPointDirective(
     const OMPCancellationPointDirective &S) {
   CGM.getOpenMPRuntime().emitCancellationPointCall(*this, S.getLocStart(),
                                                    S.getCancelRegion());
 }
 
 void CodeGenFunction::EmitOMPCancelDirective(const OMPCancelDirective &S) {
   CGM.getOpenMPRuntime().emitCancelCall(*this, S.getLocStart(),
                                         S.getCancelRegion());
 }
 
 CodeGenFunction::JumpDest
 CodeGenFunction::getOMPCancelDestination(OpenMPDirectiveKind Kind) {
   if (Kind == OMPD_parallel || Kind == OMPD_task)
     return ReturnBlock;
   else if (Kind == OMPD_for || Kind == OMPD_section || Kind == OMPD_sections)
     return BreakContinueStack.empty() ? JumpDest()
                                       : 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<CapturedStmt>(S.getAssociatedStmt());
+  EmitStmt(CS->getCapturedStmt());
+}

lib/CodeGen/CodeGenFunction.h

#endif
   void EmitOMPTaskyieldDirective(const OMPTaskyieldDirective &S);
   void EmitOMPBarrierDirective(const OMPBarrierDirective &S);
   void EmitOMPTaskwaitDirective(const OMPTaskwaitDirective &S);
   void EmitOMPTaskgroupDirective(const OMPTaskgroupDirective &S);
   void EmitOMPFlushDirective(const OMPFlushDirective &S);
   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);
   void EmitOMPCancelDirective(const OMPCancelDirective &S);
 
   /// \brief Emit inner loop of the worksharing/simd construct.
   ///
   /// \param S Directive, for which the inner loop must be emitted.
   /// \param RequiresCleanup true, if directive has some associated private
   /// variables.

lib/Parse/ParseOpenMP.cpp

 #include "RAIIObjectsForParser.h"
 #include "clang/AST/ASTConsumer.h"
 #include "clang/AST/ASTContext.h"
 #include "clang/AST/StmtOpenMP.h"
 #include "clang/Parse/ParseDiagnostic.h"
 #include "clang/Parse/Parser.h"
 #include "clang/Sema/Scope.h"
 #include "llvm/ADT/PointerIntPair.h"
+
 using namespace clang;
 
 //===----------------------------------------------------------------------===//
 // OpenMP declarative directives.
 //===----------------------------------------------------------------------===//
 
 static OpenMPDirectiveKind ParseOpenMPDirectiveKind(Parser &P) {
   // Array of foldings: F[i][0] F[i][1] ===> F[i][2].
   // E.g.: OMPD_for OMPD_simd ===> OMPD_for_simd
   // TODO: add other combined directives in topological order.
   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},
       {OMPD_parallel, OMPD_sections, OMPD_parallel_sections}};
   auto Tok = P.getCurToken();
   auto DKind =
       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) {
       TokenMatched =
           (i == 0) &&
           !P.getPreprocessor().getSpelling(Tok).compare("cancellation");
     } 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];
       }
     }
   }
   return DKind;
 }
 
 /// \brief Parsing of declarative OpenMP directives.
 ///

 ///         annot_pragma_openmp 'threadprivate' simple-variable-list
 ///         annot_pragma_openmp_end
 ///
 ///       executable-directive:
 ///         annot_pragma_openmp 'parallel' | 'simd' | 'for' | 'sections' |
 ///         'section' | 'single' | 'master' | 'critical' [ '(' <name> ')' ] |
 ///         '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
 Parser::ParseOpenMPDeclarativeOrExecutableDirective(bool StandAloneAllowed) {
   assert(Tok.is(tok::annot_pragma_openmp) && "Not an OpenMP directive!");
   ParenBraceBracketBalancer BalancerRAIIObj(*this);
   SmallVector<Expr *, 5> Identifiers;
   SmallVector<OMPClause *, 5> Clauses;
Parser::ParseOpenMPDeclarativeOrExecutableDirective(bool StandAloneAllowed) {
   case OMPD_parallel_for:
   case OMPD_parallel_for_simd:
   case OMPD_parallel_sections:
   case OMPD_task:
   case OMPD_ordered:
   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) {
       BalancedDelimiterTracker T(*this, tok::l_paren,
                                  tok::annot_pragma_openmp_end);
       if (!T.consumeOpen()) {
         if (Tok.isAnyIdentifier()) {
           DirName =

 ///
 ///    clause:
 ///       if-clause | final-clause | num_threads-clause | safelen-clause |
 ///       default-clause | private-clause | firstprivate-clause | shared-clause
 ///       | linear-clause | aligned-clause | collapse-clause |
 ///       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) {
   OMPClause *Clause = nullptr;
   bool ErrorFound = false;
   // Check if clause is allowed for the given directive.
   if (CKind != OMPC_unknown && !isAllowedClauseForDirective(DKind, CKind)) {
     Diag(Tok, diag::err_omp_unexpected_clause) << getOpenMPClauseName(CKind)

lib/Sema/SemaOpenMP.cpp

void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) {
   case OMPD_atomic: {
     Sema::CapturedParamNameType Params[] = {
         std::make_pair(StringRef(), QualType()) // __context with shared vars
     };
     ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
                              Params);
     break;
   }
+  case OMPD_target_data:
   case OMPD_target: {
     Sema::CapturedParamNameType Params[] = {
         std::make_pair(StringRef(), QualType()) // __context with shared vars
     };
     ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
                              Params);
     break;
   }
case OMPD_cancellation_point:
     break;
   case OMPD_cancel:
     assert(ClausesWithImplicit.empty() &&
            "No clauses are allowed for 'omp cancel' directive");
     assert(AStmt == nullptr &&
            "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:
     llvm_unreachable("Unknown OpenMP directive");
   }
 
   for (auto P : VarsWithInheritedDSA) {
     Diag(P.second->getExprLoc(), diag::err_omp_no_dsa_for_variable)
if (DSAStack->hasInnerTeamsRegion()) {
     }
   }
 
   getCurFunction()->setHasBranchProtectedScope();
 
   return OMPTargetDirective::Create(Context, StartLoc, EndLoc, Clauses, AStmt);
 }
 
+StmtResult Sema::ActOnOpenMPTargetDataDirective(ArrayRef<OMPClause *> Clauses,
+                                                Stmt *AStmt,
+                                                SourceLocation StartLoc,
+                                                SourceLocation EndLoc) {
+  getCurFunction()->setHasBranchProtectedScope();
+
+  return OMPTargetDataDirective::Create(Context, StartLoc, EndLoc, Clauses,
+                                        AStmt);
+}
+
 StmtResult Sema::ActOnOpenMPTeamsDirective(ArrayRef<OMPClause *> Clauses,
                                            Stmt *AStmt, SourceLocation StartLoc,
                                            SourceLocation EndLoc) {
   assert(AStmt && isa<CapturedStmt>(AStmt) && "Captured statement expected");
   CapturedStmt *CS = cast<CapturedStmt>(AStmt);
   // 1.2.2 OpenMP Language Terminology
   // Structured block - An executable statement with a single entry at the
   // top and a single exit at the bottom.
Sema::ActOnOpenMPDependClause(OpenMPDependClauseKind DepKind,
   }
 
   if (Vars.empty())
     return nullptr;
 
   return OMPDependClause::Create(Context, StartLoc, LParenLoc, EndLoc, DepKind,
                                  DepLoc, ColonLoc, Vars);
 }
-

lib/Sema/TreeTransform.h

TreeTransform<Derived>::TransformOMPTargetDirective(OMPTargetDirective *D) {
   getDerived().getSema().StartOpenMPDSABlock(OMPD_target, DirName, nullptr,
                                              D->getLocStart());
   StmtResult Res = getDerived().TransformOMPExecutableDirective(D);
   getDerived().getSema().EndOpenMPDSABlock(Res.get());
   return Res;
 }
 
 template <typename Derived>
+StmtResult TreeTransform<Derived>::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 <typename Derived>
 StmtResult
 TreeTransform<Derived>::TransformOMPTeamsDirective(OMPTeamsDirective *D) {
   DeclarationNameInfo DirName;
   getDerived().getSema().StartOpenMPDSABlock(OMPD_teams, DirName, nullptr,
                                              D->getLocStart());
   StmtResult Res = getDerived().TransformOMPExecutableDirective(D);
   getDerived().getSema().EndOpenMPDSABlock(Res.get());
   return Res;

lib/Serialization/ASTReaderStmt.cpp

 
 void ASTStmtReader::VisitOMPTargetDirective(OMPTargetDirective *D) {
   VisitStmt(D);
   // The NumClauses field was read in ReadStmtFromStream.
   ++Idx;
   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.
   ++Idx;
   VisitOMPExecutableDirective(D);
 }
 
 void ASTStmtReader::VisitOMPCancellationPointDirective(
case STMT_OMP_ATOMIC_DIRECTIVE:
           Context, Record[ASTStmtReader::NumStmtFields], Empty);
       break;
 
     case STMT_OMP_TARGET_DIRECTIVE:
       S = OMPTargetDirective::CreateEmpty(
           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);
       break;
 
     case STMT_OMP_CANCELLATION_POINT_DIRECTIVE:
       S = OMPCancellationPointDirective::CreateEmpty(Context, Empty);
       break;

lib/Serialization/ASTWriterStmt.cpp

 
 void ASTStmtWriter::VisitOMPTargetDirective(OMPTargetDirective *D) {
   VisitStmt(D);
   Record.push_back(D->getNumClauses());
   VisitOMPExecutableDirective(D);
   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);
   Code = serialization::STMT_OMP_TASKYIELD_DIRECTIVE;
 }
 
 void ASTStmtWriter::VisitOMPBarrierDirective(OMPBarrierDirective *D) {
   VisitStmt(D);

lib/StaticAnalyzer/Core/ExprEngine.cpp

switch (S->getStmtClass()) {
     case Stmt::OMPTaskyieldDirectiveClass:
     case Stmt::OMPBarrierDirectiveClass:
     case Stmt::OMPTaskwaitDirectiveClass:
     case Stmt::OMPTaskgroupDirectiveClass:
     case Stmt::OMPFlushDirectiveClass:
     case Stmt::OMPOrderedDirectiveClass:
     case Stmt::OMPAtomicDirectiveClass:
     case Stmt::OMPTargetDirectiveClass:
+    case Stmt::OMPTargetDataDirectiveClass:
     case Stmt::OMPTeamsDirectiveClass:
     case Stmt::OMPCancellationPointDirectiveClass:
     case Stmt::OMPCancelDirectiveClass:
       llvm_unreachable("Stmt should not be in analyzer evaluation loop");
 
     case Stmt::ObjCSubscriptRefExprClass:
     case Stmt::ObjCPropertyRefExprClass:
       llvm_unreachable("These are handled by PseudoObjectExpr");

test/OpenMP/target_data_ast_print.cpp

+// 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

test/OpenMP/target_data_if_messages.cpp

+// 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;
+}

test/OpenMP/target_data_messages.c

+// 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;
+}

tools/libclang/CIndex.cpp

public:
   void VisitOMPTaskgroupDirective(const OMPTaskgroupDirective *D);
   void
   VisitOMPCancellationPointDirective(const OMPCancellationPointDirective *D);
   void VisitOMPCancelDirective(const OMPCancelDirective *D);
   void VisitOMPFlushDirective(const OMPFlushDirective *D);
   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:
   void AddDeclarationNameInfo(const Stmt *S);
   void AddNestedNameSpecifierLoc(NestedNameSpecifierLoc Qualifier);
   void AddExplicitTemplateArgs(const ASTTemplateArgumentListInfo *A);
   void AddMemberRef(const FieldDecl *D, SourceLocation L);
   void AddStmt(const Stmt *S);

 void EnqueueVisitor::VisitOMPAtomicDirective(const OMPAtomicDirective *D) {
   VisitOMPExecutableDirective(D);
 }
 
 void EnqueueVisitor::VisitOMPTargetDirective(const OMPTargetDirective *D) {
   VisitOMPExecutableDirective(D);
 }
 
+void EnqueueVisitor::VisitOMPTargetDataDirective(const 
+                                                 OMPTargetDataDirective *D) {
+  VisitOMPExecutableDirective(D);
+}
+
 void EnqueueVisitor::VisitOMPTeamsDirective(const OMPTeamsDirective *D) {
   VisitOMPExecutableDirective(D);
 }
 
 void EnqueueVisitor::VisitOMPCancellationPointDirective(
     const OMPCancellationPointDirective *D) {
   VisitOMPExecutableDirective(D);
 }
CXString clang_getCursorKindSpelling(enum CXCursorKind Kind) {
   case CXCursor_OMPFlushDirective:
     return cxstring::createRef("OMPFlushDirective");
   case CXCursor_OMPOrderedDirective:
     return cxstring::createRef("OMPOrderedDirective");
   case CXCursor_OMPAtomicDirective:
     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:
     return cxstring::createRef("OMPCancellationPointDirective");
   case CXCursor_OMPCancelDirective:
     return cxstring::createRef("OMPCancelDirective");
   case CXCursor_OverloadCandidate:
       return cxstring::createRef("OverloadCandidate");

tools/libclang/CXCursor.cpp

case Stmt::OMPOrderedDirectiveClass:
     K = CXCursor_OMPOrderedDirective;
     break;
   case Stmt::OMPAtomicDirectiveClass:
     K = CXCursor_OMPAtomicDirective;
     break;
   case Stmt::OMPTargetDirectiveClass:
     K = CXCursor_OMPTargetDirective;
     break;
+  case Stmt::OMPTargetDataDirectiveClass:
+    K = CXCursor_OMPTargetDataDirective;
+    break;
   case Stmt::OMPTeamsDirectiveClass:
     K = CXCursor_OMPTeamsDirective;
     break;
   case Stmt::OMPCancellationPointDirectiveClass:
     K = CXCursor_OMPCancellationPointDirective;
     break;
   case Stmt::OMPCancelDirectiveClass:
     K = CXCursor_OMPCancelDirective;