diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index e8b8f477f91ae..5ad4c336b6c53 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -119,32 +119,6 @@ class OpenACCSeqClause : public OpenACCClause {
   }
 };
 
-// Not yet implemented, but the type name is necessary for 'seq' diagnostics, so
-// this provides a basic, do-nothing implementation. We still need to add this
-// type to the visitors/etc, as well as get it to take its proper arguments.
-class OpenACCVectorClause : public OpenACCClause {
-protected:
-  OpenACCVectorClause(SourceLocation BeginLoc, SourceLocation EndLoc)
-      : OpenACCClause(OpenACCClauseKind::Vector, BeginLoc, EndLoc) {
-    llvm_unreachable("Not yet implemented");
-  }
-
-public:
-  static bool classof(const OpenACCClause *C) {
-    return C->getClauseKind() == OpenACCClauseKind::Vector;
-  }
-
-  static OpenACCVectorClause *
-  Create(const ASTContext &Ctx, SourceLocation BeginLoc, SourceLocation EndLoc);
-
-  child_range children() {
-    return child_range(child_iterator(), child_iterator());
-  }
-  const_child_range children() const {
-    return const_child_range(const_child_iterator(), const_child_iterator());
-  }
-};
-
 /// Represents a clause that has a list of parameters.
 class OpenACCClauseWithParams : public OpenACCClause {
   /// Location of the '('.
@@ -531,6 +505,22 @@ class OpenACCWorkerClause : public OpenACCClauseWithSingleIntExpr {
                                      SourceLocation EndLoc);
 };
 
+class OpenACCVectorClause : public OpenACCClauseWithSingleIntExpr {
+protected:
+  OpenACCVectorClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
+                      Expr *IntExpr, SourceLocation EndLoc);
+
+public:
+  static bool classof(const OpenACCClause *C) {
+    return C->getClauseKind() == OpenACCClauseKind::Vector;
+  }
+
+  static OpenACCVectorClause *Create(const ASTContext &Ctx,
+                                     SourceLocation BeginLoc,
+                                     SourceLocation LParenLoc, Expr *IntExpr,
+                                     SourceLocation EndLoc);
+};
+
 class OpenACCNumWorkersClause : public OpenACCClauseWithSingleIntExpr {
   OpenACCNumWorkersClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
                           Expr *IntExpr, SourceLocation EndLoc);
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index c709795e7b21d..e78acc8dc8c57 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12702,7 +12702,7 @@ def err_acc_gang_dim_value
 def err_acc_num_arg_conflict
     : Error<"'num' argument to '%0' clause not allowed on a 'loop' construct "
             "associated with a 'kernels' construct that has a "
-            "'%select{num_gangs|num_workers}1' "
+            "'%select{num_gangs|num_workers|vector_length}1' "
             "clause">;
 def err_acc_clause_in_clause_region
     : Error<"loop with a '%0' clause may not exist in the region of a '%1' "
diff --git a/clang/include/clang/Basic/OpenACCClauses.def b/clang/include/clang/Basic/OpenACCClauses.def
index 4c0b56dc13e62..c65ebed751cf1 100644
--- a/clang/include/clang/Basic/OpenACCClauses.def
+++ b/clang/include/clang/Basic/OpenACCClauses.def
@@ -54,6 +54,7 @@ VISIT_CLAUSE(Reduction)
 VISIT_CLAUSE(Self)
 VISIT_CLAUSE(Seq)
 VISIT_CLAUSE(Tile)
+VISIT_CLAUSE(Vector)
 VISIT_CLAUSE(VectorLength)
 VISIT_CLAUSE(Wait)
 VISIT_CLAUSE(Worker)
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index e253610a84b0b..d6e56f85e2d54 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -123,6 +123,11 @@ class SemaOpenACC : public SemaBase {
   /// permits us to implement the restriction of no further 'gang' or 'worker'
   /// clauses.
   SourceLocation LoopWorkerClauseLoc;
+  /// If there is a current 'active' loop construct with a 'vector' clause on it
+  /// (on any sort of construct), this has the source location for it.  This
+  /// permits us to implement the restriction of no further 'gang', 'vector', or
+  /// 'worker' clauses.
+  SourceLocation LoopVectorClauseLoc;
 
   // Redeclaration of the version in OpenACCClause.h.
   using DeviceTypeArgument = std::pair<IdentifierInfo *, SourceLocation>;
@@ -679,6 +684,7 @@ class SemaOpenACC : public SemaBase {
     OpenACCDirectiveKind DirKind;
     SourceLocation OldLoopGangClauseOnKernelLoc;
     SourceLocation OldLoopWorkerClauseLoc;
+    SourceLocation OldLoopVectorClauseLoc;
     llvm::SmallVector<OpenACCLoopConstruct *> ParentlessLoopConstructs;
     LoopInConstructRAII LoopRAII;
 
diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp
index 638252fd811f1..1299e4f807ceb 100644
--- a/clang/lib/AST/OpenACCClause.cpp
+++ b/clang/lib/AST/OpenACCClause.cpp
@@ -44,8 +44,8 @@ bool OpenACCClauseWithCondition::classof(const OpenACCClause *C) {
 bool OpenACCClauseWithSingleIntExpr::classof(const OpenACCClause *C) {
   return OpenACCNumWorkersClause::classof(C) ||
          OpenACCVectorLengthClause::classof(C) ||
-         OpenACCWorkerClause::classof(C) || OpenACCCollapseClause::classof(C) ||
-         OpenACCAsyncClause::classof(C);
+         OpenACCVectorClause::classof(C) || OpenACCWorkerClause::classof(C) ||
+         OpenACCCollapseClause::classof(C) || OpenACCAsyncClause::classof(C);
 }
 OpenACCDefaultClause *OpenACCDefaultClause::Create(const ASTContext &C,
                                                    OpenACCDefaultClauseKind K,
@@ -424,11 +424,24 @@ OpenACCWorkerClause *OpenACCWorkerClause::Create(const ASTContext &C,
   return new (Mem) OpenACCWorkerClause(BeginLoc, LParenLoc, IntExpr, EndLoc);
 }
 
+OpenACCVectorClause::OpenACCVectorClause(SourceLocation BeginLoc,
+                                         SourceLocation LParenLoc,
+                                         Expr *IntExpr, SourceLocation EndLoc)
+    : OpenACCClauseWithSingleIntExpr(OpenACCClauseKind::Vector, BeginLoc,
+                                     LParenLoc, IntExpr, EndLoc) {
+  assert((!IntExpr || IntExpr->isInstantiationDependent() ||
+          IntExpr->getType()->isIntegerType()) &&
+         "Int expression type not scalar/dependent");
+}
+
 OpenACCVectorClause *OpenACCVectorClause::Create(const ASTContext &C,
                                                  SourceLocation BeginLoc,
+                                                 SourceLocation LParenLoc,
+                                                 Expr *IntExpr,
                                                  SourceLocation EndLoc) {
-  void *Mem = C.Allocate(sizeof(OpenACCVectorClause));
-  return new (Mem) OpenACCVectorClause(BeginLoc, EndLoc);
+  void *Mem =
+      C.Allocate(sizeof(OpenACCVectorClause), alignof(OpenACCVectorClause));
+  return new (Mem) OpenACCVectorClause(BeginLoc, LParenLoc, IntExpr, EndLoc);
 }
 
 //===----------------------------------------------------------------------===//
@@ -662,3 +675,13 @@ void OpenACCClausePrinter::VisitWorkerClause(const OpenACCWorkerClause &C) {
     OS << ")";
   }
 }
+
+void OpenACCClausePrinter::VisitVectorClause(const OpenACCVectorClause &C) {
+  OS << "vector";
+
+  if (C.hasIntExpr()) {
+    OS << "(length: ";
+    printExpr(C.getIntExpr());
+    OS << ")";
+  }
+}
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index 25b1cbb859086..01f9a30d06cd7 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -2635,6 +2635,12 @@ void OpenACCClauseProfiler::VisitWorkerClause(
     Profiler.VisitStmt(Clause.getIntExpr());
 }
 
+void OpenACCClauseProfiler::VisitVectorClause(
+    const OpenACCVectorClause &Clause) {
+  if (Clause.hasIntExpr())
+    Profiler.VisitStmt(Clause.getIntExpr());
+}
+
 void OpenACCClauseProfiler::VisitWaitClause(const OpenACCWaitClause &Clause) {
   if (Clause.hasDevNumExpr())
     Profiler.VisitStmt(Clause.getDevNumExpr());
diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp
index beccb0615f0e9..01cfb1f63f708 100644
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp
@@ -421,6 +421,7 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
     case OpenACCClauseKind::Seq:
     case OpenACCClauseKind::Tile:
     case OpenACCClauseKind::Worker:
+    case OpenACCClauseKind::Vector:
     case OpenACCClauseKind::VectorLength:
       // The condition expression will be printed as a part of the 'children',
       // but print 'clause' here so it is clear what is happening from the dump.
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 1b24331cbd87c..22aedbc70df8c 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -389,6 +389,18 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
       return false;
     }
   }
+  case OpenACCClauseKind::Vector: {
+    switch (DirectiveKind) {
+    case OpenACCDirectiveKind::Loop:
+    case OpenACCDirectiveKind::ParallelLoop:
+    case OpenACCDirectiveKind::SerialLoop:
+    case OpenACCDirectiveKind::KernelsLoop:
+    case OpenACCDirectiveKind::Routine:
+      return true;
+    default:
+      return false;
+    }
+  }
   }
 
   default:
@@ -512,14 +524,6 @@ class SemaOpenACCClauseVisitor {
 
   OpenACCClause *Visit(SemaOpenACC::OpenACCParsedClause &Clause) {
     switch (Clause.getClauseKind()) {
-    case OpenACCClauseKind::Vector: {
-      // TODO OpenACC: These are only implemented enough for the 'seq'
-      // diagnostic, otherwise treats itself as unimplemented.  When we
-      // implement these, we can remove them from here.
-      DiagIfSeqClause(Clause);
-      return isNotImplemented();
-    }
-
 #define VISIT_CLAUSE(CLAUSE_NAME)                                              \
   case OpenACCClauseKind::CLAUSE_NAME:                                         \
     return Visit##CLAUSE_NAME##Clause(Clause);
@@ -1035,6 +1039,97 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitIndependentClause(
                                           Clause.getEndLoc());
 }
 
+OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorClause(
+    SemaOpenACC::OpenACCParsedClause &Clause) {
+  if (DiagIfSeqClause(Clause))
+    return nullptr;
+  // Restrictions only properly implemented on 'loop' constructs, and it is
+  // the only construct that can do anything with this, so skip/treat as
+  // unimplemented for the combined constructs.
+  if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop)
+    return isNotImplemented();
+
+  Expr *IntExpr =
+      Clause.getNumIntExprs() != 0 ? Clause.getIntExprs()[0] : nullptr;
+  if (IntExpr) {
+    switch (SemaRef.getActiveComputeConstructInfo().Kind) {
+    case OpenACCDirectiveKind::Invalid:
+    case OpenACCDirectiveKind::Parallel:
+      // No restriction on when 'parallel' can contain an argument.
+      break;
+    case OpenACCDirectiveKind::Serial:
+      // GCC disallows this, and there is no real good reason for us to permit
+      // it, so disallow until we come up with a use case that makes sense.
+      SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_int_arg_invalid)
+          << OpenACCClauseKind::Vector << "num" << /*serial=*/3;
+      IntExpr = nullptr;
+      break;
+    case OpenACCDirectiveKind::Kernels: {
+      const auto *Itr =
+          llvm::find_if(SemaRef.getActiveComputeConstructInfo().Clauses,
+                        llvm::IsaPred<OpenACCVectorLengthClause>);
+      if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) {
+        SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict)
+            << OpenACCClauseKind::Vector << /*vector_length=*/2;
+        SemaRef.Diag((*Itr)->getBeginLoc(),
+                     diag::note_acc_previous_clause_here);
+
+        IntExpr = nullptr;
+      }
+      break;
+    }
+    default:
+      llvm_unreachable("Non compute construct in active compute construct");
+    }
+  }
+
+  // OpenACC 3.3 2.9.2: When the parent compute construct is a kernels
+  // construct, the gang clause behaves as follows. ... The region of a loop
+  // with a gang clause may not contain another loop with a gang clause unless
+  // within a nested compute region.
+  if (SemaRef.LoopGangClauseOnKernelLoc.isValid()) {
+    // This handles the 'inner loop' diagnostic, but we cannot set that we're on
+    // one of these until we get to the end of the construct.
+    SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region)
+        << OpenACCClauseKind::Vector << OpenACCClauseKind::Gang
+        << /*skip kernels construct info*/ 0;
+    SemaRef.Diag(SemaRef.LoopGangClauseOnKernelLoc,
+                 diag::note_acc_previous_clause_here);
+    return nullptr;
+  }
+
+  // OpenACC 3.3 2.9.3: The region of a loop with a 'worker' clause may not
+  // contain a loop with a gang or worker clause unless within a nested compute
+  // region.
+  if (SemaRef.LoopWorkerClauseLoc.isValid()) {
+    // This handles the 'inner loop' diagnostic, but we cannot set that we're on
+    // one of these until we get to the end of the construct.
+    SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region)
+        << OpenACCClauseKind::Vector << OpenACCClauseKind::Worker
+        << /*skip kernels construct info*/ 0;
+    SemaRef.Diag(SemaRef.LoopWorkerClauseLoc,
+                 diag::note_acc_previous_clause_here);
+    return nullptr;
+  }
+  // OpenACC 3.3 2.9.4: The region of a loop with a 'vector' clause may not
+  // contain a loop with a gang, worker, or vector clause unless within a nested
+  // compute region.
+  if (SemaRef.LoopVectorClauseLoc.isValid()) {
+    // This handles the 'inner loop' diagnostic, but we cannot set that we're on
+    // one of these until we get to the end of the construct.
+    SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region)
+        << OpenACCClauseKind::Vector << OpenACCClauseKind::Vector
+        << /*skip kernels construct info*/ 0;
+    SemaRef.Diag(SemaRef.LoopVectorClauseLoc,
+                 diag::note_acc_previous_clause_here);
+    return nullptr;
+  }
+
+  return OpenACCVectorClause::Create(Ctx, Clause.getBeginLoc(),
+                                     Clause.getLParenLoc(), IntExpr,
+                                     Clause.getEndLoc());
+}
+
 OpenACCClause *SemaOpenACCClauseVisitor::VisitWorkerClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
   if (DiagIfSeqClause(Clause))
@@ -1099,6 +1194,20 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitWorkerClause(
     return nullptr;
   }
 
+  // OpenACC 3.3 2.9.4: The region of a loop with a 'vector' clause may not
+  // contain a loop with a gang, worker, or vector clause unless within a nested
+  // compute region.
+  if (SemaRef.LoopVectorClauseLoc.isValid()) {
+    // This handles the 'inner loop' diagnostic, but we cannot set that we're on
+    // one of these until we get to the end of the construct.
+    SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region)
+        << OpenACCClauseKind::Worker << OpenACCClauseKind::Vector
+        << /*skip kernels construct info*/ 0;
+    SemaRef.Diag(SemaRef.LoopVectorClauseLoc,
+                 diag::note_acc_previous_clause_here);
+    return nullptr;
+  }
+
   return OpenACCWorkerClause::Create(Ctx, Clause.getBeginLoc(),
                                      Clause.getLParenLoc(), IntExpr,
                                      Clause.getEndLoc());
@@ -1193,6 +1302,20 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause(
     return nullptr;
   }
 
+  // OpenACC 3.3 2.9.4: The region of a loop with a 'vector' clause may not
+  // contain a loop with a gang, worker, or vector clause unless within a nested
+  // compute region.
+  if (SemaRef.LoopVectorClauseLoc.isValid()) {
+    // This handles the 'inner loop' diagnostic, but we cannot set that we're on
+    // one of these until we get to the end of the construct.
+    SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region)
+        << OpenACCClauseKind::Gang << OpenACCClauseKind::Vector
+        << /*kernels construct info*/ 1;
+    SemaRef.Diag(SemaRef.LoopVectorClauseLoc,
+                 diag::note_acc_previous_clause_here);
+    return nullptr;
+  }
+
   return OpenACCGangClause::Create(Ctx, Clause.getBeginLoc(),
                                    Clause.getLParenLoc(), GangKinds, IntExprs,
                                    Clause.getEndLoc());
@@ -1313,6 +1436,7 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
     : SemaRef(S), OldActiveComputeConstructInfo(S.ActiveComputeConstructInfo),
       DirKind(DK), OldLoopGangClauseOnKernelLoc(S.LoopGangClauseOnKernelLoc),
       OldLoopWorkerClauseLoc(S.LoopWorkerClauseLoc),
+      OldLoopVectorClauseLoc(S.LoopVectorClauseLoc),
       LoopRAII(SemaRef, /*PreserveDepth=*/false) {
   // Compute constructs end up taking their 'loop'.
   if (DirKind == OpenACCDirectiveKind::Parallel ||
@@ -1330,6 +1454,7 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
     // Implement the 'unless within a nested compute region' part.
     SemaRef.LoopGangClauseOnKernelLoc = {};
     SemaRef.LoopWorkerClauseLoc = {};
+    SemaRef.LoopVectorClauseLoc = {};
   } else if (DirKind == OpenACCDirectiveKind::Loop) {
     SetCollapseInfoBeforeAssociatedStmt(UnInstClauses, Clauses);
     SetTileInfoBeforeAssociatedStmt(UnInstClauses, Clauses);
@@ -1355,6 +1480,10 @@ SemaOpenACC::AssociatedStmtRAII::AssociatedStmtRAII(
       auto *Itr = llvm::find_if(Clauses, llvm::IsaPred<OpenACCWorkerClause>);
       if (Itr != Clauses.end())
         SemaRef.LoopWorkerClauseLoc = (*Itr)->getBeginLoc();
+
+      auto *Itr2 = llvm::find_if(Clauses, llvm::IsaPred<OpenACCVectorClause>);
+      if (Itr2 != Clauses.end())
+        SemaRef.LoopVectorClauseLoc = (*Itr2)->getBeginLoc();
     }
   }
 }
@@ -1429,6 +1558,7 @@ SemaOpenACC::AssociatedStmtRAII::~AssociatedStmtRAII() {
   SemaRef.ActiveComputeConstructInfo = OldActiveComputeConstructInfo;
   SemaRef.LoopGangClauseOnKernelLoc = OldLoopGangClauseOnKernelLoc;
   SemaRef.LoopWorkerClauseLoc = OldLoopWorkerClauseLoc;
+  SemaRef.LoopVectorClauseLoc = OldLoopVectorClauseLoc;
 
   if (DirKind == OpenACCDirectiveKind::Parallel ||
       DirKind == OpenACCDirectiveKind::Serial ||
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index 45e8b3cf6bd8f..c5a6e677ef8de 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -11815,6 +11815,33 @@ void OpenACCClauseTransform<Derived>::VisitWorkerClause(
       ParsedClause.getEndLoc());
 }
 
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitVectorClause(
+    const OpenACCVectorClause &C) {
+  if (C.hasIntExpr()) {
+    // restrictions on this expression are all "does it exist in certain
+    // situations" that are not possible to be dependent, so the only check we
+    // have is that it transforms, and is an int expression.
+    ExprResult Res = Self.TransformExpr(const_cast<Expr *>(C.getIntExpr()));
+    if (!Res.isUsable())
+      return;
+
+    Res = Self.getSema().OpenACC().ActOnIntExpr(OpenACCDirectiveKind::Invalid,
+                                                C.getClauseKind(),
+                                                C.getBeginLoc(), Res.get());
+    if (!Res.isUsable())
+      return;
+    ParsedClause.setIntExprDetails(Res.get());
+  }
+
+  NewClause = OpenACCVectorClause::Create(
+      Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
+      ParsedClause.getLParenLoc(),
+      ParsedClause.getNumIntExprs() != 0 ? ParsedClause.getIntExprs()[0]
+                                         : nullptr,
+      ParsedClause.getEndLoc());
+}
+
 template <typename Derived>
 void OpenACCClauseTransform<Derived>::VisitWaitClause(
     const OpenACCWaitClause &C) {
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index ecc5d3c59a354..1b2473f245734 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -12345,10 +12345,15 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
     return OpenACCWorkerClause::Create(getContext(), BeginLoc, LParenLoc,
                                        WorkerExpr, EndLoc);
   }
+  case OpenACCClauseKind::Vector: {
+    SourceLocation LParenLoc = readSourceLocation();
+    Expr *VectorExpr = readBool() ? readSubExpr() : nullptr;
+    return OpenACCVectorClause::Create(getContext(), BeginLoc, LParenLoc,
+                                       VectorExpr, EndLoc);
+  }
 
   case OpenACCClauseKind::Finalize:
   case OpenACCClauseKind::IfPresent:
-  case OpenACCClauseKind::Vector:
   case OpenACCClauseKind::NoHost:
   case OpenACCClauseKind::UseDevice:
   case OpenACCClauseKind::Delete:
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index 0a6e260e3e4e9..938d7b525cb95 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -8200,10 +8200,17 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
       AddStmt(const_cast<Expr *>(WC->getIntExpr()));
     return;
   }
+  case OpenACCClauseKind::Vector: {
+    const auto *VC = cast<OpenACCVectorClause>(C);
+    writeSourceLocation(VC->getLParenLoc());
+    writeBool(VC->hasIntExpr());
+    if (VC->hasIntExpr())
+      AddStmt(const_cast<Expr *>(VC->getIntExpr()));
+    return;
+  }
 
   case OpenACCClauseKind::Finalize:
   case OpenACCClauseKind::IfPresent:
-  case OpenACCClauseKind::Vector:
   case OpenACCClauseKind::NoHost:
   case OpenACCClauseKind::UseDevice:
   case OpenACCClauseKind::Delete:
diff --git a/clang/test/AST/ast-print-openacc-loop-construct.cpp b/clang/test/AST/ast-print-openacc-loop-construct.cpp
index ee11435aaa4b1..c0ca274f38dc2 100644
--- a/clang/test/AST/ast-print-openacc-loop-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-loop-construct.cpp
@@ -216,4 +216,79 @@ void foo() {
 #pragma acc kernels
 #pragma acc loop worker(num:5)
   for(;;);
+
+  // CHECK: #pragma acc loop vector
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc loop vector
+  for(;;);
+
+// CHECK: #pragma acc loop vector(length: 5)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc loop vector(5)
+  for(;;);
+
+// CHECK: #pragma acc loop vector(length: 5)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc loop vector(length:5)
+  for(;;);
+
+// CHECK: #pragma acc parallel
+// CHECK-NEXT: #pragma acc loop vector
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc parallel
+#pragma acc loop vector
+  for(;;);
+
+// CHECK: #pragma acc parallel
+// CHECK-NEXT: #pragma acc loop vector(length: 5)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc parallel
+#pragma acc loop vector(5)
+  for(;;);
+
+// CHECK: #pragma acc parallel
+// CHECK-NEXT: #pragma acc loop vector(length: 5)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc parallel
+#pragma acc loop vector(length:5)
+  for(;;);
+
+// CHECK: #pragma acc kernels
+// CHECK-NEXT: #pragma acc loop vector
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc kernels
+#pragma acc loop vector
+  for(;;);
+
+// CHECK: #pragma acc kernels
+// CHECK-NEXT: #pragma acc loop vector(length: 5)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc kernels
+#pragma acc loop vector(5)
+  for(;;);
+
+// CHECK: #pragma acc kernels
+// CHECK-NEXT: #pragma acc loop vector(length: 5)
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc kernels
+#pragma acc loop vector(length:5)
+  for(;;);
+
+// CHECK: #pragma acc serial
+// CHECK-NEXT: #pragma acc loop vector
+// CHECK-NEXT: for (;;)
+// CHECK-NEXT: ;
+#pragma acc serial
+#pragma acc loop vector
+  for(;;);
+
 }
diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c
index 81c48335cf0c4..26f39be80030b 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -898,7 +898,6 @@ void IntExprParsing() {
 #pragma acc set default_async(returns_int())
 
 
-  // expected-warning@+1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
 #pragma acc loop vector
   for(;;);
   // expected-error@+1{{expected expression}}
@@ -908,8 +907,7 @@ void IntExprParsing() {
   // expected-error@+1{{expected expression}}
 #pragma acc loop vector(invalid:)
   for(;;);
-  // expected-error@+2{{invalid tag 'invalid' on 'vector' clause}}
-  // expected-warning@+1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
+  // expected-error@+1{{invalid tag 'invalid' on 'vector' clause}}
 #pragma acc loop vector(invalid:5)
   for(;;);
   // expected-error@+1{{expected expression}}
@@ -932,20 +930,15 @@ void IntExprParsing() {
   // 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}}
+  // expected-error@+1{{invalid tag 'num' on 'vector' clause}}
 #pragma acc loop vector(num:5)
   for(;;);
-  // expected-warning@+1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
 #pragma acc loop vector(length:5)
   for(;;);
-  // expected-warning@+1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
 #pragma acc loop vector(returns_int())
   for(;;);
-  // expected-warning@+1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
 #pragma acc loop vector(length:returns_int())
   for(;;);
 
diff --git a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c b/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
index 6a975956f3ff5..ab10857e3cd85 100644
--- a/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
+++ b/clang/test/SemaOpenACC/loop-construct-auto_seq_independent-clauses.c
@@ -45,7 +45,6 @@ void uses() {
   for(;;);
 #pragma acc loop auto worker
   for(;;);
-  // expected-warning@+1{{OpenACC clause 'vector' not yet implemented}}
 #pragma acc loop auto vector
   for(;;);
   // expected-warning@+1{{OpenACC clause 'nohost' not yet implemented}}
@@ -181,7 +180,6 @@ void uses() {
   for(;;);
 #pragma acc loop worker auto
   for(;;);
-  // expected-warning@+1{{OpenACC clause 'vector' not yet implemented}}
 #pragma acc loop vector auto
   for(;;);
   // expected-warning@+1{{OpenACC clause 'nohost' not yet implemented}}
@@ -318,7 +316,6 @@ void uses() {
   for(;;);
 #pragma acc loop independent worker
   for(;;);
-  // expected-warning@+1{{OpenACC clause 'vector' not yet implemented}}
 #pragma acc loop independent vector
   for(;;);
   // expected-warning@+1{{OpenACC clause 'nohost' not yet implemented}}
@@ -454,7 +451,6 @@ void uses() {
   for(;;);
 #pragma acc loop worker independent
   for(;;);
-  // expected-warning@+1{{OpenACC clause 'vector' not yet implemented}}
 #pragma acc loop vector independent
   for(;;);
   // expected-warning@+1{{OpenACC clause 'nohost' not yet implemented}}
@@ -591,9 +587,8 @@ void uses() {
   // expected-note@+1{{previous clause is here}}
 #pragma acc loop seq worker
   for(;;);
-  // expected-error@+3{{OpenACC clause 'vector' may not appear on the same construct as a 'seq' clause on a 'loop' construct}}
-  // expected-note@+2{{previous clause is here}}
-  // expected-warning@+1{{OpenACC clause 'vector' not yet implemented}}
+  // expected-error@+2{{OpenACC clause 'vector' may not appear on the same construct as a 'seq' clause on a 'loop' construct}}
+  // expected-note@+1{{previous clause is here}}
 #pragma acc loop seq vector
   for(;;);
   // expected-warning@+1{{OpenACC clause 'finalize' not yet implemented}}
@@ -733,10 +728,8 @@ void uses() {
   // expected-note@+1{{previous clause is here}}
 #pragma acc loop worker seq
   for(;;);
-  // TODO OpenACC: when 'vector' is implemented and makes it to the AST, this should diagnose because of a conflict with 'seq'.
-  // TODOexpected-error@+3{{OpenACC clause 'vector' may not appear on the same construct as a 'seq' clause on a 'loop' construct}}
-  // TODOexpected-note@+2{{previous clause is here}}
-  // expected-warning@+1{{OpenACC clause 'vector' not yet implemented}}
+  // expected-error@+2{{OpenACC clause 'seq' may not appear on the same construct as a 'vector' clause on a 'loop' construct}}
+  // expected-note@+1{{previous clause is here}}
 #pragma acc loop vector seq
   for(;;);
   // expected-warning@+1{{OpenACC clause 'finalize' not yet implemented}}
diff --git a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
index 51da8565f4e39..f60bf35a734fe 100644
--- a/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/loop-construct-device_type-clause.c
@@ -38,7 +38,6 @@ void uses() {
   // Only 'collapse', 'gang', 'worker', 'vector', 'seq', 'independent', 'auto',
   // and 'tile'  allowed after 'device_type'.
 
-  // expected-warning@+1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
 #pragma acc loop device_type(*) vector
   for(;;);
 
diff --git a/clang/test/SemaOpenACC/loop-construct-vector-ast.cpp b/clang/test/SemaOpenACC/loop-construct-vector-ast.cpp
new file mode 100644
index 0000000000000..390497cea0dd4
--- /dev/null
+++ b/clang/test/SemaOpenACC/loop-construct-vector-ast.cpp
@@ -0,0 +1,346 @@
+// 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
+
+template<unsigned I, typename ConvertsToInt, typename Int>
+void TemplUses(ConvertsToInt CTI, Int IsI) {
+  // CHECK: FunctionTemplateDecl{{.*}}TemplUses
+  // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} 'unsigned int' depth 0 index 0 I
+  // CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 1 ConvertsToInt
+  // CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 2 Int
+  // CHECK-NEXT: FunctionDecl{{.*}}TemplUses 'void (ConvertsToInt, Int)'
+  // CHECK-NEXT: ParmVarDecl{{.*}}CTI 'ConvertsToInt'
+  // CHECK-NEXT: ParmVarDecl{{.*}}IsI 'Int'
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+  // CHECK-NEXT: vector clause
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+#pragma acc loop vector
+  for(;;);
+
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+  // CHECK-NEXT: vector clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'unsigned int' NonTypeTemplateParm{{.*}}'I' 'unsigned int'
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+#pragma acc loop vector(I)
+  for(;;);
+
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+  // CHECK-NEXT: vector clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'ConvertsToInt' lvalue ParmVar{{.*}}'CTI' 'ConvertsToInt'
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+#pragma acc loop vector(length:CTI)
+  for(;;);
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} parallel
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: vector clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Int' lvalue ParmVar{{.*}}'IsI' 'Int'
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+#pragma acc parallel
+#pragma acc loop vector(length:IsI)
+  for(;;);
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} serial
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: vector clause
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+#pragma acc serial
+#pragma acc loop vector
+  for(;;);
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} kernels
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: vector clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Int' lvalue ParmVar{{.*}}'IsI' 'Int'
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+#pragma acc kernels
+#pragma acc loop vector(length:IsI)
+  for(;;);
+
+
+  // Instantiations:
+  // CHECK-NEXT: FunctionDecl{{.*}} TemplUses 'void (Converts, int)' implicit_instantiation
+  // CHECK-NEXT: TemplateArgument integral '3U'
+  // CHECK-NEXT: TemplateArgument type 'Converts'
+  // CHECK-NEXT: RecordType{{.*}}'Converts'
+  // CHECK-NEXT: CXXRecord{{.*}}'Converts
+  // CHECK-NEXT: TemplateArgument type 'int'
+  // CHECK-NEXT: BuiltinType{{.*}}'int'
+  // CHECK-NEXT: ParmVarDecl{{.*}} CTI 'Converts'
+  // CHECK-NEXT: ParmVarDecl{{.*}} IsI 'int'
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+  // CHECK-NEXT: vector clause
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+  //
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+  // CHECK-NEXT: vector clause
+  // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}}'unsigned int'
+  // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}}'unsigned int' depth 0 index 0 I
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'unsigned int' 3
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+  //
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+  // CHECK-NEXT: vector clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <UserDefinedConversion>
+  // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int'
+  // CHECK-NEXT: MemberExpr{{.*}} .operator int
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Converts' lvalue ParmVar{{.*}}'CTI' 'Converts'
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+  //
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} parallel
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: vector clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue ParmVar{{.*}}'IsI' 'int'
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+  //
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} serial
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: vector clause
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+  //
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} kernels
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: vector clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue ParmVar{{.*}}'IsI' 'int'
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+
+
+}
+
+struct Converts{
+  operator int();
+};
+
+void uses() {
+  // CHECK: FunctionDecl{{.*}} uses
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: CallExpr
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'void (*)(Converts, int)' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'void (Converts, int)' lvalue Function{{.*}} 'TemplUses' 'void (Converts, int)'
+  // CHECK-NEXT: CXXFunctionalCastExpr{{.*}} 'Converts' functional cast to Converts <NoOp>
+  // CHECK-NEXT: InitListExpr
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+  TemplUses<3>(Converts{}, 5);
+
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}}
+  int i;
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}}
+  // CHECK-NEXT: CXXConstructExpr
+  Converts C;
+
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+  // CHECK-NEXT: vector clause
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+#pragma acc loop vector
+  for(;;);
+
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+  // CHECK-NEXT: vector clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue Var
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+#pragma acc loop vector(i)
+  for(;;);
+
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}}<orphan>
+  // CHECK-NEXT: vector clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <UserDefinedConversion>
+  // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int'
+  // CHECK-NEXT: MemberExpr{{.*}} .operator int
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Converts' lvalue Var{{.*}}'C' 'Converts'
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+#pragma acc loop vector(length:C)
+  for(;;);
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} parallel
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: vector clause
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+#pragma acc parallel
+#pragma acc loop vector
+  for(;;);
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} parallel
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: vector clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <UserDefinedConversion>
+  // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int'
+  // CHECK-NEXT: MemberExpr{{.*}} .operator int
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Converts' lvalue Var{{.*}}'C' 'Converts'
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+#pragma acc parallel
+#pragma acc loop vector(C)
+  for(;;);
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} parallel
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: vector clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue Var
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+#pragma acc parallel
+#pragma acc loop vector(length:i)
+  for(;;);
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} kernels
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: vector clause
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+#pragma acc kernels
+#pragma acc loop vector
+  for(;;);
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} kernels
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: vector clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <UserDefinedConversion>
+  // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int'
+  // CHECK-NEXT: MemberExpr{{.*}} .operator int
+  // CHECK-NEXT: DeclRefExpr{{.*}}'Converts' lvalue Var{{.*}}'C' 'Converts'
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+#pragma acc kernels
+#pragma acc loop vector(C)
+  for(;;);
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} kernels
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: vector clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}}'int' lvalue Var
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+#pragma acc kernels
+#pragma acc loop vector(length:i)
+  for(;;);
+
+  // CHECK-NEXT: OpenACCComputeConstruct 0x[[COMPUTE_ADDR:[0-9a-f]+]]{{.*}} serial
+  // CHECK-NEXT: OpenACCLoopConstruct{{.*}} parent: 0x[[COMPUTE_ADDR]]
+  // CHECK-NEXT: vector clause
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: NullStmt
+#pragma acc serial
+#pragma acc loop vector
+  for(;;);
+}
+#endif // PCH_HELPER
diff --git a/clang/test/SemaOpenACC/loop-construct-vector-clause.cpp b/clang/test/SemaOpenACC/loop-construct-vector-clause.cpp
new file mode 100644
index 0000000000000..ed2d4ee8ee0e3
--- /dev/null
+++ b/clang/test/SemaOpenACC/loop-construct-vector-clause.cpp
@@ -0,0 +1,136 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+template<typename Int, typename NotInt, typename ConvertsToInt>
+void TemplUses(Int I, NotInt NI, ConvertsToInt CTI) {
+#pragma acc loop vector(I)
+  for(;;);
+
+#pragma acc parallel
+#pragma acc loop vector(length: I)
+  for(;;);
+
+#pragma acc kernels
+#pragma acc loop vector(CTI)
+  for(;;);
+
+  // expected-error@+2{{OpenACC clause 'vector' requires expression of integer type ('NoConvert' invalid)}}
+#pragma acc kernels
+#pragma acc loop vector(length: NI)
+  for(;;);
+
+  // expected-error@+2{{'num' argument on 'vector' clause is not permitted on a 'loop' construct associated with a 'serial' compute construct}}
+#pragma acc serial
+#pragma acc loop vector(length: I)
+  for(;;);
+
+  // expected-error@+3{{'num' argument to 'vector' clause not allowed on a 'loop' construct associated with a 'kernels' construct that has a 'vector_length' clause}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc kernels vector_length(I)
+#pragma acc loop vector(length: CTI)
+  for(;;);
+
+#pragma acc loop vector
+  for(;;) {
+    for(;;);
+    // expected-error@+2{{loop with a 'vector' clause may not exist in the region of a 'vector' clause}}
+    // expected-note@-4{{previous clause is here}}
+#pragma acc loop vector
+    for(;;);
+    for(;;);
+  }
+
+#pragma acc loop vector
+  for(;;) {
+    for(;;);
+    // expected-error@+4{{loop with a 'vector' clause may not exist in the region of a 'vector' clause}}
+    // expected-error@+3{{loop with a 'worker' clause may not exist in the region of a 'vector' clause}}
+    // expected-error@+2{{loop with a 'gang' clause may not exist in the region of a 'vector' clause}}
+    // expected-note@-6 3{{previous clause is here}}
+#pragma acc loop vector, worker, gang
+    for(;;);
+    for(;;);
+  }
+
+#pragma acc loop vector
+  for(;;) {
+#pragma acc serial
+#pragma acc loop vector
+    for(;;);
+  }
+}
+
+struct NoConvert{};
+struct Converts{
+  operator int();
+};
+
+void uses() {
+  TemplUses(5, NoConvert{}, Converts{}); // expected-note{{in instantiation of function template specialization}}
+
+  unsigned i;
+  NoConvert NI;
+  Converts CTI;
+
+#pragma acc loop vector(i)
+  for(;;);
+
+#pragma acc parallel
+#pragma acc loop vector(length: i)
+  for(;;);
+
+#pragma acc kernels
+#pragma acc loop vector(CTI)
+  for(;;);
+
+  // expected-error@+2{{OpenACC clause 'vector' requires expression of integer type ('NoConvert' invalid)}}
+#pragma acc kernels
+#pragma acc loop vector(length: NI)
+  for(;;);
+
+  // expected-error@+2{{'num' argument on 'vector' clause is not permitted on a 'loop' construct associated with a 'serial' compute construct}}
+#pragma acc serial
+#pragma acc loop vector(length: i)
+  for(;;);
+
+  // expected-error@+3{{'num' argument to 'vector' clause not allowed on a 'loop' construct associated with a 'kernels' construct that has a 'vector_length' clause}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc kernels vector_length(i)
+#pragma acc loop vector(length: i)
+  for(;;);
+
+#pragma acc loop vector
+  for(;;) {
+    for(;;);
+    // expected-error@+2{{loop with a 'vector' clause may not exist in the region of a 'vector' clause}}
+    // expected-note@-4{{previous clause is here}}
+#pragma acc loop vector
+    for(;;);
+    for(;;);
+  }
+
+#pragma acc loop vector
+  for(;;) {
+#pragma acc serial
+#pragma acc loop vector
+    for(;;);
+  }
+
+#pragma acc loop vector
+  for(;;) {
+    for(;;);
+    // expected-error@+4{{loop with a 'vector' clause may not exist in the region of a 'vector' clause}}
+    // expected-error@+3{{loop with a 'worker' clause may not exist in the region of a 'vector' clause}}
+    // expected-error@+2{{loop with a 'gang' clause may not exist in the region of a 'vector' clause}}
+    // expected-note@-6 3{{previous clause is here}}
+#pragma acc loop vector, worker, gang
+    for(;;);
+    for(;;);
+  }
+
+#pragma acc loop vector
+  for(;;) {
+#pragma acc serial
+#pragma acc loop vector, worker, gang
+    for(;;);
+  }
+}
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index 4461be86ea999..55dd3d7b69114 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2893,6 +2893,12 @@ void OpenACCClauseEnqueue::VisitWorkerClause(const OpenACCWorkerClause &C) {
   if (C.hasIntExpr())
     Visitor.AddStmt(C.getIntExpr());
 }
+
+void OpenACCClauseEnqueue::VisitVectorClause(const OpenACCVectorClause &C) {
+  if (C.hasIntExpr())
+    Visitor.AddStmt(C.getIntExpr());
+}
+
 void OpenACCClauseEnqueue::VisitWaitClause(const OpenACCWaitClause &C) {
   if (const Expr *DevNumExpr = C.getDevNumExpr())
     Visitor.AddStmt(DevNumExpr);