[OpenMP] Initial parsing/sema for the 'omp target teams loop' construct

[mikerice]

Summary

Adds basic parsing/sema/serialization support for the #pragma omp target teams loop directive.

[ABataev]

LG

clang/include/clang-c/Index.h

enum CXCursorKind {
   /** OpenMP loop directive.
    */
   CXCursor_OMPGenericLoopDirective = 295,
 
   /** OpenMP teams loop directive.
    */
   CXCursor_OMPTeamsGenericLoopDirective = 296,
 
-  CXCursor_LastStmt = CXCursor_OMPTeamsGenericLoopDirective,
+  /** OpenMP target teams loop directive.
+   */
+  CXCursor_OMPTargetTeamsGenericLoopDirective = 297,
+
+  CXCursor_LastStmt = CXCursor_OMPTargetTeamsGenericLoopDirective,
 
   /**
    * 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,

clang/include/clang/AST/RecursiveASTVisitor.h

DEF_TRAVERSE_STMT(OMPMaskedDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
 DEF_TRAVERSE_STMT(OMPGenericLoopDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
 DEF_TRAVERSE_STMT(OMPTeamsGenericLoopDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
+DEF_TRAVERSE_STMT(OMPTargetTeamsGenericLoopDirective,
+                  { TRY_TO(TraverseOMPExecutableDirective(S)); })
+
 // OpenMP clauses.
 template <typename Derived>
 bool RecursiveASTVisitor<Derived>::TraverseOMPClause(OMPClause *C) {
   if (!C)
     return true;
   switch (C->getClauseKind()) {
 #define GEN_CLANG_CLAUSE_CLASS
 #define CLAUSE_CLASS(Enum, Str, Class)                                         \

clang/include/clang/AST/StmtOpenMP.h

return T->getStmtClass() == OMPSimdDirectiveClass ||
            T->getStmtClass() == OMPParallelForDirectiveClass ||
            T->getStmtClass() == OMPParallelForSimdDirectiveClass ||
            T->getStmtClass() == OMPTaskLoopDirectiveClass ||
            T->getStmtClass() == OMPTaskLoopSimdDirectiveClass ||
            T->getStmtClass() == OMPMasterTaskLoopDirectiveClass ||
            T->getStmtClass() == OMPMasterTaskLoopSimdDirectiveClass ||
            T->getStmtClass() == OMPGenericLoopDirectiveClass ||
            T->getStmtClass() == OMPTeamsGenericLoopDirectiveClass ||
+           T->getStmtClass() == OMPTargetTeamsGenericLoopDirectiveClass ||
            T->getStmtClass() == OMPParallelMasterTaskLoopDirectiveClass ||
            T->getStmtClass() == OMPParallelMasterTaskLoopSimdDirectiveClass ||
            T->getStmtClass() == OMPDistributeDirectiveClass ||
            T->getStmtClass() == OMPTargetParallelForDirectiveClass ||
            T->getStmtClass() == OMPDistributeParallelForDirectiveClass ||
            T->getStmtClass() == OMPDistributeParallelForSimdDirectiveClass ||
            T->getStmtClass() == OMPDistributeSimdDirectiveClass ||
            T->getStmtClass() == OMPTargetParallelForSimdDirectiveClass ||
static OMPTeamsGenericLoopDirective *CreateEmpty(const ASTContext &C,
                                                    unsigned CollapsedNum,
                                                    EmptyShell);
 
   static bool classof(const Stmt *T) {
     return T->getStmtClass() == OMPTeamsGenericLoopDirectiveClass;
   }
 };
 
+/// This represents '#pragma omp target teams loop' directive.
+///
+/// \code
+/// #pragma omp target teams loop private(a,b) order(concurrent)
+/// \endcode
+/// In this example directive '#pragma omp target teams loop' has
+/// clauses 'private' with the variables 'a' and 'b', and order(concurrent).
+///
+class OMPTargetTeamsGenericLoopDirective final : public OMPLoopDirective {
+  friend class ASTStmtReader;
+  friend class OMPExecutableDirective;
+  /// 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 CollapsedNum Number of collapsed nested loops.
+  ///
+  OMPTargetTeamsGenericLoopDirective(SourceLocation StartLoc,
+                                     SourceLocation EndLoc,
+                                     unsigned CollapsedNum)
+      : OMPLoopDirective(OMPTargetTeamsGenericLoopDirectiveClass,
+                         llvm::omp::OMPD_target_teams_loop, StartLoc, EndLoc,
+                         CollapsedNum) {}
+
+  /// Build an empty directive.
+  ///
+  /// \param CollapsedNum Number of collapsed nested loops.
+  ///
+  explicit OMPTargetTeamsGenericLoopDirective(unsigned CollapsedNum)
+      : OMPLoopDirective(OMPTargetTeamsGenericLoopDirectiveClass,
+                         llvm::omp::OMPD_target_teams_loop, SourceLocation(),
+                         SourceLocation(), CollapsedNum) {}
+
+public:
+  /// Creates directive with a list of \p Clauses.
+  ///
+  /// \param C AST context.
+  /// \param StartLoc Starting location of the directive kind.
+  /// \param EndLoc Ending Location of the directive.
+  /// \param CollapsedNum Number of collapsed loops.
+  /// \param Clauses List of clauses.
+  /// \param AssociatedStmt Statement, associated with the directive.
+  /// \param Exprs Helper expressions for CodeGen.
+  ///
+  static OMPTargetTeamsGenericLoopDirective *
+  Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
+         unsigned CollapsedNum, ArrayRef<OMPClause *> Clauses,
+         Stmt *AssociatedStmt, const HelperExprs &Exprs);
+
+  /// Creates an empty directive with the place
+  /// for \a NumClauses clauses.
+  ///
+  /// \param C AST context.
+  /// \param CollapsedNum Number of collapsed nested loops.
+  /// \param NumClauses Number of clauses.
+  ///
+  static OMPTargetTeamsGenericLoopDirective *CreateEmpty(const ASTContext &C,
+                                                         unsigned NumClauses,
+                                                         unsigned CollapsedNum,
+                                                         EmptyShell);
+
+  static bool classof(const Stmt *T) {
+    return T->getStmtClass() == OMPTargetTeamsGenericLoopDirectiveClass;
+  }
+};
 } // end namespace clang
 
 #endif

clang/include/clang/Basic/StmtNodes.td

 def OMPTargetTeamsDistributeParallelForDirective : StmtNode<OMPLoopDirective>;
 def OMPTargetTeamsDistributeParallelForSimdDirective : StmtNode<OMPLoopDirective>;
 def OMPTargetTeamsDistributeSimdDirective : StmtNode<OMPLoopDirective>;
 def OMPInteropDirective : StmtNode<OMPExecutableDirective>;
 def OMPDispatchDirective : StmtNode<OMPExecutableDirective>;
 def OMPMaskedDirective : StmtNode<OMPExecutableDirective>;
 def OMPGenericLoopDirective : StmtNode<OMPLoopDirective>;
 def OMPTeamsGenericLoopDirective : StmtNode<OMPLoopDirective>;
+def OMPTargetTeamsGenericLoopDirective : StmtNode<OMPLoopDirective>;

clang/include/clang/Sema/Sema.h

public:
   StmtResult ActOnOpenMPTeamsDirective(ArrayRef<OMPClause *> Clauses,
                                        Stmt *AStmt, SourceLocation StartLoc,
                                        SourceLocation EndLoc);
   /// Called on well-formed '\#pragma omp teams loop' after parsing of the
   /// associated statement.
   StmtResult ActOnOpenMPTeamsGenericLoopDirective(
       ArrayRef<OMPClause *> Clauses, Stmt *AStmt, SourceLocation StartLoc,
       SourceLocation EndLoc, VarsWithInheritedDSAType &VarsWithImplicitDSA);
+  /// Called on well-formed '\#pragma omp target teams loop' after parsing of
+  /// the associated statement.
+  StmtResult ActOnOpenMPTargetTeamsGenericLoopDirective(
+      ArrayRef<OMPClause *> Clauses, Stmt *AStmt, SourceLocation StartLoc,
+      SourceLocation EndLoc, VarsWithInheritedDSAType &VarsWithImplicitDSA);
   /// Called on well-formed '\#pragma omp cancellation point'.
   StmtResult
   ActOnOpenMPCancellationPointDirective(SourceLocation StartLoc,
                                         SourceLocation EndLoc,
                                         OpenMPDirectiveKind CancelRegion);
   /// Called on well-formed '\#pragma omp cancel'.
   StmtResult ActOnOpenMPCancelDirective(ArrayRef<OMPClause *> Clauses,
                                         SourceLocation StartLoc,

clang/include/clang/Serialization/ASTBitCodes.h

enum StmtCode {
   STMT_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_FOR_DIRECTIVE,
   STMT_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_FOR_SIMD_DIRECTIVE,
   STMT_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD_DIRECTIVE,
   STMT_OMP_INTEROP_DIRECTIVE,
   STMT_OMP_DISPATCH_DIRECTIVE,
   STMT_OMP_MASKED_DIRECTIVE,
   STMT_OMP_GENERIC_LOOP_DIRECTIVE,
   STMT_OMP_TEAMS_GENERIC_LOOP_DIRECTIVE,
+  STMT_OMP_TARGET_TEAMS_GENERIC_LOOP_DIRECTIVE,
   EXPR_OMP_ARRAY_SECTION,
   EXPR_OMP_ARRAY_SHAPING,
   EXPR_OMP_ITERATOR,
 
   // ARC
   EXPR_OBJC_BRIDGED_CAST, // ObjCBridgedCastExpr
 
   STMT_MS_DEPENDENT_EXISTS, // MSDependentExistsStmt

clang/lib/AST/StmtOpenMP.cpp

 OMPTeamsGenericLoopDirective *
 OMPTeamsGenericLoopDirective::CreateEmpty(const ASTContext &C,
                                           unsigned NumClauses,
                                           unsigned CollapsedNum, EmptyShell) {
   return createEmptyDirective<OMPTeamsGenericLoopDirective>(
       C, NumClauses, /*HasAssociatedStmt=*/true,
       numLoopChildren(CollapsedNum, OMPD_teams_loop), CollapsedNum);
 }
+
+OMPTargetTeamsGenericLoopDirective *OMPTargetTeamsGenericLoopDirective::Create(
+    const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
+    unsigned CollapsedNum, ArrayRef<OMPClause *> Clauses, Stmt *AssociatedStmt,
+    const HelperExprs &Exprs) {
+  auto *Dir = createDirective<OMPTargetTeamsGenericLoopDirective>(
+      C, Clauses, AssociatedStmt,
+      numLoopChildren(CollapsedNum, OMPD_target_teams_loop), StartLoc, EndLoc,
+      CollapsedNum);
+  Dir->setIterationVariable(Exprs.IterationVarRef);
+  Dir->setLastIteration(Exprs.LastIteration);
+  Dir->setCalcLastIteration(Exprs.CalcLastIteration);
+  Dir->setPreCond(Exprs.PreCond);
+  Dir->setCond(Exprs.Cond);
+  Dir->setInit(Exprs.Init);
+  Dir->setInc(Exprs.Inc);
+  Dir->setIsLastIterVariable(Exprs.IL);
+  Dir->setLowerBoundVariable(Exprs.LB);
+  Dir->setUpperBoundVariable(Exprs.UB);
+  Dir->setStrideVariable(Exprs.ST);
+  Dir->setEnsureUpperBound(Exprs.EUB);
+  Dir->setNextLowerBound(Exprs.NLB);
+  Dir->setNextUpperBound(Exprs.NUB);
+  Dir->setNumIterations(Exprs.NumIterations);
+  Dir->setCounters(Exprs.Counters);
+  Dir->setPrivateCounters(Exprs.PrivateCounters);
+  Dir->setInits(Exprs.Inits);
+  Dir->setUpdates(Exprs.Updates);
+  Dir->setFinals(Exprs.Finals);
+  Dir->setDependentCounters(Exprs.DependentCounters);
+  Dir->setDependentInits(Exprs.DependentInits);
+  Dir->setFinalsConditions(Exprs.FinalsConditions);
+  Dir->setPreInits(Exprs.PreInits);
+  return Dir;
+}
+
+OMPTargetTeamsGenericLoopDirective *
+OMPTargetTeamsGenericLoopDirective::CreateEmpty(const ASTContext &C,
+                                                unsigned NumClauses,
+                                                unsigned CollapsedNum,
+                                                EmptyShell) {
+  return createEmptyDirective<OMPTargetTeamsGenericLoopDirective>(
+      C, NumClauses, /*HasAssociatedStmt=*/true,
+      numLoopChildren(CollapsedNum, OMPD_target_teams_loop), CollapsedNum);
+}

clang/lib/AST/StmtPrinter.cpp

 }
 
 void StmtPrinter::VisitOMPTeamsGenericLoopDirective(
     OMPTeamsGenericLoopDirective *Node) {
   Indent() << "#pragma omp teams loop";
   PrintOMPExecutableDirective(Node);
 }
 
+void StmtPrinter::VisitOMPTargetTeamsGenericLoopDirective(
+    OMPTargetTeamsGenericLoopDirective *Node) {
+  Indent() << "#pragma omp target teams loop";
+  PrintOMPExecutableDirective(Node);
+}
+
 //===----------------------------------------------------------------------===//
 //  Expr printing methods.
 //===----------------------------------------------------------------------===//
 
 void StmtPrinter::VisitSourceLocExpr(SourceLocExpr *Node) {
   OS << Node->getBuiltinStr() << "()";
 }
 

clang/lib/AST/StmtProfile.cpp

void StmtProfiler::VisitOMPGenericLoopDirective(
   VisitOMPLoopDirective(S);
 }
 
 void StmtProfiler::VisitOMPTeamsGenericLoopDirective(
     const OMPTeamsGenericLoopDirective *S) {
   VisitOMPLoopDirective(S);
 }
 
+void StmtProfiler::VisitOMPTargetTeamsGenericLoopDirective(
+    const OMPTargetTeamsGenericLoopDirective *S) {
+  VisitOMPLoopDirective(S);
+}
+
 void StmtProfiler::VisitExpr(const Expr *S) {
   VisitStmt(S);
 }
 
 void StmtProfiler::VisitConstantExpr(const ConstantExpr *S) {
   VisitExpr(S);
 }
 

clang/lib/Basic/OpenMPKinds.cpp

return DKind == OMPD_simd || DKind == OMPD_for || DKind == OMPD_for_simd ||
          DKind == OMPD_teams_distribute ||
          DKind == OMPD_teams_distribute_simd ||
          DKind == OMPD_teams_distribute_parallel_for_simd ||
          DKind == OMPD_teams_distribute_parallel_for ||
          DKind == OMPD_target_teams_distribute ||
          DKind == OMPD_target_teams_distribute_parallel_for ||
          DKind == OMPD_target_teams_distribute_parallel_for_simd ||
          DKind == OMPD_target_teams_distribute_simd || DKind == OMPD_tile ||
-         DKind == OMPD_unroll || DKind == OMPD_loop || DKind == OMPD_teams_loop;
+         DKind == OMPD_unroll || DKind == OMPD_loop ||
+         DKind == OMPD_teams_loop || DKind == OMPD_target_teams_loop;
 }
 
 bool clang::isOpenMPWorksharingDirective(OpenMPDirectiveKind DKind) {
   return DKind == OMPD_for || DKind == OMPD_for_simd ||
          DKind == OMPD_sections || DKind == OMPD_section ||
          DKind == OMPD_single || DKind == OMPD_parallel_for ||
          DKind == OMPD_parallel_for_simd || DKind == OMPD_parallel_sections ||
          DKind == OMPD_target_parallel_for ||

 
 bool clang::isOpenMPTargetExecutionDirective(OpenMPDirectiveKind DKind) {
   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_teams || DKind == OMPD_target_teams_distribute ||
          DKind == OMPD_target_teams_distribute_parallel_for ||
          DKind == OMPD_target_teams_distribute_parallel_for_simd ||
-         DKind == OMPD_target_teams_distribute_simd;
+         DKind == OMPD_target_teams_distribute_simd ||
+         DKind == OMPD_target_teams_loop;
 }
 
 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::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 ||
          DKind == OMPD_teams_loop;
 }
 
 bool clang::isOpenMPTeamsDirective(OpenMPDirectiveKind DKind) {
-  return isOpenMPNestingTeamsDirective(DKind) ||
-         DKind == OMPD_target_teams || DKind == OMPD_target_teams_distribute ||
+  return isOpenMPNestingTeamsDirective(DKind) || DKind == OMPD_target_teams ||
+         DKind == OMPD_target_teams_distribute ||
          DKind == OMPD_target_teams_distribute_parallel_for ||
          DKind == OMPD_target_teams_distribute_parallel_for_simd ||
-         DKind == OMPD_target_teams_distribute_simd;
+         DKind == OMPD_target_teams_distribute_simd ||
+         DKind == OMPD_target_teams_loop;
 }
 
 bool clang::isOpenMPSimdDirective(OpenMPDirectiveKind DKind) {
   return DKind == OMPD_simd || DKind == OMPD_for_simd ||
          DKind == OMPD_parallel_for_simd || DKind == OMPD_taskloop_simd ||
          DKind == OMPD_master_taskloop_simd ||
          DKind == OMPD_parallel_master_taskloop_simd ||
          DKind == OMPD_distribute_parallel_for_simd ||
return isOpenMPNestingDistributeDirective(Kind) ||
          Kind == OMPD_teams_distribute_parallel_for ||
          Kind == OMPD_target_teams_distribute ||
          Kind == OMPD_target_teams_distribute_parallel_for ||
          Kind == OMPD_target_teams_distribute_parallel_for_simd ||
          Kind == OMPD_target_teams_distribute_simd;
 }
 
 bool clang::isOpenMPGenericLoopDirective(OpenMPDirectiveKind Kind) {
-  return Kind == OMPD_loop || Kind == OMPD_teams_loop;
+  return Kind == OMPD_loop || Kind == OMPD_teams_loop ||
+         Kind == OMPD_target_teams_loop;
 }
 
 bool clang::isOpenMPPrivate(OpenMPClauseKind Kind) {
   return Kind == OMPC_private || Kind == OMPC_firstprivate ||
          Kind == OMPC_lastprivate || Kind == OMPC_linear ||
          Kind == OMPC_reduction || Kind == OMPC_task_reduction ||
          Kind == OMPC_in_reduction; // TODO add next clauses like 'reduction'.
 }
void clang::getOpenMPCaptureRegions(
   case OMPD_parallel_sections:
   case OMPD_distribute_parallel_for:
   case OMPD_distribute_parallel_for_simd:
     CaptureRegions.push_back(OMPD_parallel);
     break;
   case OMPD_target_teams:
   case OMPD_target_teams_distribute:
   case OMPD_target_teams_distribute_simd:
+  case OMPD_target_teams_loop:
     CaptureRegions.push_back(OMPD_task);
     CaptureRegions.push_back(OMPD_target);
     CaptureRegions.push_back(OMPD_teams);
     break;
   case OMPD_teams:
   case OMPD_teams_distribute:
   case OMPD_teams_distribute_simd:
     CaptureRegions.push_back(OMPD_teams);

clang/lib/CodeGen/CGStmt.cpp

case Stmt::OMPMaskedDirectiveClass:
     EmitOMPMaskedDirective(cast<OMPMaskedDirective>(*S));
     break;
   case Stmt::OMPGenericLoopDirectiveClass:
     EmitOMPGenericLoopDirective(cast<OMPGenericLoopDirective>(*S));
     break;
   case Stmt::OMPTeamsGenericLoopDirectiveClass:
     llvm_unreachable("teams loop directive not supported yet.");
     break;
+  case Stmt::OMPTargetTeamsGenericLoopDirectiveClass:
+    llvm_unreachable("target teams loop directive not supported yet.");
+    break;
   }
 }
 
 bool CodeGenFunction::EmitSimpleStmt(const Stmt *S,
                                      ArrayRef<const Attr *> Attrs) {
   switch (S->getStmtClass()) {
   default:
     return false;

clang/lib/Parse/ParseOpenMP.cpp

static const OpenMPDirectiveKindExWrapper F[][3] = {
       {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, OMPD_loop, OMPD_teams_loop},
       {OMPD_target, OMPD_teams, OMPD_target_teams},
       {OMPD_target_teams, OMPD_distribute, OMPD_target_teams_distribute},
+      {OMPD_target_teams, OMPD_loop, OMPD_target_teams_loop},
       {OMPD_target_teams_distribute, OMPD_parallel,
        OMPD_target_teams_distribute_parallel},
       {OMPD_target_teams_distribute, OMPD_simd,
        OMPD_target_teams_distribute_simd},
       {OMPD_target_teams_distribute_parallel, OMPD_for,
        OMPD_target_teams_distribute_parallel_for},
       {OMPD_target_teams_distribute_parallel_for, OMPD_simd,
        OMPD_target_teams_distribute_parallel_for_simd},
Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirectiveWithExtDecl(
   case OMPD_target_teams_distribute_parallel_for:
   case OMPD_target_teams_distribute_parallel_for_simd:
   case OMPD_target_teams_distribute_simd:
   case OMPD_dispatch:
   case OMPD_masked:
   case OMPD_metadirective:
   case OMPD_loop:
   case OMPD_teams_loop:
+  case OMPD_target_teams_loop:
     Diag(Tok, diag::err_omp_unexpected_directive)
         << 1 << getOpenMPDirectiveName(DKind);
     break;
   default:
     break;
   }
   while (Tok.isNot(tok::annot_pragma_openmp_end))
     ConsumeAnyToken();
StmtResult Parser::ParseOpenMPDeclarativeOrExecutableDirective(
   case OMPD_target:
   case OMPD_teams:
   case OMPD_taskgroup:
   case OMPD_target_data:
   case OMPD_target_parallel:
   case OMPD_target_parallel_for:
   case OMPD_loop:
   case OMPD_teams_loop:
+  case OMPD_target_teams_loop:
   case OMPD_taskloop:
   case OMPD_taskloop_simd:
   case OMPD_master_taskloop:
   case OMPD_master_taskloop_simd:
   case OMPD_parallel_master_taskloop:
   case OMPD_parallel_master_taskloop_simd:
   case OMPD_distribute:
   case OMPD_distribute_parallel_for:

clang/lib/Sema/SemaExceptionSpec.cpp

CanThrowResult Sema::canThrow(const Stmt *S) {
   case Stmt::OMPTeamsDistributeParallelForSimdDirectiveClass:
   case Stmt::OMPTeamsDistributeSimdDirectiveClass:
   case Stmt::OMPInteropDirectiveClass:
   case Stmt::OMPDispatchDirectiveClass:
   case Stmt::OMPMaskedDirectiveClass:
   case Stmt::OMPMetaDirectiveClass:
   case Stmt::OMPGenericLoopDirectiveClass:
   case Stmt::OMPTeamsGenericLoopDirectiveClass:
+  case Stmt::OMPTargetTeamsGenericLoopDirectiveClass:
   case Stmt::ReturnStmtClass:
   case Stmt::SEHExceptStmtClass:
   case Stmt::SEHFinallyStmtClass:
   case Stmt::SEHLeaveStmtClass:
   case Stmt::SEHTryStmtClass:
   case Stmt::SwitchStmtClass:
   case Stmt::WhileStmtClass:
     return canSubStmtsThrow(*this, S);

clang/lib/Sema/SemaOpenMP.cpp

case OMPD_teams_distribute_simd: {
     ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
                              Params);
     break;
   }
   case OMPD_target_teams:
   case OMPD_target_parallel:
   case OMPD_target_parallel_for:
   case OMPD_target_parallel_for_simd:
+  case OMPD_target_teams_loop:
   case OMPD_target_teams_distribute:
   case OMPD_target_teams_distribute_simd: {
     QualType KmpInt32Ty = Context.getIntTypeForBitwidth(32, 1).withConst();
     QualType VoidPtrTy = Context.VoidPtrTy.withConst().withRestrict();
     QualType KmpInt32PtrTy =
         Context.getPointerType(KmpInt32Ty).withConst().withRestrict();
     QualType Args[] = {VoidPtrTy};
     FunctionProtoType::ExtProtoInfo EPI;
StmtResult Sema::ActOnOpenMPExecutableDirective(
   case OMPD_loop:
     Res = ActOnOpenMPGenericLoopDirective(ClausesWithImplicit, AStmt, StartLoc,
                                           EndLoc, VarsWithInheritedDSA);
     break;
   case OMPD_teams_loop:
     Res = ActOnOpenMPTeamsGenericLoopDirective(
         ClausesWithImplicit, AStmt, StartLoc, EndLoc, VarsWithInheritedDSA);
     break;
+  case OMPD_target_teams_loop:
+    Res = ActOnOpenMPTargetTeamsGenericLoopDirective(
+        ClausesWithImplicit, AStmt, StartLoc, EndLoc, VarsWithInheritedDSA);
+    break;
   case OMPD_declare_target:
   case OMPD_end_declare_target:
   case OMPD_threadprivate:
   case OMPD_allocate:
   case OMPD_declare_reduction:
   case OMPD_declare_mapper:
   case OMPD_declare_simd:
   case OMPD_requires:
StmtResult Sema::ActOnOpenMPTeamsGenericLoopDirective(
 
   setFunctionHasBranchProtectedScope();
   DSAStack->setParentTeamsRegionLoc(StartLoc);
 
   return OMPTeamsGenericLoopDirective::Create(
       Context, StartLoc, EndLoc, NestedLoopCount, Clauses, AStmt, B);
 }
 
+StmtResult Sema::ActOnOpenMPTargetTeamsGenericLoopDirective(
+    ArrayRef<OMPClause *> Clauses, Stmt *AStmt, SourceLocation StartLoc,
+    SourceLocation EndLoc, VarsWithInheritedDSAType &VarsWithImplicitDSA) {
+  if (!AStmt)
+    return StmtError();
+
+  // OpenMP 5.1 [2.11.7, loop construct, Restrictions]
+  // A list item may not appear in a lastprivate clause unless it is the
+  // loop iteration variable of a loop that is associated with the construct.
+  if (checkGenericLoopLastprivate(*this, Clauses, OMPD_target_teams_loop,
+                                  DSAStack))
+    return StmtError();
+
+  auto *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();
+  for (int ThisCaptureLevel = getOpenMPCaptureLevels(OMPD_target_teams_loop);
+       ThisCaptureLevel > 1; --ThisCaptureLevel) {
+    CS = cast<CapturedStmt>(CS->getCapturedStmt());
+    // 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();
+  }
+
+  OMPLoopDirective::HelperExprs B;
+  // In presence of clause 'collapse', it will define the nested loops number.
+  unsigned NestedLoopCount =
+      checkOpenMPLoop(OMPD_target_teams_loop, getCollapseNumberExpr(Clauses),
+                      /*OrderedLoopCountExpr=*/nullptr, CS, *this, *DSAStack,
+                      VarsWithImplicitDSA, B);
+  if (NestedLoopCount == 0)
+    return StmtError();
+
+  assert((CurContext->isDependentContext() || B.builtAll()) &&
+         "omp loop exprs were not built");
+
+  setFunctionHasBranchProtectedScope();
+
+  return OMPTargetTeamsGenericLoopDirective::Create(
+      Context, StartLoc, EndLoc, NestedLoopCount, Clauses, AStmt, B);
+}
+
 StmtResult Sema::ActOnOpenMPSingleDirective(ArrayRef<OMPClause *> Clauses,
                                             Stmt *AStmt,
                                             SourceLocation StartLoc,
                                             SourceLocation EndLoc) {
   if (!AStmt)
     return StmtError();
 
   assert(isa<CapturedStmt>(AStmt) && "Captured statement expected");
case OMPC_if:
     case OMPD_cancel:
     case OMPD_parallel:
     case OMPD_parallel_master:
     case OMPD_parallel_sections:
     case OMPD_parallel_for:
     case OMPD_target:
     case OMPD_target_teams:
     case OMPD_target_teams_distribute:
+    case OMPD_target_teams_loop:
     case OMPD_distribute_parallel_for:
     case OMPD_task:
     case OMPD_taskloop:
     case OMPD_master_taskloop:
     case OMPD_target_data:
     case OMPD_simd:
     case OMPD_for_simd:
     case OMPD_distribute_simd:
case OMPC_num_threads:
     case OMPD_declare_simd:
     case OMPD_declare_variant:
     case OMPD_begin_declare_variant:
     case OMPD_end_declare_variant:
     case OMPD_declare_target:
     case OMPD_end_declare_target:
     case OMPD_loop:
     case OMPD_teams_loop:
+    case OMPD_target_teams_loop:
     case OMPD_teams:
     case OMPD_simd:
     case OMPD_tile:
     case OMPD_unroll:
     case OMPD_for:
     case OMPD_for_simd:
     case OMPD_sections:
     case OMPD_section:
case OMPC_num_threads:
     break;
   case OMPC_num_teams:
     switch (DKind) {
     case OMPD_target_teams:
     case OMPD_target_teams_distribute:
     case OMPD_target_teams_distribute_simd:
     case OMPD_target_teams_distribute_parallel_for:
     case OMPD_target_teams_distribute_parallel_for_simd:
+    case OMPD_target_teams_loop:
       CaptureRegion = OMPD_target;
       break;
     case OMPD_teams_distribute_parallel_for:
     case OMPD_teams_distribute_parallel_for_simd:
     case OMPD_teams:
     case OMPD_teams_distribute:
     case OMPD_teams_distribute_simd:
     case OMPD_teams_loop:
case OMPC_num_teams:
     break;
   case OMPC_thread_limit:
     switch (DKind) {
     case OMPD_target_teams:
     case OMPD_target_teams_distribute:
     case OMPD_target_teams_distribute_simd:
     case OMPD_target_teams_distribute_parallel_for:
     case OMPD_target_teams_distribute_parallel_for_simd:
+    case OMPD_target_teams_loop:
       CaptureRegion = OMPD_target;
       break;
     case OMPD_teams_distribute_parallel_for:
     case OMPD_teams_distribute_parallel_for_simd:
     case OMPD_teams:
     case OMPD_teams_distribute:
     case OMPD_teams_distribute_simd:
     case OMPD_teams_loop:
case OMPC_schedule:
     case OMPD_declare_simd:
     case OMPD_declare_variant:
     case OMPD_begin_declare_variant:
     case OMPD_end_declare_variant:
     case OMPD_declare_target:
     case OMPD_end_declare_target:
     case OMPD_loop:
     case OMPD_teams_loop:
+    case OMPD_target_teams_loop:
     case OMPD_simd:
     case OMPD_tile:
     case OMPD_unroll:
     case OMPD_sections:
     case OMPD_section:
     case OMPD_single:
     case OMPD_master:
     case OMPD_masked:
case OMPC_dist_schedule:
     case OMPD_declare_simd:
     case OMPD_declare_variant:
     case OMPD_begin_declare_variant:
     case OMPD_end_declare_variant:
     case OMPD_declare_target:
     case OMPD_end_declare_target:
     case OMPD_loop:
     case OMPD_teams_loop:
+    case OMPD_target_teams_loop:
     case OMPD_simd:
     case OMPD_tile:
     case OMPD_unroll:
     case OMPD_for:
     case OMPD_for_simd:
     case OMPD_sections:
     case OMPD_section:
     case OMPD_single:
case OMPC_device:
     case OMPD_target_teams:
     case OMPD_target_parallel:
     case OMPD_target_teams_distribute:
     case OMPD_target_teams_distribute_simd:
     case OMPD_target_parallel_for:
     case OMPD_target_parallel_for_simd:
     case OMPD_target_teams_distribute_parallel_for:
     case OMPD_target_teams_distribute_parallel_for_simd:
+    case OMPD_target_teams_loop:
     case OMPD_dispatch:
       CaptureRegion = OMPD_task;
       break;
     case OMPD_target_data:
     case OMPD_interop:
       // Do not capture device-clause expressions.
       break;
     case OMPD_teams_distribute_parallel_for:
case OMPC_priority:
     case OMPD_declare_simd:
     case OMPD_declare_variant:
     case OMPD_begin_declare_variant:
     case OMPD_end_declare_variant:
     case OMPD_declare_target:
     case OMPD_end_declare_target:
     case OMPD_loop:
     case OMPD_teams_loop:
+    case OMPD_target_teams_loop:
     case OMPD_simd:
     case OMPD_tile:
     case OMPD_unroll:
     case OMPD_for:
     case OMPD_for_simd:
     case OMPD_sections:
     case OMPD_section:
     case OMPD_single:

clang/lib/Sema/TreeTransform.h

StmtResult TreeTransform<Derived>::TransformOMPTeamsGenericLoopDirective(
   DeclarationNameInfo DirName;
   getDerived().getSema().StartOpenMPDSABlock(OMPD_teams_loop, DirName, nullptr,
                                              D->getBeginLoc());
   StmtResult Res = getDerived().TransformOMPExecutableDirective(D);
   getDerived().getSema().EndOpenMPDSABlock(Res.get());
   return Res;
 }
 
+template <typename Derived>
+StmtResult TreeTransform<Derived>::TransformOMPTargetTeamsGenericLoopDirective(
+    OMPTargetTeamsGenericLoopDirective *D) {
+  DeclarationNameInfo DirName;
+  getDerived().getSema().StartOpenMPDSABlock(OMPD_target_teams_loop, DirName,
+                                             nullptr, D->getBeginLoc());
+  StmtResult 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())
     return nullptr;

clang/lib/Serialization/ASTReaderStmt.cpp

void ASTStmtReader::VisitOMPGenericLoopDirective(OMPGenericLoopDirective *D) {
   VisitOMPLoopDirective(D);
 }
 
 void ASTStmtReader::VisitOMPTeamsGenericLoopDirective(
     OMPTeamsGenericLoopDirective *D) {
   VisitOMPLoopDirective(D);
 }
 
+void ASTStmtReader::VisitOMPTargetTeamsGenericLoopDirective(
+    OMPTargetTeamsGenericLoopDirective *D) {
+  VisitOMPLoopDirective(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_GENERIC_LOOP_DIRECTIVE: {
       unsigned CollapsedNum = Record[ASTStmtReader::NumStmtFields];
       unsigned NumClauses = Record[ASTStmtReader::NumStmtFields + 1];
       S = OMPTeamsGenericLoopDirective::CreateEmpty(Context, NumClauses,
                                                     CollapsedNum, Empty);
       break;
     }
 
+    case STMT_OMP_TARGET_TEAMS_GENERIC_LOOP_DIRECTIVE: {
+      unsigned CollapsedNum = Record[ASTStmtReader::NumStmtFields];
+      unsigned NumClauses = Record[ASTStmtReader::NumStmtFields + 1];
+      S = OMPTargetTeamsGenericLoopDirective::CreateEmpty(Context, NumClauses,
+                                                          CollapsedNum, Empty);
+      break;
+    }
+
     case EXPR_CXX_OPERATOR_CALL:
       S = CXXOperatorCallExpr::CreateEmpty(
           Context, /*NumArgs=*/Record[ASTStmtReader::NumExprFields],
           /*HasFPFeatures=*/Record[ASTStmtReader::NumExprFields + 1], Empty);
       break;
 
     case EXPR_CXX_MEMBER_CALL:
       S = CXXMemberCallExpr::CreateEmpty(

clang/lib/Serialization/ASTWriterStmt.cpp

 }
 
 void ASTStmtWriter::VisitOMPTeamsGenericLoopDirective(
     OMPTeamsGenericLoopDirective *D) {
   VisitOMPLoopDirective(D);
   Code = serialization::STMT_OMP_TEAMS_GENERIC_LOOP_DIRECTIVE;
 }
 
+void ASTStmtWriter::VisitOMPTargetTeamsGenericLoopDirective(
+    OMPTargetTeamsGenericLoopDirective *D) {
+  VisitOMPLoopDirective(D);
+  Code = serialization::STMT_OMP_TARGET_TEAMS_GENERIC_LOOP_DIRECTIVE;
+}
+
 //===----------------------------------------------------------------------===//
 // ASTWriter Implementation
 //===----------------------------------------------------------------------===//
 
 unsigned ASTWriter::RecordSwitchCaseID(SwitchCase *S) {
   assert(SwitchCaseIDs.find(S) == SwitchCaseIDs.end() &&
          "SwitchCase recorded twice");
   unsigned NextID = SwitchCaseIDs.size();

clang/lib/StaticAnalyzer/Core/ExprEngine.cpp

switch (S->getStmtClass()) {
     case Stmt::OMPTargetTeamsDistributeParallelForSimdDirectiveClass:
     case Stmt::OMPTargetTeamsDistributeSimdDirectiveClass:
     case Stmt::OMPTileDirectiveClass:
     case Stmt::OMPInteropDirectiveClass:
     case Stmt::OMPDispatchDirectiveClass:
     case Stmt::OMPMaskedDirectiveClass:
     case Stmt::OMPGenericLoopDirectiveClass:
     case Stmt::OMPTeamsGenericLoopDirectiveClass:
+    case Stmt::OMPTargetTeamsGenericLoopDirectiveClass:
     case Stmt::CapturedStmtClass:
     case Stmt::OMPUnrollDirectiveClass:
     case Stmt::OMPMetaDirectiveClass: {
       const ExplodedNode *node = Bldr.generateSink(S, Pred, Pred->getState());
       Engine.addAbortedBlock(node, currBldrCtx->getBlock());
       break;
     }
 

clang/test/Analysis/cfg-openmp.cpp

 // CHECK-DAG:  [[#TDSB-2]]: [B3.[[#TDSB-3]]] (ImplicitCastExpr, LValueToRValue, int)
 // CHECK-DAG:  [[#TDSB-1]]: argc
 // CHECK-DAG:  [[#TDSB]]: [B3.[[#TDSB-1]]] = [B3.[[#TDSB-2]]]
 #pragma omp target
 #pragma omp teams loop firstprivate(fp) reduction(+:rd)
   for (int i = 0; i < 10; ++i)
     argc = x;
 }
+
+// CHECK-LABEL:  void targetteamsloop(int argc)
+void targetteamsloop(int argc) {
+  int x, cond, fp, rd, map;
+// CHECK-DAG:   [B3]
+// CHECK-DAG:  [[#TTDB:]]: x
+// CHECK-DAG:  [[#TTDB+1]]: [B3.[[#TTDB]]] (ImplicitCastExpr, LValueToRValue, int)
+// CHECK-DAG:  [[#TTDB+2]]: argc
+// CHECK-DAG:  [[#TTDB+3]]: [B3.[[#TTDB+2]]] = [B3.[[#TTDB+1]]]
+// CHECK-DAG:   [B1]
+// CHECK-DAG:  [[#TTD:]]: cond
+// CHECK-DAG:  [[#TTD+1]]: [B1.[[#TTD]]] (ImplicitCastExpr, LValueToRValue, int)
+// CHECK-DAG:  [[#TTD+2]]: [B1.[[#TTD+1]]] (ImplicitCastExpr, IntegralToBoolean, _Bool)
+// CHECK-DAG:  [[#TTD+3]]: fp
+// CHECK-DAG:  [[#TTD+4]]: rd
+// CHECK-DAG:  [[#TTD+5]]: [B3.[[#TTDB+2]]]
+// CHECK-DAG:  [[#TTD+6]]: [B3.[[#TTDB]]]
+// CHECK-DAG:  [[#TTD+7]]: #pragma omp target teams loop if(cond) firstprivate(fp) reduction(+: rd) map(alloc: map)
+// CHECK-DAG:    for (int i = 0;
+// CHECK-DAG:        [B3.[[#TTDB+3]]];
+#pragma omp target teams loop if(cond) firstprivate(fp) reduction(+:rd) map(alloc:map)
+  for (int i = 0; i < 10; ++i)
+    argc = x;
+}

clang/test/OpenMP/target_teams_generic_loop_ast_print.cpp

+// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -fopenmp -fopenmp-version=51 \
+// RUN:   -fsyntax-only -verify %s
+
+// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -fopenmp -fopenmp-version=51 \
+// RUN:   -ast-print %s | FileCheck %s
+
+// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -fopenmp -fopenmp-version=51 \
+// RUN:   -emit-pch -o %t %s
+
+// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -fopenmp -fopenmp-version=51 \
+// RUN:   -include-pch %t -ast-print %s | FileCheck %s
+
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+typedef void **omp_allocator_handle_t;
+extern const omp_allocator_handle_t omp_null_allocator;
+extern const omp_allocator_handle_t omp_default_mem_alloc;
+extern const omp_allocator_handle_t omp_large_cap_mem_alloc;
+extern const omp_allocator_handle_t omp_const_mem_alloc;
+extern const omp_allocator_handle_t omp_high_bw_mem_alloc;
+extern const omp_allocator_handle_t omp_low_lat_mem_alloc;
+extern const omp_allocator_handle_t omp_cgroup_mem_alloc;
+extern const omp_allocator_handle_t omp_pteam_mem_alloc;
+extern const omp_allocator_handle_t omp_thread_mem_alloc;
+
+//CHECK: template <typename T, int C, int D> void templ_foo(T t) {
+//CHECK:   T j, z;
+//CHECK:   #pragma omp target teams loop device(D) collapse(C) reduction(+: z) lastprivate(j) bind(thread) num_teams(C + 2)
+//CHECK:   for (T i = 0; i < t; ++i)
+//CHECK:       for (j = 0; j < t; ++j)
+//CHECK:           z += i + j;
+//CHECK: }
+
+//CHECK: template<> void templ_foo<int, 2, 0>(int t) {
+//CHECK:     int j, z;
+//CHECK:     #pragma omp target teams loop device(0) collapse(2) reduction(+: z) lastprivate(j) bind(thread) num_teams(2 + 2)
+//CHECK:         for (int i = 0; i < t; ++i)
+//CHECK:             for (j = 0; j < t; ++j)
+//CHECK:                 z += i + j;
+//CHECK: }
+template <typename T, int C, int D>
+void templ_foo(T t) {
+
+  T j,z;
+  #pragma omp target teams loop device(D) collapse(C) reduction(+:z) lastprivate(j) bind(thread) num_teams(C+2)
+  for (T i = 0; i<t; ++i)
+    for (j = 0; j<t; ++j)
+      z += i+j;
+}
+
+
+//CHECK: void test() {
+void test() {
+  constexpr int N = 100;
+  float MTX[N][N];
+  int aaa[1000];
+
+  //CHECK: #pragma omp target teams loop map(tofrom: MTX)
+  #pragma omp target teams loop map(MTX)
+  for (auto j = 0; j < N; ++j) {
+    MTX[0][j] = 0;
+  }
+
+  int j, z, z1;
+  //CHECK: #pragma omp target teams loop collapse(2) private(z) lastprivate(j) order(concurrent) reduction(+: z1) bind(parallel)
+  #pragma omp target teams loop collapse(2) private(z) lastprivate(j) \
+                         order(concurrent) reduction(+:z1) bind(parallel)
+  for (auto i = 0; i < N; ++i) {
+    for (j = 0; j < N; ++j) {
+      z = i+j;
+      MTX[i][j] = z;
+      z1 += z;
+    }
+  }
+
+  //CHECK: #pragma omp target teams loop bind(teams) num_teams(16) thread_limit(8) default(none)
+  #pragma omp target teams loop bind(teams) num_teams(16) thread_limit(8) default(none)
+  for (auto i = 0; i < N; ++i) { }
+
+  int pr;
+  int zzz;
+  //CHECK: #pragma omp target teams loop private(zzz) uses_allocators(omp_default_mem_alloc) allocate(omp_default_mem_alloc: zzz) if(1) device(0) map(tofrom: pr)
+  #pragma omp target teams loop private(zzz) uses_allocators(omp_default_mem_alloc) allocate(omp_default_mem_alloc:zzz) if(1) device(0) map(tofrom:pr)
+  for (int i=0; i<1000; ++i) {
+    zzz = i + 1;
+    pr = 33;
+  }
+
+  int fpr = 10;
+  int k;
+  int s = 20;
+  //CHECK: #pragma omp target teams loop bind(teams) private(pr) firstprivate(fpr) shared(s) allocate(k) reduction(+: k)
+  #pragma omp target teams loop bind(teams) private(pr) firstprivate(fpr) \
+                        shared(s) allocate(k)  reduction(+:k)
+  for (auto i = 0; i < N; ++i) {
+    pr = i + fpr + s;
+  }
+
+  short y = 3;
+  //CHECK: #pragma omp target teams loop map(tofrom: y) depend(out : y)
+  #pragma omp target teams loop map(tofrom:y) depend(out:y)
+  for (int i=0; i<10; ++i) {
+    y = 3+i;
+  }
+}
+
+//CHECK: void nobindingfunc() {
+void nobindingfunc()
+{
+  //CHECK: #pragma omp target teams loop
+  #pragma omp target teams loop
+  for (int i=0; i<10; ++i) { }
+}
+
+void bar()
+{
+  templ_foo<int,2,0>(8);
+}
+
+#endif // HEADER

clang/test/OpenMP/target_teams_generic_loop_messages.cpp

+// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -verify -fopenmp \
+// RUN:  -fopenmp-version=51 -Wuninitialized %s
+
+void foo()
+{
+  int i,j,k;
+  int z;
+
+  // expected-error@+2 {{statement after '#pragma omp target teams loop' must be a for loop}}
+  #pragma omp target teams loop bind(thread)
+  i = 0;
+
+  // OpenMP 5.1 [2.22 Nesting of regions]
+  //
+  // A barrier region may not be closely nested inside a worksharing, loop,
+  // task, taskloop, critical, ordered, atomic, or masked region.
+
+  // expected-error@+3 {{region cannot be closely nested inside 'target teams loop' region}}
+  #pragma omp target teams loop bind(thread)
+  for (i=0; i<1000; ++i) {
+    #pragma omp barrier
+  }
+
+  // A masked region may not be closely nested inside a worksharing, loop,
+  // atomic, task, or taskloop region.
+
+  // expected-error@+3 {{region cannot be closely nested inside 'target teams loop' region}}
+  #pragma omp target teams loop bind(thread)
+  for (i=0; i<1000; ++i) {
+    #pragma omp masked filter(2)
+    { }
+  }
+
+  // An ordered region that corresponds to an ordered construct without any
+  // clause or with the threads or depend clause may not be closely nested
+  // inside a critical, ordered, loop, atomic, task, or taskloop region.
+
+  // expected-error@+3 {{region cannot be closely nested inside 'target teams loop' region; perhaps you forget to enclose 'omp ordered' directive into a for or a parallel for region with 'ordered' clause?}}
+  #pragma omp target teams loop bind(thread)
+  for (i=0; i<1000; ++i) {
+    #pragma omp ordered
+    { }
+  }
+
+  // expected-error@+3 {{region cannot be closely nested inside 'target teams loop' region; perhaps you forget to enclose 'omp ordered' directive into a for or a parallel for region with 'ordered' clause?}}
+  #pragma omp target teams loop bind(thread)
+  for (i=0; i<1000; ++i) {
+    #pragma omp ordered threads
+    { }
+  }
+
+  // expected-error@+3 {{region cannot be closely nested inside 'target teams loop' region; perhaps you forget to enclose 'omp ordered' directive into a for or a parallel for region with 'ordered' clause?}}
+  #pragma omp target teams loop bind(thread)
+  for (i=0; i<1000; ++i) {
+    #pragma omp ordered depend(source)
+  }
+
+  // bind clause
+
+  // expected-error@+1 {{directive '#pragma omp target teams loop' cannot contain more than one 'bind' clause}}
+  #pragma omp target teams loop bind(thread) bind(thread)
+  for (i=0; i<1000; ++i) {
+  }
+
+  // expected-error@+1 {{expected 'teams', 'parallel' or 'thread' in OpenMP clause 'bind'}}
+  #pragma omp target teams loop bind(other)
+  for (i=0; i<1000; ++i) {
+  }
+
+  // collapse clause
+
+  // expected-error@+4 {{expected 2 for loops after '#pragma omp target teams loop', but found only 1}}
+  // expected-note@+1 {{as specified in 'collapse' clause}}
+  #pragma omp target teams loop collapse(2) bind(thread)
+  for (i=0; i<1000; ++i)
+    z = i+11;
+
+  // expected-error@+1 {{directive '#pragma omp target teams loop' cannot contain more than one 'collapse' clause}}
+  #pragma omp target teams loop collapse(2) collapse(2) bind(thread)
+  for (i=0; i<1000; ++i)
+    for (j=0; j<1000; ++j)
+      z = i+j+11;
+
+  // order clause
+
+  // expected-error@+1 {{expected 'concurrent' in OpenMP clause 'order'}}
+  #pragma omp target teams loop order(foo) bind(thread)
+  for (i=0; i<1000; ++i)
+    z = i+11;
+
+  // private clause
+
+  // expected-error@+1 {{use of undeclared identifier 'undef_var'}}
+  #pragma omp target teams loop private(undef_var) bind(thread)
+  for (i=0; i<1000; ++i)
+    z = i+11;
+
+  // lastprivate
+
+  // A list item may not appear in a lastprivate clause unless it is the loop
+  // iteration variable of a loop that is associated with the construct.
+
+  // expected-error@+1 {{only loop iteration variables are allowed in 'lastprivate' clause in 'omp target teams loop' directives}}
+  #pragma omp target teams loop lastprivate(z) bind(thread)
+  for (i=0; i<1000; ++i) {
+    z = i+11;
+  }
+
+  // expected-error@+1 {{only loop iteration variables are allowed in 'lastprivate' clause in 'omp target teams loop' directives}}
+  #pragma omp target teams loop lastprivate(k) collapse(2) bind(thread)
+  for (i=0; i<1000; ++i)
+    for (j=0; j<1000; ++j)
+      for (k=0; k<1000; ++k)
+        z = i+j+k+11;
+
+  // reduction
+
+  // expected-error@+1 {{use of undeclared identifier 'undef_var'}}
+  #pragma omp target teams loop reduction(+:undef_var) bind(thread)
+  for (i=0; i<1000; ++i)
+    z = i+11;
+
+  // nowait
+
+  // expected-error@+1 {{directive '#pragma omp target teams loop' cannot contain more than one 'nowait' clause}}
+  #pragma omp target teams loop nowait nowait
+  for (i=0; i<1000; ++i)
+    z = i+11;
+
+}
+
+template <typename T, int C>
+void templ_test(T t) {
+  T i,z;
+
+  // expected-error@+4 {{expected 2 for loops after '#pragma omp target teams loop', but found only 1}}
+  // expected-note@+1 {{as specified in 'collapse' clause}}
+  #pragma omp target teams loop collapse(C) bind(thread)
+  for (i=0; i<1000; ++i)
+    z = i+11;
+
+  // expected-error@+1 {{only loop iteration variables are allowed in 'lastprivate' clause in 'omp target teams loop' directives}}
+  #pragma omp target teams loop lastprivate(z) bind(thread)
+  for (i=0; i<1000; ++i) {
+    z = i+11;
+  }
+}
+
+void bar()
+{
+  templ_test<int, 2>(16); // expected-note {{in instantiation of function template specialization 'templ_test<int, 2>' requested here}}
+}

clang/tools/libclang/CIndex.cpp

CXString clang_getCursorKindSpelling(enum CXCursorKind Kind) {
   case CXCursor_OMPDispatchDirective:
     return cxstring::createRef("OMPDispatchDirective");
   case CXCursor_OMPMaskedDirective:
     return cxstring::createRef("OMPMaskedDirective");
   case CXCursor_OMPGenericLoopDirective:
     return cxstring::createRef("OMPGenericLoopDirective");
   case CXCursor_OMPTeamsGenericLoopDirective:
     return cxstring::createRef("OMPTeamsGenericLoopDirective");
+  case CXCursor_OMPTargetTeamsGenericLoopDirective:
+    return cxstring::createRef("OMPTargetTeamsGenericLoopDirective");
   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");

clang/tools/libclang/CXCursor.cpp

case Stmt::OMPMaskedDirectiveClass:
     K = CXCursor_OMPMaskedDirective;
     break;
   case Stmt::OMPGenericLoopDirectiveClass:
     K = CXCursor_OMPGenericLoopDirective;
     break;
   case Stmt::OMPTeamsGenericLoopDirectiveClass:
     K = CXCursor_OMPTeamsGenericLoopDirective;
     break;
+  case Stmt::OMPTargetTeamsGenericLoopDirectiveClass:
+    K = CXCursor_OMPTargetTeamsGenericLoopDirective;
+    break;
   case Stmt::BuiltinBitCastExprClass:
     K = CXCursor_BuiltinBitCastExpr;
   }
 
   CXCursor C = {K, 0, {Parent, S, TU}};
   return C;
 }
 

llvm/include/llvm/Frontend/OpenMP/OMP.td

let allowedOnceClauses = [
     VersionedClause<OMPC_Bind, 50>,
     VersionedClause<OMPC_Collapse>,
     VersionedClause<OMPC_Default>,
     VersionedClause<OMPC_NumTeams>,
     VersionedClause<OMPC_Order>,
     VersionedClause<OMPC_ThreadLimit>,
   ];
 }
+def OMP_target_teams_loop : Directive<"target teams loop"> {
+  let allowedClauses = [
+    VersionedClause<OMPC_Allocate>,
+    VersionedClause<OMPC_Depend>,
+    VersionedClause<OMPC_DefaultMap>,
+    VersionedClause<OMPC_Device>,
+    VersionedClause<OMPC_FirstPrivate>,
+    VersionedClause<OMPC_IsDevicePtr>,
+    VersionedClause<OMPC_LastPrivate>,
+    VersionedClause<OMPC_Map>,
+    VersionedClause<OMPC_Private>,
+    VersionedClause<OMPC_Reduction>,
+    VersionedClause<OMPC_Shared>,
+    VersionedClause<OMPC_UsesAllocators, 50>
+  ];
+  let allowedOnceClauses = [
+    VersionedClause<OMPC_Bind, 50>,
+    VersionedClause<OMPC_Collapse>,
+    VersionedClause<OMPC_Default>,
+    VersionedClause<OMPC_If>,
+    VersionedClause<OMPC_NoWait>,
+    VersionedClause<OMPC_NumTeams>,
+    VersionedClause<OMPC_Order>,
+    VersionedClause<OMPC_ThreadLimit>,
+  ];
+}
 def OMP_Metadirective : Directive<"metadirective"> {
   let allowedClauses = [VersionedClause<OMPC_When>];
   let allowedOnceClauses = [VersionedClause<OMPC_Default>];
 }
 def OMP_Unknown : Directive<"unknown"> {
   let isDefault = true;
 }