[OpenACC] Loop construct basic Sema and AST work (#93742)

This patch implements the 'loop' construct AST, as well as the basic appertainment rule. Additionally, it sets up the 'parent' compute construct, which is necessary for codegen/other diagnostics. A 'loop' can apply to a for or range-for loop, otherwise it has no other restrictions (though some of its clauses do).
2025-04-24 23:36:08 +00:00 · 2024-06-05 06:21:48 -07:00 · 2024-06-05 06:21:48 -07:00 · 42f4e505a3
commit 42f4e505a3
parent 2d9b83750f
47 changed files with 892 additions and 310 deletions
--- a/clang/include/clang-c/Index.h
+++ b/clang/include/clang-c/Index.h
@ -2150,7 +2150,11 @@ enum CXCursorKind {
   */
  CXCursor_OpenACCComputeConstruct = 320,

-  CXCursor_LastStmt = CXCursor_OpenACCComputeConstruct,
+  /** OpenACC Loop Construct.
+   */
+  CXCursor_OpenACCLoopConstruct = 321,
+
+  CXCursor_LastStmt = CXCursor_OpenACCLoopConstruct,

  /**
   * Cursor that represents the translation unit itself.
--- a/clang/include/clang/AST/RecursiveASTVisitor.h
+++ b/clang/include/clang/AST/RecursiveASTVisitor.h
@ -4000,6 +4000,8 @@ bool RecursiveASTVisitor<Derived>::VisitOpenACCClauseList(

 DEF_TRAVERSE_STMT(OpenACCComputeConstruct,
                  { TRY_TO(TraverseOpenACCAssociatedStmtConstruct(S)); })
+DEF_TRAVERSE_STMT(OpenACCLoopConstruct,
+                  { TRY_TO(TraverseOpenACCAssociatedStmtConstruct(S)); })

 // FIXME: look at the following tricky-seeming exprs to see if we
 // need to recurse on anything.  These are ones that have methods
--- a/clang/include/clang/AST/StmtOpenACC.h
+++ b/clang/include/clang/AST/StmtOpenACC.h
@ -113,6 +113,8 @@ public:
    return const_cast<OpenACCAssociatedStmtConstruct *>(this)->children();
  }
 };
+
+class OpenACCLoopConstruct;
 /// This class represents a compute construct, representing a 'Kind' of
 /// `parallel', 'serial', or 'kernel'. These constructs are associated with a
 /// 'structured block', defined as:
@ -165,6 +167,11 @@ class OpenACCComputeConstruct final
  }

  void setStructuredBlock(Stmt *S) { setAssociatedStmt(S); }
+  // Serialization helper function that searches the structured block for 'loop'
+  // constructs that should be associated with this, and sets their parent
+  // compute construct to this one. This isn't necessary normally, since we have
+  // the ability to record the state during parsing.
+  void findAndSetChildLoops();

 public:
  static bool classof(const Stmt *T) {
@ -176,12 +183,74 @@ public:
  static OpenACCComputeConstruct *
  Create(const ASTContext &C, OpenACCDirectiveKind K, SourceLocation BeginLoc,
         SourceLocation DirectiveLoc, SourceLocation EndLoc,
-         ArrayRef<const OpenACCClause *> Clauses, Stmt *StructuredBlock);
+         ArrayRef<const OpenACCClause *> Clauses, Stmt *StructuredBlock,
+         ArrayRef<OpenACCLoopConstruct *> AssociatedLoopConstructs);

  Stmt *getStructuredBlock() { return getAssociatedStmt(); }
  const Stmt *getStructuredBlock() const {
    return const_cast<OpenACCComputeConstruct *>(this)->getStructuredBlock();
  }
 };
+/// This class represents a 'loop' construct.  The 'loop' construct applies to a
+/// 'for' loop (or range-for loop), and is optionally associated with a Compute
+/// Construct.
+class OpenACCLoopConstruct final
+    : public OpenACCAssociatedStmtConstruct,
+      public llvm::TrailingObjects<OpenACCLoopConstruct,
+                                   const OpenACCClause *> {
+  // The compute construct this loop is associated with, or nullptr if this is
+  // an orphaned loop construct, or if it hasn't been set yet.  Because we
+  // construct the directives at the end of their statement, the 'parent'
+  // construct is not yet available at the time of construction, so this needs
+  // to be set 'later'.
+  const OpenACCComputeConstruct *ParentComputeConstruct = nullptr;
+
+  friend class ASTStmtWriter;
+  friend class ASTStmtReader;
+  friend class ASTContext;
+  friend class OpenACCComputeConstruct;
+
+  OpenACCLoopConstruct(unsigned NumClauses);
+
+  OpenACCLoopConstruct(SourceLocation Start, SourceLocation DirLoc,
+                       SourceLocation End,
+                       ArrayRef<const OpenACCClause *> Clauses, Stmt *Loop);
+  void setLoop(Stmt *Loop);
+
+  void setParentComputeConstruct(OpenACCComputeConstruct *CC) {
+    assert(!ParentComputeConstruct && "Parent already set?");
+    ParentComputeConstruct = CC;
+  }
+
+public:
+  static bool classof(const Stmt *T) {
+    return T->getStmtClass() == OpenACCLoopConstructClass;
+  }
+
+  static OpenACCLoopConstruct *CreateEmpty(const ASTContext &C,
+                                           unsigned NumClauses);
+
+  static OpenACCLoopConstruct *
+  Create(const ASTContext &C, SourceLocation BeginLoc, SourceLocation DirLoc,
+         SourceLocation EndLoc, ArrayRef<const OpenACCClause *> Clauses,
+         Stmt *Loop);
+
+  Stmt *getLoop() { return getAssociatedStmt(); }
+  const Stmt *getLoop() const {
+    return const_cast<OpenACCLoopConstruct *>(this)->getLoop();
+  }
+
+  /// OpenACC 3.3 2.9:
+  /// An orphaned loop construct is a loop construct that is not lexically
+  /// enclosed within a compute construct. The parent compute construct of a
+  /// loop construct is the nearest compute construct that lexically contains
+  /// the loop construct.
+  bool isOrphanedLoopConstruct() const {
+    return ParentComputeConstruct == nullptr;
+  }
+  const OpenACCComputeConstruct *getParentComputeConstruct() const {
+    return ParentComputeConstruct;
+  }
+};
 } // namespace clang
 #endif // LLVM_CLANG_AST_STMTOPENACC_H
--- a/clang/include/clang/AST/TextNodeDumper.h
+++ b/clang/include/clang/AST/TextNodeDumper.h
@ -408,6 +408,7 @@ public:
  VisitLifetimeExtendedTemporaryDecl(const LifetimeExtendedTemporaryDecl *D);
  void VisitHLSLBufferDecl(const HLSLBufferDecl *D);
  void VisitOpenACCConstructStmt(const OpenACCConstructStmt *S);
+  void VisitOpenACCLoopConstruct(const OpenACCLoopConstruct *S);
 };

 } // namespace clang
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@ -12413,6 +12413,9 @@ def err_acc_reduction_composite_type
 def err_acc_reduction_composite_member_type :Error<
    "OpenACC 'reduction' composite variable must not have non-scalar field">;
 def note_acc_reduction_composite_member_loc : Note<"invalid field is here">;
+def err_acc_loop_not_for_loop
+    : Error<"OpenACC 'loop' construct can only be applied to a 'for' loop">;
+def note_acc_construct_here : Note<"'%0' construct is here">;

 // AMDGCN builtins diagnostics
 def err_amdgcn_global_load_lds_size_invalid_value : Error<"invalid size value">;
--- a/clang/include/clang/Basic/StmtNodes.td
+++ b/clang/include/clang/Basic/StmtNodes.td
@ -302,3 +302,4 @@ def OpenACCConstructStmt : StmtNode<Stmt, /*abstract=*/1>;
 def OpenACCAssociatedStmtConstruct
    : StmtNode<OpenACCConstructStmt, /*abstract=*/1>;
 def OpenACCComputeConstruct : StmtNode<OpenACCAssociatedStmtConstruct>;
+def OpenACCLoopConstruct : StmtNode<OpenACCAssociatedStmtConstruct>;
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@ -15,6 +15,7 @@
 #define LLVM_CLANG_SEMA_SEMAOPENACC_H

 #include "clang/AST/DeclGroup.h"
+#include "clang/AST/StmtOpenACC.h"
 #include "clang/Basic/OpenACCKinds.h"
 #include "clang/Basic/SourceLocation.h"
 #include "clang/Sema/Ownership.h"
@ -25,6 +26,15 @@ namespace clang {
 class OpenACCClause;

 class SemaOpenACC : public SemaBase {
+private:
+  /// A collection of loop constructs in the compute construct scope that
+  /// haven't had their 'parent' compute construct set yet. Entires will only be
+  /// made to this list in the case where we know the loop isn't an orphan.
+  llvm::SmallVector<OpenACCLoopConstruct *> ParentlessLoopConstructs;
+  /// Whether we are inside of a compute construct, and should add loops to the
+  /// above collection.
+  bool InsideComputeConstruct = false;
+
 public:
  // Redeclaration of the version in OpenACCClause.h.
  using DeviceTypeArgument = std::pair<IdentifierInfo *, SourceLocation>;
@ -394,7 +404,8 @@ public:
  bool ActOnStartDeclDirective(OpenACCDirectiveKind K, SourceLocation StartLoc);
  /// Called when we encounter an associated statement for our construct, this
  /// should check legality of the statement as it appertains to this Construct.
-  StmtResult ActOnAssociatedStmt(OpenACCDirectiveKind K, StmtResult AssocStmt);
+  StmtResult ActOnAssociatedStmt(SourceLocation DirectiveLoc,
+                                 OpenACCDirectiveKind K, StmtResult AssocStmt);

  /// Called after the directive has been completely parsed, including the
  /// declaration group or associated statement.
@ -431,6 +442,20 @@ public:
                                   Expr *LowerBound,
                                   SourceLocation ColonLocFirst, Expr *Length,
                                   SourceLocation RBLoc);
+
+  /// Helper type for the registration/assignment of constructs that need to
+  /// 'know' about their parent constructs and hold a reference to them, such as
+  /// Loop needing its parent construct.
+  class AssociatedStmtRAII {
+    SemaOpenACC &SemaRef;
+    bool WasInsideComputeConstruct;
+    OpenACCDirectiveKind DirKind;
+    llvm::SmallVector<OpenACCLoopConstruct *> ParentlessLoopConstructs;
+
+  public:
+    AssociatedStmtRAII(SemaOpenACC &, OpenACCDirectiveKind);
+    ~AssociatedStmtRAII();
+  };
 };

 } // namespace clang
--- a/clang/include/clang/Serialization/ASTBitCodes.h
+++ b/clang/include/clang/Serialization/ASTBitCodes.h
@ -1946,6 +1946,7 @@ enum StmtCode {

  // OpenACC Constructs
  STMT_OPENACC_COMPUTE_CONSTRUCT,
+  STMT_OPENACC_LOOP_CONSTRUCT,
 };

 /// The kinds of designators that can occur in a
--- a/clang/lib/AST/StmtOpenACC.cpp
+++ b/clang/lib/AST/StmtOpenACC.cpp
@ -12,6 +12,8 @@

 #include "clang/AST/StmtOpenACC.h"
 #include "clang/AST/ASTContext.h"
+#include "clang/AST/RecursiveASTVisitor.h"
+#include "clang/AST/StmtCXX.h"
 using namespace clang;

 OpenACCComputeConstruct *
@ -26,11 +28,98 @@ OpenACCComputeConstruct::CreateEmpty(const ASTContext &C, unsigned NumClauses) {
 OpenACCComputeConstruct *OpenACCComputeConstruct::Create(
    const ASTContext &C, OpenACCDirectiveKind K, SourceLocation BeginLoc,
    SourceLocation DirLoc, SourceLocation EndLoc,
-    ArrayRef<const OpenACCClause *> Clauses, Stmt *StructuredBlock) {
+    ArrayRef<const OpenACCClause *> Clauses, Stmt *StructuredBlock,
+    ArrayRef<OpenACCLoopConstruct *> AssociatedLoopConstructs) {
  void *Mem = C.Allocate(
      OpenACCComputeConstruct::totalSizeToAlloc<const OpenACCClause *>(
          Clauses.size()));
  auto *Inst = new (Mem) OpenACCComputeConstruct(K, BeginLoc, DirLoc, EndLoc,
                                                 Clauses, StructuredBlock);
+
+  llvm::for_each(AssociatedLoopConstructs, [&](OpenACCLoopConstruct *C) {
+    C->setParentComputeConstruct(Inst);
+  });
+
+  return Inst;
+}
+
+void OpenACCComputeConstruct::findAndSetChildLoops() {
+  struct LoopConstructFinder : RecursiveASTVisitor<LoopConstructFinder> {
+    OpenACCComputeConstruct *Construct = nullptr;
+
+    LoopConstructFinder(OpenACCComputeConstruct *Construct)
+        : Construct(Construct) {}
+
+    bool TraverseOpenACCComputeConstruct(OpenACCComputeConstruct *C) {
+      // Stop searching if we find a compute construct.
+      return true;
+    }
+    bool TraverseOpenACCLoopConstruct(OpenACCLoopConstruct *C) {
+      // Stop searching if we find a loop construct, after taking ownership of
+      // it.
+      C->setParentComputeConstruct(Construct);
+      return true;
+    }
+  };
+
+  LoopConstructFinder f(this);
+  f.TraverseStmt(getAssociatedStmt());
+}
+
+OpenACCLoopConstruct::OpenACCLoopConstruct(unsigned NumClauses)
+    : OpenACCAssociatedStmtConstruct(
+          OpenACCLoopConstructClass, OpenACCDirectiveKind::Loop,
+          SourceLocation{}, SourceLocation{}, SourceLocation{},
+          /*AssociatedStmt=*/nullptr) {
+  std::uninitialized_value_construct(
+      getTrailingObjects<const OpenACCClause *>(),
+      getTrailingObjects<const OpenACCClause *>() + NumClauses);
+  setClauseList(
+      MutableArrayRef(getTrailingObjects<const OpenACCClause *>(), NumClauses));
+}
+
+OpenACCLoopConstruct::OpenACCLoopConstruct(
+    SourceLocation Start, SourceLocation DirLoc, SourceLocation End,
+    ArrayRef<const OpenACCClause *> Clauses, Stmt *Loop)
+    : OpenACCAssociatedStmtConstruct(OpenACCLoopConstructClass,
+                                     OpenACCDirectiveKind::Loop, Start, DirLoc,
+                                     End, Loop) {
+  // accept 'nullptr' for the loop. This is diagnosed somewhere, but this gives
+  // us some level of AST fidelity in the error case.
+  assert((Loop == nullptr || isa<ForStmt, CXXForRangeStmt>(Loop)) &&
+         "Associated Loop not a for loop?");
+  // Initialize the trailing storage.
+  std::uninitialized_copy(Clauses.begin(), Clauses.end(),
+                          getTrailingObjects<const OpenACCClause *>());
+
+  setClauseList(MutableArrayRef(getTrailingObjects<const OpenACCClause *>(),
+                                Clauses.size()));
+}
+
+void OpenACCLoopConstruct::setLoop(Stmt *Loop) {
+  assert((isa<ForStmt, CXXForRangeStmt>(Loop)) &&
+         "Associated Loop not a for loop?");
+  setAssociatedStmt(Loop);
+}
+
+OpenACCLoopConstruct *OpenACCLoopConstruct::CreateEmpty(const ASTContext &C,
+                                                        unsigned NumClauses) {
+  void *Mem =
+      C.Allocate(OpenACCLoopConstruct::totalSizeToAlloc<const OpenACCClause *>(
+          NumClauses));
+  auto *Inst = new (Mem) OpenACCLoopConstruct(NumClauses);
+  return Inst;
+}
+
+OpenACCLoopConstruct *
+OpenACCLoopConstruct::Create(const ASTContext &C, SourceLocation BeginLoc,
+                             SourceLocation DirLoc, SourceLocation EndLoc,
+                             ArrayRef<const OpenACCClause *> Clauses,
+                             Stmt *Loop) {
+  void *Mem =
+      C.Allocate(OpenACCLoopConstruct::totalSizeToAlloc<const OpenACCClause *>(
+          Clauses.size()));
+  auto *Inst =
+      new (Mem) OpenACCLoopConstruct(BeginLoc, DirLoc, EndLoc, Clauses, Loop);
  return Inst;
 }
--- a/clang/lib/AST/StmtPrinter.cpp
+++ b/clang/lib/AST/StmtPrinter.cpp
@ -1156,6 +1156,19 @@ void StmtPrinter::VisitOpenACCComputeConstruct(OpenACCComputeConstruct *S) {
  PrintStmt(S->getStructuredBlock());
 }

+void StmtPrinter::VisitOpenACCLoopConstruct(OpenACCLoopConstruct *S) {
+  Indent() << "#pragma acc loop";
+
+  if (!S->clauses().empty()) {
+    OS << ' ';
+    OpenACCClausePrinter Printer(OS, Policy);
+    Printer.VisitClauseList(S->clauses());
+  }
+  OS << '\n';
+
+  PrintStmt(S->getLoop());
+}
+
 //===----------------------------------------------------------------------===//
 //  Expr printing methods.
 //===----------------------------------------------------------------------===//
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@ -2605,6 +2605,14 @@ void StmtProfiler::VisitOpenACCComputeConstruct(
  P.VisitOpenACCClauseList(S->clauses());
 }

+void StmtProfiler::VisitOpenACCLoopConstruct(const OpenACCLoopConstruct *S) {
+  // VisitStmt handles children, so the Loop is handled.
+  VisitStmt(S);
+
+  OpenACCClauseProfiler P{*this};
+  P.VisitOpenACCClauseList(S->clauses());
+}
+
 void Stmt::Profile(llvm::FoldingSetNodeID &ID, const ASTContext &Context,
                   bool Canonical, bool ProfileLambdaExpr) const {
  StmtProfilerWithPointers Profiler(ID, Context, Canonical, ProfileLambdaExpr);
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp
@ -2869,3 +2869,10 @@ void TextNodeDumper::VisitHLSLBufferDecl(const HLSLBufferDecl *D) {
 void TextNodeDumper::VisitOpenACCConstructStmt(const OpenACCConstructStmt *S) {
  OS << " " << S->getDirectiveKind();
 }
+void TextNodeDumper::VisitOpenACCLoopConstruct(const OpenACCLoopConstruct *S) {
+
+  if (S->isOrphanedLoopConstruct())
+    OS << " <orphan>";
+  else
+    OS << " parent: " << S->getParentComputeConstruct();
+}
--- a/clang/lib/CodeGen/CGStmt.cpp
+++ b/clang/lib/CodeGen/CGStmt.cpp
@ -442,6 +442,9 @@ void CodeGenFunction::EmitStmt(const Stmt *S, ArrayRef<const Attr *> Attrs) {
  case Stmt::OpenACCComputeConstructClass:
    EmitOpenACCComputeConstruct(cast<OpenACCComputeConstruct>(*S));
    break;
+  case Stmt::OpenACCLoopConstructClass:
+    EmitOpenACCLoopConstruct(cast<OpenACCLoopConstruct>(*S));
+    break;
  }
 }

--- a/clang/lib/CodeGen/CodeGenFunction.h
+++ b/clang/lib/CodeGen/CodeGenFunction.h
@ -4062,6 +4062,13 @@ public:
    EmitStmt(S.getStructuredBlock());
  }

+  void EmitOpenACCLoopConstruct(const OpenACCLoopConstruct &S) {
+    // TODO OpenACC: Implement this.  It is currently implemented as a 'no-op',
+    // simply emitting its loop, but in the future we will implement
+    // some sort of IR.
+    EmitStmt(S.getLoop());
+  }
+
  //===--------------------------------------------------------------------===//
  //                         LValue Expression Emission
  //===--------------------------------------------------------------------===//
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@ -571,6 +571,7 @@ bool doesDirectiveHaveAssociatedStmt(OpenACCDirectiveKind DirKind) {
  case OpenACCDirectiveKind::Parallel:
  case OpenACCDirectiveKind::Serial:
  case OpenACCDirectiveKind::Kernels:
+  case OpenACCDirectiveKind::Loop:
    return true;
  }
  llvm_unreachable("Unhandled directive->assoc stmt");
@ -1447,13 +1448,14 @@ StmtResult Parser::ParseOpenACCDirectiveStmt() {
    return StmtError();

  StmtResult AssocStmt;
-
+  SemaOpenACC::AssociatedStmtRAII AssocStmtRAII(getActions().OpenACC(),
+                                                DirInfo.DirKind);
  if (doesDirectiveHaveAssociatedStmt(DirInfo.DirKind)) {
    ParsingOpenACCDirectiveRAII DirScope(*this, /*Value=*/false);
    ParseScope ACCScope(this, getOpenACCScopeFlags(DirInfo.DirKind));

-    AssocStmt = getActions().OpenACC().ActOnAssociatedStmt(DirInfo.DirKind,
-                                                           ParseStatement());
+    AssocStmt = getActions().OpenACC().ActOnAssociatedStmt(
+        DirInfo.StartLoc, DirInfo.DirKind, ParseStatement());
  }

  return getActions().OpenACC().ActOnEndStmtDirective(
--- a/clang/lib/Sema/SemaExceptionSpec.cpp
+++ b/clang/lib/Sema/SemaExceptionSpec.cpp
@ -1425,6 +1425,7 @@ CanThrowResult Sema::canThrow(const Stmt *S) {

    // Most statements can throw if any substatement can throw.
  case Stmt::OpenACCComputeConstructClass:
+  case Stmt::OpenACCLoopConstructClass:
  case Stmt::AttributedStmtClass:
  case Stmt::BreakStmtClass:
  case Stmt::CapturedStmtClass:
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@ -33,6 +33,7 @@ bool diagnoseConstructAppertainment(SemaOpenACC &S, OpenACCDirectiveKind K,
  case OpenACCDirectiveKind::Parallel:
  case OpenACCDirectiveKind::Serial:
  case OpenACCDirectiveKind::Kernels:
+  case OpenACCDirectiveKind::Loop:
    if (!IsStmt)
      return S.Diag(StartLoc, diag::err_acc_construct_appertainment) << K;
    break;
@ -370,6 +371,30 @@ bool checkValidAfterDeviceType(

 SemaOpenACC::SemaOpenACC(Sema &S) : SemaBase(S) {}

+SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(SemaOpenACC &S,
+                                                    OpenACCDirectiveKind DK)
+    : SemaRef(S), WasInsideComputeConstruct(S.InsideComputeConstruct),
+      DirKind(DK) {
+  // Compute constructs end up taking their 'loop'.
+  if (DirKind == OpenACCDirectiveKind::Parallel ||
+      DirKind == OpenACCDirectiveKind::Serial ||
+      DirKind == OpenACCDirectiveKind::Kernels) {
+    SemaRef.InsideComputeConstruct = true;
+    SemaRef.ParentlessLoopConstructs.swap(ParentlessLoopConstructs);
+  }
+}
+
+SemaOpenACC::AssociatedStmtRAII::~AssociatedStmtRAII() {
+  SemaRef.InsideComputeConstruct = WasInsideComputeConstruct;
+  if (DirKind == OpenACCDirectiveKind::Parallel ||
+      DirKind == OpenACCDirectiveKind::Serial ||
+      DirKind == OpenACCDirectiveKind::Kernels) {
+    assert(SemaRef.ParentlessLoopConstructs.empty() &&
+           "Didn't consume loop construct list?");
+    SemaRef.ParentlessLoopConstructs.swap(ParentlessLoopConstructs);
+  }
+}
+
 OpenACCClause *
 SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses,
                         OpenACCParsedClause &Clause) {
@ -927,6 +952,7 @@ void SemaOpenACC::ActOnConstruct(OpenACCDirectiveKind K,
  case OpenACCDirectiveKind::Parallel:
  case OpenACCDirectiveKind::Serial:
  case OpenACCDirectiveKind::Kernels:
+  case OpenACCDirectiveKind::Loop:
    // Nothing to do here, there is no real legalization that needs to happen
    // here as these constructs do not take any arguments.
    break;
@ -1348,16 +1374,34 @@ StmtResult SemaOpenACC::ActOnEndStmtDirective(OpenACCDirectiveKind K,
    return StmtError();
  case OpenACCDirectiveKind::Parallel:
  case OpenACCDirectiveKind::Serial:
-  case OpenACCDirectiveKind::Kernels:
-    // TODO OpenACC: Add clauses to the construct here.
-    return OpenACCComputeConstruct::Create(
+  case OpenACCDirectiveKind::Kernels: {
+    auto *ComputeConstruct = OpenACCComputeConstruct::Create(
        getASTContext(), K, StartLoc, DirLoc, EndLoc, Clauses,
+        AssocStmt.isUsable() ? AssocStmt.get() : nullptr,
+        ParentlessLoopConstructs);
+
+    ParentlessLoopConstructs.clear();
+    return ComputeConstruct;
+  }
+  case OpenACCDirectiveKind::Loop: {
+    auto *LoopConstruct = OpenACCLoopConstruct::Create(
+        getASTContext(), StartLoc, DirLoc, EndLoc, Clauses,
        AssocStmt.isUsable() ? AssocStmt.get() : nullptr);
+
+    // If we are in the scope of a compute construct, add this to the list of
+    // loop constructs that need assigning to the next closing compute
+    // construct.
+    if (InsideComputeConstruct)
+      ParentlessLoopConstructs.push_back(LoopConstruct);
+
+    return LoopConstruct;
+  }
  }
  llvm_unreachable("Unhandled case in directive handling?");
 }

-StmtResult SemaOpenACC::ActOnAssociatedStmt(OpenACCDirectiveKind K,
+StmtResult SemaOpenACC::ActOnAssociatedStmt(SourceLocation DirectiveLoc,
+                                            OpenACCDirectiveKind K,
                                            StmtResult AssocStmt) {
  switch (K) {
  default:
@ -1375,6 +1419,14 @@ StmtResult SemaOpenACC::ActOnAssociatedStmt(OpenACCDirectiveKind K,
    // an interpretation of it is to allow this and treat the initializer as
    // the 'structured block'.
    return AssocStmt;
+  case OpenACCDirectiveKind::Loop:
+    if (AssocStmt.isUsable() &&
+        !isa<CXXForRangeStmt, ForStmt>(AssocStmt.get())) {
+      Diag(AssocStmt.get()->getBeginLoc(), diag::err_acc_loop_not_for_loop);
+      Diag(DirectiveLoc, diag::note_acc_construct_here) << K;
+      return StmtError();
+    }
+    return AssocStmt;
  }
  llvm_unreachable("Invalid associated statement application");
 }
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@ -4041,6 +4041,15 @@ public:
                                                     EndLoc, Clauses, StrBlock);
  }

+  StmtResult RebuildOpenACCLoopConstruct(SourceLocation BeginLoc,
+                                         SourceLocation DirLoc,
+                                         SourceLocation EndLoc,
+                                         ArrayRef<OpenACCClause *> Clauses,
+                                         StmtResult Loop) {
+    return getSema().OpenACC().ActOnEndStmtDirective(
+        OpenACCDirectiveKind::Loop, BeginLoc, DirLoc, EndLoc, Clauses, Loop);
+  }
+
 private:
  TypeLoc TransformTypeInObjectScope(TypeLoc TL,
                                     QualType ObjectType,
@ -11541,8 +11550,31 @@ template <typename Derived>
 StmtResult TreeTransform<Derived>::TransformOpenACCComputeConstruct(
    OpenACCComputeConstruct *C) {
  getSema().OpenACC().ActOnConstruct(C->getDirectiveKind(), C->getBeginLoc());
-  // FIXME: When implementing this for constructs that can take arguments, we
-  // should do Sema for them here.
+
+  if (getSema().OpenACC().ActOnStartStmtDirective(C->getDirectiveKind(),
+                                                  C->getBeginLoc()))
+    return StmtError();
+
+  llvm::SmallVector<OpenACCClause *> TransformedClauses =
+      getDerived().TransformOpenACCClauseList(C->getDirectiveKind(),
+                                              C->clauses());
+  // Transform Structured Block.
+  SemaOpenACC::AssociatedStmtRAII AssocStmtRAII(getSema().OpenACC(),
+                                                C->getDirectiveKind());
+  StmtResult StrBlock = getDerived().TransformStmt(C->getStructuredBlock());
+  StrBlock = getSema().OpenACC().ActOnAssociatedStmt(
+      C->getBeginLoc(), C->getDirectiveKind(), StrBlock);
+
+  return getDerived().RebuildOpenACCComputeConstruct(
+      C->getDirectiveKind(), C->getBeginLoc(), C->getDirectiveLoc(),
+      C->getEndLoc(), TransformedClauses, StrBlock);
+}
+
+template <typename Derived>
+StmtResult
+TreeTransform<Derived>::TransformOpenACCLoopConstruct(OpenACCLoopConstruct *C) {
+
+  getSema().OpenACC().ActOnConstruct(C->getDirectiveKind(), C->getBeginLoc());

  if (getSema().OpenACC().ActOnStartStmtDirective(C->getDirectiveKind(),
                                                  C->getBeginLoc()))
@ -11552,14 +11584,16 @@ StmtResult TreeTransform<Derived>::TransformOpenACCComputeConstruct(
      getDerived().TransformOpenACCClauseList(C->getDirectiveKind(),
                                              C->clauses());

-  // Transform Structured Block.
-  StmtResult StrBlock = getDerived().TransformStmt(C->getStructuredBlock());
-  StrBlock =
-      getSema().OpenACC().ActOnAssociatedStmt(C->getDirectiveKind(), StrBlock);
+  // Transform Loop.
+  SemaOpenACC::AssociatedStmtRAII AssocStmtRAII(getSema().OpenACC(),
+                                                C->getDirectiveKind());
+  StmtResult Loop = getDerived().TransformStmt(C->getLoop());
+  Loop = getSema().OpenACC().ActOnAssociatedStmt(C->getBeginLoc(),
+                                                 C->getDirectiveKind(), Loop);

-  return getDerived().RebuildOpenACCComputeConstruct(
-      C->getDirectiveKind(), C->getBeginLoc(), C->getDirectiveLoc(),
-      C->getEndLoc(), TransformedClauses, StrBlock);
+  return getDerived().RebuildOpenACCLoopConstruct(
+      C->getBeginLoc(), C->getDirectiveLoc(), C->getEndLoc(),
+      TransformedClauses, Loop);
 }

 //===----------------------------------------------------------------------===//
--- a/clang/lib/Serialization/ASTReaderStmt.cpp
+++ b/clang/lib/Serialization/ASTReaderStmt.cpp
@ -2810,6 +2810,12 @@ void ASTStmtReader::VisitOpenACCAssociatedStmtConstruct(
 void ASTStmtReader::VisitOpenACCComputeConstruct(OpenACCComputeConstruct *S) {
  VisitStmt(S);
  VisitOpenACCAssociatedStmtConstruct(S);
+  S->findAndSetChildLoops();
+}
+
+void ASTStmtReader::VisitOpenACCLoopConstruct(OpenACCLoopConstruct *S) {
+  VisitStmt(S);
+  VisitOpenACCAssociatedStmtConstruct(S);
 }

 //===----------------------------------------------------------------------===//
@ -4235,6 +4241,11 @@ Stmt *ASTReader::ReadStmtFromStream(ModuleFile &F) {
      S = OpenACCComputeConstruct::CreateEmpty(Context, NumClauses);
      break;
    }
+    case STMT_OPENACC_LOOP_CONSTRUCT: {
+      unsigned NumClauses = Record[ASTStmtReader::NumStmtFields];
+      S = OpenACCLoopConstruct::CreateEmpty(Context, NumClauses);
+      break;
+    }
    case EXPR_REQUIRES:
      unsigned numLocalParameters = Record[ASTStmtReader::NumExprFields];
      unsigned numRequirement = Record[ASTStmtReader::NumExprFields + 1];
--- a/clang/lib/Serialization/ASTWriterStmt.cpp
+++ b/clang/lib/Serialization/ASTWriterStmt.cpp
@ -2863,6 +2863,12 @@ void ASTStmtWriter::VisitOpenACCComputeConstruct(OpenACCComputeConstruct *S) {
  Code = serialization::STMT_OPENACC_COMPUTE_CONSTRUCT;
 }

+void ASTStmtWriter::VisitOpenACCLoopConstruct(OpenACCLoopConstruct *S) {
+  VisitStmt(S);
+  VisitOpenACCAssociatedStmtConstruct(S);
+  Code = serialization::STMT_OPENACC_LOOP_CONSTRUCT;
+}
+
 //===----------------------------------------------------------------------===//
 // ASTWriter Implementation
 //===----------------------------------------------------------------------===//
--- a/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
+++ b/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp
@ -1822,6 +1822,7 @@ void ExprEngine::Visit(const Stmt *S, ExplodedNode *Pred,
    case Stmt::OMPTargetParallelGenericLoopDirectiveClass:
    case Stmt::CapturedStmtClass:
    case Stmt::OpenACCComputeConstructClass:
+    case Stmt::OpenACCLoopConstructClass:
    case Stmt::OMPUnrollDirectiveClass:
    case Stmt::OMPMetaDirectiveClass: {
      const ExplodedNode *node = Bldr.generateSink(S, Pred, Pred->getState());
--- a/clang/test/AST/ast-print-openacc-loop-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-loop-construct.cpp
@ -0,0 +1,9 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-deprecated-clause-alias -ast-print %s -o - | FileCheck %s
+
+void foo() {
+// CHECK: #pragma acc loop
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc loop
+  for(;;);
+}
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@ -37,23 +37,23 @@ void func() {
  // expected-warning@+1{{OpenACC construct 'host_data' not yet implemented, pragma ignored}}
 #pragma acc host_data if_present, if_present

-  // expected-warning@+4{{OpenACC clause 'seq' not yet implemented, clause ignored}}
-  // expected-warning@+3{{OpenACC clause 'independent' not yet implemented, clause ignored}}
-  // expected-warning@+2{{OpenACC clause 'auto' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-warning@+3{{OpenACC clause 'seq' not yet implemented, clause ignored}}
+  // expected-warning@+2{{OpenACC clause 'independent' not yet implemented, clause ignored}}
+  // expected-warning@+1{{OpenACC clause 'auto' not yet implemented, clause ignored}}
 #pragma acc loop seq independent auto
+  for(;;){}

-  // expected-warning@+4{{OpenACC clause 'seq' not yet implemented, clause ignored}}
-  // expected-warning@+3{{OpenACC clause 'independent' not yet implemented, clause ignored}}
-  // expected-warning@+2{{OpenACC clause 'auto' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-warning@+3{{OpenACC clause 'seq' not yet implemented, clause ignored}}
+  // expected-warning@+2{{OpenACC clause 'independent' not yet implemented, clause ignored}}
+  // expected-warning@+1{{OpenACC clause 'auto' not yet implemented, clause ignored}}
 #pragma acc loop seq, independent auto
+  for(;;){}

-  // expected-warning@+4{{OpenACC clause 'seq' not yet implemented, clause ignored}}
-  // expected-warning@+3{{OpenACC clause 'independent' not yet implemented, clause ignored}}
-  // expected-warning@+2{{OpenACC clause 'auto' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-warning@+3{{OpenACC clause 'seq' not yet implemented, clause ignored}}
+  // expected-warning@+2{{OpenACC clause 'independent' not yet implemented, clause ignored}}
+  // expected-warning@+1{{OpenACC clause 'auto' not yet implemented, clause ignored}}
 #pragma acc loop seq independent, auto
+  for(;;){}

  // expected-warning@+4{{OpenACC clause 'seq' not yet implemented, clause ignored}}
  // expected-warning@+3{{OpenACC clause 'independent' not yet implemented, clause ignored}}
@ -67,65 +67,57 @@ void func() {
  // expected-warning@+2{{OpenACC clause 'auto' not yet implemented, clause ignored}}
  // expected-warning@+1{{OpenACC construct 'serial loop' not yet implemented, pragma ignored}}
 #pragma acc serial loop seq, independent auto
-  {}
+  for(;;){}

  // expected-warning@+4{{OpenACC clause 'seq' not yet implemented, clause ignored}}
  // expected-warning@+3{{OpenACC clause 'independent' not yet implemented, clause ignored}}
  // expected-warning@+2{{OpenACC clause 'auto' not yet implemented, clause ignored}}
  // expected-warning@+1{{OpenACC construct 'parallel loop' not yet implemented, pragma ignored}}
 #pragma acc parallel loop seq independent, auto
-  {}
+  for(;;){}


+  // expected-error@+1{{expected identifier}}
+#pragma acc loop , seq
+  for(;;){}
+
  // expected-error@+2{{expected identifier}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
-#pragma acc loop , seq
-
-  // expected-error@+3{{expected identifier}}
-  // expected-warning@+2{{OpenACC clause 'seq' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-warning@+1{{OpenACC clause 'seq' not yet implemented, clause ignored}}
 #pragma acc loop seq,
+  for(;;){}

-  // expected-error@+2{{expected '('}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+1{{expected '('}}
 #pragma acc loop collapse
  for(;;){}

-  // expected-error@+2{{expected expression}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+1{{expected expression}}
 #pragma acc loop collapse()
  for(;;){}

-  // expected-error@+3{{invalid tag 'unknown' on 'collapse' clause}}
-  // expected-error@+2{{expected expression}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+2{{invalid tag 'unknown' on 'collapse' clause}}
+  // expected-error@+1{{expected expression}}
 #pragma acc loop collapse(unknown:)
  for(;;){}

-  // expected-error@+2{{expected expression}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+1{{expected expression}}
 #pragma acc loop collapse(force:)
  for(;;){}

-  // expected-error@+3{{invalid tag 'unknown' on 'collapse' clause}}
-  // expected-warning@+2{{OpenACC clause 'collapse' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+2{{invalid tag 'unknown' on 'collapse' clause}}
+  // expected-warning@+1{{OpenACC clause 'collapse' not yet implemented, clause ignored}}
 #pragma acc loop collapse(unknown:5)
  for(;;){}

-  // expected-warning@+2{{OpenACC clause 'collapse' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-warning@+1{{OpenACC clause 'collapse' not yet implemented, clause ignored}}
 #pragma acc loop collapse(force:5)
  for(;;){}

-  // expected-warning@+2{{OpenACC clause 'collapse' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-warning@+1{{OpenACC clause 'collapse' not yet implemented, clause ignored}}
 #pragma acc loop collapse(5)
  for(;;){}

-  // expected-error@+3{{expected ')'}}
-  // expected-note@+2{{to match this '('}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
 #pragma acc loop collapse(5, 6)
  for(;;){}
 }
@ -989,108 +981,108 @@ void IntExprParsing() {
 #pragma acc set default_async(returns_int())


-  // expected-warning@+2{{OpenACC clause 'vector' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-warning@+1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
 #pragma acc loop vector
-  // expected-error@+2{{expected expression}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  for(;;);
+  // expected-error@+1{{expected expression}}
 #pragma acc loop vector()
-  // expected-error@+3{{invalid tag 'invalid' on 'vector' clause}}
-  // expected-error@+2{{expected expression}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  for(;;);
+  // expected-error@+2{{invalid tag 'invalid' on 'vector' clause}}
+  // expected-error@+1{{expected expression}}
 #pragma acc loop vector(invalid:)
-  // expected-error@+3{{invalid tag 'invalid' on 'vector' clause}}
-  // expected-warning@+2{{OpenACC clause 'vector' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  for(;;);
+  // expected-error@+2{{invalid tag 'invalid' on 'vector' clause}}
+  // expected-warning@+1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
 #pragma acc loop vector(invalid:5)
-  // expected-error@+2{{expected expression}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  for(;;);
+  // expected-error@+1{{expected expression}}
 #pragma acc loop vector(length:)
-  // expected-error@+3{{invalid tag 'num' on 'vector' clause}}
-  // expected-error@+2{{expected expression}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  for(;;);
+  // expected-error@+2{{invalid tag 'num' on 'vector' clause}}
+  // expected-error@+1{{expected expression}}
 #pragma acc loop vector(num:)
-  // expected-error@+3{{expected ')'}}
-  // expected-note@+2{{to match this '('}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  for(;;);
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
 #pragma acc loop vector(5, 4)
-  // expected-error@+3{{expected ')'}}
-  // expected-note@+2{{to match this '('}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  for(;;);
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
 #pragma acc loop vector(length:6,4)
-  // expected-error@+4{{invalid tag 'num' on 'vector' clause}}
-  // expected-error@+3{{expected ')'}}
-  // expected-note@+2{{to match this '('}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
-#pragma acc loop vector(num:6,4)
-  // expected-warning@+2{{OpenACC clause 'vector' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
-#pragma acc loop vector(5)
+  for(;;);
  // expected-error@+3{{invalid tag 'num' on 'vector' clause}}
-  // expected-warning@+2{{OpenACC clause 'vector' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
+#pragma acc loop vector(num:6,4)
+  for(;;);
+  // expected-warning@+1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
+#pragma acc loop vector(5)
+  for(;;);
+  // expected-error@+2{{invalid tag 'num' on 'vector' clause}}
+  // expected-warning@+1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
 #pragma acc loop vector(num:5)
-  // expected-warning@+2{{OpenACC clause 'vector' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  for(;;);
+  // expected-warning@+1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
 #pragma acc loop vector(length:5)
-  // expected-warning@+2{{OpenACC clause 'vector' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  for(;;);
+  // expected-warning@+1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
 #pragma acc loop vector(returns_int())
-  // expected-warning@+2{{OpenACC clause 'vector' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  for(;;);
+  // expected-warning@+1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
 #pragma acc loop vector(length:returns_int())
+  for(;;);

-  // expected-warning@+2{{OpenACC clause 'worker' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-warning@+1{{OpenACC clause 'worker' not yet implemented, clause ignored}}
 #pragma acc loop worker
-  // expected-error@+2{{expected expression}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  for(;;);
+  // expected-error@+1{{expected expression}}
 #pragma acc loop worker()
-  // expected-error@+3{{invalid tag 'invalid' on 'worker' clause}}
-  // expected-error@+2{{expected expression}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  for(;;);
+  // expected-error@+2{{invalid tag 'invalid' on 'worker' clause}}
+  // expected-error@+1{{expected expression}}
 #pragma acc loop worker(invalid:)
-  // expected-error@+3{{invalid tag 'invalid' on 'worker' clause}}
-  // expected-warning@+2{{OpenACC clause 'worker' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  for(;;);
+  // expected-error@+2{{invalid tag 'invalid' on 'worker' clause}}
+  // expected-warning@+1{{OpenACC clause 'worker' not yet implemented, clause ignored}}
 #pragma acc loop worker(invalid:5)
-  // expected-error@+2{{expected expression}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  for(;;);
+  // expected-error@+1{{expected expression}}
 #pragma acc loop worker(num:)
-  // expected-error@+3{{invalid tag 'length' on 'worker' clause}}
-  // expected-error@+2{{expected expression}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  for(;;);
+  // expected-error@+2{{invalid tag 'length' on 'worker' clause}}
+  // expected-error@+1{{expected expression}}
 #pragma acc loop worker(length:)
-  // expected-error@+3{{expected ')'}}
-  // expected-note@+2{{to match this '('}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  for(;;);
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
 #pragma acc loop worker(5, 4)
-  // expected-error@+3{{expected ')'}}
-  // expected-note@+2{{to match this '('}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  for(;;);
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
 #pragma acc loop worker(num:6,4)
-  // expected-error@+4{{invalid tag 'length' on 'worker' clause}}
-  // expected-error@+3{{expected ')'}}
-  // expected-note@+2{{to match this '('}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  for(;;);
+  // expected-error@+3{{invalid tag 'length' on 'worker' clause}}
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
 #pragma acc loop worker(length:6,4)
-  // expected-warning@+2{{OpenACC clause 'worker' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  for(;;);
+  // expected-warning@+1{{OpenACC clause 'worker' not yet implemented, clause ignored}}
 #pragma acc loop worker(5)
-  // expected-error@+3{{invalid tag 'length' on 'worker' clause}}
-  // expected-warning@+2{{OpenACC clause 'worker' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  for(;;);
+  // expected-error@+2{{invalid tag 'length' on 'worker' clause}}
+  // expected-warning@+1{{OpenACC clause 'worker' not yet implemented, clause ignored}}
 #pragma acc loop worker(length:5)
-  // expected-warning@+2{{OpenACC clause 'worker' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  for(;;);
+  // expected-warning@+1{{OpenACC clause 'worker' not yet implemented, clause ignored}}
 #pragma acc loop worker(num:5)
-  // expected-warning@+2{{OpenACC clause 'worker' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  for(;;);
+  // expected-warning@+1{{OpenACC clause 'worker' not yet implemented, clause ignored}}
 #pragma acc loop worker(returns_int())
-  // expected-error@+3{{invalid tag 'length' on 'worker' clause}}
-  // expected-warning@+2{{OpenACC clause 'worker' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  for(;;);
+  // expected-error@+2{{invalid tag 'length' on 'worker' clause}}
+  // expected-warning@+1{{OpenACC clause 'worker' not yet implemented, clause ignored}}
 #pragma acc loop worker(length:returns_int())
+  for(;;);
 }

 void device_type() {
@ -1236,238 +1228,196 @@ void AsyncArgument() {
 void Tile() {

  int* Foo;
-  // expected-error@+2{{expected '('}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+1{{expected '('}}
 #pragma acc loop tile
  for(;;){}
-  // expected-error@+4{{expected expression}}
-  // expected-error@+3{{expected ')'}}
-  // expected-note@+2{{to match this '('}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+3{{expected expression}}
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
 #pragma acc loop tile(
  for(;;){}
-  // expected-error@+3{{expected expression}}
-  // expected-warning@+2{{OpenACC clause 'tile' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+2{{expected expression}}
+  // expected-warning@+1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
 #pragma acc loop tile()
  for(;;){}
-  // expected-error@+4{{expected expression}}
-  // expected-error@+3{{expected ')'}}
-  // expected-note@+2{{to match this '('}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+3{{expected expression}}
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
 #pragma acc loop tile(,
  for(;;){}
-  // expected-error@+3{{expected expression}}
-  // expected-warning@+2{{OpenACC clause 'tile' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+2{{expected expression}}
+  // expected-warning@+1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
 #pragma acc loop tile(,)
  for(;;){}
-  // expected-error@+3{{use of undeclared identifier 'invalid'}}
-  // expected-warning@+2{{OpenACC clause 'tile' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+2{{use of undeclared identifier 'invalid'}}
+  // expected-warning@+1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
 #pragma acc loop tile(returns_int(), *, invalid, *)
  for(;;){}

-  // expected-error@+3{{expected expression}}
-  // expected-warning@+2{{OpenACC clause 'tile' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+2{{expected expression}}
+  // expected-warning@+1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
 #pragma acc loop tile(returns_int() *, Foo, *)
  for(;;){}

-  // expected-error@+3{{indirection requires pointer operand ('int' invalid)}}
-  // expected-warning@+2{{OpenACC clause 'tile' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+2{{indirection requires pointer operand ('int' invalid)}}
+  // expected-warning@+1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
 #pragma acc loop tile(* returns_int() , *)
  for(;;){}

-  // expected-warning@+2{{OpenACC clause 'tile' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-warning@+1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
 #pragma acc loop tile(*)
  for(;;){}
-  // expected-warning@+2{{OpenACC clause 'tile' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-warning@+1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
 #pragma acc loop tile(*Foo, *Foo)
  for(;;){}
-  // expected-warning@+2{{OpenACC clause 'tile' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-warning@+1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
 #pragma acc loop tile(5)
  for(;;){}
-  // expected-warning@+2{{OpenACC clause 'tile' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-warning@+1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
 #pragma acc loop tile(*, 5)
  for(;;){}
-  // expected-warning@+2{{OpenACC clause 'tile' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-warning@+1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
 #pragma acc loop tile(5, *)
  for(;;){}
-  // expected-warning@+2{{OpenACC clause 'tile' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-warning@+1{{OpenACC clause 'tile' not yet implemented, clause ignored}}
 #pragma acc loop tile(5, *, 3, *)
  for(;;){}
 }

 void Gang() {
-  // expected-warning@+2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
 #pragma acc loop gang
  for(;;){}
-  // expected-error@+4{{expected expression}}
-  // expected-error@+3{{expected ')'}}
-  // expected-note@+2{{to match this '('}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+3{{expected expression}}
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
 #pragma acc loop gang(
  for(;;){}
-  // expected-error@+3{{expected expression}}
-  // expected-warning@+2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+2{{expected expression}}
+  // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
 #pragma acc loop gang()
  for(;;){}

-  // expected-error@+3{{expected expression}}
-  // expected-warning@+2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+2{{expected expression}}
+  // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
 #pragma acc loop gang(5, *)
  for(;;){}

-  // expected-error@+3{{expected expression}}
-  // expected-warning@+2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+2{{expected expression}}
+  // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
 #pragma acc loop gang(*)
  for(;;){}

-  // expected-error@+3{{expected expression}}
-  // expected-warning@+2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+2{{expected expression}}
+  // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
 #pragma acc loop gang(5, num:*)
  for(;;){}

-  // expected-error@+3{{expected expression}}
-  // expected-warning@+2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+2{{expected expression}}
+  // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
 #pragma acc loop gang(num:5, *)
  for(;;){}

-  // expected-error@+3{{expected expression}}
-  // expected-warning@+2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+2{{expected expression}}
+  // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
 #pragma acc loop gang(num:5, num:*)
  for(;;){}

-  // expected-error@+3{{expected expression}}
-  // expected-warning@+2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+2{{expected expression}}
+  // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
 #pragma acc loop gang(num:*)
  for(;;){}

-  // expected-warning@+2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
 #pragma acc loop gang(dim:5)
  for(;;){}

-  // expected-error@+3{{expected expression}}
-  // expected-warning@+2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+2{{expected expression}}
+  // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
 #pragma acc loop gang(dim:5, dim:*)
  for(;;){}

-  // expected-error@+3{{expected expression}}
-  // expected-warning@+2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+2{{expected expression}}
+  // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
 #pragma acc loop gang(dim:*)
  for(;;){}

-  // expected-warning@+2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
 #pragma acc loop gang(static:*)
  for(;;){}

-  // expected-warning@+2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
 #pragma acc loop gang(static:*, static:5)
  for(;;){}

-  // expected-warning@+2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
 #pragma acc loop gang(static:*, 5)
  for(;;){}

-  // expected-warning@+2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
 #pragma acc loop gang(static:45, 5)
  for(;;){}

-  // expected-error@+4{{expected expression}}
-  // expected-error@+3{{expected ')'}}
-  // expected-note@+2{{to match this '('}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+3{{expected expression}}
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
 #pragma acc loop gang(static:45,
  for(;;){}

-  // expected-error@+3{{expected ')'}}
-  // expected-note@+2{{to match this '('}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
 #pragma acc loop gang(static:45
  for(;;){}

-  // expected-error@+4{{expected expression}}
-  // expected-error@+3{{expected ')'}}
-  // expected-note@+2{{to match this '('}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+3{{expected expression}}
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
 #pragma acc loop gang(static:*,
  for(;;){}

-  // expected-error@+3{{expected ')'}}
-  // expected-note@+2{{to match this '('}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
 #pragma acc loop gang(static:*
  for(;;){}

-  // expected-error@+4{{expected expression}}
-  // expected-error@+3{{expected ')'}}
-  // expected-note@+2{{to match this '('}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+3{{expected expression}}
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
 #pragma acc loop gang(45,
  for(;;){}

-  // expected-error@+3{{expected ')'}}
-  // expected-note@+2{{to match this '('}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
 #pragma acc loop gang(45
  for(;;){}

-  // expected-error@+4{{expected expression}}
-  // expected-error@+3{{expected ')'}}
-  // expected-note@+2{{to match this '('}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+3{{expected expression}}
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
 #pragma acc loop gang(num:45,
  for(;;){}

-  // expected-error@+3{{expected ')'}}
-  // expected-note@+2{{to match this '('}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
 #pragma acc loop gang(num:45
  for(;;){}

-  // expected-error@+4{{expected expression}}
-  // expected-error@+3{{expected ')'}}
-  // expected-note@+2{{to match this '('}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+3{{expected expression}}
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
 #pragma acc loop gang(dim:45,
  for(;;){}

-  // expected-error@+3{{expected ')'}}
-  // expected-note@+2{{to match this '('}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
 #pragma acc loop gang(dim:45
  for(;;){}

-  // expected-warning@+2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
 #pragma acc loop gang(static:*, dim:returns_int(), 5)
  for(;;){}

-  // expected-warning@+2{{OpenACC clause 'gang' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-warning@+1{{OpenACC clause 'gang' not yet implemented, clause ignored}}
 #pragma acc loop gang(num: 32, static:*, dim:returns_int(), 5)
  for(;;){}

--- a/clang/test/ParserOpenACC/parse-clauses.cpp
+++ b/clang/test/ParserOpenACC/parse-clauses.cpp
@ -2,13 +2,11 @@

 template<unsigned I, typename T>
 void templ() {
-  // expected-warning@+2{{OpenACC clause 'collapse' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-warning@+1{{OpenACC clause 'collapse' not yet implemented, clause ignored}}
 #pragma acc loop collapse(I)
  for(;;){}

-  // expected-warning@+2{{OpenACC clause 'collapse' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-warning@+1{{OpenACC clause 'collapse' not yet implemented, clause ignored}}
 #pragma acc loop collapse(T::value)
  for(;;){}

--- a/clang/test/ParserOpenACC/parse-constructs.c
+++ b/clang/test/ParserOpenACC/parse-constructs.c
@ -82,8 +82,7 @@ void func() {
  // expected-warning@+1{{OpenACC construct 'host_data' not yet implemented, pragma ignored}}
 #pragma acc host_data clause list
  for(;;){}
-  // expected-error@+2{{invalid OpenACC clause 'clause'}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented, pragma ignored}}
+  // expected-error@+1{{invalid OpenACC clause 'clause'}}
 #pragma acc loop clause list
  for(;;){}
  // expected-error@+1{{invalid OpenACC clause 'invalid'}}
--- a/clang/test/SemaOpenACC/compute-construct-async-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-async-clause.c
@ -39,8 +39,7 @@ void Test() {
 #pragma acc kernels async(SomeE)
  while(1);

-  // expected-error@+2{{OpenACC 'async' clause is not valid on 'loop' directive}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented}}
+  // expected-error@+1{{OpenACC 'async' clause is not valid on 'loop' directive}}
 #pragma acc loop async(1)
  for(;;);
 }
--- a/clang/test/SemaOpenACC/compute-construct-attach-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-attach-clause.c
@ -59,8 +59,7 @@ void uses() {
 #pragma acc parallel attach(s.PtrMem)
  while (1);

-  // expected-error@+2{{OpenACC 'attach' clause is not valid on 'loop' directive}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented}}
+  // expected-error@+1{{OpenACC 'attach' clause is not valid on 'loop' directive}}
 #pragma acc loop attach(LocalInt)
  for(;;);
 }
--- a/clang/test/SemaOpenACC/compute-construct-copy-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-copy-clause.c
@ -60,16 +60,13 @@ void uses(int IntParam, short *PointerParam, float ArrayParam[5], Complete Compo
 #pragma acc parallel copy((float)ArrayParam[2])
  while(1);

-  // expected-error@+2{{OpenACC 'copy' clause is not valid on 'loop' directive}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented}}
+  // expected-error@+1{{OpenACC 'copy' clause is not valid on 'loop' directive}}
 #pragma acc loop copy(LocalInt)
  for(;;);
-  // expected-error@+2{{OpenACC 'pcopy' clause is not valid on 'loop' directive}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented}}
+  // expected-error@+1{{OpenACC 'pcopy' clause is not valid on 'loop' directive}}
 #pragma acc loop pcopy(LocalInt)
  for(;;);
-  // expected-error@+2{{OpenACC 'present_or_copy' clause is not valid on 'loop' directive}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented}}
+  // expected-error@+1{{OpenACC 'present_or_copy' clause is not valid on 'loop' directive}}
 #pragma acc loop present_or_copy(LocalInt)
  for(;;);
 }
--- a/clang/test/SemaOpenACC/compute-construct-copyin-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-copyin-clause.c
@ -66,16 +66,13 @@ void uses(int IntParam, short *PointerParam, float ArrayParam[5], Complete Compo
 #pragma acc parallel copyin(invalid:(float)ArrayParam[2])
  while(1);

-  // expected-error@+2{{OpenACC 'copyin' clause is not valid on 'loop' directive}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented}}
+  // expected-error@+1{{OpenACC 'copyin' clause is not valid on 'loop' directive}}
 #pragma acc loop copyin(LocalInt)
  for(;;);
-  // expected-error@+2{{OpenACC 'pcopyin' clause is not valid on 'loop' directive}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented}}
+  // expected-error@+1{{OpenACC 'pcopyin' clause is not valid on 'loop' directive}}
 #pragma acc loop pcopyin(LocalInt)
  for(;;);
-  // expected-error@+2{{OpenACC 'present_or_copyin' clause is not valid on 'loop' directive}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented}}
+  // expected-error@+1{{OpenACC 'present_or_copyin' clause is not valid on 'loop' directive}}
 #pragma acc loop present_or_copyin(LocalInt)
  for(;;);
 }
--- a/clang/test/SemaOpenACC/compute-construct-copyout-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-copyout-clause.c
@ -66,16 +66,13 @@ void uses(int IntParam, short *PointerParam, float ArrayParam[5], Complete Compo
 #pragma acc parallel copyout(invalid:(float)ArrayParam[2])
  while(1);

-  // expected-error@+2{{OpenACC 'copyout' clause is not valid on 'loop' directive}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented}}
+  // expected-error@+1{{OpenACC 'copyout' clause is not valid on 'loop' directive}}
 #pragma acc loop copyout(LocalInt)
  for(;;);
-  // expected-error@+2{{OpenACC 'pcopyout' clause is not valid on 'loop' directive}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented}}
+  // expected-error@+1{{OpenACC 'pcopyout' clause is not valid on 'loop' directive}}
 #pragma acc loop pcopyout(LocalInt)
  for(;;);
-  // expected-error@+2{{OpenACC 'present_or_copyout' clause is not valid on 'loop' directive}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented}}
+  // expected-error@+1{{OpenACC 'present_or_copyout' clause is not valid on 'loop' directive}}
 #pragma acc loop present_or_copyout(LocalInt)
  for(;;);
 }
--- a/clang/test/SemaOpenACC/compute-construct-create-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-create-clause.c
@ -67,16 +67,13 @@ void uses(int IntParam, short *PointerParam, float ArrayParam[5], Complete Compo
 #pragma acc parallel create(invalid:(float)ArrayParam[2])
  while(1);

-  // expected-error@+2{{OpenACC 'create' clause is not valid on 'loop' directive}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented}}
+  // expected-error@+1{{OpenACC 'create' clause is not valid on 'loop' directive}}
 #pragma acc loop create(LocalInt)
  for(;;);
-  // expected-error@+2{{OpenACC 'pcreate' clause is not valid on 'loop' directive}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented}}
+  // expected-error@+1{{OpenACC 'pcreate' clause is not valid on 'loop' directive}}
 #pragma acc loop pcreate(LocalInt)
  for(;;);
-  // expected-error@+2{{OpenACC 'present_or_create' clause is not valid on 'loop' directive}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented}}
+  // expected-error@+1{{OpenACC 'present_or_create' clause is not valid on 'loop' directive}}
 #pragma acc loop present_or_create(LocalInt)
  for(;;);
 }
--- a/clang/test/SemaOpenACC/compute-construct-default-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-default-clause.c
@ -43,18 +43,16 @@ void SingleOnly() {
  #pragma acc data default(none)
  while(0);

-  // expected-warning@+2{{OpenACC construct 'loop' not yet implemented}}
  // expected-error@+1{{OpenACC 'default' clause is not valid on 'loop' directive}}
  #pragma acc loop default(none)
-  while(0);
+  for(;;);

  // expected-warning@+2{{OpenACC construct 'wait' not yet implemented}}
  // expected-error@+1{{OpenACC 'default' clause is not valid on 'wait' directive}}
  #pragma acc wait default(none)
  while(0);

-  // expected-error@+2{{OpenACC 'default' clause is not valid on 'loop' directive}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented}}
+  // expected-error@+1{{OpenACC 'default' clause is not valid on 'loop' directive}}
 #pragma acc loop default(present)
  for(;;);
 }
--- a/clang/test/SemaOpenACC/compute-construct-deviceptr-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-deviceptr-clause.c
@ -59,8 +59,7 @@ void uses() {
 #pragma acc parallel deviceptr(s.PtrMem)
  while (1);

-  // expected-error@+2{{OpenACC 'deviceptr' clause is not valid on 'loop' directive}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented}}
+  // expected-error@+1{{OpenACC 'deviceptr' clause is not valid on 'loop' directive}}
 #pragma acc loop deviceptr(LocalInt)
  for(;;);
 }
--- a/clang/test/SemaOpenACC/compute-construct-firstprivate-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-firstprivate-clause.c
@ -53,8 +53,7 @@ void uses(int IntParam, short *PointerParam, float ArrayParam[5], Complete Compo
 #pragma acc parallel firstprivate((float)ArrayParam[2])
  while(1);

-  // expected-error@+2{{OpenACC 'firstprivate' clause is not valid on 'loop' directive}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented}}
+  // expected-error@+1{{OpenACC 'firstprivate' clause is not valid on 'loop' directive}}
 #pragma acc loop firstprivate(LocalInt)
  for(;;);
 }
--- a/clang/test/SemaOpenACC/compute-construct-if-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-if-clause.c
@ -60,8 +60,7 @@ void BoolExpr(int *I, float *F) {
 #pragma acc kernels loop if (*I < *F)
  while(0);

-  // expected-error@+2{{OpenACC 'if' clause is not valid on 'loop' directive}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented}}
+  // expected-error@+1{{OpenACC 'if' clause is not valid on 'loop' directive}}
 #pragma acc loop if(I)
  for(;;);
 }
--- a/clang/test/SemaOpenACC/compute-construct-no_create-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-no_create-clause.c
@ -52,8 +52,7 @@ void uses(int IntParam, short *PointerParam, float ArrayParam[5], Complete Compo
 #pragma acc parallel no_create((float)ArrayParam[2])
  while(1);

-  // expected-error@+2{{OpenACC 'no_create' clause is not valid on 'loop' directive}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented}}
+  // expected-error@+1{{OpenACC 'no_create' clause is not valid on 'loop' directive}}
 #pragma acc loop no_create(LocalInt)
  for(;;);
 }
--- a/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-num_gangs-clause.c
@ -52,8 +52,7 @@ void Test() {
 #pragma acc parallel num_gangs(getS(), 1, getS(), 1)
  while(1);

-  // expected-error@+2{{OpenACC 'num_gangs' clause is not valid on 'loop' directive}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented}}
+  // expected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'loop' directive}}
 #pragma acc loop num_gangs(1)
  for(;;);
 }
--- a/clang/test/SemaOpenACC/compute-construct-num_workers-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-num_workers-clause.c
@ -31,8 +31,7 @@ void Test() {
 #pragma acc kernels num_workers(SomeE)
  while(1);

-  // expected-error@+2{{OpenACC 'num_workers' clause is not valid on 'loop' directive}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented}}
+  // expected-error@+1{{OpenACC 'num_workers' clause is not valid on 'loop' directive}}
 #pragma acc loop num_workers(1)
  for(;;);
 }
--- a/clang/test/SemaOpenACC/compute-construct-present-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-present-clause.c
@ -52,8 +52,7 @@ void uses(int IntParam, short *PointerParam, float ArrayParam[5], Complete Compo
 #pragma acc parallel present((float)ArrayParam[2])
  while(1);

-  // expected-error@+2{{OpenACC 'present' clause is not valid on 'loop' directive}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented}}
+  // expected-error@+1{{OpenACC 'present' clause is not valid on 'loop' directive}}
 #pragma acc loop present(LocalInt)
  for(;;);
 }
--- a/clang/test/SemaOpenACC/compute-construct-self-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-self-clause.c
@ -80,8 +80,7 @@ void WarnMaybeNotUsed(int val1, int val2) {
 #pragma acc parallel if(invalid) self(val1)
  while(0);

-  // expected-error@+2{{OpenACC 'self' clause is not valid on 'loop' directive}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented}}
+  // expected-error@+1{{OpenACC 'self' clause is not valid on 'loop' directive}}
 #pragma acc loop self
  for(;;);
 }
--- a/clang/test/SemaOpenACC/compute-construct-vector_length-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-vector_length-clause.c
@ -31,8 +31,7 @@ void Test() {
 #pragma acc kernels vector_length(SomeE)
  while(1);

-  // expected-error@+2{{OpenACC 'vector_length' clause is not valid on 'loop' directive}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented}}
+  // expected-error@+1{{OpenACC 'vector_length' clause is not valid on 'loop' directive}}
 #pragma acc loop vector_length(1)
  for(;;);
 }
--- a/clang/test/SemaOpenACC/compute-construct-wait-clause.c
+++ b/clang/test/SemaOpenACC/compute-construct-wait-clause.c
@ -36,8 +36,7 @@ void uses() {
 #pragma acc parallel wait(devnum:arr : queues: arr, NC, 5)
  while(1);

-  // expected-error@+2{{OpenACC 'wait' clause is not valid on 'loop' directive}}
-  // expected-warning@+1{{OpenACC construct 'loop' not yet implemented}}
+  // expected-error@+1{{OpenACC 'wait' clause is not valid on 'loop' directive}}
 #pragma acc loop wait
  for(;;);
 }
--- a/clang/test/SemaOpenACC/loop-ast.cpp
+++ b/clang/test/SemaOpenACC/loop-ast.cpp
@ -0,0 +1,182 @@
+
+// RUN: %clang_cc1 %s -fopenacc -ast-dump | FileCheck %s
+
+// Test this with PCH.
+// RUN: %clang_cc1 %s -fopenacc -emit-pch -o %t %s
+// RUN: %clang_cc1 %s -fopenacc -include-pch %t -ast-dump-all | FileCheck %s
+
+#ifndef PCH_HELPER
+#define PCH_HELPER
+
+void NormalFunc() {
+  // CHECK-LABEL: NormalFunc
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc loop
+  for(;;);
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} <orphan>
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+
+  int array[5];
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl
+#pragma acc loop
+  for(auto x : array){}
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} <orphan>
+  // CHECK-NEXT: CXXForRangeStmt
+  // CHECK: CompoundStmt
+
+#pragma acc parallel
+  // CHECK-NEXT: OpenACCComputeConstruct {{.*}}parallel
+  // CHECK-NEXT: CompoundStmt
+  {
+#pragma acc parallel
+    // CHECK-NEXT: OpenACCComputeConstruct [[PAR_ADDR:[0-9a-fx]+]] {{.*}}parallel
+    // CHECK-NEXT: CompoundStmt
+    {
+#pragma acc loop
+      for(;;);
+    // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: [[PAR_ADDR]]
+    // CHECK-NEXT: ForStmt
+    // CHECK-NEXT: <<<NULL>>>
+    // CHECK-NEXT: <<<NULL>>>
+    // CHECK-NEXT: <<<NULL>>>
+    // CHECK-NEXT: <<<NULL>>>
+    // CHECK-NEXT: NullStmt
+    }
+  }
+}
+
+template<typename T>
+void TemplFunc() {
+  // CHECK-LABEL: FunctionTemplateDecl {{.*}}TemplFunc
+  // CHECK-NEXT: TemplateTypeParmDecl
+  // CHECK-NEXT: FunctionDecl{{.*}}TemplFunc
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc loop
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} <orphan>
+  for(typename T::type t = 0; t < 5;++t) {
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} referenced t 'typename T::type'
+  // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 0
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: BinaryOperator{{.*}} '<dependent type>' '<'
+  // CHECK-NEXT: DeclRefExpr {{.*}} 'typename T::type' lvalue Var
+  // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 5
+  // CHECK-NEXT: UnaryOperator{{.*}} '<dependent type>' lvalue prefix '++'
+  // CHECK-NEXT: DeclRefExpr {{.*}} 'typename T::type' lvalue Var
+  // CHECK-NEXT: CompoundStmt
+    typename T::type I;
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} I 'typename T::type'
+
+  }
+
+#pragma acc parallel
+  {
+    // CHECK-NEXT: OpenACCComputeConstruct {{.*}}parallel
+    // CHECK-NEXT: CompoundStmt
+#pragma acc parallel
+    {
+    // CHECK-NEXT: OpenACCComputeConstruct [[PAR_ADDR_UNINST:[0-9a-fx]+]] {{.*}}parallel
+    // CHECK-NEXT: CompoundStmt
+#pragma acc loop
+    // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: [[PAR_ADDR_UNINST]]
+    // CHECK-NEXT: ForStmt
+    // CHECK-NEXT: <<<NULL>>>
+    // CHECK-NEXT: <<<NULL>>>
+    // CHECK-NEXT: <<<NULL>>>
+    // CHECK-NEXT: <<<NULL>>>
+    // CHECK-NEXT: NullStmt
+      for(;;);
+
+#pragma acc loop
+    // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: [[PAR_ADDR_UNINST]]
+    // CHECK-NEXT: ForStmt
+    // CHECK-NEXT: <<<NULL>>>
+    // CHECK-NEXT: <<<NULL>>>
+    // CHECK-NEXT: <<<NULL>>>
+    // CHECK-NEXT: <<<NULL>>>
+    // CHECK-NEXT: NullStmt
+      for(;;);
+    }
+  }
+
+  typename T::type array[5];
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl
+
+#pragma acc loop
+  for(auto x : array){}
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} <orphan>
+  // CHECK-NEXT: CXXForRangeStmt
+  // CHECK: CompoundStmt
+
+  // Instantiation:
+  // CHECK-NEXT: FunctionDecl{{.*}} TemplFunc 'void ()' implicit_instantiation
+  // CHECK-NEXT: TemplateArgument type 'S'
+  // CHECK-NEXT: RecordType{{.*}} 'S'
+  // CHECK-NEXT: CXXRecord{{.*}} 'S'
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} <orphan>
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} used t 'typename S::type':'int'
+  // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 0
+  // CHECK-NEXT: <<<NULL>>
+  // CHECK-NEXT: BinaryOperator{{.*}} 'bool' '<'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'typename S::type':'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr {{.*}} 'typename S::type':'int' lvalue Var
+  // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 5
+  // CHECK-NEXT: UnaryOperator{{.*}} 'typename S::type':'int' lvalue prefix '++'
+  // CHECK-NEXT: DeclRefExpr {{.*}} 'typename S::type':'int' lvalue Var
+  // CHECK-NEXT: CompoundStmt
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} I 'typename S::type':'int'
+
+  // CHECK-NEXT: OpenACCComputeConstruct {{.*}}parallel
+  // CHECK-NEXT: CompoundStmt
+  //
+  // CHECK-NEXT: OpenACCComputeConstruct [[PAR_ADDR_INST:[0-9a-fx]+]] {{.*}}parallel
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: [[PAR_ADDR_INST]]
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: [[PAR_ADDR_INST]]
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} <orphan>
+  // CHECK-NEXT: CXXForRangeStmt
+  // CHECK: CompoundStmt
+}
+
+struct S {
+  using type = int;
+};
+
+void use() {
+  TemplFunc<S>();
+}
+#endif
+
--- a/clang/test/SemaOpenACC/loop-loc-and-stmt.c
+++ b/clang/test/SemaOpenACC/loop-loc-and-stmt.c
@ -0,0 +1,38 @@
+// RUN: %clang_cc1 %s -verify -fopenacc
+
+// expected-error@+1{{OpenACC construct 'loop' cannot be used here; it can only be used in a statement context}}
+#pragma acc loop
+
+// expected-error@+1{{OpenACC construct 'loop' cannot be used here; it can only be used in a statement context}}
+#pragma acc loop
+int foo;
+
+struct S {
+// expected-error@+1{{OpenACC construct 'loop' cannot be used here; it can only be used in a statement context}}
+#pragma acc loop
+  int i;
+};
+
+void func() {
+  // expected-error@+2{{expected expression}}
+#pragma acc loop
+  int foo;
+
+  // expected-error@+3{{OpenACC 'loop' construct can only be applied to a 'for' loop}}
+  // expected-note@+1{{'loop' construct is here}}
+#pragma acc loop
+  while(0);
+
+  // expected-error@+3{{OpenACC 'loop' construct can only be applied to a 'for' loop}}
+  // expected-note@+1{{'loop' construct is here}}
+#pragma acc loop
+  do{}while(0);
+
+  // expected-error@+3{{OpenACC 'loop' construct can only be applied to a 'for' loop}}
+  // expected-note@+1{{'loop' construct is here}}
+#pragma acc loop
+  {}
+
+#pragma acc loop
+  for(;;);
+}
--- a/clang/test/SemaOpenACC/loop-loc-and-stmt.cpp
+++ b/clang/test/SemaOpenACC/loop-loc-and-stmt.cpp
@ -0,0 +1,80 @@
+// RUN: %clang_cc1 %s -verify -fopenacc
+//
+// expected-error@+1{{OpenACC construct 'loop' cannot be used here; it can only be used in a statement context}}
+#pragma acc loop
+
+// expected-error@+1{{OpenACC construct 'loop' cannot be used here; it can only be used in a statement context}}
+#pragma acc loop
+int foo;
+
+struct S {
+// expected-error@+1{{OpenACC construct 'loop' cannot be used here; it can only be used in a statement context}}
+#pragma acc loop
+  int i;
+
+  void mem_func() {
+  // expected-error@+3{{OpenACC 'loop' construct can only be applied to a 'for' loop}}
+  // expected-note@+1{{'loop' construct is here}}
+#pragma acc loop
+    int foo;
+
+  // expected-error@+3{{OpenACC 'loop' construct can only be applied to a 'for' loop}}
+  // expected-note@+1{{'loop' construct is here}}
+#pragma acc loop
+    while(0);
+
+  // expected-error@+3{{OpenACC 'loop' construct can only be applied to a 'for' loop}}
+  // expected-note@+1{{'loop' construct is here}}
+#pragma acc loop
+    do{}while(0);
+
+  // expected-error@+3{{OpenACC 'loop' construct can only be applied to a 'for' loop}}
+  // expected-note@+1{{'loop' construct is here}}
+#pragma acc loop
+    {}
+
+#pragma acc loop
+    for(;;);
+
+    int array[5];
+
+#pragma acc loop
+    for(auto X : array){}
+}
+};
+
+template<typename T>
+void templ_func() {
+  // expected-error@+3{{OpenACC 'loop' construct can only be applied to a 'for' loop}}
+  // expected-note@+1{{'loop' construct is here}}
+#pragma acc loop
+  int foo;
+
+  // expected-error@+3{{OpenACC 'loop' construct can only be applied to a 'for' loop}}
+  // expected-note@+1{{'loop' construct is here}}
+#pragma acc loop
+  while(T{});
+
+  // expected-error@+3{{OpenACC 'loop' construct can only be applied to a 'for' loop}}
+  // expected-note@+1{{'loop' construct is here}}
+#pragma acc loop
+  do{}while(0);
+
+  // expected-error@+3{{OpenACC 'loop' construct can only be applied to a 'for' loop}}
+  // expected-note@+1{{'loop' construct is here}}
+#pragma acc loop
+  {}
+
+#pragma acc loop
+  for(T i;;);
+
+  T array[5];
+
+#pragma acc loop
+  for(auto X : array){}
+}
+
+void use() {
+  templ_func<int>();
+}
+
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@ -2170,6 +2170,7 @@ public:
  void VisitRequiresExpr(const RequiresExpr *E);
  void VisitCXXParenListInitExpr(const CXXParenListInitExpr *E);
  void VisitOpenACCComputeConstruct(const OpenACCComputeConstruct *D);
+  void VisitOpenACCLoopConstruct(const OpenACCLoopConstruct *D);
  void VisitOMPExecutableDirective(const OMPExecutableDirective *D);
  void VisitOMPLoopBasedDirective(const OMPLoopBasedDirective *D);
  void VisitOMPLoopDirective(const OMPLoopDirective *D);
@ -3496,6 +3497,12 @@ void EnqueueVisitor::VisitOpenACCComputeConstruct(
    EnqueueChildren(Clause);
 }

+void EnqueueVisitor::VisitOpenACCLoopConstruct(const OpenACCLoopConstruct *C) {
+  EnqueueChildren(C);
+  for (auto *Clause : C->clauses())
+    EnqueueChildren(Clause);
+}
+
 void EnqueueVisitor::VisitAnnotateAttr(const AnnotateAttr *A) {
  EnqueueChildren(A);
 }
@ -6234,6 +6241,8 @@ CXString clang_getCursorKindSpelling(enum CXCursorKind Kind) {
    return cxstring::createRef("ConceptDecl");
  case CXCursor_OpenACCComputeConstruct:
    return cxstring::createRef("OpenACCComputeConstruct");
+  case CXCursor_OpenACCLoopConstruct:
+    return cxstring::createRef("OpenACCLoopConstruct");
  }

  llvm_unreachable("Unhandled CXCursorKind");
--- a/clang/tools/libclang/CXCursor.cpp
+++ b/clang/tools/libclang/CXCursor.cpp
@ -873,6 +873,9 @@ CXCursor cxcursor::MakeCXCursor(const Stmt *S, const Decl *Parent,
  case Stmt::OpenACCComputeConstructClass:
    K = CXCursor_OpenACCComputeConstruct;
    break;
+  case Stmt::OpenACCLoopConstructClass:
+    K = CXCursor_OpenACCLoopConstruct;
+    break;
  case Stmt::OMPTargetParallelGenericLoopDirectiveClass:
    K = CXCursor_OMPTargetParallelGenericLoopDirective;
    break;