[OpenMP] Sema and parsing for 'target teams' pragma

[kkwli0]

Summary

This patch is to add sema and parsing support for 'target teams' pragma.

[ABataev]

LG

include/clang-c/Index.h

enum CXCursorKind {
   /** \brief OpenMP teams distribute parallel for simd directive.
    */
   CXCursor_OMPTeamsDistributeParallelForSimdDirective = 273,
 
   /** \brief OpenMP teams distribute parallel for directive.
    */
   CXCursor_OMPTeamsDistributeParallelForDirective = 274,
 
-  CXCursor_LastStmt = CXCursor_OMPTeamsDistributeParallelForDirective,
+  /** \brief OpenMP target teams directive.
+   */
+  CXCursor_OMPTargetTeamsDirective = 275,
+
+  CXCursor_LastStmt = CXCursor_OMPTargetTeamsDirective,
 
   /**
    * \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/RecursiveASTVisitor.h

DEF_TRAVERSE_STMT(OMPTeamsDistributeSimdDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
 DEF_TRAVERSE_STMT(OMPTeamsDistributeParallelForSimdDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
 DEF_TRAVERSE_STMT(OMPTeamsDistributeParallelForDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
+DEF_TRAVERSE_STMT(OMPTargetTeamsDirective,
+                  { TRY_TO(TraverseOMPExecutableDirective(S)); })
+
 // OpenMP clauses.
 template <typename Derived>
 bool RecursiveASTVisitor<Derived>::TraverseOMPClause(OMPClause *C) {
   if (!C)
     return true;
   switch (C->getClauseKind()) {
 #define OPENMP_CLAUSE(Name, Class)                                             \
   case OMPC_##Name:                                                            \

include/clang/AST/StmtOpenMP.h

public:
   CreateEmpty(const ASTContext &C, unsigned NumClauses, unsigned CollapsedNum,
               EmptyShell);
 
   static bool classof(const Stmt *T) {
     return T->getStmtClass() == OMPTeamsDistributeParallelForDirectiveClass;
   }
 };
 
+/// This represents '#pragma omp target teams' directive.
+///
+/// \code
+/// #pragma omp target teams if(a>0)
+/// \endcode
+/// In this example directive '#pragma omp target teams' has clause 'if' with
+/// condition 'a>0'.
+///
+class OMPTargetTeamsDirective final : public OMPExecutableDirective {
+  friend class ASTStmtReader;
+  /// 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 Number of clauses.
+  ///
+  OMPTargetTeamsDirective(SourceLocation StartLoc, SourceLocation EndLoc,
+                          unsigned NumClauses)
+      : OMPExecutableDirective(this, OMPTargetTeamsDirectiveClass,
+                               OMPD_target_teams, StartLoc, EndLoc, NumClauses,
+                               1) {}
+
+  /// Build an empty directive.
+  ///
+  /// \param NumClauses Number of clauses.
+  ///
+  explicit OMPTargetTeamsDirective(unsigned NumClauses)
+      : OMPExecutableDirective(this, OMPTargetTeamsDirectiveClass,
+                               OMPD_target_teams, SourceLocation(),
+                               SourceLocation(), NumClauses, 1) {}
+
+public:
+  /// 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 OMPTargetTeamsDirective *Create(const ASTContext &C,
+                                         SourceLocation StartLoc,
+                                         SourceLocation EndLoc,
+                                         ArrayRef<OMPClause *> Clauses,
+                                         Stmt *AssociatedStmt);
+
+  /// Creates an empty directive with the place for \a NumClauses clauses.
+  ///
+  /// \param C AST context.
+  /// \param NumClauses Number of clauses.
+  ///
+  static OMPTargetTeamsDirective *CreateEmpty(const ASTContext &C,
+                                              unsigned NumClauses, EmptyShell);
+
+  static bool classof(const Stmt *T) {
+    return T->getStmtClass() == OMPTargetTeamsDirectiveClass;
+  }
+};
+
 } // end namespace clang
 
 #endif

include/clang/Basic/OpenMPKinds.h

 /// \brief Checks if the specified directive is a target data offload directive.
 /// \param DKind Specified directive.
 /// \return true - the directive is a target data offload directive like
 /// 'omp target data', 'omp target update', 'omp target enter data',
 /// 'omp target exit data'
 /// otherwise - false.
 bool isOpenMPTargetDataManagementDirective(OpenMPDirectiveKind DKind);
 
-/// \brief Checks if the specified directive is a teams-kind directive.
+/// Checks if the specified composite/combined directive constitutes a teams
+/// directive in the outermost nest.  For example
+/// 'omp teams distribute' or 'omp teams distribute parallel for'.
+/// \param DKind Specified directive.
+/// \return true - the directive has teams on the outermost nest, otherwise -
+/// false.
+bool isOpenMPNestingTeamsDirective(OpenMPDirectiveKind DKind);
+
+/// Checks if the specified directive is a teams-kind directive.  For example,
+/// 'omp teams distribute' or 'omp target teams'.
 /// \param DKind Specified directive.
-/// \return true - the directive is a teams-like directive like 'omp teams',
-/// otherwise - false.
+/// \return true - the directive is a teams-like directive, otherwise - false.
 bool isOpenMPTeamsDirective(OpenMPDirectiveKind DKind);
 
 /// \brief Checks if the specified directive is a simd directive.
 /// \param DKind Specified directive.
 /// \return true - the directive is a simd directive like 'omp simd',
 /// otherwise - false.
 bool isOpenMPSimdDirective(OpenMPDirectiveKind DKind);
 

include/clang/Basic/OpenMPKinds.def

 #define OPENMP_TEAMS_DISTRIBUTE_SIMD_CLAUSE(Name)
 #endif
 #ifndef OPENMP_TEAMS_DISTRIBUTE_PARALLEL_FOR_SIMD_CLAUSE
 #define OPENMP_TEAMS_DISTRIBUTE_PARALLEL_FOR_SIMD_CLAUSE(Name)
 #endif
 #ifndef OPENMP_TEAMS_DISTRIBUTE_PARALLEL_FOR_CLAUSE
 #define OPENMP_TEAMS_DISTRIBUTE_PARALLEL_FOR_CLAUSE(Name)
 #endif
+#ifndef OPENMP_TARGET_TEAMS_CLAUSE
+#define OPENMP_TARGET_TEAMS_CLAUSE(Name)
+#endif
 
 // OpenMP directives.
 OPENMP_DIRECTIVE(threadprivate)
 OPENMP_DIRECTIVE(parallel)
 OPENMP_DIRECTIVE(task)
 OPENMP_DIRECTIVE(simd)
 OPENMP_DIRECTIVE(for)
 OPENMP_DIRECTIVE(sections)

 OPENMP_DIRECTIVE_EXT(distribute_parallel_for_simd, "distribute parallel for simd")
 OPENMP_DIRECTIVE_EXT(distribute_simd, "distribute simd")
 OPENMP_DIRECTIVE_EXT(target_parallel_for_simd, "target parallel for simd")
 OPENMP_DIRECTIVE_EXT(target_simd, "target simd")
 OPENMP_DIRECTIVE_EXT(teams_distribute, "teams distribute")
 OPENMP_DIRECTIVE_EXT(teams_distribute_simd, "teams distribute simd")
 OPENMP_DIRECTIVE_EXT(teams_distribute_parallel_for_simd, "teams distribute parallel for simd")
 OPENMP_DIRECTIVE_EXT(teams_distribute_parallel_for, "teams distribute parallel for")
+OPENMP_DIRECTIVE_EXT(target_teams, "target teams")
 
 // OpenMP clauses.
 OPENMP_CLAUSE(if, OMPIfClause)
 OPENMP_CLAUSE(final, OMPFinalClause)
 OPENMP_CLAUSE(num_threads, OMPNumThreadsClause)
 OPENMP_CLAUSE(safelen, OMPSafelenClause)
 OPENMP_CLAUSE(simdlen, OMPSimdlenClause)
 OPENMP_CLAUSE(collapse, OMPCollapseClause)

 OPENMP_TEAMS_DISTRIBUTE_PARALLEL_FOR_CLAUSE(private)
 OPENMP_TEAMS_DISTRIBUTE_PARALLEL_FOR_CLAUSE(shared)
 OPENMP_TEAMS_DISTRIBUTE_PARALLEL_FOR_CLAUSE(reduction)
 OPENMP_TEAMS_DISTRIBUTE_PARALLEL_FOR_CLAUSE(schedule)
 OPENMP_TEAMS_DISTRIBUTE_PARALLEL_FOR_CLAUSE(linear)
 OPENMP_TEAMS_DISTRIBUTE_PARALLEL_FOR_CLAUSE(num_teams)
 OPENMP_TEAMS_DISTRIBUTE_PARALLEL_FOR_CLAUSE(thread_limit)
 
+// Clauses allowed for OpenMP directive 'target teams'.
+OPENMP_TARGET_TEAMS_CLAUSE(if)
+OPENMP_TARGET_TEAMS_CLAUSE(device)
+OPENMP_TARGET_TEAMS_CLAUSE(map)
+OPENMP_TARGET_TEAMS_CLAUSE(private)
+OPENMP_TARGET_TEAMS_CLAUSE(nowait)
+OPENMP_TARGET_TEAMS_CLAUSE(depend)
+OPENMP_TARGET_TEAMS_CLAUSE(defaultmap)
+OPENMP_TARGET_TEAMS_CLAUSE(firstprivate)
+OPENMP_TARGET_TEAMS_CLAUSE(is_device_ptr)
+OPENMP_TARGET_TEAMS_CLAUSE(default)
+OPENMP_TARGET_TEAMS_CLAUSE(shared)
+OPENMP_TARGET_TEAMS_CLAUSE(reduction)
+OPENMP_TARGET_TEAMS_CLAUSE(num_teams)
+OPENMP_TARGET_TEAMS_CLAUSE(thread_limit)
+
 #undef OPENMP_TASKLOOP_SIMD_CLAUSE
 #undef OPENMP_TASKLOOP_CLAUSE
 #undef OPENMP_LINEAR_KIND
 #undef OPENMP_DEPEND_KIND
 #undef OPENMP_SCHEDULE_MODIFIER
 #undef OPENMP_SCHEDULE_KIND
 #undef OPENMP_PROC_BIND_KIND
 #undef OPENMP_DEFAULT_KIND

 #undef OPENMP_DISTRIBUTE_PARALLEL_FOR_SIMD_CLAUSE
 #undef OPENMP_DISTRIBUTE_SIMD_CLAUSE
 #undef OPENMP_TARGET_PARALLEL_FOR_SIMD_CLAUSE
 #undef OPENMP_TARGET_SIMD_CLAUSE
 #undef OPENMP_TEAMS_DISTRIBUTE_CLAUSE
 #undef OPENMP_TEAMS_DISTRIBUTE_SIMD_CLAUSE
 #undef OPENMP_TEAMS_DISTRIBUTE_PARALLEL_FOR_SIMD_CLAUSE
 #undef OPENMP_TEAMS_DISTRIBUTE_PARALLEL_FOR_CLAUSE
+#undef OPENMP_TARGET_TEAMS_CLAUSE

include/clang/Basic/StmtNodes.td

 def OMPDistributeParallelForSimdDirective : DStmt<OMPLoopDirective>;
 def OMPDistributeSimdDirective : DStmt<OMPLoopDirective>;
 def OMPTargetParallelForSimdDirective : DStmt<OMPLoopDirective>;
 def OMPTargetSimdDirective : DStmt<OMPLoopDirective>;
 def OMPTeamsDistributeDirective : DStmt<OMPLoopDirective>;
 def OMPTeamsDistributeSimdDirective : DStmt<OMPLoopDirective>;
 def OMPTeamsDistributeParallelForSimdDirective : DStmt<OMPLoopDirective>;
 def OMPTeamsDistributeParallelForDirective : DStmt<OMPLoopDirective>;
+def OMPTargetTeamsDirective : DStmt<OMPExecutableDirective>;

include/clang/Sema/Sema.h

StmtResult ActOnOpenMPTeamsDistributeParallelForSimdDirective(
       SourceLocation EndLoc,
       llvm::DenseMap<ValueDecl *, Expr *> &VarsWithImplicitDSA);
   /// Called on well-formed '\#pragma omp teams distribute parallel for'
   /// after parsing of the associated statement.
   StmtResult ActOnOpenMPTeamsDistributeParallelForDirective(
       ArrayRef<OMPClause *> Clauses, Stmt *AStmt, SourceLocation StartLoc,
       SourceLocation EndLoc,
       llvm::DenseMap<ValueDecl *, Expr *> &VarsWithImplicitDSA);
+  /// Called on well-formed '\#pragma omp target teams' after parsing of the
+  /// associated statement.
+  StmtResult ActOnOpenMPTargetTeamsDirective(ArrayRef<OMPClause *> Clauses,
+                                             Stmt *AStmt,
+                                             SourceLocation StartLoc,
+                                             SourceLocation EndLoc);
 
   /// Checks correctness of linear modifiers.
   bool CheckOpenMPLinearModifier(OpenMPLinearClauseKind LinKind,
                                  SourceLocation LinLoc);
   /// Checks that the specified declaration matches requirements for the linear
   /// decls.
   bool CheckOpenMPLinearDecl(ValueDecl *D, SourceLocation ELoc,
                              OpenMPLinearClauseKind LinKind, QualType Type);

include/clang/Serialization/ASTBitCodes.h

enum StmtCode {
       STMT_OMP_DISTRIBUTE_PARALLEL_FOR_SIMD_DIRECTIVE,
       STMT_OMP_DISTRIBUTE_SIMD_DIRECTIVE,
       STMT_OMP_TARGET_PARALLEL_FOR_SIMD_DIRECTIVE,
       STMT_OMP_TARGET_SIMD_DIRECTIVE,
       STMT_OMP_TEAMS_DISTRIBUTE_DIRECTIVE,
       STMT_OMP_TEAMS_DISTRIBUTE_SIMD_DIRECTIVE,
       STMT_OMP_TEAMS_DISTRIBUTE_PARALLEL_FOR_SIMD_DIRECTIVE,
       STMT_OMP_TEAMS_DISTRIBUTE_PARALLEL_FOR_DIRECTIVE,
+      STMT_OMP_TARGET_TEAMS_DIRECTIVE,
       EXPR_OMP_ARRAY_SECTION,
 
       // ARC
       EXPR_OBJC_BRIDGED_CAST,     // ObjCBridgedCastExpr
       
       STMT_MS_DEPENDENT_EXISTS,   // MSDependentExistsStmt
       EXPR_LAMBDA                 // LambdaExpr
     };

lib/AST/StmtOpenMP.cpp

OMPTeamsDistributeParallelForDirective::CreateEmpty(const ASTContext &C,
   void *Mem = C.Allocate(
       Size + sizeof(OMPClause *) * NumClauses +
       sizeof(Stmt *) *
           numLoopChildren(CollapsedNum, OMPD_teams_distribute_parallel_for));
   return new (Mem)
       OMPTeamsDistributeParallelForDirective(CollapsedNum, NumClauses);
 }
 
+OMPTargetTeamsDirective *OMPTargetTeamsDirective::Create(
+    const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
+    ArrayRef<OMPClause *> Clauses, Stmt *AssociatedStmt) {
+  auto Size =
+      llvm::alignTo(sizeof(OMPTargetTeamsDirective), alignof(OMPClause *));
+  void *Mem =
+      C.Allocate(Size + sizeof(OMPClause *) * Clauses.size() + sizeof(Stmt *));
+  OMPTargetTeamsDirective *Dir =
+      new (Mem) OMPTargetTeamsDirective(StartLoc, EndLoc, Clauses.size());
+  Dir->setClauses(Clauses);
+  Dir->setAssociatedStmt(AssociatedStmt);
+  return Dir;
+}
+
+OMPTargetTeamsDirective *
+OMPTargetTeamsDirective::CreateEmpty(const ASTContext &C, unsigned NumClauses,
+                                     EmptyShell) {
+  auto Size =
+      llvm::alignTo(sizeof(OMPTargetTeamsDirective), alignof(OMPClause *));
+  void *Mem =
+      C.Allocate(Size + sizeof(OMPClause *) * NumClauses + sizeof(Stmt *));
+  return new (Mem) OMPTargetTeamsDirective(NumClauses);
+}

lib/AST/StmtPrinter.cpp

 }
 
 void StmtPrinter::VisitOMPTeamsDistributeParallelForDirective(
     OMPTeamsDistributeParallelForDirective *Node) {
   Indent() << "#pragma omp teams distribute parallel for ";
   PrintOMPExecutableDirective(Node);
 }
 
+void StmtPrinter::VisitOMPTargetTeamsDirective(OMPTargetTeamsDirective *Node) {
+  Indent() << "#pragma omp target teams ";
+  PrintOMPExecutableDirective(Node);
+}
+
 //===----------------------------------------------------------------------===//
 //  Expr printing methods.
 //===----------------------------------------------------------------------===//
 
 void StmtPrinter::VisitDeclRefExpr(DeclRefExpr *Node) {
   if (auto *OCED = dyn_cast<OMPCapturedExprDecl>(Node->getDecl())) {
     OCED->getInit()->IgnoreImpCasts()->printPretty(OS, nullptr, Policy);
     return;

lib/AST/StmtProfile.cpp

void StmtProfiler::VisitOMPTeamsDistributeParallelForSimdDirective(
   VisitOMPLoopDirective(S);
 }
 
 void StmtProfiler::VisitOMPTeamsDistributeParallelForDirective(
     const OMPTeamsDistributeParallelForDirective *S) {
   VisitOMPLoopDirective(S);
 }
 
+void StmtProfiler::VisitOMPTargetTeamsDirective(
+    const OMPTargetTeamsDirective *S) {
+  VisitOMPExecutableDirective(S);
+}
+
 void StmtProfiler::VisitExpr(const Expr *S) {
   VisitStmt(S);
 }
 
 void StmtProfiler::VisitDeclRefExpr(const DeclRefExpr *S) {
   VisitExpr(S);
   if (!Canonical)
     VisitNestedNameSpecifier(S->getQualifier());

lib/Basic/OpenMPKinds.cpp

 #define OPENMP_TEAMS_DISTRIBUTE_PARALLEL_FOR_CLAUSE(Name)                      \
   case OMPC_##Name:                                                            \
     return true;
 #include "clang/Basic/OpenMPKinds.def"
     default:
       break;
     }
     break;
+  case OMPD_target_teams:
+    switch (CKind) {
+#define OPENMP_TARGET_TEAMS_CLAUSE(Name)                                       \
+  case OMPC_##Name:                                                            \
+    return true;
+#include "clang/Basic/OpenMPKinds.def"
+    default:
+      break;
+    }
+    break;
   case OMPD_declare_target:
   case OMPD_end_declare_target:
   case OMPD_unknown:
   case OMPD_threadprivate:
   case OMPD_section:
   case OMPD_master:
   case OMPD_taskyield:
   case OMPD_barrier:
return DKind == OMPD_parallel || DKind == OMPD_parallel_for ||
          DKind == OMPD_distribute_parallel_for_simd ||
          DKind == OMPD_target_parallel_for_simd ||
          DKind == OMPD_teams_distribute_parallel_for ||
          DKind == OMPD_teams_distribute_parallel_for_simd;
   // TODO add next directives.
 }
 
 bool clang::isOpenMPTargetExecutionDirective(OpenMPDirectiveKind DKind) {
-  // TODO add next directives.
   return DKind == OMPD_target || DKind == OMPD_target_parallel ||
          DKind == OMPD_target_parallel_for || 
-         DKind == OMPD_target_parallel_for_simd || DKind == OMPD_target_simd;
+         DKind == OMPD_target_parallel_for_simd || DKind == OMPD_target_simd ||
+         DKind == OMPD_target_teams;
 }
 
 bool clang::isOpenMPTargetDataManagementDirective(OpenMPDirectiveKind DKind) {
   return DKind == OMPD_target_data || DKind == OMPD_target_enter_data ||
          DKind == OMPD_target_exit_data || DKind == OMPD_target_update;
 }
 
-bool clang::isOpenMPTeamsDirective(OpenMPDirectiveKind DKind) {
+bool clang::isOpenMPNestingTeamsDirective(OpenMPDirectiveKind DKind) {
   return DKind == OMPD_teams || DKind == OMPD_teams_distribute ||
          DKind == OMPD_teams_distribute_simd ||
          DKind == OMPD_teams_distribute_parallel_for_simd ||
          DKind == OMPD_teams_distribute_parallel_for;
-  // TODO add next directives.
+}
+
+bool clang::isOpenMPTeamsDirective(OpenMPDirectiveKind DKind) {
+  return isOpenMPNestingTeamsDirective(DKind) ||
+         DKind == OMPD_target_teams;
 }
 
 bool clang::isOpenMPSimdDirective(OpenMPDirectiveKind DKind) {
   return DKind == OMPD_simd || DKind == OMPD_for_simd ||
          DKind == OMPD_parallel_for_simd || DKind == OMPD_taskloop_simd ||
          DKind == OMPD_distribute_parallel_for_simd ||
          DKind == OMPD_distribute_simd || DKind == OMPD_target_simd ||
          DKind == OMPD_teams_distribute_simd ||

lib/CodeGen/CGStmt.cpp

#include "clang/AST/StmtNodes.inc"
   case Stmt::OMPTeamsDistributeParallelForSimdDirectiveClass:
     EmitOMPTeamsDistributeParallelForSimdDirective(
         cast<OMPTeamsDistributeParallelForSimdDirective>(*S));
     break;
   case Stmt::OMPTeamsDistributeParallelForDirectiveClass:
     EmitOMPTeamsDistributeParallelForDirective(
         cast<OMPTeamsDistributeParallelForDirective>(*S));
     break;
+  case Stmt::OMPTargetTeamsDirectiveClass:
+    EmitOMPTargetTeamsDirective(cast<OMPTargetTeamsDirective>(*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

CGM.getOpenMPRuntime().emitInlinedDirective(
       *this, OMPD_teams_distribute_parallel_for,
       [&S](CodeGenFunction &CGF, PrePostActionTy &) {
         OMPLoopScope PreInitScope(CGF, S);
         CGF.EmitStmt(
             cast<CapturedStmt>(S.getAssociatedStmt())->getCapturedStmt());
       });
 }
 
+void CodeGenFunction::EmitOMPTargetTeamsDirective(
+    const OMPTargetTeamsDirective &S) {
+  CGM.getOpenMPRuntime().emitInlinedDirective(*this, OMPD_target_teams,
+      [&S](CodeGenFunction &CGF, PrePostActionTy &) {
+        CGF.EmitStmt(cast<CapturedStmt>(S.getAssociatedStmt())->getCapturedStmt());
+      });
+}
+
 /// \brief Emit a helper variable and return corresponding lvalue.
 static LValue EmitOMPHelperVar(CodeGenFunction &CGF,
                                const DeclRefExpr *Helper) {
   auto VDecl = cast<VarDecl>(Helper->getDecl());
   CGF.EmitVarDecl(*VDecl);
   return CGF.EmitLValue(Helper);
 }
 

lib/CodeGen/CodeGenFunction.h

#endif
   void EmitOMPTargetSimdDirective(const OMPTargetSimdDirective &S);
   void EmitOMPTeamsDistributeDirective(const OMPTeamsDistributeDirective &S);
   void
   EmitOMPTeamsDistributeSimdDirective(const OMPTeamsDistributeSimdDirective &S);
   void EmitOMPTeamsDistributeParallelForSimdDirective(
       const OMPTeamsDistributeParallelForSimdDirective &S);
   void EmitOMPTeamsDistributeParallelForDirective(
       const OMPTeamsDistributeParallelForDirective &S);
+  void EmitOMPTargetTeamsDirective(const OMPTargetTeamsDirective &S);
 
   /// Emit outlined function for the target directive.
   static std::pair<llvm::Function * /*OutlinedFn*/,
                    llvm::Constant * /*OutlinedFnID*/>
   EmitOMPTargetDirectiveOutlinedFunction(CodeGenModule &CGM,
                                          const OMPTargetDirective &S,
                                          StringRef ParentName,
                                          bool IsOffloadEntry);

lib/Parse/ParseOpenMP.cpp

static const unsigned F[][3] = {
     { OMPD_target, OMPD_parallel, OMPD_target_parallel },
     { OMPD_target, OMPD_simd, OMPD_target_simd },
     { OMPD_target_parallel, OMPD_for, OMPD_target_parallel_for },
     { OMPD_target_parallel_for, OMPD_simd, OMPD_target_parallel_for_simd },
     { OMPD_teams, OMPD_distribute, OMPD_teams_distribute },
     { OMPD_teams_distribute, OMPD_simd, OMPD_teams_distribute_simd },
     { OMPD_teams_distribute, OMPD_parallel, OMPD_teams_distribute_parallel },
     { OMPD_teams_distribute_parallel, OMPD_for, OMPD_teams_distribute_parallel_for },
-    { OMPD_teams_distribute_parallel_for, OMPD_simd, OMPD_teams_distribute_parallel_for_simd }
+    { OMPD_teams_distribute_parallel_for, OMPD_simd, OMPD_teams_distribute_parallel_for_simd },
+    { OMPD_target, OMPD_teams, OMPD_target_teams }
   };
   enum { CancellationPoint = 0, DeclareReduction = 1, TargetData = 2 };
   auto Tok = P.getCurToken();
   unsigned DKind =
       Tok.isAnnotation()
           ? static_cast<unsigned>(OMPD_unknown)
           : getOpenMPDirectiveKindEx(P.getPreprocessor().getSpelling(Tok));
   if (DKind == OMPD_unknown)
Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirectiveWithExtDecl(
   case OMPD_distribute_parallel_for_simd:
   case OMPD_distribute_simd:
   case OMPD_target_parallel_for_simd:
   case OMPD_target_simd:
   case OMPD_teams_distribute:
   case OMPD_teams_distribute_simd:
   case OMPD_teams_distribute_parallel_for_simd:
   case OMPD_teams_distribute_parallel_for:
+  case OMPD_target_teams:
     Diag(Tok, diag::err_omp_unexpected_directive)
         << getOpenMPDirectiveName(DKind);
     break;
   }
   while (Tok.isNot(tok::annot_pragma_openmp_end))
     ConsumeAnyToken();
   ConsumeAnyToken();
   return nullptr;

 ///         'taskgroup' | 'teams' | 'taskloop' | 'taskloop simd' |
 ///         'distribute' | 'target enter data' | 'target exit data' |
 ///         'target parallel' | 'target parallel for' |
 ///         'target update' | 'distribute parallel for' |
 ///         'distribute paralle for simd' | 'distribute simd' |
 ///         'target parallel for simd' | 'target simd' |
 ///         'teams distribute' | 'teams distribute simd' |
 ///         'teams distribute parallel for simd' |
-///         'teams distribute parallel for' {clause}
+///         'teams distribute parallel for' | 'target teams' {clause}
 ///         annot_pragma_openmp_end
 ///
 StmtResult Parser::ParseOpenMPDeclarativeOrExecutableDirective(
     AllowedContsructsKind Allowed) {
   assert(Tok.is(tok::annot_pragma_openmp) && "Not an OpenMP directive!");
   ParenBraceBracketBalancer BalancerRAIIObj(*this);
   SmallVector<OMPClause *, 5> Clauses;
   SmallVector<llvm::PointerIntPair<OMPClause *, 1, bool>, OMPC_unknown + 1>
StmtResult Parser::ParseOpenMPDeclarativeOrExecutableDirective(
   case OMPD_distribute_parallel_for:
   case OMPD_distribute_parallel_for_simd:
   case OMPD_distribute_simd:
   case OMPD_target_parallel_for_simd:
   case OMPD_target_simd:
   case OMPD_teams_distribute:
   case OMPD_teams_distribute_simd:
   case OMPD_teams_distribute_parallel_for_simd:
-  case OMPD_teams_distribute_parallel_for: {
+  case OMPD_teams_distribute_parallel_for:
+  case OMPD_target_teams: {
     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 =

lib/Sema/SemaOpenMP.cpp

 } // namespace
 
 void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) {
   switch (DKind) {
   case OMPD_parallel:
   case OMPD_parallel_for:
   case OMPD_parallel_for_simd:
   case OMPD_parallel_sections:
-  case OMPD_teams: {
+  case OMPD_teams:
+  case OMPD_target_teams: {
     QualType KmpInt32Ty = Context.getIntTypeForBitwidth(32, 1);
     QualType KmpInt32PtrTy =
         Context.getPointerType(KmpInt32Ty).withConst().withRestrict();
     Sema::CapturedParamNameType Params[] = {
         std::make_pair(".global_tid.", KmpInt32PtrTy),
         std::make_pair(".bound_tid.", KmpInt32PtrTy),
         std::make_pair(StringRef(), QualType()) // __context with shared vars
     };
if (CurrentRegion == OMPD_section) {
             << getOpenMPDirectiveName(ParentRegion);
         return true;
       }
       return false;
     }
     // Allow some constructs (except teams) to be orphaned (they could be
     // used in functions, called from OpenMP regions with the required
     // preconditions).
-    if (ParentRegion == OMPD_unknown && !isOpenMPTeamsDirective(CurrentRegion))
+    if (ParentRegion == OMPD_unknown &&
+        !isOpenMPNestingTeamsDirective(CurrentRegion))
       return false;
     if (CurrentRegion == OMPD_cancellation_point ||
         CurrentRegion == OMPD_cancel) {
       // OpenMP [2.16, Nesting of Regions]
       // A cancellation point construct for which construct-type-clause is
       // taskgroup must be nested inside a task construct. A cancellation
       // point construct for which construct-type-clause is not taskgroup must
       // be closely nested inside an OpenMP construct that matches the type
if (CurrentRegion == OMPD_cancellation_point ||
       // OpenMP [2.8.1,simd Construct, Restrictions]
       // An ordered construct with the simd clause is the only OpenMP construct
       // that can appear in the simd region.
       NestingProhibited = ParentRegion == OMPD_critical ||
                           isOpenMPTaskingDirective(ParentRegion) ||
                           !(isOpenMPSimdDirective(ParentRegion) ||
                             Stack->isParentOrderedRegion());
       Recommend = ShouldBeInOrderedRegion;
-    } else if (isOpenMPTeamsDirective(CurrentRegion)) {
+    } else if (isOpenMPNestingTeamsDirective(CurrentRegion)) {
       // OpenMP [2.16, Nesting of Regions]
       // If specified, a teams construct must be contained within a target
       // construct.
       NestingProhibited = ParentRegion != OMPD_target;
       OrphanSeen = ParentRegion == OMPD_unknown;
       Recommend = ShouldBeInTargetRegion;
       Stack->setParentTeamsRegionLoc(Stack->getConstructLoc());
     }
-    if (!NestingProhibited && ParentRegion == OMPD_teams) {
+    if (!NestingProhibited &&
+        !isOpenMPTargetExecutionDirective(CurrentRegion) &&
+        !isOpenMPTargetDataManagementDirective(CurrentRegion) &&
+        (ParentRegion == OMPD_teams || ParentRegion == OMPD_target_teams)) {
       // OpenMP [2.16, Nesting of Regions]
       // distribute, parallel, parallel sections, parallel workshare, and the
       // parallel loop and parallel loop SIMD constructs are the only OpenMP
       // constructs that can be closely nested in the teams region.
       NestingProhibited = !isOpenMPParallelDirective(CurrentRegion) &&
                           !isOpenMPDistributeDirective(CurrentRegion);
       Recommend = ShouldBeInParallelRegion;
     }
     if (!NestingProhibited &&
         isOpenMPNestingDistributeDirective(CurrentRegion)) {
       // OpenMP 4.5 [2.17 Nesting of Regions]
       // The region associated with the distribute construct must be strictly
       // nested inside a teams region
-      NestingProhibited = ParentRegion != OMPD_teams;
+      NestingProhibited =
+          (ParentRegion != OMPD_teams && ParentRegion != OMPD_target_teams);
       Recommend = ShouldBeInTeamsRegion;
     }
     if (!NestingProhibited &&
         (isOpenMPTargetExecutionDirective(CurrentRegion) ||
          isOpenMPTargetDataManagementDirective(CurrentRegion))) {
       // OpenMP 4.5 [2.17 Nesting of Regions]
       // If a target, target update, target data, target enter data, or
       // target exit data construct is encountered during execution of a
Res = ActOnOpenMPTeamsDistributeParallelForSimdDirective(
         ClausesWithImplicit, AStmt, StartLoc, EndLoc, VarsWithInheritedDSA);
     AllowedNameModifiers.push_back(OMPD_parallel);
     break;
   case OMPD_teams_distribute_parallel_for:
     Res = ActOnOpenMPTeamsDistributeParallelForDirective(
         ClausesWithImplicit, AStmt, StartLoc, EndLoc, VarsWithInheritedDSA);
     AllowedNameModifiers.push_back(OMPD_parallel);
     break;
+  case OMPD_target_teams:
+    Res = ActOnOpenMPTargetTeamsDirective(ClausesWithImplicit, AStmt, StartLoc,
+                                          EndLoc);
+    AllowedNameModifiers.push_back(OMPD_target);
+    break;
   case OMPD_declare_target:
   case OMPD_end_declare_target:
   case OMPD_threadprivate:
   case OMPD_declare_reduction:
   case OMPD_declare_simd:
     llvm_unreachable("OpenMP Directive is not allowed");
   case OMPD_unknown:
     llvm_unreachable("Unknown OpenMP directive");
if (!CurContext->isDependentContext()) {
     }
   }
 
   getCurFunction()->setHasBranchProtectedScope();
   return OMPTeamsDistributeParallelForDirective::Create(
       Context, StartLoc, EndLoc, NestedLoopCount, Clauses, AStmt, B);
 }
 
+StmtResult Sema::ActOnOpenMPTargetTeamsDirective(ArrayRef<OMPClause *> Clauses,
+                                                 Stmt *AStmt,
+                                                 SourceLocation StartLoc,
+                                                 SourceLocation EndLoc) {
+  if (!AStmt)
+    return StmtError();
+
+  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.
+  // The point of exit cannot be a branch out of the structured block.
+  // longjmp() and throw() must not violate the entry/exit criteria.
+  CS->getCapturedDecl()->setNothrow();
+
+  getCurFunction()->setHasBranchProtectedScope();
+
+  return OMPTargetTeamsDirective::Create(Context, StartLoc, EndLoc, Clauses,
+                                         AStmt);
+}
+
 OMPClause *Sema::ActOnOpenMPSingleExprClause(OpenMPClauseKind Kind, Expr *Expr,
                                              SourceLocation StartLoc,
                                              SourceLocation LParenLoc,
                                              SourceLocation EndLoc) {
   OMPClause *Res = nullptr;
   switch (Kind) {
   case OMPC_final:
     Res = ActOnOpenMPFinalClause(Expr, StartLoc, LParenLoc, EndLoc);
for (auto &RefExpr : VarList) {
     DSAStackTy::DSAVarData DVar = DSAStack->getTopDSA(D, false);
     if (DVar.CKind != OMPC_unknown && DVar.CKind != OMPC_private) {
       Diag(ELoc, diag::err_omp_wrong_dsa) << getOpenMPClauseName(DVar.CKind)
                                           << getOpenMPClauseName(OMPC_private);
       ReportOriginalDSA(*this, DSAStack, D, DVar);
       continue;
     }
 
+    auto CurrDir = DSAStack->getCurrentDirective();
     // Variably modified types are not supported for tasks.
     if (!Type->isAnyPointerType() && Type->isVariablyModifiedType() &&
-        isOpenMPTaskingDirective(DSAStack->getCurrentDirective())) {
+        isOpenMPTaskingDirective(CurrDir)) {
       Diag(ELoc, diag::err_omp_variably_modified_type_not_supported)
           << getOpenMPClauseName(OMPC_private) << Type
-          << getOpenMPDirectiveName(DSAStack->getCurrentDirective());
+          << getOpenMPDirectiveName(CurrDir);
       bool IsDecl =
           !VD ||
           VD->isThisDeclarationADefinition(Context) == VarDecl::DeclarationOnly;
       Diag(D->getLocation(),
            IsDecl ? diag::note_previous_decl : diag::note_defined_here)
           << D;
       continue;
     }
 
     // OpenMP 4.5 [2.15.5.1, Restrictions, p.3]
     // A list item cannot appear in both a map clause and a data-sharing
     // attribute clause on the same construct
-    if (DSAStack->getCurrentDirective() == OMPD_target) {
+    if (CurrDir == OMPD_target || CurrDir == OMPD_target_teams) {
       OpenMPClauseKind ConflictKind;
       if (DSAStack->checkMappableExprComponentListsForDecl(
               VD, /*CurrentRegionOnly=*/true,
               [&](OMPClauseMappableExprCommon::MappableExprComponentListRef,
                   OpenMPClauseKind WhereFoundClauseKind) -> bool {
                 ConflictKind = WhereFoundClauseKind;
                 return true;
               })) {
         Diag(ELoc, diag::err_omp_variable_in_given_clause_and_dsa)
             << getOpenMPClauseName(OMPC_private)
             << getOpenMPClauseName(ConflictKind)
-            << getOpenMPDirectiveName(DSAStack->getCurrentDirective());
+            << getOpenMPDirectiveName(CurrDir);
         ReportOriginalDSA(*this, DSAStack, D, DVar);
         continue;
       }
     }
 
     // OpenMP [2.9.3.3, Restrictions, C/C++, p.1]
     //  A variable of class type (or array thereof) that appears in a private
     //  clause requires an accessible, unambiguous default constructor for the
if (!IsImplicitClause) {
           Diag(ELoc, diag::err_omp_firstprivate_and_lastprivate_in_distribute);
           ReportOriginalDSA(*this, DSAStack, D, DVar);
           continue;
         }
       }
       // OpenMP 4.5 [2.15.5.1, Restrictions, p.3]
       // A list item cannot appear in both a map clause and a data-sharing
       // attribute clause on the same construct
-      if (CurrDir == OMPD_target) {
+      if (CurrDir == OMPD_target || CurrDir == OMPD_target_teams) {
         OpenMPClauseKind ConflictKind;
         if (DSAStack->checkMappableExprComponentListsForDecl(
                 VD, /*CurrentRegionOnly=*/true,
                 [&](OMPClauseMappableExprCommon::MappableExprComponentListRef,
                     OpenMPClauseKind WhereFoundClauseKind) -> bool {
                   ConflictKind = WhereFoundClauseKind;
                   return true;
                 })) {
if (CKind == OMPC_map) {
             << getOpenMPSimpleClauseTypeName(OMPC_map, MapType)
             << getOpenMPDirectiveName(DKind);
         continue;
       }
 
       // OpenMP 4.5 [2.15.5.1, Restrictions, p.3]
       // A list item cannot appear in both a map clause and a data-sharing
       // attribute clause on the same construct
-      if (DKind == OMPD_target && VD) {
+      if ((DKind == OMPD_target || DKind == OMPD_target_teams) && VD) {
         auto DVar = DSAS->getTopDSA(VD, false);
         if (isOpenMPPrivate(DVar.CKind)) {
           SemaRef.Diag(ELoc, diag::err_omp_variable_in_given_clause_and_dsa)
               << getOpenMPClauseName(DVar.CKind)
               << getOpenMPClauseName(OMPC_map)
               << getOpenMPDirectiveName(DSAS->getCurrentDirective());
           ReportOriginalDSA(SemaRef, DSAS, CurDeclaration, DVar);
           continue;

lib/Sema/TreeTransform.h

StmtResult TreeTransform<Derived>::TransformOMPTeamsDistributeParallelForDirective(
   DeclarationNameInfo DirName;
   getDerived().getSema().StartOpenMPDSABlock(OMPD_teams_distribute_parallel_for,
       DirName, nullptr, D->getLocStart());
   StmtResult Res = getDerived().TransformOMPExecutableDirective(D);
   getDerived().getSema().EndOpenMPDSABlock(Res.get());
   return Res;
 }
 
+template <typename Derived>
+StmtResult TreeTransform<Derived>::TransformOMPTargetTeamsDirective(
+    OMPTargetTeamsDirective *D) {
+  DeclarationNameInfo DirName;
+  getDerived().getSema().StartOpenMPDSABlock(OMPD_target_teams, DirName,
+                                             nullptr, D->getLocStart());
+  auto Res = getDerived().TransformOMPExecutableDirective(D);
+  getDerived().getSema().EndOpenMPDSABlock(Res.get());
+  return Res;
+}
 
 //===----------------------------------------------------------------------===//
 // OpenMP clause transformation
 //===----------------------------------------------------------------------===//
 template <typename Derived>
 OMPClause *TreeTransform<Derived>::TransformOMPIfClause(OMPIfClause *C) {
   ExprResult Cond = getDerived().TransformExpr(C->getCondition());
   if (Cond.isInvalid())

lib/Serialization/ASTReaderStmt.cpp

void ASTStmtReader::VisitOMPTeamsDistributeParallelForSimdDirective(
   VisitOMPLoopDirective(D);
 }
 
 void ASTStmtReader::VisitOMPTeamsDistributeParallelForDirective(
     OMPTeamsDistributeParallelForDirective *D) {
   VisitOMPLoopDirective(D);
 }
 
+void ASTStmtReader::VisitOMPTargetTeamsDirective(OMPTargetTeamsDirective *D) {
+  VisitStmt(D);
+  // The NumClauses field was read in ReadStmtFromStream.
+  ++Idx;
+  VisitOMPExecutableDirective(D);
+}
+
 //===----------------------------------------------------------------------===//
 // ASTReader Implementation
 //===----------------------------------------------------------------------===//
 
 Stmt *ASTReader::ReadStmt(ModuleFile &F) {
   switch (ReadingKind) {
   case Read_None:
     llvm_unreachable("should not call this when not reading anything");
while (true) {
     case STMT_OMP_TEAMS_DISTRIBUTE_PARALLEL_FOR_DIRECTIVE: {
       auto NumClauses = Record[ASTStmtReader::NumStmtFields];
       auto CollapsedNum = Record[ASTStmtReader::NumStmtFields + 1];
       S = OMPTeamsDistributeParallelForDirective::CreateEmpty(
           Context, NumClauses, CollapsedNum, Empty);
       break;
     }
 
+    case STMT_OMP_TARGET_TEAMS_DIRECTIVE: {
+      S = OMPTargetTeamsDirective::CreateEmpty(
+          Context, Record[ASTStmtReader::NumStmtFields], Empty);
+      break;
+    }
+ 
     case EXPR_CXX_OPERATOR_CALL:
       S = new (Context) CXXOperatorCallExpr(Context, Empty);
       break;
 
     case EXPR_CXX_MEMBER_CALL:
       S = new (Context) CXXMemberCallExpr(Context, Empty);
       break;
 

lib/Serialization/ASTWriterStmt.cpp

 }
 
 void ASTStmtWriter::VisitOMPTeamsDistributeParallelForDirective(
     OMPTeamsDistributeParallelForDirective *D) {
   VisitOMPLoopDirective(D);
   Code = serialization::STMT_OMP_TEAMS_DISTRIBUTE_PARALLEL_FOR_DIRECTIVE;
 }
 
+void ASTStmtWriter::VisitOMPTargetTeamsDirective(OMPTargetTeamsDirective *D) {
+  VisitStmt(D);
+  Record.push_back(D->getNumClauses());
+  VisitOMPExecutableDirective(D);
+  Code = serialization::STMT_OMP_TARGET_TEAMS_DIRECTIVE;
+}
+
 //===----------------------------------------------------------------------===//
 // ASTWriter Implementation
 //===----------------------------------------------------------------------===//
 
 unsigned ASTWriter::RecordSwitchCaseID(SwitchCase *S) {
   assert(SwitchCaseIDs.find(S) == SwitchCaseIDs.end() &&
          "SwitchCase recorded twice");
   unsigned NextID = SwitchCaseIDs.size();

lib/StaticAnalyzer/Core/ExprEngine.cpp

switch (S->getStmtClass()) {
     case Stmt::OMPDistributeParallelForSimdDirectiveClass:
     case Stmt::OMPDistributeSimdDirectiveClass:
     case Stmt::OMPTargetParallelForSimdDirectiveClass:
     case Stmt::OMPTargetSimdDirectiveClass:
     case Stmt::OMPTeamsDistributeDirectiveClass:
     case Stmt::OMPTeamsDistributeSimdDirectiveClass:
     case Stmt::OMPTeamsDistributeParallelForSimdDirectiveClass:
     case Stmt::OMPTeamsDistributeParallelForDirectiveClass:
+    case Stmt::OMPTargetTeamsDirectiveClass:
       llvm_unreachable("Stmt should not be in analyzer evaluation loop");
 
     case Stmt::ObjCSubscriptRefExprClass:
     case Stmt::ObjCPropertyRefExprClass:
       llvm_unreachable("These are handled by PseudoObjectExpr");
 
     case Stmt::GNUNullExprClass: {
       // GNU __null is a pointer-width integer, not an actual pointer.

test/OpenMP/nesting_of_regions.cpp

for (int i = 0; i < 10; ++i)
       ;
   }
 #pragma omp parallel
   {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'parallel' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp parallel
+  {
+#pragma omp target teams // OK
+    ++a;
+  }
 
 // SIMD DIRECTIVE
 #pragma omp simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int j = 0; j < 10; ++j)
       ;
   }
 #pragma omp simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int j = 0; j < 10; ++j)
       ;
   }
+#pragma omp simd
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // expected-error {{OpenMP constructs may not be nested inside a simd region}}
+    ++a;
+  }
 
 // FOR DIRECTIVE
 #pragma omp for
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'for' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int j = 0; j < 10; ++j)
       ;
   }
 #pragma omp for
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'for' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int j = 0; j < 10; ++j)
       ;
   }
+#pragma omp for
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // OK
+    ++a;
+  }
 
 // FOR SIMD DIRECTIVE
 #pragma omp for simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int j = 0; j < 10; ++j)
       ;
   }
 #pragma omp for simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int j = 0; j < 10; ++j)
       ;
   }
+#pragma omp for simd
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // expected-error {{OpenMP constructs may not be nested inside a simd region}}
+    a++;
+  }
 
 // SECTIONS DIRECTIVE
 #pragma omp sections
   {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'sections' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
       ;
   }
 #pragma omp sections
   {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'sections' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp sections
+  {
+#pragma omp target teams // OK
+    ++a;
+  }
 
 // SECTION DIRECTIVE
 #pragma omp section // expected-error {{orphaned 'omp section' directives are prohibited, it must be closely nested to a sections region}}
   {
     bar();
   }
 #pragma omp sections
   {
#pragma omp teams distribute parallel for simd // expected-error {{region cannot be closely nested inside 'section' region; perhaps you forget to enclose 'omp teams distribute parallel for simd' directive into a target region?}}
   }
 #pragma omp sections
   {
 #pragma omp section
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'section' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp sections
+  {
+#pragma omp section
+#pragma omp target teams // OK
+    ++a;
+  }
 
 // SINGLE DIRECTIVE
 #pragma omp single
   {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'single' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
       ;
   }
 #pragma omp single
   {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'single' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp single
+  {
+#pragma omp target teams // OK
+    ++a;
+  }
 
 // MASTER DIRECTIVE
 #pragma omp master
   {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'master' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
       ;
   }
 #pragma omp master
   {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'master' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp master
+  {
+#pragma omp target teams // OK
+    a++;
+  }
 
 // CRITICAL DIRECTIVE
 #pragma omp critical
   {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'critical' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
       ;
   }
 #pragma omp critical
   {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'critical' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp critical
+  {
+#pragma omp target teams // OK
+    a++;
+  }
 
 // PARALLEL FOR DIRECTIVE
 #pragma omp parallel for
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'parallel for' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int j = 0; j < 10; ++j)
       ;
   }
 #pragma omp parallel for
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'parallel for' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int j = 0; j < 10; ++j)
       ;
   }
+#pragma omp parallel for
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // OK
+    a++;
+  }
 
 // PARALLEL FOR SIMD DIRECTIVE
 #pragma omp parallel for simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int j = 0; j < 10; ++j)
       ;
   }
 #pragma omp parallel for simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int j = 0; j < 10; ++j)
       ;
   }
+#pragma omp parallel for simd
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // expected-error {{OpenMP constructs may not be nested inside a simd region}}
+    a++;
+  }
 
 // PARALLEL SECTIONS DIRECTIVE
 #pragma omp parallel sections
   {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'parallel sections' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
       ;
   }
 #pragma omp parallel sections
   {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'parallel sections' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp parallel sections
+  {
+#pragma omp target teams // OK
+    a++;
+  }
 
 // TASK DIRECTIVE
 #pragma omp task
 #pragma omp for // expected-error {{region cannot be closely nested inside 'task' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
   for (int i = 0; i < 10; ++i)
     ;
 #pragma omp task
 #pragma omp simd
for (int i = 0; i < 10; ++i)
       ;
   }
 #pragma omp task
   {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'task' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp task
+  {
+#pragma omp target teams // OK
+    a++;
+  }
 
 // ORDERED DIRECTIVE
 #pragma omp ordered
   {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'ordered' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
       ;
   }
 #pragma omp ordered
   {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'ordered' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp ordered
+  {
+#pragma omp target teams // OK
+    a++;
+  }
 
 // ATOMIC DIRECTIVE
 #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 for // expected-error {{OpenMP constructs may not be nested inside an atomic region}}
     for (int i = 0; i < 10; ++i)

 #pragma omp atomic
   // expected-error@+2 {{the statement for 'atomic' must be an expression statement of form '++x;', '--x;', 'x++;', 'x--;', 'x binop= expr;', 'x = x binop expr' or 'x = expr binop x', where x is an l-value expression with scalar type}}
   // expected-note@+1 {{expected an expression statement}}
   {
 #pragma omp teams distribute parallel for // expected-error {{OpenMP constructs may not be nested inside an atomic region}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp atomic
+  // expected-error@+2 {{the statement for 'atomic' must be an expression statement of form '++x;', '--x;', 'x++;', 'x--;', 'x binop= expr;', 'x = x binop expr' or 'x = expr binop x', where x is an l-value expression with scalar type}}
+  // expected-note@+1 {{expected an expression statement}}
+  {
+#pragma omp target teams // expected-error {{OpenMP constructs may not be nested inside an atomic region}}
+    a++;
+  }
 
 // TARGET DIRECTIVE
 #pragma omp target
 #pragma omp parallel
   bar();
 #pragma omp target
 #pragma omp for
   for (int i = 0; i < 10; ++i)
for (int i = 0; i < 10; ++i)
       ;
   }
 #pragma omp target
   { 
 #pragma omp teams distribute parallel for // OK
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp target
+  { 
+#pragma omp target teams // expected-error {{region cannot be nested inside 'target' region}}
+    for (int i = 0; i < 10; ++i)
+      ;
+  }
 
 // TARGET PARALLEL DIRECTIVE
 #pragma omp target parallel
 #pragma omp parallel
   bar();
 #pragma omp target parallel
 #pragma omp for
   for (int i = 0; i < 10; ++i)
for (int i = 0; i < 10; ++i)
       ;
   }
 #pragma omp target parallel
   { 
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'target parallel' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp target parallel
+  { 
+#pragma omp target teams // expected-error {{region cannot be nested inside 'target parallel' region}}
+    a++;
+  }
 
 // TARGET PARALLEL FOR DIRECTIVE
 #pragma omp target parallel for
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'target parallel for' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int j = 0; j < 10; ++j)
       ;
   }    
 #pragma omp target parallel for
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'target parallel for' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int j = 0; j < 10; ++j)
       ;
   }    
+#pragma omp target parallel for
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // expected-error {{region cannot be nested inside 'target parallel for' region}}
+    a++;
+  }    
 
 // TEAMS DIRECTIVE
 #pragma omp teams // expected-error {{orphaned 'omp teams' directives are prohibited; perhaps you forget to enclose the directive into a target region?}}
   bar();
 #pragma omp target
 #pragma omp teams
 #pragma omp parallel
   bar();

 #pragma omp teams
   {
 #pragma omp atomic // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp atomic' directive into a parallel region?}}
     ++a;
   }
 #pragma omp target
 #pragma omp teams
   {
-#pragma omp target // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp target' directive into a parallel region?}}
+#pragma omp target // expected-error {{region cannot be nested inside 'target' region}}
     ++a;
   }
 #pragma omp target
 #pragma omp teams
   {
 #pragma omp target parallel // expected-error {{region cannot be nested inside 'target' region}}
     ++a;
   }
 #pragma omp target
 #pragma omp teams
 #pragma omp target parallel for // expected-error {{region cannot be nested inside 'target' region}}
   for (int i = 0; i < 10; ++i)
     ;
 #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 enter data map(to: a) // expected-error {{region cannot be nested inside 'target' region}}
     ++a;
   }
 #pragma omp target
 #pragma omp teams
   {
-#pragma omp target exit data map(from: a) // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp target exit data' directive into a parallel region?}}
+#pragma omp target exit data map(from: a) // expected-error {{region cannot be nested inside 'target' region}}
     ++a;
   }
 #pragma omp target
 #pragma omp teams
   {
 #pragma omp teams // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
     ++a;
   }
#pragma omp distribute
   for (int i = 0; i < 10; ++i)
     ;
 #pragma omp distribute
   for (int j = 0; j < 10; ++j)
     ;        
 #pragma omp target
 #pragma omp teams
   {
-#pragma omp target update to(a) // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp target update' directive into a parallel region?}}
+#pragma omp target update to(a) // expected-error {{region cannot be nested inside 'target' region}}
   }
 #pragma omp target
 #pragma omp teams
   {
 #pragma omp distribute parallel for
   for (int i = 0; i < 10; ++i)
     ;
 #pragma omp distribute parallel for
for (int i = 0; i < 10; ++i)
     ;
 #pragma omp distribute simd
   for (int j = 0; j < 10; ++j)
     ;  
   }
 #pragma omp target
 #pragma omp teams
   {
-#pragma omp target simd // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp target simd' directive into a parallel region?}}
+#pragma omp target simd // expected-error {{region cannot be nested inside 'target' region}}
   for (int i = 0; i < 10; ++i)
     ;
   }
 #pragma omp target
 #pragma omp teams
   {
 #pragma omp teams distribute // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp teams distribute' directive into a target region?}}
   for (int i = 0; i < 10; ++i)
#pragma omp teams distribute parallel for simd // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp teams distribute parallel for simd' directive into a target region?}}
   }
 #pragma omp target
 #pragma omp teams
   {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
   for (int i = 0; i < 10; ++i)
     ;
   }
+#pragma omp target
+#pragma omp teams
+  {
+#pragma omp target teams // expected-error {{region cannot be nested inside 'target' region}}
+    a++;
+  }
 
 // TASKLOOP DIRECTIVE
 #pragma omp taskloop
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'taskloop' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int j = 0; j < 10; ++j)
     ++a;
   }
 #pragma omp taskloop
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'taskloop' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
   for (int j = 0; j < 10; ++j)
     ++a;
   }
+#pragma omp taskloop
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // OK
+    ++a;
+  }
 
 // DISTRIBUTE DIRECTIVE
 #pragma omp target
 #pragma omp teams
 #pragma omp distribute
   for (int i = 0; i < 10; ++i) {
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'distribute' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int i = 0; i < 10; ++i)

 #pragma omp target
 #pragma omp teams
 #pragma omp distribute
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'distribute' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int j = 0; j < 10; ++j)
       ;
   }
+#pragma omp target
+#pragma omp teams
+#pragma omp distribute
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // expected-error {{region cannot be nested inside 'target' region}}
+    a++;
+  }
 
 // DISTRIBUTE PARALLEL FOR DIRECTIVE
 #pragma omp target
 #pragma omp teams
 #pragma omp distribute parallel for
   for (int i = 0; i < 10; ++i) {
 #pragma omp distribute parallel for // expected-error {{region cannot be closely nested inside 'distribute parallel for' region; perhaps you forget to enclose 'omp distribute parallel for' directive into a teams region?}}
     for (int i = 0; i < 10; ++i)

 #pragma omp target
 #pragma omp teams
 #pragma omp distribute parallel for
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'distribute parallel for' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp target
+#pragma omp teams
+#pragma omp distribute parallel for
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // expected-error {{region cannot be nested inside 'target' region}}
+    a++;
+  }
 
 // DISTRIBUTE PARALLEL FOR SIMD DIRECTIVE
 #pragma omp target
 #pragma omp teams
 #pragma omp distribute parallel for simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp distribute parallel for simd // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int i = 0; i < 10; ++i)

 #pragma omp target
 #pragma omp teams
 #pragma omp distribute parallel for simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int i = 0; i < 10; ++i)
       ++a;
   }
+#pragma omp target
+#pragma omp teams
+#pragma omp distribute parallel for simd
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // expected-error {{OpenMP constructs may not be nested inside a simd region}}
+    ++a;
+  }
 
 // TARGET SIMD DIRECTIVE
 #pragma omp target simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp distribute parallel for simd // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int i = 0; i < 10; ++i)
       ;
   }
#pragma omp teams distribute parallel for simd // expected-error {{region cannot be closely nested inside 'teams distribute' region; perhaps you forget to enclose 'omp teams distribute parallel for simd' directive into a target region?}}
   }
 #pragma omp target
 #pragma omp teams distribute
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'teams distribute' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int i = 0; i < 10; ++i)
       ++a;
   }
+#pragma omp target
+#pragma omp teams distribute
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // expected-error {{region cannot be nested inside 'target' region}}
+    ++a;
+  }
 
 // TEAMS DISTRIBUTE SIMD DIRECTIVE
 #pragma omp teams distribute simd // expected-error {{orphaned 'omp teams distribute simd' directives are prohibited; perhaps you forget to enclose the directive into a target region?}}
   for (int i = 0; i < 10; ++i)
     ;
 #pragma omp target
 #pragma omp teams distribute simd
   for (int i = 0; i < 10; ++i) {
#pragma omp teams distribute parallel for simd // expected-error {{OpenMP constructs may not be nested inside a simd region}}
   }
 #pragma omp target
 #pragma omp teams distribute simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int i = 0; i < 10; ++i)
       ++a;
   }
+#pragma omp target
+#pragma omp teams distribute simd
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // expected-error {{OpenMP constructs may not be nested inside a simd region}}
+    ++a;
+  }
 
 // TEAMS DISTRIBUTE PARALLEL FOR SIMD DIRECTIVE
 #pragma omp teams distribute parallel for simd // expected-error {{orphaned 'omp teams distribute parallel for simd' directives are prohibited; perhaps you forget to enclose the directive into a target region?}}
   for (int i = 0; i < 10; ++i)
     ;
 #pragma omp target
 #pragma omp teams distribute parallel for simd
   for (int i = 0; i < 10; ++i) {
#pragma omp teams distribute parallel for simd // expected-error {{OpenMP constructs may not be nested inside a simd region}}
   }
 #pragma omp target
 #pragma omp teams distribute parallel for simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int i = 0; i < 10; ++i)
       ++a;
   }
+#pragma omp target
+#pragma omp teams distribute parallel for simd
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // expected-error {{OpenMP constructs may not be nested inside a simd region}}
+    ++a;
+  }
 
 // TEAMS DISTRIBUTE PARALLEL FOR DIRECTIVE
 #pragma omp teams distribute parallel for // expected-error {{orphaned 'omp teams distribute parallel for' directives are prohibited; perhaps you forget to enclose the directive into a target region?}}
   for (int i = 0; i < 10; ++i)
     ;
 #pragma omp target
 #pragma omp teams distribute parallel for
   for (int i = 0; i < 10; ++i) {
#pragma omp teams distribute parallel for simd // expected-error {{region cannot be closely nested inside 'teams distribute parallel for' region; perhaps you forget to enclose 'omp teams distribute parallel for simd' directive into a target region?}}
   }
 #pragma omp target
 #pragma omp teams distribute parallel for
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'teams distribute parallel for' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int i = 0; i < 10; ++i)
       ++a;
   }
+#pragma omp target
+#pragma omp teams distribute parallel for
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // expected-error {{region cannot be nested inside 'target' region}}
+    ++a;
+  }
+
+// TARGET TEAMS DIRECTIVE
+#pragma omp target teams
+#pragma omp parallel
+  bar();
+#pragma omp target teams
+#pragma omp for // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
+  for (int i = 0; i < 10; ++i)
+    ;
+#pragma omp target teams
+#pragma omp simd // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp simd' directive into a parallel region?}}
+  for (int i = 0; i < 10; ++i)
+    ;
+#pragma omp target teams
+#pragma omp for simd // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp for simd' directive into a parallel region?}}
+  for (int i = 0; i < 10; ++i)
+    ;
+#pragma omp target teams
+#pragma omp sections // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp sections' directive into a parallel region?}}
+  {
+    bar();
+  }
+#pragma omp target teams
+#pragma omp section // expected-error {{'omp section' directive must be closely nested to a sections region, not a target teams region}}
+  {
+    bar();
+  }
+#pragma omp target teams
+#pragma omp single // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp single' directive into a parallel region?}}
+  bar();
+
+#pragma omp target teams
+#pragma omp master // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp master' directive into a parallel region?}}
+  {
+    bar();
+  }
+#pragma omp target teams
+#pragma omp critical // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp critical' directive into a parallel region?}}
+  {
+    bar();
+  }
+#pragma omp target teams
+#pragma omp parallel for
+  for (int i = 0; i < 10; ++i)
+    ;
+#pragma omp target teams
+#pragma omp parallel for simd
+  for (int i = 0; i < 10; ++i)
+    ;
+#pragma omp target teams
+#pragma omp parallel sections
+  {
+    bar();
+  }
+#pragma omp target teams
+#pragma omp task // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp task' directive into a parallel region?}}
+  {
+    bar();
+  }
+#pragma omp target teams
+  {
+#pragma omp taskyield // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp taskyield' directive into a parallel region?}}
+    bar();
+  }
+#pragma omp target teams
+  {
+#pragma omp barrier // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp barrier' directive into a parallel region?}}
+    bar();
+  }
+#pragma omp target teams
+  {
+#pragma omp taskwait // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp taskwait' directive into a parallel region?}}
+    bar();
+  }
+#pragma omp target teams
+  {
+#pragma omp flush // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp flush' directive into a parallel region?}}
+    bar();
+  }
+#pragma omp target teams
+  {
+#pragma omp ordered // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp ordered' directive into a for or a parallel for region with 'ordered' clause?}}
+    bar();
+  }
+#pragma omp target teams
+  {
+#pragma omp atomic // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp atomic' directive into a parallel region?}}
+    ++a;
+  }
+#pragma omp target teams
+  {
+#pragma omp target // expected-error {{region cannot be nested inside 'target teams' region}}
+    ++a;
+  }
+#pragma omp target teams
+  {
+#pragma omp target parallel // expected-error {{region cannot be nested inside 'target teams' region}}
+    ++a;
+  }
+#pragma omp target teams
+#pragma omp target parallel for // expected-error {{region cannot be nested inside 'target teams' region}}
+  for (int i = 0; i < 10; ++i)
+    ;
+#pragma omp target teams
+  {
+#pragma omp target enter data map(to: a) // expected-error {{region cannot be nested inside 'target teams' region}}
+  }
+#pragma omp target teams
+  {
+#pragma omp target exit data map(from: a) // expected-error {{region cannot be nested inside 'target teams' region}}
+  }
+#pragma omp target teams
+  {
+#pragma omp teams // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
+    ++a;
+  }
+#pragma omp target teams
+  {
+#pragma omp taskloop // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp taskloop' directive into a parallel region?}}
+  for (int i = 0; i < 10; ++i)
+    ++a;
+  }
+#pragma omp target teams
+#pragma omp distribute
+  for (int i = 0; i < 10; ++i)
+    ;
+#pragma omp target teams
+  {
+#pragma omp target update to(a) // expected-error {{region cannot be nested inside 'target teams' region}}
+    ++a;
+  }
+#pragma omp target teams
+#pragma omp distribute parallel for
+  for (int i = 0; i < 10; ++i)
+    ;
+#pragma omp target teams
+  {
+#pragma omp distribute parallel for
+  for (int i = 0; i < 10; ++i)
+    ;
+#pragma omp distribute parallel for
+  for (int j = 0; j < 10; ++j)
+    ;  
+  }
+#pragma omp target teams
+  {
+#pragma omp distribute parallel for simd
+  for (int i = 0; i < 10; ++i)
+    ;
+#pragma omp distribute parallel for simd
+  for (int j = 0; j < 10; ++j)
+    ;  
+  }
+#pragma omp target teams
+  {
+#pragma omp distribute simd
+  for (int i = 0; i < 10; ++i)
+    ;
+#pragma omp distribute simd
+  for (int j = 0; j < 10; ++j)
+    ;  
+  }
+#pragma omp target teams
+  {
+#pragma omp target simd // expected-error {{region cannot be nested inside 'target teams' region}}
+  for (int i = 0; i < 10; ++i)
+    ;
+  }
+#pragma omp target teams
+  {
+#pragma omp teams distribute // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp teams distribute' directive into a target region?}}
+  for (int i = 0; i < 10; ++i)
+    ;
+  }
+#pragma omp target teams
+  {
+#pragma omp teams distribute simd // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp teams distribute simd' directive into a target region?}}
+  for (int i = 0; i < 10; ++i)
+    ;
+  }
+#pragma omp target teams
+  {
+#pragma omp teams distribute parallel for simd // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp teams distribute parallel for simd' directive into a target region?}}
+  for (int i = 0; i < 10; ++i)
+    ;
+  }
+#pragma omp target teams
+  {
+#pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
+  for (int i = 0; i < 10; ++i)
+    ;
+  }
+#pragma omp target teams
+  {
+#pragma omp target teams // expected-error {{region cannot be nested inside 'target teams' region}}
+    a++;
+  }
+
 }
 
 void foo() {
   int a = 0;
 // PARALLEL DIRECTIVE
 #pragma omp parallel
 #pragma omp for
   for (int i = 0; i < 10; ++i)
for (int i = 0; i < 10; ++i)
       ;
   }
 #pragma omp parallel
   {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'parallel' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp parallel
+  {
+#pragma omp target teams // OK
+    a++;
+  }
 
 // SIMD DIRECTIVE
 #pragma omp simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int j = 0; j < 10; ++j)
       ;
   }
 #pragma omp simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int j = 0; j < 10; ++j)
       ;
   }
+#pragma omp simd
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // expected-error {{OpenMP constructs may not be nested inside a simd region}}
+    a++;
+  }
 
 // FOR DIRECTIVE
 #pragma omp for
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'for' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int j = 0; j < 10; ++j)
       ;
   }
 #pragma omp for
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'for' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int j = 0; j < 10; ++j)
       ;
   }
+#pragma omp for
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // OK
+    a++;
+  }
 
 // FOR SIMD DIRECTIVE
 #pragma omp for simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int j = 0; j < 10; ++j)
       ;
   }
 #pragma omp for simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int j = 0; j < 10; ++j)
       ;
   }
+#pragma omp for simd
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // expected-error {{OpenMP constructs may not be nested inside a simd region}}
+    a++;
+  }
 
 // SECTIONS DIRECTIVE
 #pragma omp sections
   {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'sections' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
       ;
   }
 #pragma omp sections
   {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'sections' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp sections
+  {
+#pragma omp target teams // OK
+    a++;
+  }
 
 // SECTION DIRECTIVE
 #pragma omp section // expected-error {{orphaned 'omp section' directives are prohibited, it must be closely nested to a sections region}}
   {
     bar();
   }
 #pragma omp sections
   {
#pragma omp teams distribute parallel for simd // expected-error {{region cannot be closely nested inside 'section' region; perhaps you forget to enclose 'omp teams distribute parallel for simd' directive into a target region?}}
   }
 #pragma omp sections
   {
 #pragma omp section
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'section' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp sections
+  {
+#pragma omp section
+#pragma omp target teams // OK
+    a++;
+  }
 
 // SINGLE DIRECTIVE
 #pragma omp single
   {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'single' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
       ;
   }
 #pragma omp single
   {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'single' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp single
+  {
+#pragma omp target teams // oK
+    a++;
+  }
 
 // MASTER DIRECTIVE
 #pragma omp master
   {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'master' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
       ;
   }
 #pragma omp master
   {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'master' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp master
+  {
+#pragma omp target teams // OK
+    a++;
+  }
 
 // CRITICAL DIRECTIVE
 #pragma omp critical
   {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'critical' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
       ;
   }
 #pragma omp critical
   {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'critical' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp critical
+  {
+#pragma omp target teams // OK
+    a++;
+  }
 
 // PARALLEL FOR DIRECTIVE
 #pragma omp parallel for
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'parallel for' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int j = 0; j < 10; ++j)
       ;
   }
 #pragma omp parallel for
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'parallel for' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int j = 0; j < 10; ++j)
       ;
   }
+#pragma omp parallel for
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // OK
+    a++;
+  }
 
 // PARALLEL FOR SIMD DIRECTIVE
 #pragma omp parallel for simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int j = 0; j < 10; ++j)
       ;
   }
 #pragma omp parallel for simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int j = 0; j < 10; ++j)
       ;
   }
+#pragma omp parallel for simd
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // expected-error {{OpenMP constructs may not be nested inside a simd region}}
+    a++;
+  }
 
 // PARALLEL SECTIONS DIRECTIVE
 #pragma omp parallel sections
   {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'parallel sections' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
       ;
   }
 #pragma omp parallel sections
   {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'parallel sections' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp parallel sections
+  {
+#pragma omp target teams // OK
+    a++;
+  }
 
 // TASK DIRECTIVE
 #pragma omp task
 #pragma omp for // expected-error {{region cannot be closely nested inside 'task' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
   for (int i = 0; i < 10; ++i)
     ;
 #pragma omp task
 #pragma omp simd
for (int i = 0; i < 10; ++i)
       ;
   }
 #pragma omp task
   {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'task' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp task
+  {
+#pragma omp target teams // OK
+    a++;
+  }
 
 // ATOMIC DIRECTIVE
 #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 for // expected-error {{OpenMP constructs may not be nested inside an atomic region}}
     for (int i = 0; i < 10; ++i)

 #pragma omp atomic
   // expected-error@+2 {{the statement for 'atomic' must be an expression statement of form '++x;', '--x;', 'x++;', 'x--;', 'x binop= expr;', 'x = x binop expr' or 'x = expr binop x', where x is an l-value expression with scalar type}}
   // expected-note@+1 {{expected an expression statement}}
   {
 #pragma omp teams distribute parallel for // expected-error {{OpenMP constructs may not be nested inside an atomic region}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp atomic
+  // expected-error@+2 {{the statement for 'atomic' must be an expression statement of form '++x;', '--x;', 'x++;', 'x--;', 'x binop= expr;', 'x = x binop expr' or 'x = expr binop x', where x is an l-value expression with scalar type}}
+  // expected-note@+1 {{expected an expression statement}}
+  {
+#pragma omp target teams // expected-error {{OpenMP constructs may not be nested inside an atomic region}}
+    a++;
+  }
 
 // TARGET DIRECTIVE
 #pragma omp target
 #pragma omp parallel
   bar();
 #pragma omp target
 #pragma omp for
   for (int i = 0; i < 10; ++i)
for (int i = 0; i < 10; ++i)
       ;
   }
 #pragma omp target
   { 
 #pragma omp teams distribute parallel for // OK
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp target
+  { 
+#pragma omp target teams // expected-error {{region cannot be nested inside 'target' region}}
+    a++;
+  }
 
 // TARGET PARALLEL DIRECTIVE
 #pragma omp target parallel
 #pragma omp parallel
   bar();
 #pragma omp target parallel
 #pragma omp for
   for (int i = 0; i < 10; ++i)
for (int i = 0; i < 10; ++i)
       ;
   }
 #pragma omp target parallel
   { 
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'target parallel' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp target parallel
+  { 
+#pragma omp target teams // expected-error {{region cannot be nested inside 'target parallel' region}}
+    a++;
+  }
 
 // TARGET PARALLEL FOR DIRECTIVE
 #pragma omp target parallel for
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'target parallel for' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int j = 0; j < 10; ++j)
       ;
   }
 #pragma omp target parallel for
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'target parallel for' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int j = 0; j < 10; ++j)
       ;
   }
+#pragma omp target parallel for
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // expected-error {{region cannot be nested inside 'target parallel for' region}}
+    a++;
+  }
 
 // TEAMS DIRECTIVE
 #pragma omp target
 #pragma omp teams
 #pragma omp parallel
   bar();
 #pragma omp target
 #pragma omp teams

 #pragma omp teams
   {
 #pragma omp atomic // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp atomic' directive into a parallel region?}}
     ++a;
   }
 #pragma omp target
 #pragma omp teams
   {
-#pragma omp target // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp target' directive into a parallel region?}}
+#pragma omp target // expected-error {{region cannot be nested inside 'target' region}}
     ++a;
   }
 #pragma omp target
 #pragma omp teams
   {
 #pragma omp target parallel // expected-error {{region cannot be nested inside 'target' region}}
     ++a;
   }
 #pragma omp target
 #pragma omp teams
 #pragma omp target parallel for // expected-error {{region cannot be nested inside 'target' region}}
   for (int i = 0; i < 10; ++i)
     ;
 #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 enter data map(to: a) // expected-error {{region cannot be nested inside 'target' region}}
   }
 #pragma omp target
 #pragma omp teams
   {
-#pragma omp target exit data map(from: a) // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp target exit data' directive into a parallel region?}}
+#pragma omp target exit data map(from: a) // expected-error {{region cannot be nested inside 'target' region}}
   }
 #pragma omp target
 #pragma omp teams
   {
 #pragma omp teams // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
     ++a;
   }
 #pragma omp target
#pragma omp distribute
   for (int i = 0; i < 10; ++i)
     ;
 #pragma omp distribute
   for (int j = 0; j < 10; ++j)
     ;
 #pragma omp target
 #pragma omp teams
   {
-#pragma omp target update to(a) // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp target update' directive into a parallel region?}}
+#pragma omp target update to(a) // expected-error {{region cannot be nested inside 'target' region}}
     ++a;
   }
 #pragma omp target
 #pragma omp teams
 #pragma omp distribute parallel for
   for (int i = 0; i < 10; ++i)
     ;
 #pragma omp target
for (int i = 0; i < 10; ++i)
     ;
 #pragma omp distribute simd
   for (int j = 0; j < 10; ++j)
     ;  
   }
 #pragma omp target
 #pragma omp teams
   {
-#pragma omp target simd // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp target simd' directive into a parallel region?}}
+#pragma omp target simd // expected-error {{region cannot be nested inside 'target' region}}
   for (int i = 0; i < 10; ++i)
     ;
   }
 #pragma omp target
 #pragma omp teams
   {
 #pragma omp teams distribute // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp teams distribute' directive into a target region?}}
   for (int i = 0; i < 10; ++i)
#pragma omp teams distribute parallel for simd // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp teams distribute parallel for simd' directive into a target region?}}
   }
 #pragma omp target
 #pragma omp teams
   {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'teams' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
   for (int i = 0; i < 10; ++i)
     ;
   }
+#pragma omp target
+#pragma omp teams
+  {
+#pragma omp target teams // expected-error {{region cannot be nested inside 'target' region}}
+    a++;
+  }
 
 // TASKLOOP DIRECTIVE
 #pragma omp taskloop
   for (int i = 0; i < 10; ++i) {
 #pragma omp for // expected-error {{region cannot be closely nested inside 'taskloop' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
       ;
   }
 #pragma omp taskloop
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'taskloop' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp taskloop
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // OK
+    a++;
+  }
 
 // DISTRIBUTE DIRECTIVE
 #pragma omp target
 #pragma omp teams
 #pragma omp distribute
   for (int i = 0; i < 10; ++i) {
 #pragma omp distribute // expected-error {{region cannot be closely nested inside 'distribute' region; perhaps you forget to enclose 'omp distribute' directive into a teams region?}}
     for (int i = 0; i < 10; ++i)

 #pragma omp target
 #pragma omp teams
 #pragma omp distribute
   for (int i = 0; i < 10; ++i) {
 #pragma omp target parallel for // expected-error {{region cannot be nested inside 'target' region}}
   for (int i = 0; i < 10; ++i)
     ++a;
   }
+#pragma omp target
+#pragma omp teams
+#pragma omp distribute
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // expected-error {{region cannot be nested inside 'target' region}}
+  for (int i = 0; i < 10; ++i)
+    ++a;
+  }
 
 // DISTRIBUTE PARALLEL FOR DIRECTIVE
 #pragma omp target
 #pragma omp teams
 #pragma omp distribute parallel for
   for (int i = 0; i < 10; ++i) {
 #pragma omp distribute parallel for // expected-error {{region cannot be closely nested inside 'distribute parallel for' region; perhaps you forget to enclose 'omp distribute parallel for' directive into a teams region?}}
     for (int i = 0; i < 10; ++i)

 #pragma omp target
 #pragma omp teams
 #pragma omp distribute parallel for
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'distribute parallel for' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
   for (int i = 0; i < 10; ++i)
     ++a;
   }
+#pragma omp target
+#pragma omp teams
+#pragma omp distribute parallel for
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // expected-error {{region cannot be nested inside 'target' region}}
+    ++a;
+  }
 
 // DISTRIBUTE PARALLEL FOR SIMD DIRECTIVE
 #pragma omp target
 #pragma omp teams
 #pragma omp distribute parallel for simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp distribute parallel for simd // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int i = 0; i < 10; ++i)

 #pragma omp target
 #pragma omp teams
 #pragma omp distribute parallel for simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp target
+#pragma omp teams
+#pragma omp distribute parallel for simd
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // expected-error {{OpenMP constructs may not be nested inside a simd region}}
+    a++;
+  }
 
 // DISTRIBUTE SIMD DIRECTIVE
 #pragma omp target
 #pragma omp teams
 #pragma omp distribute simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp distribute parallel for simd // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int i = 0; i < 10; ++i)

 #pragma omp target
 #pragma omp teams
 #pragma omp distribute simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp target
+#pragma omp teams
+#pragma omp distribute simd
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // expected-error {{OpenMP constructs may not be nested inside a simd region}}
+    a++;
+  }
 
 // TARGET SIMD DIRECTIVE
 #pragma omp target simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp distribute parallel for simd // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int i = 0; i < 10; ++i)
       ;
   }
for (int i = 0; i < 10; ++i)
       ;
   }
 #pragma omp target simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int i = 0; i < 10; ++i)
       ;
   }
+#pragma omp target simd
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // expected-error {{OpenMP constructs may not be nested inside a simd region}}
+    a++;
+  }
 
 // TEAMS DISTRIBUTE DIRECTIVE
 #pragma omp teams distribute // expected-error {{orphaned 'omp teams distribute' directives are prohibited; perhaps you forget to enclose the directive into a target region?}}
   for (int i = 0; i < 10; ++i)
     ;
 #pragma omp target
 #pragma omp teams distribute
   for (int i = 0; i < 10; ++i) {
#pragma omp teams distribute parallel for simd // expected-error {{region cannot be closely nested inside 'teams distribute' region; perhaps you forget to enclose 'omp teams distribute parallel for simd' directive into a target region?}}
   }
 #pragma omp target
 #pragma omp teams distribute
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'teams distribute' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int i = 0; i < 10; ++i)
       ++a;
   }
+#pragma omp target
+#pragma omp teams distribute
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // expected-error {{region cannot be nested inside 'target' region}}
+    ++a;
+  }
 
 // TEAMS DISTRIBUTE SIMD DIRECTIVE
 #pragma omp teams distribute simd // expected-error {{orphaned 'omp teams distribute simd' directives are prohibited; perhaps you forget to enclose the directive into a target region?}}
   for (int i = 0; i < 10; ++i)
     ;
 #pragma omp target
 #pragma omp teams distribute simd
   for (int i = 0; i < 10; ++i) {
#pragma omp teams distribute parallel for simd // expected-error {{OpenMP constructs may not be nested inside a simd region}}
   }
 #pragma omp target
 #pragma omp teams distribute simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int i = 0; i < 10; ++i)
       ++a;
   }
+#pragma omp target
+#pragma omp teams distribute simd
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // expected-error {{OpenMP constructs may not be nested inside a simd region}}
+    ++a;
+  }
 
 // TEAMS DISTRIBUTE PARALLEL FOR SIMD DIRECTIVE
 #pragma omp teams distribute parallel for simd // expected-error {{orphaned 'omp teams distribute parallel for simd' directives are prohibited; perhaps you forget to enclose the directive into a target region?}}
   for (int i = 0; i < 10; ++i)
     ;
 #pragma omp target
 #pragma omp teams distribute parallel for simd
   for (int i = 0; i < 10; ++i) {
#pragma omp teams distribute parallel for simd // expected-error {{OpenMP constructs may not be nested inside a simd region}}
   }
 #pragma omp target
 #pragma omp teams distribute parallel for simd
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{OpenMP constructs may not be nested inside a simd region}}
     for (int i = 0; i < 10; ++i)
       ++a;
   }
-
-//--------------------
+#pragma omp target
+#pragma omp teams distribute parallel for simd
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // expected-error {{OpenMP constructs may not be nested inside a simd region}}
+    ++a;
+  }
 
 // TEAMS DISTRIBUTE PARALLEL FOR DIRECTIVE
 #pragma omp teams distribute parallel for // expected-error {{orphaned 'omp teams distribute parallel for' directives are prohibited; perhaps you forget to enclose the directive into a target region?}}
   for (int i = 0; i < 10; ++i)
     ;
 #pragma omp target
 #pragma omp teams distribute parallel for
   for (int i = 0; i < 10; ++i) {
#pragma omp teams distribute parallel for simd // expected-error {{region cannot be closely nested inside 'teams distribute parallel for' region; perhaps you forget to enclose 'omp teams distribute parallel for simd' directive into a target region?}}
   }
 #pragma omp target
 #pragma omp teams distribute parallel for
   for (int i = 0; i < 10; ++i) {
 #pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'teams distribute parallel for' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
     for (int i = 0; i < 10; ++i)
       ++a;
   }
+#pragma omp target
+#pragma omp teams distribute parallel for
+  for (int i = 0; i < 10; ++i) {
+#pragma omp target teams // expected-error {{region cannot be nested inside 'target' region}}
+    ++a;
+  }
+
+// TARGET TEAMS DIRECTIVE
+#pragma omp target teams
+#pragma omp parallel
+  bar();
+#pragma omp target teams
+#pragma omp for // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp for' directive into a parallel region?}}
+  for (int i = 0; i < 10; ++i)
+    ;
+#pragma omp target teams
+#pragma omp simd // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp simd' directive into a parallel region?}}
+  for (int i = 0; i < 10; ++i)
+    ;
+#pragma omp target teams
+#pragma omp for simd // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp for simd' directive into a parallel region?}}
+  for (int i = 0; i < 10; ++i)
+    ;
+#pragma omp target teams
+#pragma omp sections // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp sections' directive into a parallel region?}}
+  {
+    bar();
+  }
+#pragma omp target teams
+#pragma omp section // expected-error {{'omp section' directive must be closely nested to a sections region, not a target teams region}}
+  {
+    bar();
+  }
+#pragma omp target teams
+#pragma omp single // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp single' directive into a parallel region?}}
+  bar();
+
+#pragma omp target teams
+#pragma omp master // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp master' directive into a parallel region?}}
+  {
+    bar();
+  }
+#pragma omp target teams
+#pragma omp critical // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp critical' directive into a parallel region?}}
+  {
+    bar();
+  }
+#pragma omp target teams
+#pragma omp parallel for
+  for (int i = 0; i < 10; ++i)
+    ;
+#pragma omp target teams
+#pragma omp parallel for simd
+  for (int i = 0; i < 10; ++i)
+    ;
+#pragma omp target teams
+#pragma omp parallel sections
+  {
+    bar();
+  }
+#pragma omp target teams
+#pragma omp task // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp task' directive into a parallel region?}}
+  {
+    bar();
+  }
+#pragma omp target teams
+  {
+#pragma omp taskyield // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp taskyield' directive into a parallel region?}}
+    bar();
+  }
+#pragma omp target teams
+  {
+#pragma omp barrier // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp barrier' directive into a parallel region?}}
+    bar();
+  }
+#pragma omp target teams
+  {
+#pragma omp taskwait // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp taskwait' directive into a parallel region?}}
+    bar();
+  }
+#pragma omp target teams
+  {
+#pragma omp flush // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp flush' directive into a parallel region?}}
+    bar();
+  }
+#pragma omp target teams
+  {
+#pragma omp ordered // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp ordered' directive into a for or a parallel for region with 'ordered' clause?}}
+    bar();
+  }
+#pragma omp target teams
+  {
+#pragma omp atomic // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp atomic' directive into a parallel region?}}
+    ++a;
+  }
+#pragma omp target teams
+  {
+#pragma omp target // expected-error {{region cannot be nested inside 'target teams' region}}
+    ++a;
+  }
+#pragma omp target teams
+  {
+#pragma omp target parallel // expected-error {{region cannot be nested inside 'target teams' region}}
+    ++a;
+  }
+#pragma omp target teams
+#pragma omp target parallel for // expected-error {{region cannot be nested inside 'target teams' region}}
+  for (int i = 0; i < 10; ++i)
+    ;
+#pragma omp target teams
+  {
+#pragma omp target enter data map(to: a) // expected-error {{region cannot be nested inside 'target teams' region}}
+  }
+#pragma omp target teams
+  {
+#pragma omp target exit data map(from: a) // expected-error {{region cannot be nested inside 'target teams' region}}
+  }
+#pragma omp target teams
+  {
+#pragma omp teams // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp teams' directive into a target region?}}
+    ++a;
+  }
+#pragma omp target teams
+  {
+#pragma omp taskloop // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp taskloop' directive into a parallel region?}}
+  for (int i = 0; i < 10; ++i)
+    ++a;
+  }
+#pragma omp target teams
+#pragma omp distribute
+  for (int i = 0; i < 10; ++i)
+    ;
+#pragma omp target teams
+  {
+#pragma omp target update to(a) // expected-error {{region cannot be nested inside 'target teams' region}}
+    ++a;
+  }
+#pragma omp target teams
+#pragma omp distribute parallel for
+  for (int i = 0; i < 10; ++i)
+    ;
+#pragma omp target teams
+  {
+#pragma omp distribute parallel for
+  for (int i = 0; i < 10; ++i)
+    ;
+#pragma omp distribute parallel for
+  for (int j = 0; j < 10; ++j)
+    ;  
+  }
+#pragma omp target teams
+  {
+#pragma omp distribute parallel for simd
+  for (int i = 0; i < 10; ++i)
+    ;
+#pragma omp distribute parallel for simd
+  for (int j = 0; j < 10; ++j)
+    ;  
+  }
+#pragma omp target teams
+  {
+#pragma omp distribute simd
+  for (int i = 0; i < 10; ++i)
+    ;
+#pragma omp distribute simd
+  for (int j = 0; j < 10; ++j)
+    ;  
+  }
+#pragma omp target teams
+  {
+#pragma omp target simd // expected-error {{region cannot be nested inside 'target teams' region}}
+  for (int i = 0; i < 10; ++i)
+    ;
+  }
+#pragma omp target teams
+  {
+#pragma omp teams distribute // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp teams distribute' directive into a target region?}}
+  for (int i = 0; i < 10; ++i)
+    ;
+  }
+#pragma omp target teams
+  {
+#pragma omp teams distribute simd // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp teams distribute simd' directive into a target region?}}
+  for (int i = 0; i < 10; ++i)
+    ;
+  }
+#pragma omp target teams
+  {
+#pragma omp teams distribute parallel for simd // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp teams distribute parallel for simd' directive into a target region?}}
+  for (int i = 0; i < 10; ++i)
+    ;
+  }
+#pragma omp target teams
+  {
+#pragma omp teams distribute parallel for // expected-error {{region cannot be closely nested inside 'target teams' region; perhaps you forget to enclose 'omp teams distribute parallel for' directive into a target region?}}
+  for (int i = 0; i < 10; ++i)
+    ;
+  }
+#pragma omp target teams
+  {
+#pragma omp target teams // expected-error {{region cannot be nested inside 'target teams' region}}
+    a++;
+  }
 
   return foo<int>();
 }
 

test/OpenMP/target_teams_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() {}
+
+template <class T>
+struct S {
+  operator T() {return T();}
+  static T TS;
+  #pragma omp threadprivate(TS)
+};
+
+// CHECK:      template <class T> struct S {
+// CHECK:        static T TS;
+// CHECK-NEXT:   #pragma omp threadprivate(S::TS)
+// CHECK:      };
+// CHECK:      template<> struct S<int> {
+// CHECK:        static int TS;
+// CHECK-NEXT:   #pragma omp threadprivate(S<int>::TS)
+// CHECK-NEXT: }
+// CHECK:      template<> struct S<long> {
+// CHECK:        static long TS;
+// CHECK-NEXT:   #pragma omp threadprivate(S<long>::TS)
+// CHECK-NEXT: }
+
+template <typename T, int C>
+T tmain(T argc, T *argv) {
+  T b = argc, c, d, e, f, g;
+  static T a;
+  S<T> s;
+#pragma omp target teams
+  a=2;
+#pragma omp target teams default(none), private(argc,b) firstprivate(argv) shared (d) reduction(+:c) reduction(max:e) num_teams(C) thread_limit(d*C)
+  foo();
+#pragma omp target teams reduction(^:e, f) reduction(&& : g)
+  foo();
+  return 0;
+}
+
+// CHECK: template <typename T, int C> T tmain(T argc, T *argv) {
+// CHECK-NEXT: T b = argc, c, d, e, f, g;
+// CHECK-NEXT: static T a;
+// CHECK-NEXT: S<T> s;
+// CHECK-NEXT: #pragma omp target teams
+// CHECK-NEXT: a = 2;
+// CHECK-NEXT: #pragma omp target teams default(none) private(argc,b) firstprivate(argv) shared(d) reduction(+: c) reduction(max: e) num_teams(C) thread_limit(d * C)
+// CHECK-NEXT: foo()
+// CHECK-NEXT: #pragma omp target teams reduction(^: e,f) reduction(&&: g)
+// CHECK-NEXT: foo()
+// CHECK: template<> int tmain<int, 5>(int argc, int *argv) {
+// CHECK-NEXT: int b = argc, c, d, e, f, g;
+// CHECK-NEXT: static int a;
+// CHECK-NEXT: S<int> s;
+// CHECK-NEXT: #pragma omp target teams
+// CHECK-NEXT: a = 2;
+// CHECK-NEXT: #pragma omp target teams default(none) private(argc,b) firstprivate(argv) shared(d) reduction(+: c) reduction(max: e) num_teams(5) thread_limit(d * 5)
+// CHECK-NEXT: foo()
+// CHECK-NEXT: #pragma omp target teams reduction(^: e,f) reduction(&&: g)
+// CHECK-NEXT: foo()
+// CHECK: template<> long tmain<long, 1>(long argc, long *argv) {
+// CHECK-NEXT: long b = argc, c, d, e, f, g;
+// CHECK-NEXT: static long a;
+// CHECK-NEXT: S<long> s;
+// CHECK-NEXT: #pragma omp target teams
+// CHECK-NEXT: a = 2;
+// CHECK-NEXT: #pragma omp target teams default(none) private(argc,b) firstprivate(argv) shared(d) reduction(+: c) reduction(max: e) num_teams(1) thread_limit(d * 1)
+// CHECK-NEXT: foo()
+// CHECK-NEXT: #pragma omp target teams reduction(^: e,f) reduction(&&: g)
+// CHECK-NEXT: foo()
+
+enum Enum { };
+
+int main (int argc, char **argv) {
+  long x;
+  int b = argc, c, d, e, f, g;
+  static int a;
+  #pragma omp threadprivate(a)
+  Enum ee;
+// CHECK: Enum ee;
+#pragma omp target teams
+// CHECK-NEXT: #pragma omp target teams
+  a=2;
+// CHECK-NEXT: a = 2;
+#pragma omp target teams default(none), private(argc,b) num_teams(f) firstprivate(argv) reduction(| : c, d) reduction(* : e) thread_limit(f+g)
+// CHECK-NEXT: #pragma omp target teams default(none) private(argc,b) num_teams(f) firstprivate(argv) reduction(|: c,d) reduction(*: e) thread_limit(f + g)
+  foo();
+// CHECK-NEXT: foo();
+  return tmain<int, 5>(b, &b) + tmain<long, 1>(x, &x);
+}
+
+extern template int S<int>::TS;
+extern template long S<long>::TS;
+#endif

test/OpenMP/target_teams_default_messages.cpp

+// RUN: %clang_cc1 -verify -fopenmp -o - %s
+
+void foo();
+
+int main(int argc, char **argv) {
+#pragma omp target teams default // expected-error {{expected '(' after 'default'}}
+  foo();
+#pragma omp target teams default ( // expected-error {{expected 'none' or 'shared' in OpenMP clause 'default'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams default () // expected-error {{expected 'none' or 'shared' in OpenMP clause 'default'}}
+  foo();
+#pragma omp target teams default (none // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams default (shared), default(shared) // expected-error {{directive '#pragma omp target teams' cannot contain more than one 'default' clause}}
+  foo();
+#pragma omp target teams default (x) // expected-error {{expected 'none' or 'shared' in OpenMP clause 'default'}}
+  foo();
+
+#pragma omp target teams default(none)
+  ++argc; // expected-error {{variable 'argc' must have explicitly specified data sharing attributes}}
+
+#pragma omp target teams default(none)
+#pragma omp parallel default(shared)
+  ++argc;
+  return 0;
+}

test/OpenMP/target_teams_defaultmap_messages.cpp

+// RUN: %clang_cc1 -verify -fopenmp %s
+
+void foo() {
+}
+
+template <class T, typename S, int N, int ST>
+T tmain(T argc, S **argv) {
+  #pragma omp target teams defaultmap // expected-error {{expected '(' after 'defaultmap'}}
+  foo();
+  #pragma omp target teams defaultmap ( // expected-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+  #pragma omp target teams defaultmap () // expected-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}}
+  foo();
+  #pragma omp target teams defaultmap (tofrom // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-warning {{missing ':' after defaultmap modifier - ignoring}} expected-error {{expected 'scalar' in OpenMP clause 'defaultmap'}}
+  foo();
+  #pragma omp target teams defaultmap (tofrom: // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected 'scalar' in OpenMP clause 'defaultmap'}}
+  foo();
+  #pragma omp target teams defaultmap (tofrom) // expected-warning {{missing ':' after defaultmap modifier - ignoring}} expected-error {{expected 'scalar' in OpenMP clause 'defaultmap'}}
+  foo();
+  #pragma omp target teams defaultmap (tofrom scalar) // expected-warning {{missing ':' after defaultmap modifier - ignoring}}
+  foo();
+  #pragma omp target teams defaultmap (tofrom, // expected-error {{expected ')'}} expected-error {{expected 'scalar' in OpenMP clause 'defaultmap'}} expected-warning {{missing ':' after defaultmap modifier - ignoring}} expected-note {{to match this '('}}
+  foo();
+  #pragma omp target teams defaultmap (scalar: // expected-error {{expected ')'}} expected-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}} expected-note {{to match this '('}}
+  foo();
+  #pragma omp target teams defaultmap (tofrom, scalar // expected-error {{expected ')'}} expected-warning {{missing ':' after defaultmap modifier - ignoring}} expected-error {{expected 'scalar' in OpenMP clause 'defaultmap'}} expected-note {{to match this '('}}
+  foo();
+
+  return argc;
+}
+
+int main(int argc, char **argv) {
+  #pragma omp target teams defaultmap // expected-error {{expected '(' after 'defaultmap'}}
+  foo();
+  #pragma omp target teams defaultmap ( // expected-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+  #pragma omp target teams defaultmap () // expected-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}}
+  foo();
+  #pragma omp target teams defaultmap (tofrom // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-warning {{missing ':' after defaultmap modifier - ignoring}} expected-error {{expected 'scalar' in OpenMP clause 'defaultmap'}}
+  foo();
+  #pragma omp target teams defaultmap (tofrom: // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected 'scalar' in OpenMP clause 'defaultmap'}}
+  foo();
+  #pragma omp target teams defaultmap (tofrom) // expected-warning {{missing ':' after defaultmap modifier - ignoring}} expected-error {{expected 'scalar' in OpenMP clause 'defaultmap'}}
+  foo();
+  #pragma omp target teams defaultmap (tofrom scalar) // expected-warning {{missing ':' after defaultmap modifier - ignoring}}
+  foo();
+  #pragma omp target teams defaultmap (tofrom, // expected-error {{expected ')'}} expected-error {{expected 'scalar' in OpenMP clause 'defaultmap'}} expected-warning {{missing ':' after defaultmap modifier - ignoring}} expected-note {{to match this '('}}
+  foo();
+  #pragma omp target teams defaultmap (scalar: // expected-error {{expected ')'}} expected-error {{expected 'tofrom' in OpenMP clause 'defaultmap'}} expected-note {{to match this '('}}
+  foo();
+  #pragma omp target teams defaultmap (tofrom, scalar // expected-error {{expected ')'}} expected-warning {{missing ':' after defaultmap modifier - ignoring}} expected-error {{expected 'scalar' in OpenMP clause 'defaultmap'}} expected-note {{to match this '('}}
+  foo();
+
+  return tmain<int, char, 1, 0>(argc, argv);
+}
+

test/OpenMP/target_teams_depend_messages.cpp

+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 100 -o - -std=c++11 %s
+
+void foo() {
+}
+
+bool foobool(int argc) {
+  return argc;
+}
+
+struct S1; // expected-note {{declared here}}
+
+class vector {
+  public:
+    int operator[](int index) { return 0; }
+};
+
+int main(int argc, char **argv, char *env[]) {
+  vector vec;
+  typedef float V __attribute__((vector_size(16)));
+  V a;
+  auto arr = x; // expected-error {{use of undeclared identifier 'x'}}
+
+#pragma omp target teams depend // expected-error {{expected '(' after 'depend'}}
+  foo();
+#pragma omp target teams depend ( // expected-error {{expected 'in', 'out' or 'inout' in OpenMP clause 'depend'}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-warning {{missing ':' after dependency type - ignoring}}
+  foo();
+#pragma omp target teams depend () // expected-error {{expected 'in', 'out' or 'inout' in OpenMP clause 'depend'}} expected-warning {{missing ':' after dependency type - ignoring}}
+  foo();
+#pragma omp target teams depend (argc // expected-error {{expected 'in', 'out' or 'inout' in OpenMP clause 'depend'}} expected-warning {{missing ':' after dependency type - ignoring}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams depend (source : argc) // expected-error {{expected 'in', 'out' or 'inout' in OpenMP clause 'depend'}}
+  foo();
+#pragma omp target teams depend (source) // expected-error {{expected expression}} expected-warning {{missing ':' after dependency type - ignoring}}
+  foo();
+#pragma omp target teams depend (in : argc)) // expected-warning {{extra tokens at the end of '#pragma omp target teams' are ignored}}
+  foo();
+#pragma omp target teams depend (out: ) // expected-error {{expected expression}}
+  foo();
+#pragma omp target teams depend (inout : foobool(argc)), depend (in, argc) // expected-error {{expected variable name, array element or array section}} expected-warning {{missing ':' after dependency type - ignoring}} expected-error {{expected expression}}
+  foo();
+#pragma omp target teams depend (out :S1) // expected-error {{'S1' does not refer to a value}}
+  foo();
+#pragma omp target teams depend(in : argv[1][1] = '2') // expected-error {{expected variable name, array element or array section}}
+  foo();
+#pragma omp target teams depend (in : vec[1]) // expected-error {{expected variable name, array element or array section}}
+  foo();
+#pragma omp target teams depend (in : argv[0])
+  foo();
+#pragma omp target teams depend (in : ) // expected-error {{expected expression}}
+  foo();
+#pragma omp target teams depend (in : main) // expected-error {{expected variable name, array element or array section}}
+  foo();
+#pragma omp target teams depend(in : a[0]) // expected-error{{expected variable name, array element or array section}}
+  foo();
+#pragma omp target teams depend (in : vec[1:2]) // expected-error {{ value is not an array or pointer}}
+  foo();
+#pragma omp target teams depend (in : argv[ // expected-error {{expected expression}} expected-error {{expected ']'}} expected-error {{expected ')'}} expected-note {{to match this '['}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams depend (in : argv[: // expected-error {{expected expression}} expected-error {{expected ']'}} expected-error {{expected ')'}} expected-note {{to match this '['}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams depend (in : argv[:] // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams depend (in : argv[argc: // expected-error {{expected expression}} expected-error {{expected ']'}} expected-error {{expected ')'}} expected-note {{to match this '['}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams depend (in : argv[argc:argc] // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams depend (in : argv[0:-1]) // expected-error {{section length is evaluated to a negative value -1}}
+  foo();
+#pragma omp target teams depend (in : argv[-1:0])
+  foo();
+#pragma omp target teams depend (in : argv[:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+  foo();
+#pragma omp target teams depend (in : argv[3:4:1]) // expected-error {{expected ']'}} expected-note {{to match this '['}}
+  foo();
+#pragma omp target teams depend(in:a[0:1]) // expected-error {{subscripted value is not an array or pointer}}
+  foo();
+#pragma omp target teams depend(in:argv[argv[:2]:1]) // expected-error {{OpenMP array section is not allowed here}}
+  foo();
+#pragma omp target teams depend(in:argv[0:][:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+  foo();
+#pragma omp target teams depend(in:env[0:][:]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is an array of unknown bound}}
+  foo();
+#pragma omp target teams depend(in : argv[ : argc][1 : argc - 1])
+  foo();
+#pragma omp target teams depend(in : arr[0])
+  foo();
+
+  return 0;
+}

test/OpenMP/target_teams_device_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 teams device // expected-error {{expected '(' after 'device'}}
+  foo();
+#pragma omp target teams device ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams device () // expected-error {{expected expression}}
+  foo();
+#pragma omp target teams device (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams device (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target teams' are ignored}}
+  foo();
+#pragma omp target teams device (argc > 0 ? argv[1] : argv[2]) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}}
+  foo();
+#pragma omp target teams device (argc + argc)
+  foo();
+#pragma omp target teams device (argc), device (argc+1) // expected-error {{directive '#pragma omp target teams' cannot contain more than one 'device' clause}}
+  foo();
+#pragma omp target teams device (S1) // expected-error {{'S1' does not refer to a value}}
+  foo();
+#pragma omp target teams device (-2) // expected-error {{argument to 'device' clause must be a non-negative integer value}}
+  foo();
+#pragma omp target teams device (-10u)
+  foo();
+#pragma omp target teams device (3.14) // expected-error {{expression must have integral or unscoped enumeration type, not 'double'}}
+  foo();
+
+  return 0;
+}

test/OpenMP/target_teams_firstprivate_messages.cpp

+// RUN: %clang_cc1 -verify -fopenmp %s
+
+void foo() {
+}
+
+bool foobool(int argc) {
+  return argc;
+}
+
+struct S1; // expected-note {{declared here}} expected-note{{forward declaration of 'S1'}}
+extern S1 a;
+class S2 {
+  mutable int a;
+
+public:
+  S2() : a(0) {}
+  S2(const S2 &s2) : a(s2.a) {}
+  static float S2s;
+  static const float S2sc;
+};
+const float S2::S2sc = 0;
+const S2 b;
+const S2 ba[5];
+class S3 {
+  int a;
+
+public:
+  S3() : a(0) {}
+  S3(const S3 &s3) : a(s3.a) {}
+};
+const S3 c;
+const S3 ca[5];
+extern const int f;
+class S4 {
+  int a;
+  S4();
+  S4(const S4 &s4); // expected-note {{implicitly declared private here}}
+public:
+  S4(int v) : a(v) {}
+};
+class S5 {
+  int a;
+  S5() : a(0) {}
+  S5(const S5 &s5) : a(s5.a) {} // expected-note {{implicitly declared private here}}
+public:
+  S5(int v) : a(v) {}
+};
+
+S3 h;
+#pragma omp threadprivate(h) // expected-note {{defined as threadprivate or thread local}}
+
+namespace A {
+double x;
+#pragma omp threadprivate(x) // expected-note {{defined as threadprivate or thread local}}
+}
+namespace B {
+using A::x;
+}
+
+int main(int argc, char **argv) {
+  const int d = 5;
+  const int da[5] = {0};
+  S4 e(4);
+  S5 g(5);
+  int i;
+  int &j = i;
+#pragma omp target teams firstprivate // expected-error {{expected '(' after 'firstprivate'}}
+  foo();
+#pragma omp target teams firstprivate( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams firstprivate() // expected-error {{expected expression}}
+  foo();
+#pragma omp target teams firstprivate(argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams firstprivate(argc, // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams firstprivate(argc > 0 ? argv[1] : argv[2]) // expected-error {{expected variable name}}
+  foo();
+#pragma omp target teams firstprivate(argc)
+  foo();
+#pragma omp target teams firstprivate(S1) // expected-error {{'S1' does not refer to a value}}
+  foo();
+#pragma omp target teams firstprivate(a, b, c, d, f) // expected-error {{firstprivate variable with incomplete type 'S1'}}
+  foo();
+#pragma omp target teams firstprivate(argv[1]) // expected-error {{expected variable name}}
+  foo();
+#pragma omp target teams firstprivate(ba)
+  foo();
+#pragma omp target teams firstprivate(ca)
+  foo();
+#pragma omp target teams firstprivate(da)
+  foo();
+#pragma omp target teams firstprivate(S2::S2s)
+  foo();
+#pragma omp target teams firstprivate(S2::S2sc)
+  foo();
+#pragma omp target teams firstprivate(e, g) // expected-error {{calling a private constructor of class 'S4'}} expected-error {{calling a private constructor of class 'S5'}}
+  foo();
+#pragma omp target teams firstprivate(h, B::x) // expected-error 2 {{threadprivate or thread local variable cannot be firstprivate}}
+  foo();
+#pragma omp target teams private(i), firstprivate(i) // expected-error {{private variable cannot be firstprivate}} expected-note{{defined as private}}
+  foo();
+#pragma omp target teams shared(i)
+  foo();
+#pragma omp target teams firstprivate(i)
+  foo();
+#pragma omp target teams firstprivate(j)
+  foo();
+  static int m;
+#pragma omp target teams firstprivate(m) // OK
+  foo();
+
+  return 0;
+}

test/OpenMP/target_teams_if_messages.cpp

+// RUN: %clang_cc1 -verify -fopenmp %s
+
+void foo() {
+}
+
+bool foobool(int argc) {
+  return argc;
+}
+
+struct S1; // expected-note {{declared here}}
+
+template <class T, class S> // expected-note {{declared here}}
+int tmain(T argc, S **argv) {
+#pragma omp target teams if // expected-error {{expected '(' after 'if'}}
+  foo();
+#pragma omp target teams if ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams if () // expected-error {{expected expression}}
+  foo();
+#pragma omp target teams if (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams if (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target teams' are ignored}}
+  foo();
+#pragma omp target teams if (argc > 0 ? argv[1] : argv[2])
+  foo();
+#pragma omp target teams if (foobool(argc)), if (true) // expected-error {{directive '#pragma omp target teams' cannot contain more than one 'if' clause}}
+  foo();
+#pragma omp target teams if (S) // expected-error {{'S' does not refer to a value}}
+  foo();
+#pragma omp target teams if (argv[1]=2) // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams if (argc argc) // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams if(argc)
+  foo();
+#pragma omp target teams if(target : // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams if(target : argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams if(target : argc)
+  foo();
+#pragma omp target teams if(target : argc) if (for:argc) // expected-error {{directive name modifier 'for' is not allowed for '#pragma omp target teams'}}
+  foo();
+#pragma omp target teams if(target : argc) if (target:argc) // expected-error {{directive '#pragma omp target teams' cannot contain more than one 'if' clause with 'target' name modifier}}
+  foo();
+#pragma omp target teams if(target : argc) if (argc) // expected-error {{no more 'if' clause is allowed}} expected-note {{previous clause with directive name modifier specified here}}
+  foo();
+
+  return 0;
+}
+
+int main(int argc, char **argv) {
+#pragma omp target teams if // expected-error {{expected '(' after 'if'}}
+  foo();
+#pragma omp target teams if ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams if () // expected-error {{expected expression}}
+  foo();
+#pragma omp target teams if (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams if (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target teams' are ignored}}
+  foo();
+#pragma omp target teams if (argc > 0 ? argv[1] : argv[2])
+  foo();
+#pragma omp target teams if (foobool(argc)), if (true) // expected-error {{directive '#pragma omp target teams' cannot contain more than one 'if' clause}}
+  foo();
+#pragma omp target teams if (S1) // expected-error {{'S1' does not refer to a value}}
+  foo();
+#pragma omp target teams if (argv[1]=2) // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams if (argc argc) // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams if (1 0) // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams if(if(tmain(argc, argv) // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams if(target : // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams if(target : argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams if(target : argc)
+  foo();
+#pragma omp target teams if(target : argc) if (for:argc) // expected-error {{directive name modifier 'for' is not allowed for '#pragma omp target teams'}}
+  foo();
+#pragma omp target teams if(target : argc) if (target:argc) // expected-error {{directive '#pragma omp target teams' cannot contain more than one 'if' clause with 'target' name modifier}}
+  foo();
+#pragma omp target teams if(target : argc) if (argc) // expected-error {{no more 'if' clause is allowed}} expected-note {{previous clause with directive name modifier specified here}}
+  foo();
+
+  return tmain(argc, argv);
+}

test/OpenMP/target_teams_is_device_ptr_ast_print.cpp

+// RUN: %clang_cc1 -verify -fopenmp -std=c++11 -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
+struct ST {
+  int *a;
+};
+typedef int arr[10];
+typedef ST STarr[10];
+struct SA {
+  const int da[5] = { 0 };
+  ST g[10];
+  STarr &rg = g;
+  int i;
+  int &j = i;
+  int *k = &j;
+  int *&z = k;
+  int aa[10];
+  arr &raa = aa;
+  void func(int arg) {
+#pragma omp target teams is_device_ptr(k)
+    {}
+#pragma omp target teams is_device_ptr(z)
+    {}
+#pragma omp target teams is_device_ptr(aa) // OK
+    {}
+#pragma omp target teams is_device_ptr(raa) // OK
+    {}
+#pragma omp target teams is_device_ptr(g) // OK
+    {}
+#pragma omp target teams is_device_ptr(rg) // OK
+    {}
+#pragma omp target teams is_device_ptr(da) // OK
+    {}
+  return;
+ }
+};
+// CHECK: struct SA
+// CHECK-NEXT: const int da[5] = {0};
+// CHECK-NEXT: ST g[10];
+// CHECK-NEXT: STarr &rg = this->g;
+// CHECK-NEXT: int i;
+// CHECK-NEXT: int &j = this->i;
+// CHECK-NEXT: int *k = &this->j;
+// CHECK-NEXT: int *&z = this->k;
+// CHECK-NEXT: int aa[10];
+// CHECK-NEXT: arr &raa = this->aa;
+// CHECK-NEXT: func(
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(this->k)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(this->z)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(this->aa)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(this->raa)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(this->g)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(this->rg)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(this->da)
+
+struct SB {
+  unsigned A;
+  unsigned B;
+  float Arr[100];
+  float *Ptr;
+  float *foo() {
+    return &Arr[0];
+  }
+};
+
+struct SC {
+  unsigned A : 2;
+  unsigned B : 3;
+  unsigned C;
+  unsigned D;
+  float Arr[100];
+  SB S;
+  SB ArrS[100];
+  SB *PtrS;
+  SB *&RPtrS;
+  float *Ptr;
+
+  SC(SB *&_RPtrS) : RPtrS(_RPtrS) {}
+};
+
+union SD {
+  unsigned A;
+  float B;
+};
+
+struct S1;
+extern S1 a;
+class S2 {
+  mutable int a;
+public:
+  S2():a(0) { }
+  S2(S2 &s2):a(s2.a) { }
+  static float S2s;
+  static const float S2sc;
+};
+const float S2::S2sc = 0;
+const S2 b;
+const S2 ba[5];
+class S3 {
+  int a;
+public:
+  S3():a(0) { }
+  S3(S3 &s3):a(s3.a) { }
+};
+const S3 c;
+const S3 ca[5];
+extern const int f;
+class S4 {
+  int a;
+  S4();
+  S4(const S4 &s4);
+public:
+  S4(int v):a(v) { }
+};
+class S5 {
+  int a;
+  S5():a(0) {}
+  S5(const S5 &s5):a(s5.a) { }
+public:
+  S5(int v):a(v) { }
+};
+
+S3 h;
+#pragma omp threadprivate(h)
+
+typedef struct {
+  int a;
+} S6;
+
+template <typename T>
+T tmain(T argc) {
+  const T da[5] = { 0 };
+  S6 h[10];
+  auto &rh = h;
+  T i;
+  T &j = i;
+  T *k = &j;
+  T *&z = k;
+  T aa[10];
+  auto &raa = aa;
+#pragma omp target teams is_device_ptr(k)
+  {}
+#pragma omp target teams is_device_ptr(z)
+  {}
+#pragma omp target teams is_device_ptr(aa)
+  {}
+#pragma omp target teams is_device_ptr(raa)
+  {}
+#pragma omp target teams is_device_ptr(h)
+  {}
+#pragma omp target teams is_device_ptr(rh)
+  {}
+#pragma omp target teams is_device_ptr(da)
+  {}
+  return 0;
+}
+
+// CHECK: template<> int tmain<int>(int argc) {
+// CHECK-NEXT: const int da[5] = {0};
+// CHECK-NEXT: S6 h[10];
+// CHECK-NEXT: auto &rh = h;
+// CHECK-NEXT: int i;
+// CHECK-NEXT: int &j = i;
+// CHECK-NEXT: int *k = &j;
+// CHECK-NEXT: int *&z = k;
+// CHECK-NEXT: int aa[10];
+// CHECK-NEXT: auto &raa = aa;
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(k)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(z)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(aa)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(raa)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(h)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(rh)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(da)
+
+// CHECK: template<> int *tmain<int *>(int *argc) {
+// CHECK-NEXT: int *const da[5] = {0};
+// CHECK-NEXT: S6 h[10];
+// CHECK-NEXT: auto &rh = h;
+// CHECK-NEXT: int *i;
+// CHECK-NEXT: int *&j = i;
+// CHECK-NEXT: int **k = &j;
+// CHECK-NEXT: int **&z = k;
+// CHECK-NEXT: int *aa[10];
+// CHECK-NEXT: auto &raa = aa;
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(k)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(z)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(aa)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(raa)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(h)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(rh)
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(da)
+
+// CHECK-LABEL: int main(int argc, char **argv) {
+int main(int argc, char **argv) {
+  const int da[5] = { 0 };
+  S6 h[10];
+  auto &rh = h;
+  int i;
+  int &j = i;
+  int *k = &j;
+  int *&z = k;
+  int aa[10];
+  auto &raa = aa;
+// CHECK-NEXT: const int da[5] = {0};
+// CHECK-NEXT: S6 h[10];
+// CHECK-NEXT: auto &rh = h;
+// CHECK-NEXT: int i;
+// CHECK-NEXT: int &j = i;
+// CHECK-NEXT: int *k = &j;
+// CHECK-NEXT: int *&z = k;
+// CHECK-NEXT: int aa[10];
+// CHECK-NEXT: auto &raa = aa;
+#pragma omp target teams is_device_ptr(k)
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(k)
+  {}
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+#pragma omp target teams is_device_ptr(z)
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(z)
+  {}
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+#pragma omp target teams is_device_ptr(aa)
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(aa)
+  {}
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+#pragma omp target teams is_device_ptr(raa)
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(raa)
+  {}
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+#pragma omp target teams is_device_ptr(h)
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(h)
+  {}
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+#pragma omp target teams is_device_ptr(rh)
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(rh)
+  {}
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+#pragma omp target teams is_device_ptr(da)
+// CHECK-NEXT: #pragma omp target teams is_device_ptr(da)
+  {}
+// CHECK-NEXT: {
+// CHECK-NEXT: }
+  return tmain<int>(argc) + *tmain<int *>(&argc);
+}
+#endif

test/OpenMP/target_teams_is_device_ptr_messages.cpp

+// RUN: %clang_cc1 -std=c++11 -verify -fopenmp -ferror-limit 200 %s
+struct ST {
+  int *a;
+};
+typedef int arr[10];
+typedef ST STarr[10];
+struct SA {
+  const int d = 5;
+  const int da[5] = { 0 };
+  ST e;
+  ST g[10];
+  STarr &rg = g;
+  int i;
+  int &j = i;
+  int *k = &j;
+  int *&z = k;
+  int aa[10];
+  arr &raa = aa;
+  void func(int arg) {
+#pragma omp target teams is_device_ptr // expected-error {{expected '(' after 'is_device_ptr'}}
+    {}
+#pragma omp target teams is_device_ptr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+    {}
+#pragma omp target teams is_device_ptr() // expected-error {{expected expression}}
+    {}
+#pragma omp target teams is_device_ptr(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
+    {}
+#pragma omp target teams is_device_ptr(arg // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected pointer, array, reference to pointer, or reference to array in 'is_device_ptr clause'}}
+    {}
+#pragma omp target teams is_device_ptr(i) // expected-error {{expected pointer, array, reference to pointer, or reference to array in 'is_device_ptr clause'}}
+    {}
+#pragma omp target teams is_device_ptr(j) // expected-error {{expected pointer, array, reference to pointer, or reference to array in 'is_device_ptr clause'}}
+    {}
+#pragma omp target teams is_device_ptr(k) // OK
+    {}
+#pragma omp target teams is_device_ptr(z) // OK
+    {}
+#pragma omp target teams is_device_ptr(aa) // OK
+    {}
+#pragma omp target teams is_device_ptr(raa) // OK
+    {}    
+#pragma omp target teams is_device_ptr(e) // expected-error{{expected pointer, array, reference to pointer, or reference to array in 'is_device_ptr clause'}}
+    {}
+#pragma omp target teams is_device_ptr(g) // OK
+    {}
+#pragma omp target teams is_device_ptr(rg) // OK
+    {}
+#pragma omp target teams is_device_ptr(k,i,j) // expected-error2 {{expected pointer, array, reference to pointer, or reference to array in 'is_device_ptr clause'}}
+    {}
+#pragma omp target teams is_device_ptr(d) // expected-error{{expected pointer, array, reference to pointer, or reference to array in 'is_device_ptr clause'}}
+    {}
+#pragma omp target teams is_device_ptr(da) // OK
+    {}
+  return;
+ }
+};
+struct SB {
+  unsigned A;
+  unsigned B;
+  float Arr[100];
+  float *Ptr;
+  float *foo() {
+    return &Arr[0];
+  }
+};
+
+struct SC {
+  unsigned A : 2;
+  unsigned B : 3;
+  unsigned C;
+  unsigned D;
+  float Arr[100];
+  SB S;
+  SB ArrS[100];
+  SB *PtrS;
+  SB *&RPtrS;
+  float *Ptr;
+
+  SC(SB *&_RPtrS) : RPtrS(_RPtrS) {}
+};
+
+union SD {
+  unsigned A;
+  float B;
+};
+
+struct S1;
+extern S1 a;
+class S2 {
+  mutable int a;
+public:
+  S2():a(0) { }
+  S2(S2 &s2):a(s2.a) { }
+  static float S2s;
+  static const float S2sc;
+};
+const float S2::S2sc = 0;
+const S2 b;
+const S2 ba[5];
+class S3 {
+  int a;
+public:
+  S3():a(0) { }
+  S3(S3 &s3):a(s3.a) { }
+};
+const S3 c;
+const S3 ca[5];
+extern const int f;
+class S4 {
+  int a;
+  S4();
+  S4(const S4 &s4);
+public:
+  S4(int v):a(v) { }
+};
+class S5 {
+  int a;
+  S5():a(0) {}
+  S5(const S5 &s5):a(s5.a) { }
+public:
+  S5(int v):a(v) { }
+};
+
+S3 h;
+#pragma omp threadprivate(h)
+
+typedef struct {
+  int a;
+} S6;
+
+template <typename T, int I>
+T tmain(T argc) {
+  const T d = 5;
+  const T da[5] = { 0 };
+  S4 e(4);
+  S5 g(5);
+  S6 h[10];
+  auto &rh = h;
+  T i;
+  T &j = i;
+  T *k = &j;
+  T *&z = k;
+  T aa[10];
+  auto &raa = aa;
+  S6 *ps;
+#pragma omp target teams is_device_ptr // expected-error {{expected '(' after 'is_device_ptr'}}
+  {}
+#pragma omp target teams is_device_ptr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+  {}
+#pragma omp target teams is_device_ptr() // expected-error {{expected expression}}
+  {}
+#pragma omp target teams is_device_ptr(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
+  {}
+#pragma omp target teams is_device_ptr(argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error{{expected pointer, array, reference to pointer, or reference to array in 'is_device_ptr clause'}}
+  {}
+#pragma omp target teams is_device_ptr(i) // expected-error {{expected pointer, array, reference to pointer, or reference to array in 'is_device_ptr clause'}}
+  {}
+#pragma omp target teams is_device_ptr(j) // expected-error {{expected pointer, array, reference to pointer, or reference to array in 'is_device_ptr clause'}}
+  {}
+#pragma omp target teams is_device_ptr(k) // OK
+  {}
+#pragma omp target teams is_device_ptr(z) // OK
+  {}
+#pragma omp target teams is_device_ptr(aa) // OK
+  {}
+#pragma omp target teams is_device_ptr(raa) // OK
+  {}
+#pragma omp target teams is_device_ptr(e) // expected-error{{expected pointer, array, reference to pointer, or reference to array in 'is_device_ptr clause'}}
+  {}
+#pragma omp target teams is_device_ptr(g) // expected-error{{expected pointer, array, reference to pointer, or reference to array in 'is_device_ptr clause'}}
+  {}
+#pragma omp target teams is_device_ptr(h) // OK
+  {}
+#pragma omp target teams is_device_ptr(rh) // OK
+  {}
+#pragma omp target teams is_device_ptr(k,i,j) // expected-error2 {{expected pointer, array, reference to pointer, or reference to array in 'is_device_ptr clause'}}
+  {}
+#pragma omp target teams is_device_ptr(d) // expected-error{{expected pointer, array, reference to pointer, or reference to array in 'is_device_ptr clause'}}
+  {}
+#pragma omp target teams is_device_ptr(da) // OK
+  {}
+#pragma omp target teams map(ps) is_device_ptr(ps) // expected-error{{variable already marked as mapped in current construct}} expected-note{{used here}}
+  {}
+#pragma omp target teams is_device_ptr(ps) map(ps) // expected-error{{variable already marked as mapped in current construct}} expected-note{{used here}}
+  {}
+#pragma omp target teams map(ps->a) is_device_ptr(ps) // expected-error{{variable already marked as mapped in current construct}} expected-note{{used here}}
+  {}
+#pragma omp target teams is_device_ptr(ps) map(ps->a) // expected-error{{pointer cannot be mapped along with a section derived from itself}} expected-note{{used here}}
+  {}
+#pragma omp target teams is_device_ptr(ps) firstprivate(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target teams' directive}}
+  {}
+#pragma omp target teams firstprivate(ps) is_device_ptr(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target teams' directive}} expected-note{{defined as firstprivate}}
+  {}
+#pragma omp target teams is_device_ptr(ps) private(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target teams' directive}}
+  {}
+#pragma omp target teams private(ps) is_device_ptr(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target teams' directive}} expected-note{{defined as private}}
+  {}
+  return 0;
+}
+
+int main(int argc, char **argv) {
+  const int d = 5;
+  const int da[5] = { 0 };
+  S4 e(4);
+  S5 g(5);
+  S6 h[10];
+  auto &rh = h;
+  int i;
+  int &j = i;
+  int *k = &j;
+  int *&z = k;
+  int aa[10];
+  auto &raa = aa;
+  S6 *ps;
+#pragma omp target teams is_device_ptr // expected-error {{expected '(' after 'is_device_ptr'}}
+  {}
+#pragma omp target teams is_device_ptr( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+  {}
+#pragma omp target teams is_device_ptr() // expected-error {{expected expression}}
+  {}
+#pragma omp target teams is_device_ptr(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
+  {}
+#pragma omp target teams is_device_ptr(argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected pointer, array, reference to pointer, or reference to array in 'is_device_ptr clause'}}
+  {}
+#pragma omp target teams is_device_ptr(i) // expected-error {{expected pointer, array, reference to pointer, or reference to array in 'is_device_ptr clause'}}
+  {}
+#pragma omp target teams is_device_ptr(j) // expected-error {{expected pointer, array, reference to pointer, or reference to array in 'is_device_ptr clause'}}
+  {}
+#pragma omp target teams is_device_ptr(k) // OK
+  {}
+#pragma omp target teams is_device_ptr(z) // OK
+  {}
+#pragma omp target teams is_device_ptr(aa) // OK
+  {}
+#pragma omp target teams is_device_ptr(raa) // OK
+  {}
+#pragma omp target teams is_device_ptr(e) // expected-error{{expected pointer, array, reference to pointer, or reference to array in 'is_device_ptr clause'}}
+  {}
+#pragma omp target teams is_device_ptr(g) // expected-error{{expected pointer, array, reference to pointer, or reference to array in 'is_device_ptr clause'}}
+  {}
+#pragma omp target teams is_device_ptr(h) // OK
+  {}
+#pragma omp target teams is_device_ptr(rh) // OK
+  {}
+#pragma omp target teams is_device_ptr(k,i,j) // expected-error2 {{expected pointer, array, reference to pointer, or reference to array in 'is_device_ptr clause'}}
+  {}
+#pragma omp target teams is_device_ptr(d) // expected-error{{expected pointer, array, reference to pointer, or reference to array in 'is_device_ptr clause'}}
+  {}
+#pragma omp target teams is_device_ptr(da) // OK
+  {}
+#pragma omp target teams map(ps) is_device_ptr(ps) // expected-error{{variable already marked as mapped in current construct}} expected-note{{used here}}
+  {}
+#pragma omp target teams is_device_ptr(ps) map(ps) // expected-error{{variable already marked as mapped in current construct}} expected-note{{used here}}
+  {}
+#pragma omp target teams map(ps->a) is_device_ptr(ps) // expected-error{{variable already marked as mapped in current construct}} expected-note{{used here}}
+  {}
+#pragma omp target teams is_device_ptr(ps) map(ps->a) // expected-error{{pointer cannot be mapped along with a section derived from itself}} expected-note{{used here}}
+  {}
+#pragma omp target teams is_device_ptr(ps) firstprivate(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target teams' directive}}
+  {}
+#pragma omp target teams firstprivate(ps) is_device_ptr(ps) // expected-error{{firstprivate variable cannot be in a is_device_ptr clause in '#pragma omp target teams' directive}} expected-note{{defined as firstprivate}}
+  {}
+#pragma omp target teams is_device_ptr(ps) private(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target teams' directive}}
+  {}
+#pragma omp target teams private(ps) is_device_ptr(ps) // expected-error{{private variable cannot be in a is_device_ptr clause in '#pragma omp target teams' directive}} expected-note{{defined as private}}
+  {}
+  return tmain<int, 3>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}}
+}

test/OpenMP/target_teams_map_messages.cpp

+// RUN: %clang_cc1 -verify -fopenmp -ferror-limit 200 %s
+// rUN: %clang_cc1 -DCCODE -verify -fopenmp -ferror-limit 200 -x c %s
+#ifdef CCODE
+void foo(int arg) {
+  const int n = 0;
+
+  double marr[10][10][10];
+
+  #pragma omp target teams map(marr[2][0:2][0:2]) // expected-error {{array section does not specify contiguous storage}}
+  {}
+  #pragma omp target teams map(marr[:][0:][:])
+  {}
+  #pragma omp target teams map(marr[:][1:][:]) // expected-error {{array section does not specify contiguous storage}}
+  {}
+  #pragma omp target teams map(marr[:][n:][:])
+  {}
+}
+#else
+template <typename T, int I>
+struct SA {
+  static int ss;
+  #pragma omp threadprivate(ss) // expected-note {{defined as threadprivate or thread local}}
+  float a;
+  int b[12];
+  float *c;
+  T d;
+  float e[I];
+  T *f;
+  void func(int arg) {
+    #pragma omp target teams map(arg,a,d)
+    {}
+    #pragma omp target teams map(arg[2:2],a,d) // expected-error {{subscripted value is not an array or pointer}}
+    {}
+    #pragma omp target teams map(arg,a*2) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+    {}
+    #pragma omp target teams map(arg,(c+1)[2]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+    {}
+    #pragma omp target teams map(arg,a[:2],d) // expected-error {{subscripted value is not an array or pointer}}
+    {}
+    #pragma omp target teams map(arg,a,d[:2]) // expected-error {{subscripted value is not an array or pointer}}
+    {}
+
+    #pragma omp target teams map(to:ss) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
+    {}
+
+    #pragma omp target teams map(to:b,e)
+    {}
+    #pragma omp target teams map(to:b,e) map(to:b) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+    {}
+    #pragma omp target teams map(to:b[:2],e)
+    {}
+    #pragma omp target teams map(to:b,e[:])
+    {}
+    #pragma omp target teams map(b[-1:]) // expected-error {{array section must be a subset of the original array}}
+    {}
+    #pragma omp target teams map(b[:-1]) // expected-error {{section length is evaluated to a negative value -1}}
+    {}
+
+    #pragma omp target teams map(always, tofrom: c,f)
+    {}
+    #pragma omp target teams map(always, tofrom: c[1:2],f)
+    {}
+    #pragma omp target teams map(always, tofrom: c,f[1:2])
+    {}
+    #pragma omp target teams map(always, tofrom: c[:],f)   // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+    {}
+    #pragma omp target teams map(always, tofrom: c,f[:])   // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+    {}
+    return;
+  }
+};
+
+struct SB {
+  unsigned A;
+  unsigned B;
+  float Arr[100];
+  float *Ptr;
+  float *foo() {
+    return &Arr[0];
+  }
+};
+
+struct SC {
+  unsigned A : 2;
+  unsigned B : 3;
+  unsigned C;
+  unsigned D;
+  float Arr[100];
+  SB S;
+  SB ArrS[100];
+  SB *PtrS;
+  SB *&RPtrS;
+  float *Ptr;
+
+  SC(SB *&_RPtrS) : RPtrS(_RPtrS) {}
+};
+
+union SD {
+  unsigned A;
+  float B;
+};
+
+void SAclient(int arg) {
+  SA<int,123> s;
+  s.func(arg); // expected-note {{in instantiation of member function}}
+  double marr[10][10][10];
+  double marr2[5][10][1];
+  double mvla[5][arg][10];
+  double ***mptr;
+  const int n = 0;
+  const int m = 1;
+  double mvla2[5][arg][m+n+10];
+
+  SB *p;
+
+  SD u;
+  SC r(p),t(p);
+  #pragma omp target teams map(r)
+  {}
+  #pragma omp target teams map(marr[2][0:2][0:2]) // expected-error {{array section does not specify contiguous storage}}
+  {}
+  #pragma omp target teams map(marr[:][0:2][0:2]) // expected-error {{array section does not specify contiguous storage}}
+  {}
+  #pragma omp target teams map(marr[2][3][0:2])
+  {}
+  #pragma omp target teams map(marr[:][:][:])
+  {}
+  #pragma omp target teams map(marr[:2][:][:])
+  {}
+  #pragma omp target teams map(marr[arg:][:][:])
+  {}
+  #pragma omp target teams map(marr[arg:])
+  {}
+  #pragma omp target teams map(marr[arg:][:arg][:]) // correct if arg is the size of dimension 2
+  {}
+  #pragma omp target teams map(marr[:arg][:])
+  {}
+  #pragma omp target teams map(marr[:arg][n:])
+  {}
+  #pragma omp target teams map(marr[:][:arg][n:]) // correct if arg is the size of  dimension 2
+  {}
+  #pragma omp target teams map(marr[:][:m][n:]) // expected-error {{array section does not specify contiguous storage}}
+  {}
+  #pragma omp target teams map(marr[n:m][:arg][n:])
+  {}
+  #pragma omp target teams map(marr[:2][:1][:]) // expected-error {{array section does not specify contiguous storage}}
+  {}
+  #pragma omp target teams map(marr[:2][1:][:]) // expected-error {{array section does not specify contiguous storage}}
+  {}
+  #pragma omp target teams map(marr[:2][:][:1]) // expected-error {{array section does not specify contiguous storage}}
+  {}
+  #pragma omp target teams map(marr[:2][:][1:]) // expected-error {{array section does not specify contiguous storage}}
+  {}
+  #pragma omp target teams map(marr[:1][:2][:])
+  {}
+  #pragma omp target teams map(marr[:1][0][:])
+  {}
+  #pragma omp target teams map(marr[:arg][:2][:]) // correct if arg is 1
+  {}
+  #pragma omp target teams map(marr[:1][3:1][:2])
+  {}
+  #pragma omp target teams map(marr[:1][3:arg][:2]) // correct if arg is 1
+  {}
+  #pragma omp target teams map(marr[:1][3:2][:2]) // expected-error {{array section does not specify contiguous storage}}
+  {}
+  #pragma omp target teams map(marr[:2][:10][:])
+  {}
+  #pragma omp target teams map(marr[:2][:][:5+5])
+  {}
+  #pragma omp target teams map(marr[:2][2+2-4:][0:5+5])
+  {}
+
+  #pragma omp target teams map(marr[:1][:2][0]) // expected-error {{array section does not specify contiguous storage}}
+  {}
+  #pragma omp target teams map(marr2[:1][:2][0])
+  {}
+
+  #pragma omp target teams map(mvla[:1][:][0]) // correct if the size of dimension 2 is 1.
+  {}
+  #pragma omp target teams map(mvla[:2][:arg][:]) // correct if arg is the size of dimension 2.
+  {}
+  #pragma omp target teams map(mvla[:1][:2][0]) // expected-error {{array section does not specify contiguous storage}}
+   {}
+  #pragma omp target teams map(mvla[1][2:arg][:])
+  {}
+  #pragma omp target teams map(mvla[:1][:][:])
+  {}
+  #pragma omp target teams map(mvla2[:1][:2][:11])
+  {}
+  #pragma omp target teams map(mvla2[:1][:2][:10]) // expected-error {{array section does not specify contiguous storage}}
+  {}
+
+  #pragma omp target teams map(mptr[:2][2+2-4:1][0:5+5]) // expected-error {{array section does not specify contiguous storage}}
+  {}
+  #pragma omp target teams map(mptr[:1][:2-1][2:4-3])
+  {}
+  #pragma omp target teams map(mptr[:1][:arg][2:4-3]) // correct if arg is 1.
+  {}
+  #pragma omp target teams map(mptr[:1][:2-1][0:2])
+  {}
+  #pragma omp target teams map(mptr[:1][:2][0:2]) // expected-error {{array section does not specify contiguous storage}}
+  {}
+  #pragma omp target teams map(mptr[:1][:][0:2]) // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+  {}
+  #pragma omp target teams map(mptr[:2][:1][0:2]) // expected-error {{array section does not specify contiguous storage}}
+  {}
+
+  #pragma omp target teams map(r.ArrS[0].B)
+  {}
+  #pragma omp target teams map(r.ArrS[:1].B) // expected-error {{OpenMP array section is not allowed here}}
+  {}
+  #pragma omp target teams map(r.ArrS[:arg].B) // expected-error {{OpenMP array section is not allowed here}}
+  {}
+  #pragma omp target teams map(r.ArrS[0].Arr[1:23])
+  {}
+  #pragma omp target teams map(r.ArrS[0].Arr[1:arg])
+  {}
+  #pragma omp target teams map(r.ArrS[0].Arr[arg:23])
+  {}
+  #pragma omp target teams map(r.ArrS[0].Error) // expected-error {{no member named 'Error' in 'SB'}}
+  {}
+  #pragma omp target teams map(r.ArrS[0].A, r.ArrS[1].A) // expected-error {{multiple array elements associated with the same variable are not allowed in map clauses of the same construct}} expected-note {{used here}}
+  {}
+  #pragma omp target teams map(r.ArrS[0].A, t.ArrS[1].A)
+  {}
+  #pragma omp target teams map(r.PtrS[0], r.PtrS->B) // expected-error {{same pointer derreferenced in multiple different ways in map clause expressions}} expected-note {{used here}}
+  {}
+  #pragma omp target teams map(r.RPtrS[0], r.RPtrS->B) // expected-error {{same pointer derreferenced in multiple different ways in map clause expressions}} expected-note {{used here}}
+  {}
+  #pragma omp target teams map(r.S.Arr[:12])
+  {}
+  #pragma omp target teams map(r.S.foo()[:12]) // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+  {}
+  #pragma omp target teams map(r.C, r.D)
+  {}
+  #pragma omp target teams map(r.C, r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+  {}
+  #pragma omp target teams map(r.C) map(r.C) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+  {}
+  #pragma omp target teams map(r.C, r.S)  // this would be an error only caught at runtime - Sema would have to make sure there is not way for the missing data between fields to be mapped somewhere else.
+  {}
+  #pragma omp target teams map(r, r.S)  // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+  {}
+  #pragma omp target teams map(r.C, t.C)
+  {}
+  #pragma omp target teams map(r.A)   // expected-error {{bit fields cannot be used to specify storage in a 'map' clause}}
+  {}
+  #pragma omp target teams map(r.Arr)
+  {}
+  #pragma omp target teams map(r.Arr[3:5])
+  {}
+  #pragma omp target teams map(r.Ptr[3:5])
+  {}
+  #pragma omp target teams map(r.ArrS[3:5].A)   // expected-error {{OpenMP array section is not allowed here}}
+  {}
+  #pragma omp target teams map(r.ArrS[3:5].Arr[6:7])   // expected-error {{OpenMP array section is not allowed here}}
+  {}
+  #pragma omp target teams map(r.ArrS[3].Arr[6:7])
+  {}
+  #pragma omp target teams map(r.S.Arr[4:5])
+  {}
+  #pragma omp target teams map(r.S.Ptr[4:5])
+  {}
+  #pragma omp target teams map(r.S.Ptr[:])  // expected-error {{section length is unspecified and cannot be inferred because subscripted value is not an array}}
+  {}
+  #pragma omp target teams map((p+1)->A)  // expected-error {{expected expression containing only member accesses and/or array sections based on named variables}}
+  {}
+  #pragma omp target teams map(u.B)  // expected-error {{mapped storage cannot be derived from a union}}
+  {}
+
+  #pragma omp target data map(to: r.C) //expected-note {{used here}}
+  {
+    #pragma omp target teams map(r.D)  // expected-error {{original storage of expression in data environment is shared but data environment do not fully contain mapped expression storage}}
+    {}
+  }
+
+  #pragma omp target data map(to: t.Ptr) //expected-note {{used here}}
+  {
+    #pragma omp target teams map(t.Ptr[:23])  // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+    {}
+  }
+
+  #pragma omp target data map(to: t.C, t.D)
+  {
+  #pragma omp target data map(to: t.C)
+  {
+    #pragma omp target teams map(t.D)
+    {}
+  }
+  }
+
+  #pragma omp target data map(to: t)
+  {
+  #pragma omp target data map(to: t.C)
+  {
+    #pragma omp target teams map(t.D)
+    {}
+  }
+  }
+}
+void foo() {
+}
+
+bool foobool(int argc) {
+  return argc;
+}
+
+struct S1; // expected-note 2 {{declared here}}
+extern S1 a;
+class S2 {
+  mutable int a;
+public:
+  S2():a(0) { }
+  S2(S2 &s2):a(s2.a) { }
+  static float S2s; // expected-note 4 {{mappable type cannot contain static members}}
+  static const float S2sc; // expected-note 4 {{mappable type cannot contain static members}}
+};
+const float S2::S2sc = 0;
+const S2 b;
+const S2 ba[5];
+class S3 {
+  int a;
+public:
+  S3():a(0) { }
+  S3(S3 &s3):a(s3.a) { }
+};
+const S3 c;
+const S3 ca[5];
+extern const int f;
+class S4 {
+  int a;
+  S4();
+  S4(const S4 &s4);
+public:
+  S4(int v):a(v) { }
+};
+class S5 {
+  int a;
+  S5():a(0) {}
+  S5(const S5 &s5):a(s5.a) { }
+public:
+  S5(int v):a(v) { }
+};
+
+template <class T>
+struct S6;
+
+template<>
+struct S6<int>  // expected-note {{mappable type cannot be polymorphic}}
+{
+   virtual void foo();
+};
+
+S3 h;
+#pragma omp threadprivate(h) // expected-note 2 {{defined as threadprivate or thread local}}
+
+typedef int from;
+
+template <typename T, int I> // expected-note {{declared here}}
+T tmain(T argc) {
+  const T d = 5;
+  const T da[5] = { 0 };
+  S4 e(4);
+  S5 g(5);
+  T i, t[20];
+  T &j = i;
+  T *k = &j;
+  T x;
+  T y;
+  T to, tofrom, always;
+  const T (&l)[5] = da;
+#pragma omp target teams map // expected-error {{expected '(' after 'map'}}
+  {}
+#pragma omp target teams map( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+  {}
+#pragma omp target teams map() // expected-error {{expected expression}}
+  {}
+#pragma omp target teams map(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
+  {}
+#pragma omp target teams map(to argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected ',' or ')' in 'map' clause}}
+  {}
+#pragma omp target teams map(to:) // expected-error {{expected expression}}
+  {}
+#pragma omp target teams map(from: argc, // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  {}
+#pragma omp target teams map(x: y) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
+  {}
+#pragma omp target teams map(x)
+  foo();
+#pragma omp target teams map(tofrom: t[:I])
+  foo();
+#pragma omp target teams map(T: a) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}} expected-error {{incomplete type 'S1' where a complete type is required}}
+  foo();
+#pragma omp target teams map(T) // expected-error {{'T' does not refer to a value}}
+  foo();
+#pragma omp target teams map(I) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+  foo();
+#pragma omp target teams map(S2::S2s)
+  foo();
+#pragma omp target teams map(S2::S2sc)
+  foo();
+#pragma omp target teams map(x)
+  foo();
+#pragma omp target teams map(to: x)
+  foo();
+#pragma omp target teams map(to: to)
+  foo();
+#pragma omp target teams map(to)
+  foo();
+#pragma omp target teams map(to, x)
+  foo();
+
+#pragma omp target data map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
+#pragma omp target data map(tofrom: argc > 0 ? x : y) // expected-error 2 {{expected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target data map(argc)
+#pragma omp target data map(S1) // expected-error {{'S1' does not refer to a value}}
+#pragma omp target data map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} expected-error 2 {{type 'S2' is not mappable to target}}
+#pragma omp target data map(ba) // expected-error 2 {{type 'S2' is not mappable to target}}
+#pragma omp target data map(ca)
+#pragma omp target data map(da)
+#pragma omp target data map(S2::S2s)
+#pragma omp target data map(S2::S2sc)
+#pragma omp target data map(e, g)
+#pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
+#pragma omp target data map(k) map(k) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+#pragma omp target teams map(k), map(k[:5]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}} expected-note 2 {{used here}}
+  foo();
+
+#pragma omp target data map(da)
+#pragma omp target teams map(da[:4])
+  foo();
+
+#pragma omp target data map(k, j, l) // expected-note 2 {{used here}}
+#pragma omp target data map(k[:4]) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(j)
+#pragma omp target teams map(l) map(l[:5]) // expected-error 2 {{variable already marked as mapped in current construct}} expected-note 2 {{used here}}
+  foo();
+
+#pragma omp target data map(k[:4], j, l[:5]) // expected-note 2 {{used here}}
+#pragma omp target data map(k) // expected-error 2 {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(j)
+#pragma omp target teams map(l)
+  foo();
+
+#pragma omp target data map(always, tofrom: x)
+#pragma omp target data map(always: x) // expected-error {{missing map type}}
+#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always'}} expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
+#pragma omp target data map(always, tofrom: always, tofrom, x)
+#pragma omp target teams map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
+  foo();
+  return 0;
+}
+
+int main(int argc, char **argv) {
+  const int d = 5;
+  const int da[5] = { 0 };
+  S4 e(4);
+  S5 g(5);
+  int i;
+  int &j = i;
+  int *k = &j;
+  S6<int> m;
+  int x;
+  int y;
+  int to, tofrom, always;
+  const int (&l)[5] = da;
+#pragma omp target data map // expected-error {{expected '(' after 'map'}} expected-error {{expected at least one map clause for '#pragma omp target data'}}
+#pragma omp target data map( // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected expression}}
+#pragma omp target data map() // expected-error {{expected expression}}
+#pragma omp target data map(alloc) // expected-error {{use of undeclared identifier 'alloc'}}
+#pragma omp target data map(to argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{expected ',' or ')' in 'map' clause}}
+#pragma omp target data map(to:) // expected-error {{expected expression}}
+#pragma omp target data map(from: argc, // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+#pragma omp target data map(x: y) // expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
+#pragma omp target teams map(x)
+  foo();
+
+#pragma omp target teams map(to: x)
+  foo();
+#pragma omp target teams map(to: to)
+  foo();
+#pragma omp target teams map(to)
+  foo();
+#pragma omp target teams map(to, x)
+  foo();
+
+#pragma omp target data map(to x) // expected-error {{expected ',' or ')' in 'map' clause}}
+#pragma omp target data map(tofrom: argc > 0 ? argv[1] : argv[2]) // expected-error {{xpected expression containing only member accesses and/or array sections based on named variables}}
+#pragma omp target data map(argc)
+#pragma omp target data map(S1) // expected-error {{'S1' does not refer to a value}}
+#pragma omp target data map(a, b, c, d, f) // expected-error {{incomplete type 'S1' where a complete type is required}} expected-error 2 {{type 'S2' is not mappable to target}}
+#pragma omp target data map(argv[1])
+#pragma omp target data map(ba) // expected-error 2 {{type 'S2' is not mappable to target}}
+#pragma omp target data map(ca)
+#pragma omp target data map(da)
+#pragma omp target data map(S2::S2s)
+#pragma omp target data map(S2::S2sc)
+#pragma omp target data map(e, g)
+#pragma omp target data map(h) // expected-error {{threadprivate variables are not allowed in 'map' clause}}
+#pragma omp target data map(k), map(k) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+#pragma omp target teams map(k), map(k[:5]) // expected-error {{pointer cannot be mapped along with a section derived from itself}} expected-note {{used here}}
+  foo();
+
+#pragma omp target data map(da)
+#pragma omp target teams map(da[:4])
+  foo();
+
+#pragma omp target data map(k, j, l) // expected-note {{used here}}
+#pragma omp target data map(k[:4]) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(j)
+#pragma omp target teams map(l) map(l[:5]) // expected-error {{variable already marked as mapped in current construct}} expected-note {{used here}}
+  foo();
+
+#pragma omp target data map(k[:4], j, l[:5]) // expected-note {{used here}}
+#pragma omp target data map(k) // expected-error {{pointer cannot be mapped along with a section derived from itself}}
+#pragma omp target data map(j)
+#pragma omp target teams map(l)
+  foo();
+
+#pragma omp target data map(always, tofrom: x)
+#pragma omp target data map(always: x) // expected-error {{missing map type}}
+#pragma omp target data map(tofrom, always: x) // expected-error {{incorrect map type modifier, expected 'always'}} expected-error {{incorrect map type, expected one of 'to', 'from', 'tofrom', 'alloc', 'release', or 'delete'}}
+#pragma omp target data map(always, tofrom: always, tofrom, x)
+#pragma omp target teams map(tofrom j) // expected-error {{expected ',' or ')' in 'map' clause}}
+  foo();
+
+#pragma omp target teams private(j) map(j) // expected-error {{private variable cannot be in a map clause in '#pragma omp target teams' directive}}  expected-note {{defined as private}}
+  {}
+
+#pragma omp target teams firstprivate(j) map(j)  // expected-error {{firstprivate variable cannot be in a map clause in '#pragma omp target teams' directive}} expected-note {{defined as firstprivate}}
+  {}
+
+#pragma omp target teams map(m) // expected-error {{type 'S6<int>' is not mappable to target}}
+  {}
+
+  return tmain<int, 3>(argc)+tmain<from, 4>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 3>' requested here}} expected-note {{in instantiation of function template specialization 'tmain<int, 4>' requested here}}
+}
+#endif

test/OpenMP/target_teams_messages.cpp

+// RUN: %clang_cc1 -verify -fopenmp -std=c++11 -o - %s
+
+void foo() {
+}
+
+#pragma omp target teams // expected-error {{unexpected OpenMP directive '#pragma omp target teams'}}
+
+int main(int argc, char **argv) {
+#pragma omp target teams { // expected-warning {{extra tokens at the end of '#pragma omp target teams' are ignored}}
+  foo();
+#pragma omp target teams ( // expected-warning {{extra tokens at the end of '#pragma omp target teams' are ignored}}
+  foo();
+#pragma omp target teams [ // expected-warning {{extra tokens at the end of '#pragma omp target teams' are ignored}}
+  foo();
+#pragma omp target teams ] // expected-warning {{extra tokens at the end of '#pragma omp target teams' are ignored}}
+  foo();
+#pragma omp target teams ) // expected-warning {{extra tokens at the end of '#pragma omp target teams' are ignored}}
+  foo();
+#pragma omp target teams } // expected-warning {{extra tokens at the end of '#pragma omp target teams' are ignored}}
+  foo();
+#pragma omp target teams
+  foo();
+#pragma omp target teams unknown() // expected-warning {{extra tokens at the end of '#pragma omp target teams' are ignored}}
+  foo();
+  L1:
+    foo();
+#pragma omp target teams
+  ;
+#pragma omp target teams
+  {
+    goto L1; // expected-error {{use of undeclared label 'L1'}}
+    argc++;
+  }
+
+  for (int i = 0; i < 10; ++i) {
+    switch(argc) {
+     case (0):
+      #pragma omp target teams
+      {
+        foo();
+        break; // expected-error {{'break' statement not in loop or switch statement}}
+        continue; // expected-error {{'continue' statement not in loop statement}}
+      }
+      default:
+       break;
+    }
+  }
+#pragma omp target teams default(none)
+  ++argc; // expected-error {{variable 'argc' must have explicitly specified data sharing attributes}}
+
+  goto L2; // expected-error {{use of undeclared label 'L2'}}
+#pragma omp target teams
+  L2:
+  foo();
+#pragma omp target teams
+  {
+    return 1; // expected-error {{cannot return from OpenMP region}}
+  }
+
+  [[]] // expected-error {{an attribute list cannot appear here}}
+#pragma omp target teams
+  for (int n = 0; n < 100; ++n) {}
+
+  return 0;
+}

test/OpenMP/target_teams_nowait_messages.cpp

+// RUN: %clang_cc1 -verify -fopenmp -o - %s
+
+void foo() {
+}
+
+int main(int argc, char **argv) {
+#pragma omp target teams nowait( // expected-warning {{extra tokens at the end of '#pragma omp target teams' are ignored}}
+  foo();
+#pragma omp target teams nowait (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target teams' are ignored}}
+  foo();
+#pragma omp target teams nowait device (-10u)
+  foo();
+#pragma omp target teams nowait (3.14) device (-10u) // expected-warning {{extra tokens at the end of '#pragma omp target teams' are ignored}}
+  foo();
+
+  return 0;
+}

test/OpenMP/target_teams_num_teams_messages.cpp

+// RUN: %clang_cc1 -verify -fopenmp -std=c++11 -ferror-limit 100 -o - %s
+
+void foo() {
+}
+
+bool foobool(int argc) {
+  return argc;
+}
+
+struct S1; // expected-note 2 {{declared here}}
+
+template <typename T, int C> // expected-note {{declared here}}
+T tmain(T argc) {
+  char **a;
+#pragma omp target teams num_teams(C)
+  foo();
+#pragma omp target teams num_teams(T) // expected-error {{'T' does not refer to a value}}
+  foo();
+#pragma omp target teams num_teams // expected-error {{expected '(' after 'num_teams'}}
+  foo();
+#pragma omp target teams num_teams( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams num_teams() // expected-error {{expected expression}}
+  foo();
+#pragma omp target teams num_teams(argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams num_teams(argc)) // expected-warning {{extra tokens at the end of '#pragma omp target teams' are ignored}}
+  foo();
+#pragma omp target teams num_teams(argc > 0 ? a[1] : a[2]) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}}
+  foo();
+#pragma omp target teams num_teams(argc + argc)
+  foo();
+#pragma omp target teams num_teams(argc), num_teams (argc+1) // expected-error {{directive '#pragma omp target teams' cannot contain more than one 'num_teams' clause}}
+  foo();
+#pragma omp target teams num_teams(S1) // expected-error {{'S1' does not refer to a value}}
+  foo();
+#pragma omp target teams num_teams(-2) // expected-error {{argument to 'num_teams' clause must be a strictly positive integer value}}
+  foo();
+#pragma omp target teams num_teams(-10u)
+  foo();
+#pragma omp target teams num_teams(3.14) // expected-error 2 {{expression must have integral or unscoped enumeration type, not 'double'}}
+  foo();
+
+  return 0;
+}
+
+int main(int argc, char **argv) {
+#pragma omp target teams num_teams // expected-error {{expected '(' after 'num_teams'}}
+  foo();
+
+#pragma omp target teams num_teams ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+
+#pragma omp target teams num_teams () // expected-error {{expected expression}}
+  foo();
+
+#pragma omp target teams num_teams (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+
+#pragma omp target teams num_teams (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target teams' are ignored}}
+  foo();
+
+#pragma omp target teams num_teams (argc > 0 ? argv[1] : argv[2]) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}}
+  foo();
+
+#pragma omp target teams num_teams (argc + argc)
+  foo();
+
+#pragma omp target teams num_teams (argc), num_teams (argc+1) // expected-error {{directive '#pragma omp target teams' cannot contain more than one 'num_teams' clause}}
+  foo();
+
+#pragma omp target teams num_teams (S1) // expected-error {{'S1' does not refer to a value}}
+  foo();
+
+#pragma omp target teams num_teams (-2) // expected-error {{argument to 'num_teams' clause must be a strictly positive integer value}}
+  foo();
+
+#pragma omp target teams num_teams (-10u)
+  foo();
+
+#pragma omp target teams num_teams (3.14) // expected-error {{expression must have integral or unscoped enumeration type, not 'double'}}
+  foo();
+
+  return tmain<int, 10>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 10>' requested here}}
+}

test/OpenMP/target_teams_private_messages.cpp

+// RUN: %clang_cc1 -verify -fopenmp %s
+
+void foo() {
+}
+
+bool foobool(int argc) {
+  return argc;
+}
+
+struct S1; // expected-note {{declared here}} expected-note{{forward declaration of 'S1'}}
+extern S1 a;
+class S2 {
+  mutable int a;
+public:
+  S2():a(0) { }
+  static float S2s; // expected-note {{static data member is predetermined as shared}}
+};
+const S2 b;
+const S2 ba[5];
+class S3 {
+  int a;
+public:
+  S3():a(0) { }
+};
+const S3 c; // expected-note {{global variable is predetermined as shared}}
+const S3 ca[5]; // expected-note {{global variable is predetermined as shared}}
+extern const int f; // expected-note {{global variable is predetermined as shared}}
+class S4 {
+  int a;
+  S4(); // expected-note {{implicitly declared private here}}
+public:
+  S4(int v):a(v) { }
+};
+class S5 {
+  int a;
+  S5():a(0) {} // expected-note {{implicitly declared private here}}
+public:
+  S5(int v):a(v) { }
+};
+
+int threadvar;
+#pragma omp threadprivate(threadvar) // expected-note {{defined as threadprivate or thread local}}
+
+namespace A {
+double x;
+#pragma omp threadprivate(x) // expected-note {{defined as threadprivate or thread local}}
+}
+namespace B {
+using A::x;
+}
+
+int main(int argc, char **argv) {
+  const int d = 5; // expected-note {{constant variable is predetermined as shared}}
+  const int da[5] = { 0 }; // expected-note {{constant variable is predetermined as shared}}
+  S4 e(4);
+  S5 g(5);
+  int i;
+  int &j = i;
+#pragma omp target teams private // expected-error {{expected '(' after 'private'}}
+  foo();
+#pragma omp target teams private ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams private () // expected-error {{expected expression}}
+  foo();
+#pragma omp target teams private (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams private (argc, // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams private (argc > 0 ? argv[1] : argv[2]) // expected-error {{expected variable name}}
+  foo();
+#pragma omp target teams private (argc argv) // expected-error {{expected ',' or ')' in 'private' clause}}
+  foo();
+#pragma omp target teams private (S1) // expected-error {{'S1' does not refer to a value}}
+  foo();
+#pragma omp target teams private (a, b, c, d, f) // expected-error {{a private variable with incomplete type 'S1'}} expected-error 3 {{shared variable cannot be private}}
+  foo();
+#pragma omp target teams private (argv[1]) // expected-error {{expected variable name}}
+  foo();
+#pragma omp target teams private(ba)
+  foo();
+#pragma omp target teams private(ca) // expected-error {{shared variable cannot be private}}
+  foo();
+#pragma omp target teams private(da) // expected-error {{shared variable cannot be private}}
+  foo();
+#pragma omp target teams private(S2::S2s) // expected-error {{shared variable cannot be private}}
+  foo();
+#pragma omp target teams private(e, g) // expected-error {{calling a private constructor of class 'S4'}} expected-error {{calling a private constructor of class 'S5'}}
+  foo();
+#pragma omp target teams private(threadvar, B::x) // expected-error 2 {{threadprivate or thread local variable cannot be private}}
+  foo();
+#pragma omp target teams shared(i), private(i) // expected-error {{shared variable cannot be private}} expected-note {{defined as shared}}
+  foo();
+#pragma omp target teams firstprivate(i) private(i) // expected-error {{firstprivate variable cannot be private}} expected-note {{defined as firstprivate}}
+  foo();
+#pragma omp target teams private(i)
+  foo();
+#pragma omp target teams private(j)
+  foo();
+#pragma omp target teams firstprivate(i)
+  for (int k = 0; k < 10; ++k) {
+#pragma omp parallel private(i)
+    foo();
+  }
+  static int m;
+#pragma omp target teams private(m) // OK
+  foo();
+
+  return 0;
+}

test/OpenMP/target_teams_reduction_messages.cpp

+// RUN: %clang_cc1 -verify -fopenmp -o - %s
+// RUN: %clang_cc1 -verify -fopenmp -std=c++98 -o - %s
+// RUN: %clang_cc1 -verify -fopenmp -std=c++11 -o - %s
+
+void foo() {
+}
+
+bool foobool(int argc) {
+  return argc;
+}
+
+void foobar(int &ref) {
+#pragma omp target teams reduction(+:ref)
+  foo();
+}
+
+struct S1; // expected-note {{declared here}} expected-note 4 {{forward declaration of 'S1'}}
+extern S1 a;
+class S2 {
+  mutable int a;
+  S2 &operator+(const S2 &arg) { return (*this); } // expected-note 3 {{implicitly declared private here}}
+
+public:
+  S2() : a(0) {}
+  S2(S2 &s2) : a(s2.a) {}
+  static float S2s; // expected-note 2 {{static data member is predetermined as shared}}
+  static const float S2sc;
+};
+const float S2::S2sc = 0; // expected-note 2 {{'S2sc' defined here}}
+S2 b;                     // expected-note 3 {{'b' defined here}}
+const S2 ba[5];           // expected-note 2 {{'ba' defined here}}
+class S3 {
+  int a;
+
+public:
+  int b;
+  S3() : a(0) {}
+  S3(const S3 &s3) : a(s3.a) {}
+  S3 operator+(const S3 &arg1) { return arg1; }
+};
+int operator+(const S3 &arg1, const S3 &arg2) { return 5; }
+S3 c;               // expected-note 3 {{'c' defined here}}
+const S3 ca[5];     // expected-note 2 {{'ca' defined here}}
+extern const int f; // expected-note 4 {{'f' declared here}}
+class S4 {
+  int a;
+  S4(); // expected-note {{implicitly declared private here}}
+  S4(const S4 &s4);
+  S4 &operator+(const S4 &arg) { return (*this); }
+
+public:
+  S4(int v) : a(v) {}
+};
+S4 &operator&=(S4 &arg1, S4 &arg2) { return arg1; }
+class S5 {
+  int a;
+  S5() : a(0) {} // expected-note {{implicitly declared private here}}
+  S5(const S5 &s5) : a(s5.a) {}
+  S5 &operator+(const S5 &arg);
+
+public:
+  S5(int v) : a(v) {}
+};
+class S6 { // expected-note 3 {{candidate function (the implicit copy assignment operator) not viable: no known conversion from 'int' to 'const S6' for 1st argument}}
+#if __cplusplus >= 201103L // C++11 or later
+// expected-note@-2 3 {{candidate function (the implicit move assignment operator) not viable}}
+#endif
+  int a;
+
+public:
+  S6() : a(6) {}
+  operator int() { return 6; }
+} o;
+
+S3 h, k;
+#pragma omp threadprivate(h) // expected-note 2 {{defined as threadprivate or thread local}}
+
+template <class T>       // expected-note {{declared here}}
+T tmain(T argc) {
+  const T d = T();       // expected-note 4 {{'d' defined here}}
+  const T da[5] = {T()}; // expected-note 2 {{'da' defined here}}
+  T qa[5] = {T()};
+  T i;
+  T &j = i;                    // expected-note 4 {{'j' defined here}}
+  S3 &p = k;                   // expected-note 2 {{'p' defined here}}
+  const T &r = da[(int)i];     // expected-note 2 {{'r' defined here}}
+  T &q = qa[(int)i];           // expected-note 2 {{'q' defined here}}
+  T fl;
+#pragma omp target teams reduction // expected-error {{expected '(' after 'reduction'}}
+  foo();
+#pragma omp target teams reduction + // expected-error {{expected '(' after 'reduction'}} expected-warning {{extra tokens at the end of '#pragma omp target teams' are ignored}}
+  foo();
+#pragma omp target teams reduction( // expected-error {{expected unqualified-id}} expected-warning {{missing ':' after reduction identifier - ignoring}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams reduction(- // expected-warning {{missing ':' after reduction identifier - ignoring}} expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams reduction() // expected-error {{expected unqualified-id}} expected-warning {{missing ':' after reduction identifier - ignoring}}
+  foo();
+#pragma omp target teams reduction(*) // expected-warning {{missing ':' after reduction identifier - ignoring}} expected-error {{expected expression}}
+  foo();
+#pragma omp target teams reduction(\) // expected-error {{expected unqualified-id}} expected-warning {{missing ':' after reduction identifier - ignoring}}
+  foo();
+#pragma omp target teams reduction(& : argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{invalid operands to binary expression ('float' and 'float')}}
+  foo();
+#pragma omp target teams reduction(| : argc, // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{invalid operands to binary expression ('float' and 'float')}}
+  foo();
+#pragma omp target teams reduction(|| : argc ? i : argc) // expected-error 2 {{expected variable name, array element or array section}}
+  foo();
+#pragma omp target teams reduction(foo : argc) //expected-error {{incorrect reduction identifier, expected one of '+', '-', '*', '&', '|', '^', '&&', '||', 'min' or 'max' or declare reduction for type 'float'}} expected-error {{incorrect reduction identifier, expected one of '+', '-', '*', '&', '|', '^', '&&', '||', 'min' or 'max' or declare reduction for type 'int'}}
+  foo();
+#pragma omp target teams reduction(&& : argc)
+  foo();
+#pragma omp target teams reduction(^ : T) // expected-error {{'T' does not refer to a value}}
+  foo();
+#pragma omp target teams reduction(+ : a, b, c, d, f) // expected-error {{a reduction list item with incomplete type 'S1'}} expected-error 3 {{const-qualified list item cannot be reduction}} expected-error 2 {{'operator+' is a private member of 'S2'}}
+  foo();
+#pragma omp target teams reduction(min : a, b, c, d, f) // expected-error {{a reduction list item with incomplete type 'S1'}} expected-error 4 {{arguments of OpenMP clause 'reduction' for 'min' or 'max' must be of arithmetic type}} expected-error 3 {{const-qualified list item cannot be reduction}}
+  foo();
+#pragma omp target teams reduction(max : h.b) // expected-error {{expected variable name, array element or array section}}
+  foo();
+#pragma omp target teams reduction(+ : ba) // expected-error {{const-qualified list item cannot be reduction}}
+  foo();
+#pragma omp target teams reduction(* : ca) // expected-error {{const-qualified list item cannot be reduction}}
+  foo();
+#pragma omp target teams reduction(- : da) // expected-error {{const-qualified list item cannot be reduction}} expected-error {{const-qualified list item cannot be reduction}}
+  foo();
+#pragma omp target teams reduction(^ : fl) // expected-error {{invalid operands to binary expression ('float' and 'float')}}
+  foo();
+#pragma omp target teams reduction(&& : S2::S2s) // expected-error {{shared variable cannot be reduction}}
+  foo();
+#pragma omp target teams reduction(&& : S2::S2sc) // expected-error {{const-qualified list item cannot be reduction}}
+  foo();
+#pragma omp target teams reduction(+ : h, k) // expected-error {{threadprivate or thread local variable cannot be reduction}}
+  foo();
+#pragma omp target teams reduction(+ : o) // expected-error 2 {{no viable overloaded '='}}
+  foo();
+#pragma omp target teams private(i), reduction(+ : j), reduction(+ : q) // expected-error 4 {{argument of OpenMP clause 'reduction' must reference the same object in all threads}}
+  foo();
+#pragma omp parallel private(k)
+#pragma omp target teams reduction(+ : p), reduction(+ : p) // expected-error 2 {{argument of OpenMP clause 'reduction' must reference the same object in all threads}}
+  foo();
+#pragma omp target teams reduction(+ : p), reduction(+ : p) // expected-error 2 {{variable can appear only once in OpenMP 'reduction' clause}} expected-note 2 {{previously referenced here}}
+  foo();
+#pragma omp target teams reduction(+ : r) // expected-error 2 {{const-qualified list item cannot be reduction}}
+  foo();
+#pragma omp parallel shared(i)
+#pragma omp parallel reduction(min : i)
+#pragma omp target teams reduction(max : j) // expected-error 2 {{argument of OpenMP clause 'reduction' must reference the same object in all threads}}
+  foo();
+#pragma omp target teams
+#pragma omp parallel for private(fl)
+  for (int i = 0; i < 10; ++i)
+  {}
+#pragma omp target teams reduction(+ : fl)
+    foo();
+#pragma omp target teams
+#pragma omp parallel for reduction(- : fl)
+  for (int i = 0; i < 10; ++i)
+  {}
+#pragma omp target teams reduction(+ : fl)
+    foo();
+
+  return T();
+}
+
+namespace A {
+double x;
+#pragma omp threadprivate(x) // expected-note {{defined as threadprivate or thread local}}
+}
+namespace B {
+using A::x;
+}
+
+int main(int argc, char **argv) {
+  const int d = 5;       // expected-note 2 {{'d' defined here}}
+  const int da[5] = {0}; // expected-note {{'da' defined here}}
+  int qa[5] = {0};
+  S4 e(4);
+  S5 g(5);
+  int i;
+  int &j = i;                  // expected-note 2 {{'j' defined here}}
+  S3 &p = k;                   // expected-note 2 {{'p' defined here}}
+  const int &r = da[i];        // expected-note {{'r' defined here}}
+  int &q = qa[i];              // expected-note {{'q' defined here}}
+  float fl;
+#pragma omp target teams reduction // expected-error {{expected '(' after 'reduction'}}
+  foo();
+#pragma omp target teams reduction + // expected-error {{expected '(' after 'reduction'}} expected-warning {{extra tokens at the end of '#pragma omp target teams' are ignored}}
+  foo();
+#pragma omp target teams reduction( // expected-error {{expected unqualified-id}} expected-warning {{missing ':' after reduction identifier - ignoring}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams reduction(- // expected-warning {{missing ':' after reduction identifier - ignoring}} expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams reduction() // expected-error {{expected unqualified-id}} expected-warning {{missing ':' after reduction identifier - ignoring}}
+  foo();
+#pragma omp target teams reduction(*) // expected-warning {{missing ':' after reduction identifier - ignoring}} expected-error {{expected expression}}
+  foo();
+#pragma omp target teams reduction(\) // expected-error {{expected unqualified-id}} expected-warning {{missing ':' after reduction identifier - ignoring}}
+  foo();
+#pragma omp target teams reduction(foo : argc // expected-error {{expected ')'}} expected-note {{to match this '('}} expected-error {{incorrect reduction identifier, expected one of '+', '-', '*', '&', '|', '^', '&&', '||', 'min' or 'max'}}
+  foo();
+#pragma omp target teams reduction(| : argc, // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams reduction(|| : argc > 0 ? argv[1] : argv[2]) // expected-error {{expected variable name, array element or array section}}
+  foo();
+#pragma omp target teams reduction(~ : argc) // expected-error {{expected unqualified-id}}
+  foo();
+#pragma omp target teams reduction(&& : argc)
+  foo();
+#pragma omp target teams reduction(^ : S1) // expected-error {{'S1' does not refer to a value}}
+  foo();
+#pragma omp target teams reduction(+ : a, b, c, d, f) // expected-error {{a reduction list item with incomplete type 'S1'}} expected-error 2 {{const-qualified list item cannot be reduction}} expected-error {{'operator+' is a private member of 'S2'}}
+  foo();
+#pragma omp target teams reduction(min : a, b, c, d, f) // expected-error {{a reduction list item with incomplete type 'S1'}} expected-error 2 {{arguments of OpenMP clause 'reduction' for 'min' or 'max' must be of arithmetic type}} expected-error 2 {{const-qualified list item cannot be reduction}}
+  foo();
+#pragma omp target teams reduction(max : h.b) // expected-error {{expected variable name, array element or array section}}
+  foo();
+#pragma omp target teams reduction(+ : ba) // expected-error {{const-qualified list item cannot be reduction}}
+  foo();
+#pragma omp target teams reduction(* : ca) // expected-error {{const-qualified list item cannot be reduction}}
+  foo();
+#pragma omp target teams reduction(- : da) // expected-error {{const-qualified list item cannot be reduction}}
+  foo();
+#pragma omp target teams reduction(^ : fl) // expected-error {{invalid operands to binary expression ('float' and 'float')}}
+  foo();
+#pragma omp target teams reduction(&& : S2::S2s) // expected-error {{shared variable cannot be reduction}}
+  foo();
+#pragma omp target teams reduction(&& : S2::S2sc) // expected-error {{const-qualified list item cannot be reduction}}
+  foo();
+#pragma omp target teams reduction(& : e, g) // expected-error {{calling a private constructor of class 'S4'}} expected-error {{invalid operands to binary expression ('S4' and 'S4')}} expected-error {{calling a private constructor of class 'S5'}} expected-error {{invalid operands to binary expression ('S5' and 'S5')}}
+  foo();
+#pragma omp target teams reduction(+ : h, k, B::x) // expected-error 2 {{threadprivate or thread local variable cannot be reduction}}
+  foo();
+#pragma omp target teams reduction(+ : o) // expected-error {{no viable overloaded '='}}
+  foo();
+#pragma omp target teams private(i), reduction(+ : j), reduction(+ : q) // expected-error 2 {{argument of OpenMP clause 'reduction' must reference the same object in all threads}}
+  foo();
+#pragma omp parallel private(k)
+#pragma omp target teams reduction(+ : p), reduction(+ : p) // expected-error 2 {{argument of OpenMP clause 'reduction' must reference the same object in all threads}}
+  foo();
+#pragma omp target teams reduction(+ : p), reduction(+ : p) // expected-error {{variable can appear only once in OpenMP 'reduction' clause}} expected-note {{previously referenced here}}
+  foo();
+#pragma omp target teams reduction(+ : r) // expected-error {{const-qualified list item cannot be reduction}}
+  foo();
+#pragma omp parallel shared(i)
+#pragma omp parallel reduction(min : i)
+#pragma omp target teams reduction(max : j) // expected-error {{argument of OpenMP clause 'reduction' must reference the same object in all threads}}
+  foo();
+#pragma omp target teams
+#pragma omp parallel for private(fl)
+  for (int i = 0; i < 10; ++i)
+  {}
+#pragma omp target teams reduction(+ : fl)
+    foo();
+#pragma omp target teams
+#pragma omp parallel for reduction(- : fl)
+  for (int i = 0; i < 10; ++i)
+  {}
+#pragma omp target teams reduction(+ : fl)
+    foo();
+  static int m;
+#pragma omp target teams reduction(+ : m) // OK
+  foo();
+
+  return tmain(argc) + tmain(fl); // expected-note {{in instantiation of function template specialization 'tmain<int>' requested here}} expected-note {{in instantiation of function template specialization 'tmain<float>' requested here}}
+}

test/OpenMP/target_teams_shared_messages.cpp

+// RUN: %clang_cc1 -verify -fopenmp %s
+
+void foo() {
+}
+
+bool foobool(int argc) {
+  return argc;
+}
+
+struct S1; // expected-note {{declared here}}
+extern S1 a;
+class S2 {
+  mutable int a;
+public:
+  S2():a(0) { }
+  S2(S2 &s2):a(s2.a) { }
+};
+const S2 b;
+const S2 ba[5];
+class S3 {
+  int a;
+public:
+  S3():a(0) { }
+  S3(S3 &s3):a(s3.a) { }
+};
+const S3 c;
+const S3 ca[5];
+extern const int f;
+class S4 {
+  int a;
+  S4();
+  S4(const S4 &s4);
+public:
+  S4(int v):a(v) { }
+};
+class S5 {
+  int a;
+  S5():a(0) {}
+  S5(const S5 &s5):a(s5.a) { }
+public:
+  S5(int v):a(v) { }
+};
+
+S3 h;
+#pragma omp threadprivate(h) // expected-note {{defined as threadprivate or thread local}}
+
+namespace A {
+double x;
+#pragma omp threadprivate(x) // expected-note {{defined as threadprivate or thread local}}
+}
+namespace B {
+using A::x;
+}
+
+int main(int argc, char **argv) {
+  const int d = 5;
+  const int da[5] = { 0 };
+  S4 e(4);
+  S5 g(5);
+  int i;
+  int &j = i;
+#pragma omp target teams shared // expected-error {{expected '(' after 'shared'}}
+  foo();
+#pragma omp target teams shared ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams shared () // expected-error {{expected expression}}
+  foo();
+#pragma omp target teams shared (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams shared (argc, // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams shared (argc > 0 ? argv[1] : argv[2]) // expected-error {{expected variable name}}
+  foo();
+#pragma omp target teams shared (argc)
+ foo();
+#pragma omp target teams shared (S1) // expected-error {{'S1' does not refer to a value}}
+  foo();
+#pragma omp target teams shared (a, b, c, d, f)
+  foo();
+#pragma omp target teams shared (argv[1]) // expected-error {{expected variable name}}
+  foo();
+#pragma omp target teams shared(ba)
+  foo();
+#pragma omp target teams shared(ca)
+  foo();
+#pragma omp target teams shared(da)
+  foo();
+#pragma omp target teams shared(e, g)
+  foo();
+#pragma omp target teams shared(h, B::x) // expected-error 2 {{threadprivate or thread local variable cannot be shared}}
+  foo();
+#pragma omp target teams private(i), shared(i) // expected-error {{private variable cannot be shared}} expected-note {{defined as private}}
+  foo();
+#pragma omp target teams firstprivate(i), shared(i) // expected-error {{firstprivate variable cannot be shared}} expected-note {{defined as firstprivate}}
+  foo();
+#pragma omp target teams private(i)
+  foo();
+#pragma omp target teams shared(i)
+  foo();
+#pragma omp target teams shared(j)
+  foo();
+#pragma omp target teams firstprivate(i)
+  foo();
+#pragma omp target teams shared(i)
+  foo();
+#pragma omp target teams shared(j)
+  foo();
+
+  return 0;
+}

test/OpenMP/target_teams_thread_limit_messages.cpp

+// RUN: %clang_cc1 -verify -fopenmp -std=c++11 -ferror-limit 100 -o - %s
+
+void foo() {
+}
+
+bool foobool(int argc) {
+  return argc;
+}
+
+struct S1; // expected-note 2 {{declared here}}
+
+template <typename T, int C> // expected-note {{declared here}}
+T tmain(T argc) {
+  char **a;
+#pragma omp target teams thread_limit(C)
+  foo();
+#pragma omp target teams thread_limit(T) // expected-error {{'T' does not refer to a value}}
+  foo();
+#pragma omp target teams thread_limit // expected-error {{expected '(' after 'thread_limit'}}
+  foo();
+#pragma omp target teams thread_limit( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams thread_limit() // expected-error {{expected expression}}
+  foo();
+#pragma omp target teams thread_limit(argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+#pragma omp target teams thread_limit(argc)) // expected-warning {{extra tokens at the end of '#pragma omp target teams' are ignored}}
+  foo();
+#pragma omp target teams thread_limit(argc > 0 ? a[1] : a[2]) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}}
+  foo();
+#pragma omp target teams thread_limit(argc + argc)
+  foo();
+#pragma omp target teams thread_limit(argc), thread_limit (argc+1) // expected-error {{directive '#pragma omp target teams' cannot contain more than one 'thread_limit' clause}}
+  foo();
+#pragma omp target teams thread_limit(S1) // expected-error {{'S1' does not refer to a value}}
+  foo();
+#pragma omp target teams thread_limit(-2) // expected-error {{argument to 'thread_limit' clause must be a strictly positive integer value}}
+  foo();
+#pragma omp target teams thread_limit(-10u)
+  foo();
+#pragma omp target teams thread_limit(3.14) // expected-error 2 {{expression must have integral or unscoped enumeration type, not 'double'}}
+  foo();
+
+  return 0;
+}
+
+int main(int argc, char **argv) {
+#pragma omp target teams thread_limit // expected-error {{expected '(' after 'thread_limit'}}
+  foo();
+
+#pragma omp target teams thread_limit ( // expected-error {{expected expression}} expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+
+#pragma omp target teams thread_limit () // expected-error {{expected expression}}
+  foo();
+
+#pragma omp target teams thread_limit (argc // expected-error {{expected ')'}} expected-note {{to match this '('}}
+  foo();
+
+#pragma omp target teams thread_limit (argc)) // expected-warning {{extra tokens at the end of '#pragma omp target teams' are ignored}}
+  foo();
+
+#pragma omp target teams thread_limit (argc > 0 ? argv[1] : argv[2]) // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}}
+  foo();
+
+#pragma omp target teams thread_limit (argc + argc)
+  foo();
+
+#pragma omp target teams thread_limit (argc), thread_limit (argc+1) // expected-error {{directive '#pragma omp target teams' cannot contain more than one 'thread_limit' clause}}
+  foo();
+
+#pragma omp target teams thread_limit (S1) // expected-error {{'S1' does not refer to a value}}
+  foo();
+
+#pragma omp target teams thread_limit (-2) // expected-error {{argument to 'thread_limit' clause must be a strictly positive integer value}}
+  foo();
+
+#pragma omp target teams thread_limit (-10u)
+  foo();
+
+#pragma omp target teams thread_limit (3.14) // expected-error {{expression must have integral or unscoped enumeration type, not 'double'}}
+  foo();
+
+  return tmain<int, 10>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 10>' requested here}}
+}

tools/libclang/CIndex.cpp

public:
   void VisitOMPTargetSimdDirective(const OMPTargetSimdDirective *D);
   void VisitOMPTeamsDistributeDirective(const OMPTeamsDistributeDirective *D);
   void VisitOMPTeamsDistributeSimdDirective(
       const OMPTeamsDistributeSimdDirective *D);
   void VisitOMPTeamsDistributeParallelForSimdDirective(
       const OMPTeamsDistributeParallelForSimdDirective *D);
   void VisitOMPTeamsDistributeParallelForDirective(
       const OMPTeamsDistributeParallelForDirective *D);
+  void VisitOMPTargetTeamsDirective(const OMPTargetTeamsDirective *D);
 
 private:
   void AddDeclarationNameInfo(const Stmt *S);
   void AddNestedNameSpecifierLoc(NestedNameSpecifierLoc Qualifier);
   void AddExplicitTemplateArgs(const TemplateArgumentLoc *A,
                                unsigned NumTemplateArgs);
   void AddMemberRef(const FieldDecl *D, SourceLocation L);
   void AddStmt(const Stmt *S);
void EnqueueVisitor::VisitOMPTeamsDistributeParallelForSimdDirective(
   VisitOMPLoopDirective(D);
 }
 
 void EnqueueVisitor::VisitOMPTeamsDistributeParallelForDirective(
     const OMPTeamsDistributeParallelForDirective *D) {
   VisitOMPLoopDirective(D);
 }
 
+void EnqueueVisitor::VisitOMPTargetTeamsDirective(
+    const OMPTargetTeamsDirective *D) {
+  VisitOMPExecutableDirective(D);
+}
+
 void CursorVisitor::EnqueueWorkList(VisitorWorkList &WL, const Stmt *S) {
   EnqueueVisitor(WL, MakeCXCursor(S, StmtParent, TU,RegionOfInterest)).Visit(S);
 }
 
 bool CursorVisitor::IsInRegionOfInterest(CXCursor C) {
   if (RegionOfInterest.isValid()) {
     SourceRange Range = getRawCursorExtent(C);
     if (Range.isInvalid() || CompareRegionOfInterest(Range))
CXString clang_getCursorKindSpelling(enum CXCursorKind Kind) {
   case CXCursor_OMPTeamsDistributeDirective:
     return cxstring::createRef("OMPTeamsDistributeDirective");
   case CXCursor_OMPTeamsDistributeSimdDirective:
     return cxstring::createRef("OMPTeamsDistributeSimdDirective");
   case CXCursor_OMPTeamsDistributeParallelForSimdDirective:
     return cxstring::createRef("OMPTeamsDistributeParallelForSimdDirective");
   case CXCursor_OMPTeamsDistributeParallelForDirective:
     return cxstring::createRef("OMPTeamsDistributeParallelForDirective");
+  case CXCursor_OMPTargetTeamsDirective:
+    return cxstring::createRef("OMPTargetTeamsDirective");
   case CXCursor_OverloadCandidate:
       return cxstring::createRef("OverloadCandidate");
   case CXCursor_TypeAliasTemplateDecl:
       return cxstring::createRef("TypeAliasTemplateDecl");
   case CXCursor_StaticAssert:
       return cxstring::createRef("StaticAssert");
   case CXCursor_FriendDecl:
     return cxstring::createRef("FriendDecl");

tools/libclang/CXCursor.cpp

case Stmt::OMPTeamsDistributeSimdDirectiveClass:
     K = CXCursor_OMPTeamsDistributeSimdDirective;
     break;
   case Stmt::OMPTeamsDistributeParallelForSimdDirectiveClass:
     K = CXCursor_OMPTeamsDistributeParallelForSimdDirective;
     break;
   case Stmt::OMPTeamsDistributeParallelForDirectiveClass:
     K = CXCursor_OMPTeamsDistributeParallelForDirective;
     break;
+  case Stmt::OMPTargetTeamsDirectiveClass:
+    K = CXCursor_OMPTargetTeamsDirective;
+    break;
   }
 
   CXCursor C = { K, 0, { Parent, S, TU } };
   return C;
 }
 
 CXCursor cxcursor::MakeCursorObjCSuperClassRef(ObjCInterfaceDecl *Super, 
                                                SourceLocation Loc,