diff --git a/docs/analyzer/checkers.rst b/docs/analyzer/checkers.rst
index e5b60bbf15b..f50fff918cd 100644
--- a/docs/analyzer/checkers.rst
+++ b/docs/analyzer/checkers.rst
@@ -242,10 +242,35 @@ C++ Checkers.
 
 .. _cplusplus-InnerPointer:
 
-cplusplus.InnerPointer
-""""""""""""""""""""""
+cplusplus.InnerPointer (C++)
+""""""""""""""""""""""""""""
 Check for inner pointers of C++ containers used after re/deallocation.
 
+Many container methods in the C++ standard library are known to invalidate
+"references" (including actual references, iterators and raw pointers) to
+elements of the container. Using such references after they are invalidated
+causes undefined behavior, which is a common source of memory errors in C++ that
+this checker is capable of finding.
+
+The checker is currently limited to ``std::string`` objects and doesn't
+recognize some of the more sophisticated approaches to passing unowned pointers
+around, such as ``std::string_view``.
+
+.. code-block:: cpp
+
+ void deref_after_assignment() {
+   std::string s = "llvm";
+   const char *c = s.data(); // note: pointer to inner buffer of 'std::string' obtained here
+   s = "clang"; // note: inner buffer of 'std::string' reallocated by call to 'operator='
+   consume(c); // warn: inner pointer of container used after re/deallocation
+ }
+
+ const char *return_temp(int x) {
+   return std::to_string(x).c_str(); // warn: inner pointer of container used after re/deallocation
+   // note: pointer to inner buffer of 'std::string' obtained here
+   // note: inner buffer of 'std::string' deallocated by call to destructor
+ }
+
 .. _cplusplus-NewDelete:
 
 cplusplus.NewDelete (C++)
diff --git a/include/clang/AST/StmtOpenMP.h b/include/clang/AST/StmtOpenMP.h
index e37f5b1e000..c9efe323871 100644
--- a/include/clang/AST/StmtOpenMP.h
+++ b/include/clang/AST/StmtOpenMP.h
@@ -448,7 +448,8 @@ class OMPLoopDirective : public OMPExecutableDirective {
     PreInitsOffset = 8,
     // The '...End' enumerators do not correspond to child expressions - they
     // specify the offset to the end (and start of the following counters/
-    // updates/finals arrays).
+    // updates/finals/dependent_counters/dependent_inits/finals_conditions
+    // arrays).
     DefaultEnd = 9,
     // The following 8 exprs are used by worksharing and distribute loops only.
     IsLastIterVariableOffset = 9,
@@ -474,7 +475,8 @@ class OMPLoopDirective : public OMPExecutableDirective {
     CombinedNextUpperBoundOffset = 27,
     CombinedDistConditionOffset = 28,
     CombinedParForInDistConditionOffset = 29,
-    // Offset to the end (and start of the following counters/updates/finals
+    // Offset to the end (and start of the following
+    // counters/updates/finals/dependent_counters/dependent_inits/finals_conditions
     // arrays) for combined distribute loop directives.
     CombinedDistributeEnd = 30,
   };
@@ -517,6 +519,30 @@ class OMPLoopDirective : public OMPExecutableDirective {
     return MutableArrayRef<Expr *>(Storage, CollapsedNum);
   }
 
+  /// Get the dependent counters storage.
+  MutableArrayRef<Expr *> getDependentCounters() {
+    Expr **Storage = reinterpret_cast<Expr **>(
+        &*std::next(child_begin(),
+                    getArraysOffset(getDirectiveKind()) + 5 * CollapsedNum));
+    return MutableArrayRef<Expr *>(Storage, CollapsedNum);
+  }
+
+  /// Get the dependent inits storage.
+  MutableArrayRef<Expr *> getDependentInits() {
+    Expr **Storage = reinterpret_cast<Expr **>(
+        &*std::next(child_begin(),
+                    getArraysOffset(getDirectiveKind()) + 6 * CollapsedNum));
+    return MutableArrayRef<Expr *>(Storage, CollapsedNum);
+  }
+
+  /// Get the finals conditions storage.
+  MutableArrayRef<Expr *> getFinalsConditions() {
+    Expr **Storage = reinterpret_cast<Expr **>(
+        &*std::next(child_begin(),
+                    getArraysOffset(getDirectiveKind()) + 7 * CollapsedNum));
+    return MutableArrayRef<Expr *>(Storage, CollapsedNum);
+  }
+
 protected:
   /// Build instance of loop directive of class \a Kind.
   ///
@@ -551,9 +577,10 @@ class OMPLoopDirective : public OMPExecutableDirective {
   /// Children number.
   static unsigned numLoopChildren(unsigned CollapsedNum,
                                   OpenMPDirectiveKind Kind) {
-    return getArraysOffset(Kind) + 5 * CollapsedNum; // Counters,
-                                                     // PrivateCounters, Inits,
-                                                     // Updates and Finals
+    return getArraysOffset(Kind) +
+           8 * CollapsedNum; // Counters, PrivateCounters, Inits,
+                             // Updates, Finals, DependentCounters,
+                             // DependentInits, FinalsConditions.
   }
 
   void setIterationVariable(Expr *IV) {
@@ -703,6 +730,9 @@ class OMPLoopDirective : public OMPExecutableDirective {
   void setInits(ArrayRef<Expr *> A);
   void setUpdates(ArrayRef<Expr *> A);
   void setFinals(ArrayRef<Expr *> A);
+  void setDependentCounters(ArrayRef<Expr *> A);
+  void setDependentInits(ArrayRef<Expr *> A);
+  void setFinalsConditions(ArrayRef<Expr *> A);
 
 public:
   /// The expressions built to support OpenMP loops in combined/composite
@@ -798,6 +828,15 @@ class OMPLoopDirective : public OMPExecutableDirective {
     SmallVector<Expr *, 4> Updates;
     /// Final loop counter values for GodeGen.
     SmallVector<Expr *, 4> Finals;
+    /// List of counters required for the generation of the non-rectangular
+    /// loops.
+    SmallVector<Expr *, 4> DependentCounters;
+    /// List of initializers required for the generation of the non-rectangular
+    /// loops.
+    SmallVector<Expr *, 4> DependentInits;
+    /// List of final conditions required for the generation of the
+    /// non-rectangular loops.
+    SmallVector<Expr *, 4> FinalsConditions;
     /// Init statement for all captured expressions.
     Stmt *PreInits;
 
@@ -813,7 +852,9 @@ class OMPLoopDirective : public OMPExecutableDirective {
     }
 
     /// Initialize all the fields to null.
-    /// \param Size Number of elements in the counters/finals/updates arrays.
+    /// \param Size Number of elements in the
+    /// counters/finals/updates/dependent_counters/dependent_inits/finals_conditions
+    /// arrays.
     void clear(unsigned Size) {
       IterationVarRef = nullptr;
       LastIteration = nullptr;
@@ -839,12 +880,18 @@ class OMPLoopDirective : public OMPExecutableDirective {
       Inits.resize(Size);
       Updates.resize(Size);
       Finals.resize(Size);
+      DependentCounters.resize(Size);
+      DependentInits.resize(Size);
+      FinalsConditions.resize(Size);
       for (unsigned i = 0; i < Size; ++i) {
         Counters[i] = nullptr;
         PrivateCounters[i] = nullptr;
         Inits[i] = nullptr;
         Updates[i] = nullptr;
         Finals[i] = nullptr;
+        DependentCounters[i] = nullptr;
+        DependentInits[i] = nullptr;
+        FinalsConditions[i] = nullptr;
       }
       PreInits = nullptr;
       DistCombinedFields.LB = nullptr;
@@ -1078,6 +1125,24 @@ class OMPLoopDirective : public OMPExecutableDirective {
     return const_cast<OMPLoopDirective *>(this)->getFinals();
   }
 
+  ArrayRef<Expr *> dependent_counters() { return getDependentCounters(); }
+
+  ArrayRef<Expr *> dependent_counters() const {
+    return const_cast<OMPLoopDirective *>(this)->getDependentCounters();
+  }
+
+  ArrayRef<Expr *> dependent_inits() { return getDependentInits(); }
+
+  ArrayRef<Expr *> dependent_inits() const {
+    return const_cast<OMPLoopDirective *>(this)->getDependentInits();
+  }
+
+  ArrayRef<Expr *> finals_conditions() { return getFinalsConditions(); }
+
+  ArrayRef<Expr *> finals_conditions() const {
+    return const_cast<OMPLoopDirective *>(this)->getFinalsConditions();
+  }
+
   static bool classof(const Stmt *T) {
     return T->getStmtClass() == OMPSimdDirectiveClass ||
            T->getStmtClass() == OMPForDirectiveClass ||
diff --git a/include/clang/Driver/CC1Options.td b/include/clang/Driver/CC1Options.td
index 4335207fc4b..e2029d181d2 100644
--- a/include/clang/Driver/CC1Options.td
+++ b/include/clang/Driver/CC1Options.td
@@ -140,7 +140,8 @@ def analyzer_checker_help_developer : Flag<["-"], "analyzer-checker-help-develop
            "and debug checkers">;
 
 def analyzer_config_help : Flag<["-"], "analyzer-config-help">,
-  HelpText<"Display the list of -analyzer-config options">;
+  HelpText<"Display the list of -analyzer-config options. These are meant for "
+           "development purposes only!">;
 
 def analyzer_list_enabled_checkers : Flag<["-"], "analyzer-list-enabled-checkers">,
   HelpText<"Display the list of enabled analyzer checkers">;
diff --git a/include/clang/StaticAnalyzer/Core/BugReporter/BugReporter.h b/include/clang/StaticAnalyzer/Core/BugReporter/BugReporter.h
index c3d7ba3120f..c2a5bb7f629 100644
--- a/include/clang/StaticAnalyzer/Core/BugReporter/BugReporter.h
+++ b/include/clang/StaticAnalyzer/Core/BugReporter/BugReporter.h
@@ -180,10 +180,12 @@ class BugReport : public llvm::ilist_node<BugReport> {
   /// to the user. This method allows to rest the location which should be used
   /// for uniquing reports. For example, memory leaks checker, could set this to
   /// the allocation site, rather then the location where the bug is reported.
-  BugReport(BugType& bt, StringRef desc, const ExplodedNode *errornode,
+  BugReport(BugType &bt, StringRef desc, const ExplodedNode *errornode,
             PathDiagnosticLocation LocationToUnique, const Decl *DeclToUnique)
       : BT(bt), Description(desc), UniqueingLocation(LocationToUnique),
-        UniqueingDecl(DeclToUnique), ErrorNode(errornode) {}
+        UniqueingDecl(DeclToUnique), ErrorNode(errornode),
+        ErrorNodeRange(getStmt() ? getStmt()->getSourceRange()
+                                 : SourceRange()) {}
 
   virtual ~BugReport() = default;
 
diff --git a/include/clang/StaticAnalyzer/Core/BugReporter/BugReporterVisitors.h b/include/clang/StaticAnalyzer/Core/BugReporter/BugReporterVisitors.h
index 9069270a38e..8488a889763 100644
--- a/include/clang/StaticAnalyzer/Core/BugReporter/BugReporterVisitors.h
+++ b/include/clang/StaticAnalyzer/Core/BugReporter/BugReporterVisitors.h
@@ -143,8 +143,6 @@ class FindLastStoreBRVisitor final : public BugReporterVisitor {
 
   /// \param V We're searching for the store where \c R received this value.
   /// \param R The region we're tracking.
-  /// \param EnableNullFPSuppression Whether we should employ false positive
-  ///         suppression (inlined defensive checks, returned null).
   /// \param TKind May limit the amount of notes added to the bug report.
   /// \param OriginSFC Only adds notes when the last store happened in a
   ///        different stackframe to this one. Disregarded if the tracking kind
diff --git a/lib/AST/StmtOpenMP.cpp b/lib/AST/StmtOpenMP.cpp
index 4e829897ceb..7fda574bae0 100644
--- a/lib/AST/StmtOpenMP.cpp
+++ b/lib/AST/StmtOpenMP.cpp
@@ -72,6 +72,25 @@ void OMPLoopDirective::setFinals(ArrayRef<Expr *> A) {
   std::copy(A.begin(), A.end(), getFinals().begin());
 }
 
+void OMPLoopDirective::setDependentCounters(ArrayRef<Expr *> A) {
+  assert(
+      A.size() == getCollapsedNumber() &&
+      "Number of dependent counters is not the same as the collapsed number");
+  llvm::copy(A, getDependentCounters().begin());
+}
+
+void OMPLoopDirective::setDependentInits(ArrayRef<Expr *> A) {
+  assert(A.size() == getCollapsedNumber() &&
+         "Number of dependent inits is not the same as the collapsed number");
+  llvm::copy(A, getDependentInits().begin());
+}
+
+void OMPLoopDirective::setFinalsConditions(ArrayRef<Expr *> A) {
+  assert(A.size() == getCollapsedNumber() &&
+         "Number of finals conditions is not the same as the collapsed number");
+  llvm::copy(A, getFinalsConditions().begin());
+}
+
 OMPParallelDirective *OMPParallelDirective::Create(
     const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
     ArrayRef<OMPClause *> Clauses, Stmt *AssociatedStmt, bool HasCancel) {
@@ -122,6 +141,9 @@ OMPSimdDirective::Create(const ASTContext &C, SourceLocation StartLoc,
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
+  Dir->setDependentCounters(Exprs.DependentCounters);
+  Dir->setDependentInits(Exprs.DependentInits);
+  Dir->setFinalsConditions(Exprs.FinalsConditions);
   Dir->setPreInits(Exprs.PreInits);
   return Dir;
 }
@@ -170,6 +192,9 @@ OMPForDirective::Create(const ASTContext &C, SourceLocation StartLoc,
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
+  Dir->setDependentCounters(Exprs.DependentCounters);
+  Dir->setDependentInits(Exprs.DependentInits);
+  Dir->setFinalsConditions(Exprs.FinalsConditions);
   Dir->setPreInits(Exprs.PreInits);
   Dir->setHasCancel(HasCancel);
   return Dir;
@@ -220,6 +245,9 @@ OMPForSimdDirective::Create(const ASTContext &C, SourceLocation StartLoc,
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
+  Dir->setDependentCounters(Exprs.DependentCounters);
+  Dir->setDependentInits(Exprs.DependentInits);
+  Dir->setFinalsConditions(Exprs.FinalsConditions);
   Dir->setPreInits(Exprs.PreInits);
   return Dir;
 }
@@ -383,6 +411,9 @@ OMPParallelForDirective *OMPParallelForDirective::Create(
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
+  Dir->setDependentCounters(Exprs.DependentCounters);
+  Dir->setDependentInits(Exprs.DependentInits);
+  Dir->setFinalsConditions(Exprs.FinalsConditions);
   Dir->setPreInits(Exprs.PreInits);
   Dir->setHasCancel(HasCancel);
   return Dir;
@@ -432,6 +463,9 @@ OMPParallelForSimdDirective *OMPParallelForSimdDirective::Create(
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
+  Dir->setDependentCounters(Exprs.DependentCounters);
+  Dir->setDependentInits(Exprs.DependentInits);
+  Dir->setFinalsConditions(Exprs.FinalsConditions);
   Dir->setPreInits(Exprs.PreInits);
   return Dir;
 }
@@ -772,6 +806,9 @@ OMPTargetParallelForDirective *OMPTargetParallelForDirective::Create(
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
+  Dir->setDependentCounters(Exprs.DependentCounters);
+  Dir->setDependentInits(Exprs.DependentInits);
+  Dir->setFinalsConditions(Exprs.FinalsConditions);
   Dir->setPreInits(Exprs.PreInits);
   Dir->setHasCancel(HasCancel);
   return Dir;
@@ -914,6 +951,9 @@ OMPTaskLoopDirective *OMPTaskLoopDirective::Create(
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
+  Dir->setDependentCounters(Exprs.DependentCounters);
+  Dir->setDependentInits(Exprs.DependentInits);
+  Dir->setFinalsConditions(Exprs.FinalsConditions);
   Dir->setPreInits(Exprs.PreInits);
   return Dir;
 }
@@ -963,6 +1003,9 @@ OMPTaskLoopSimdDirective *OMPTaskLoopSimdDirective::Create(
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
+  Dir->setDependentCounters(Exprs.DependentCounters);
+  Dir->setDependentInits(Exprs.DependentInits);
+  Dir->setFinalsConditions(Exprs.FinalsConditions);
   Dir->setPreInits(Exprs.PreInits);
   return Dir;
 }
@@ -1011,6 +1054,9 @@ OMPDistributeDirective *OMPDistributeDirective::Create(
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
+  Dir->setDependentCounters(Exprs.DependentCounters);
+  Dir->setDependentInits(Exprs.DependentInits);
+  Dir->setFinalsConditions(Exprs.FinalsConditions);
   Dir->setPreInits(Exprs.PreInits);
   return Dir;
 }
@@ -1089,6 +1135,9 @@ OMPDistributeParallelForDirective *OMPDistributeParallelForDirective::Create(
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
+  Dir->setDependentCounters(Exprs.DependentCounters);
+  Dir->setDependentInits(Exprs.DependentInits);
+  Dir->setFinalsConditions(Exprs.FinalsConditions);
   Dir->setPreInits(Exprs.PreInits);
   Dir->setCombinedLowerBoundVariable(Exprs.DistCombinedFields.LB);
   Dir->setCombinedUpperBoundVariable(Exprs.DistCombinedFields.UB);
@@ -1157,6 +1206,9 @@ OMPDistributeParallelForSimdDirective::Create(
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
+  Dir->setDependentCounters(Exprs.DependentCounters);
+  Dir->setDependentInits(Exprs.DependentInits);
+  Dir->setFinalsConditions(Exprs.FinalsConditions);
   Dir->setPreInits(Exprs.PreInits);
   Dir->setCombinedLowerBoundVariable(Exprs.DistCombinedFields.LB);
   Dir->setCombinedUpperBoundVariable(Exprs.DistCombinedFields.UB);
@@ -1219,6 +1271,9 @@ OMPDistributeSimdDirective *OMPDistributeSimdDirective::Create(
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
+  Dir->setDependentCounters(Exprs.DependentCounters);
+  Dir->setDependentInits(Exprs.DependentInits);
+  Dir->setFinalsConditions(Exprs.FinalsConditions);
   Dir->setPreInits(Exprs.PreInits);
   return Dir;
 }
@@ -1271,6 +1326,9 @@ OMPTargetParallelForSimdDirective *OMPTargetParallelForSimdDirective::Create(
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
+  Dir->setDependentCounters(Exprs.DependentCounters);
+  Dir->setDependentInits(Exprs.DependentInits);
+  Dir->setFinalsConditions(Exprs.FinalsConditions);
   Dir->setPreInits(Exprs.PreInits);
   return Dir;
 }
@@ -1315,6 +1373,9 @@ OMPTargetSimdDirective::Create(const ASTContext &C, SourceLocation StartLoc,
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
+  Dir->setDependentCounters(Exprs.DependentCounters);
+  Dir->setDependentInits(Exprs.DependentInits);
+  Dir->setFinalsConditions(Exprs.FinalsConditions);
   Dir->setPreInits(Exprs.PreInits);
   return Dir;
 }
@@ -1363,6 +1424,9 @@ OMPTeamsDistributeDirective *OMPTeamsDistributeDirective::Create(
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
+  Dir->setDependentCounters(Exprs.DependentCounters);
+  Dir->setDependentInits(Exprs.DependentInits);
+  Dir->setFinalsConditions(Exprs.FinalsConditions);
   Dir->setPreInits(Exprs.PreInits);
   return Dir;
 }
@@ -1414,6 +1478,9 @@ OMPTeamsDistributeSimdDirective *OMPTeamsDistributeSimdDirective::Create(
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
+  Dir->setDependentCounters(Exprs.DependentCounters);
+  Dir->setDependentInits(Exprs.DependentInits);
+  Dir->setFinalsConditions(Exprs.FinalsConditions);
   Dir->setPreInits(Exprs.PreInits);
   return Dir;
 }
@@ -1471,6 +1538,9 @@ OMPTeamsDistributeParallelForSimdDirective::Create(
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
+  Dir->setDependentCounters(Exprs.DependentCounters);
+  Dir->setDependentInits(Exprs.DependentInits);
+  Dir->setFinalsConditions(Exprs.FinalsConditions);
   Dir->setPreInits(Exprs.PreInits);
   Dir->setCombinedLowerBoundVariable(Exprs.DistCombinedFields.LB);
   Dir->setCombinedUpperBoundVariable(Exprs.DistCombinedFields.UB);
@@ -1540,6 +1610,9 @@ OMPTeamsDistributeParallelForDirective::Create(
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
+  Dir->setDependentCounters(Exprs.DependentCounters);
+  Dir->setDependentInits(Exprs.DependentInits);
+  Dir->setFinalsConditions(Exprs.FinalsConditions);
   Dir->setPreInits(Exprs.PreInits);
   Dir->setCombinedLowerBoundVariable(Exprs.DistCombinedFields.LB);
   Dir->setCombinedUpperBoundVariable(Exprs.DistCombinedFields.UB);
@@ -1628,6 +1701,9 @@ OMPTargetTeamsDistributeDirective *OMPTargetTeamsDistributeDirective::Create(
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
+  Dir->setDependentCounters(Exprs.DependentCounters);
+  Dir->setDependentInits(Exprs.DependentInits);
+  Dir->setFinalsConditions(Exprs.FinalsConditions);
   Dir->setPreInits(Exprs.PreInits);
   return Dir;
 }
@@ -1688,6 +1764,9 @@ OMPTargetTeamsDistributeParallelForDirective::Create(
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
+  Dir->setDependentCounters(Exprs.DependentCounters);
+  Dir->setDependentInits(Exprs.DependentInits);
+  Dir->setFinalsConditions(Exprs.FinalsConditions);
   Dir->setPreInits(Exprs.PreInits);
   Dir->setCombinedLowerBoundVariable(Exprs.DistCombinedFields.LB);
   Dir->setCombinedUpperBoundVariable(Exprs.DistCombinedFields.UB);
@@ -1761,6 +1840,9 @@ OMPTargetTeamsDistributeParallelForSimdDirective::Create(
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
+  Dir->setDependentCounters(Exprs.DependentCounters);
+  Dir->setDependentInits(Exprs.DependentInits);
+  Dir->setFinalsConditions(Exprs.FinalsConditions);
   Dir->setPreInits(Exprs.PreInits);
   Dir->setCombinedLowerBoundVariable(Exprs.DistCombinedFields.LB);
   Dir->setCombinedUpperBoundVariable(Exprs.DistCombinedFields.UB);
@@ -1826,6 +1908,9 @@ OMPTargetTeamsDistributeSimdDirective::Create(
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
+  Dir->setDependentCounters(Exprs.DependentCounters);
+  Dir->setDependentInits(Exprs.DependentInits);
+  Dir->setFinalsConditions(Exprs.FinalsConditions);
   Dir->setPreInits(Exprs.PreInits);
   return Dir;
 }
diff --git a/lib/Basic/Targets/AMDGPU.cpp b/lib/Basic/Targets/AMDGPU.cpp
index f2d49fdc7e1..cc9791db499 100644
--- a/lib/Basic/Targets/AMDGPU.cpp
+++ b/lib/Basic/Targets/AMDGPU.cpp
@@ -138,9 +138,6 @@ bool AMDGPUTargetInfo::initFeatureMap(
 
   // XXX - What does the member GPU mean if device name string passed here?
   if (isAMDGCN(getTriple())) {
-    if (CPU.empty())
-      CPU = "gfx600";
-
     switch (llvm::AMDGPU::parseArchAMDGCN(CPU)) {
     case GK_GFX1012:
     case GK_GFX1011:
@@ -196,7 +193,7 @@ bool AMDGPUTargetInfo::initFeatureMap(
     case GK_GFX600:
       break;
     case GK_NONE:
-      return false;
+      break;
     default:
       llvm_unreachable("Unhandled GPU!");
     }
diff --git a/lib/CodeGen/CGLoopInfo.cpp b/lib/CodeGen/CGLoopInfo.cpp
index 8e4a0a2a9b5..c51efdc5d11 100644
--- a/lib/CodeGen/CGLoopInfo.cpp
+++ b/lib/CodeGen/CGLoopInfo.cpp
@@ -253,12 +253,18 @@ LoopInfo::createLoopVectorizeMetadata(const LoopAttributes &Attrs,
   Args.append(LoopProperties.begin(), LoopProperties.end());
 
   // Setting vectorize.predicate
-  if (Attrs.VectorizePredicateEnable != LoopAttributes::Unspecified) {
+  bool IsVectorPredicateEnabled = false;
+  if (Attrs.VectorizePredicateEnable != LoopAttributes::Unspecified &&
+      Attrs.VectorizeEnable != LoopAttributes::Disable &&
+      Attrs.VectorizeWidth < 1) {
+
+    IsVectorPredicateEnabled =
+        (Attrs.VectorizePredicateEnable == LoopAttributes::Enable);
+
     Metadata *Vals[] = {
         MDString::get(Ctx, "llvm.loop.vectorize.predicate.enable"),
-        ConstantAsMetadata::get(ConstantInt::get(
-            llvm::Type::getInt1Ty(Ctx),
-            (Attrs.VectorizePredicateEnable == LoopAttributes::Enable)))};
+        ConstantAsMetadata::get(ConstantInt::get(llvm::Type::getInt1Ty(Ctx),
+                                                 IsVectorPredicateEnabled))};
     Args.push_back(MDNode::get(Ctx, Vals));
   }
 
@@ -281,12 +287,15 @@ LoopInfo::createLoopVectorizeMetadata(const LoopAttributes &Attrs,
   }
 
   // Setting vectorize.enable
-  if (Attrs.VectorizeEnable != LoopAttributes::Unspecified) {
+  if (Attrs.VectorizeEnable != LoopAttributes::Unspecified ||
+      IsVectorPredicateEnabled) {
     Metadata *Vals[] = {
         MDString::get(Ctx, "llvm.loop.vectorize.enable"),
         ConstantAsMetadata::get(ConstantInt::get(
             llvm::Type::getInt1Ty(Ctx),
-            (Attrs.VectorizeEnable == LoopAttributes::Enable)))};
+            IsVectorPredicateEnabled
+                ? true
+                : (Attrs.VectorizeEnable == LoopAttributes::Enable)))};
     Args.push_back(MDNode::get(Ctx, Vals));
   }
 
diff --git a/lib/CodeGen/CGStmtOpenMP.cpp b/lib/CodeGen/CGStmtOpenMP.cpp
index e8fbca5108a..2ca7ebc6b0d 100644
--- a/lib/CodeGen/CGStmtOpenMP.cpp
+++ b/lib/CodeGen/CGStmtOpenMP.cpp
@@ -1324,6 +1324,16 @@ void CodeGenFunction::EmitOMPLoopBody(const OMPLoopDirective &D,
   // On a continue in the body, jump to the end.
   JumpDest Continue = getJumpDestInCurrentScope("omp.body.continue");
   BreakContinueStack.push_back(BreakContinue(LoopExit, Continue));
+  for (const Expr *E : D.finals_conditions()) {
+    if (!E)
+      continue;
+    // Check that loop counter in non-rectangular nest fits into the iteration
+    // space.
+    llvm::BasicBlock *NextBB = createBasicBlock("omp.body.next");
+    EmitBranchOnBoolExpr(E, NextBB, Continue.getBlock(),
+                         getProfileCount(D.getBody()));
+    EmitBlock(NextBB);
+  }
   // Emit loop body.
   EmitStmt(D.getBody());
   // The end (updates/cleanups).
@@ -1553,8 +1563,28 @@ static void emitPreCond(CodeGenFunction &CGF, const OMPLoopDirective &S,
       CGF.EmitIgnoredExpr(I);
     }
   }
+  // Create temp loop control variables with their init values to support
+  // non-rectangular loops.
+  CodeGenFunction::OMPMapVars PreCondVars;
+  for (const Expr * E: S.dependent_counters()) {
+    if (!E)
+      continue;
+    assert(!E->getType().getNonReferenceType()->isRecordType() &&
+           "dependent counter must not be an iterator.");
+    const auto *VD = cast<VarDecl>(cast<DeclRefExpr>(E)->getDecl());
+    Address CounterAddr =
+        CGF.CreateMemTemp(VD->getType().getNonReferenceType());
+    (void)PreCondVars.setVarAddr(CGF, VD, CounterAddr);
+  }
+  (void)PreCondVars.apply(CGF);
+  for (const Expr *E : S.dependent_inits()) {
+    if (!E)
+      continue;
+    CGF.EmitIgnoredExpr(E);
+  }
   // Check that loop is executed at least one time.
   CGF.EmitBranchOnBoolExpr(Cond, TrueBlock, FalseBlock, TrueCount);
+  PreCondVars.restore(CGF);
 }
 
 void CodeGenFunction::EmitOMPLinearClause(
diff --git a/lib/Parse/ParsePragma.cpp b/lib/Parse/ParsePragma.cpp
index f63ab3a744f..a386a4d71c3 100644
--- a/lib/Parse/ParsePragma.cpp
+++ b/lib/Parse/ParsePragma.cpp
@@ -1006,18 +1006,13 @@ struct PragmaLoopHintInfo {
 } // end anonymous namespace
 
 static std::string PragmaLoopHintString(Token PragmaName, Token Option) {
-  std::string PragmaString;
-  if (PragmaName.getIdentifierInfo()->getName() == "loop") {
-    PragmaString = "clang loop ";
-    PragmaString += Option.getIdentifierInfo()->getName();
-  } else if (PragmaName.getIdentifierInfo()->getName() == "unroll_and_jam") {
-    PragmaString = "unroll_and_jam";
-  } else {
-    assert(PragmaName.getIdentifierInfo()->getName() == "unroll" &&
-           "Unexpected pragma name");
-    PragmaString = "unroll";
-  }
-  return PragmaString;
+  StringRef Str = PragmaName.getIdentifierInfo()->getName();
+  std::string ClangLoopStr = (llvm::Twine("clang loop ") + Str).str();
+  return llvm::StringSwitch<StringRef>(Str)
+      .Case("loop", ClangLoopStr)
+      .Case("unroll_and_jam", Str)
+      .Case("unroll", Str)
+      .Default("");
 }
 
 bool Parser::HandlePragmaLoopHint(LoopHint &Hint) {
@@ -1041,12 +1036,12 @@ bool Parser::HandlePragmaLoopHint(LoopHint &Hint) {
 
   // Return a valid hint if pragma unroll or nounroll were specified
   // without an argument.
-  bool PragmaUnroll = PragmaNameInfo->getName() == "unroll";
-  bool PragmaNoUnroll = PragmaNameInfo->getName() == "nounroll";
-  bool PragmaUnrollAndJam = PragmaNameInfo->getName() == "unroll_and_jam";
-  bool PragmaNoUnrollAndJam = PragmaNameInfo->getName() == "nounroll_and_jam";
-  if (Toks.empty() && (PragmaUnroll || PragmaNoUnroll || PragmaUnrollAndJam ||
-                       PragmaNoUnrollAndJam)) {
+  auto IsLoopHint = llvm::StringSwitch<bool>(PragmaNameInfo->getName())
+                        .Cases("unroll", "nounroll", "unroll_and_jam",
+                               "nounroll_and_jam", true)
+                        .Default(false);
+
+  if (Toks.empty() && IsLoopHint) {
     ConsumeAnnotationToken();
     Hint.Range = Info->PragmaName.getLocation();
     return true;
diff --git a/lib/Sema/SemaExprCXX.cpp b/lib/Sema/SemaExprCXX.cpp
index 355459cd514..fbae19a2490 100644
--- a/lib/Sema/SemaExprCXX.cpp
+++ b/lib/Sema/SemaExprCXX.cpp
@@ -6800,14 +6800,10 @@ ExprResult Sema::ActOnStartCXXMemberReference(Scope *S, Expr *Base,
   // it's legal for the type to be incomplete if this is a pseudo-destructor
   // call.  We'll do more incomplete-type checks later in the lookup process,
   // so just skip this check for ObjC types.
-  if (BaseType->isObjCObjectOrInterfaceType()) {
+  if (!BaseType->isRecordType()) {
     ObjectType = ParsedType::make(BaseType);
     MayBePseudoDestructor = true;
     return Base;
-  } else if (!BaseType->isRecordType()) {
-    ObjectType = nullptr;
-    MayBePseudoDestructor = true;
-    return Base;
   }
 
   // The object type must be complete (or dependent), or
diff --git a/lib/Sema/SemaInit.cpp b/lib/Sema/SemaInit.cpp
index c7c7c4dfdb7..dea061a30ee 100644
--- a/lib/Sema/SemaInit.cpp
+++ b/lib/Sema/SemaInit.cpp
@@ -7360,7 +7360,7 @@ static bool shouldTrackImplicitObjectArg(const CXXMethodDecl *Callee) {
              OO == OverloadedOperatorKind::OO_Star;
     }
     return llvm::StringSwitch<bool>(Callee->getName())
-        .Cases("front", "back", "at", true)
+        .Cases("front", "back", "at", "top", "value", true)
         .Default(false);
   }
   return false;
diff --git a/lib/Sema/SemaOpenMP.cpp b/lib/Sema/SemaOpenMP.cpp
index 49c7c3aca14..3c23e43a20a 100644
--- a/lib/Sema/SemaOpenMP.cpp
+++ b/lib/Sema/SemaOpenMP.cpp
@@ -4709,6 +4709,54 @@ StmtResult Sema::ActOnOpenMPParallelDirective(ArrayRef<OMPClause *> Clauses,
 }
 
 namespace {
+/// Iteration space of a single for loop.
+struct LoopIterationSpace final {
+  /// True if the condition operator is the strict compare operator (<, > or
+  /// !=).
+  bool IsStrictCompare = false;
+  /// Condition of the loop.
+  Expr *PreCond = nullptr;
+  /// This expression calculates the number of iterations in the loop.
+  /// It is always possible to calculate it before starting the loop.
+  Expr *NumIterations = nullptr;
+  /// The loop counter variable.
+  Expr *CounterVar = nullptr;
+  /// Private loop counter variable.
+  Expr *PrivateCounterVar = nullptr;
+  /// This is initializer for the initial value of #CounterVar.
+  Expr *CounterInit = nullptr;
+  /// This is step for the #CounterVar used to generate its update:
+  /// #CounterVar = #CounterInit + #CounterStep * CurrentIteration.
+  Expr *CounterStep = nullptr;
+  /// Should step be subtracted?
+  bool Subtract = false;
+  /// Source range of the loop init.
+  SourceRange InitSrcRange;
+  /// Source range of the loop condition.
+  SourceRange CondSrcRange;
+  /// Source range of the loop increment.
+  SourceRange IncSrcRange;
+  /// Minimum value that can have the loop control variable. Used to support
+  /// non-rectangular loops. Applied only for LCV with the non-iterator types,
+  /// since only such variables can be used in non-loop invariant expressions.
+  Expr *MinValue = nullptr;
+  /// Maximum value that can have the loop control variable. Used to support
+  /// non-rectangular loops. Applied only for LCV with the non-iterator type,
+  /// since only such variables can be used in non-loop invariant expressions.
+  Expr *MaxValue = nullptr;
+  /// true, if the lower bound depends on the outer loop control var.
+  bool IsNonRectangularLB = false;
+  /// true, if the upper bound depends on the outer loop control var.
+  bool IsNonRectangularUB = false;
+  /// Index of the loop this loop depends on and forms non-rectangular loop
+  /// nest.
+  unsigned LoopDependentIdx = 0;
+  /// Final condition for the non-rectangular loop nest support. It is used to
+  /// check that the number of iterations for this particular counter must be
+  /// finished.
+  Expr *FinalCondition = nullptr;
+};
+
 /// Helper class for checking canonical form of the OpenMP loops and
 /// extracting iteration space of each loop in the loop nest, that will be used
 /// for IR generation.
@@ -4758,6 +4806,9 @@ class OpenMPIterationSpaceChecker {
   Optional<unsigned> CondDependOnLC;
   /// Checks if the provide statement depends on the loop counter.
   Optional<unsigned> doesDependOnLoopCounter(const Stmt *S, bool IsInitializer);
+  /// Original condition required for checking of the exit condition for
+  /// non-rectangular loop.
+  Expr *Condition = nullptr;
 
 public:
   OpenMPIterationSpaceChecker(Sema &SemaRef, DSAStackTy &Stack,
@@ -4789,7 +4840,7 @@ class OpenMPIterationSpaceChecker {
   bool isStrictTestOp() const { return TestIsStrictOp; }
   /// Build the expression to calculate the number of iterations.
   Expr *buildNumIterations(
-      Scope *S, const bool LimitedType,
+      Scope *S, ArrayRef<LoopIterationSpace> ResultIterSpaces, bool LimitedType,
       llvm::MapVector<const Expr *, DeclRefExpr *> &Captures) const;
   /// Build the precondition expression for the loops.
   Expr *
@@ -4813,8 +4864,21 @@ class OpenMPIterationSpaceChecker {
                        llvm::MapVector<const Expr *, DeclRefExpr *> &Captures,
                        SourceLocation Loc, Expr *Inc = nullptr,
                        OverloadedOperatorKind OOK = OO_Amp);
+  /// Builds the minimum value for the loop counter.
+  std::pair<Expr *, Expr *> buildMinMaxValues(
+      Scope *S, llvm::MapVector<const Expr *, DeclRefExpr *> &Captures) const;
+  /// Builds final condition for the non-rectangular loops.
+  Expr *buildFinalCondition(Scope *S) const;
   /// Return true if any expression is dependent.
   bool dependent() const;
+  /// Returns true if the initializer forms non-rectangular loop.
+  bool doesInitDependOnLC() const { return InitDependOnLC.hasValue(); }
+  /// Returns true if the condition forms non-rectangular loop.
+  bool doesCondDependOnLC() const { return CondDependOnLC.hasValue(); }
+  /// Returns index of the loop we depend on (starting from 1), or 0 otherwise.
+  unsigned getLoopDependentIdx() const {
+    return InitDependOnLC.getValueOr(CondDependOnLC.getValueOr(0));
+  }
 
 private:
   /// Check the right-hand side of an assignment in the increment
@@ -5013,9 +5077,9 @@ class LoopCounterRefChecker final
     return false;
   }
   bool VisitStmt(const Stmt *S) {
-    bool Res = true;
+    bool Res = false;
     for (const Stmt *Child : S->children())
-      Res = Child && Visit(Child) && Res;
+      Res = (Child && Visit(Child)) || Res;
     return Res;
   }
   explicit LoopCounterRefChecker(Sema &SemaRef, DSAStackTy &Stack,
@@ -5165,6 +5229,7 @@ bool OpenMPIterationSpaceChecker::checkAndSetCond(Expr *S) {
     SemaRef.Diag(DefaultLoc, diag::err_omp_loop_not_canonical_cond) << LCDecl;
     return true;
   }
+  Condition = S;
   S = getExprAsWritten(S);
   SourceLocation CondLoc = S->getBeginLoc();
   if (auto *BO = dyn_cast<BinaryOperator>(S)) {
@@ -5351,15 +5416,177 @@ tryBuildCapture(Sema &SemaRef, Expr *Capture,
 
 /// Build the expression to calculate the number of iterations.
 Expr *OpenMPIterationSpaceChecker::buildNumIterations(
-    Scope *S, const bool LimitedType,
+    Scope *S, ArrayRef<LoopIterationSpace> ResultIterSpaces, bool LimitedType,
     llvm::MapVector<const Expr *, DeclRefExpr *> &Captures) const {
   ExprResult Diff;
   QualType VarType = LCDecl->getType().getNonReferenceType();
   if (VarType->isIntegerType() || VarType->isPointerType() ||
       SemaRef.getLangOpts().CPlusPlus) {
+    Expr *LBVal = LB;
+    Expr *UBVal = UB;
+    // LB = TestIsLessOp.getValue() ? min(LB(MinVal), LB(MaxVal)) :
+    // max(LB(MinVal), LB(MaxVal))
+    if (InitDependOnLC) {
+      const LoopIterationSpace &IS =
+          ResultIterSpaces[ResultIterSpaces.size() - 1 -
+                           InitDependOnLC.getValueOr(
+                               CondDependOnLC.getValueOr(0))];
+      if (!IS.MinValue || !IS.MaxValue)
+        return nullptr;
+      // OuterVar = Min
+      ExprResult MinValue =
+          SemaRef.ActOnParenExpr(DefaultLoc, DefaultLoc, IS.MinValue);
+      if (!MinValue.isUsable())
+        return nullptr;
+
+      ExprResult LBMinVal = SemaRef.BuildBinOp(S, DefaultLoc, BO_Assign,
+                                               IS.CounterVar, MinValue.get());
+      if (!LBMinVal.isUsable())
+        return nullptr;
+      // OuterVar = Min, LBVal
+      LBMinVal =
+          SemaRef.BuildBinOp(S, DefaultLoc, BO_Comma, LBMinVal.get(), LBVal);
+      if (!LBMinVal.isUsable())
+        return nullptr;
+      // (OuterVar = Min, LBVal)
+      LBMinVal = SemaRef.ActOnParenExpr(DefaultLoc, DefaultLoc, LBMinVal.get());
+      if (!LBMinVal.isUsable())
+        return nullptr;
+
+      // OuterVar = Max
+      ExprResult MaxValue =
+          SemaRef.ActOnParenExpr(DefaultLoc, DefaultLoc, IS.MaxValue);
+      if (!MaxValue.isUsable())
+        return nullptr;
+
+      ExprResult LBMaxVal = SemaRef.BuildBinOp(S, DefaultLoc, BO_Assign,
+                                               IS.CounterVar, MaxValue.get());
+      if (!LBMaxVal.isUsable())
+        return nullptr;
+      // OuterVar = Max, LBVal
+      LBMaxVal =
+          SemaRef.BuildBinOp(S, DefaultLoc, BO_Comma, LBMaxVal.get(), LBVal);
+      if (!LBMaxVal.isUsable())
+        return nullptr;
+      // (OuterVar = Max, LBVal)
+      LBMaxVal = SemaRef.ActOnParenExpr(DefaultLoc, DefaultLoc, LBMaxVal.get());
+      if (!LBMaxVal.isUsable())
+        return nullptr;
+
+      Expr *LBMin = tryBuildCapture(SemaRef, LBMinVal.get(), Captures).get();
+      Expr *LBMax = tryBuildCapture(SemaRef, LBMaxVal.get(), Captures).get();
+      if (!LBMin || !LBMax)
+        return nullptr;
+      // LB(MinVal) < LB(MaxVal)
+      ExprResult MinLessMaxRes =
+          SemaRef.BuildBinOp(S, DefaultLoc, BO_LT, LBMin, LBMax);
+      if (!MinLessMaxRes.isUsable())
+        return nullptr;
+      Expr *MinLessMax =
+          tryBuildCapture(SemaRef, MinLessMaxRes.get(), Captures).get();
+      if (!MinLessMax)
+        return nullptr;
+      if (TestIsLessOp.getValue()) {
+        // LB(MinVal) < LB(MaxVal) ? LB(MinVal) : LB(MaxVal) - min(LB(MinVal),
+        // LB(MaxVal))
+        ExprResult MinLB = SemaRef.ActOnConditionalOp(DefaultLoc, DefaultLoc,
+                                                      MinLessMax, LBMin, LBMax);
+        if (!MinLB.isUsable())
+          return nullptr;
+        LBVal = MinLB.get();
+      } else {
+        // LB(MinVal) < LB(MaxVal) ? LB(MaxVal) : LB(MinVal) - max(LB(MinVal),
+        // LB(MaxVal))
+        ExprResult MaxLB = SemaRef.ActOnConditionalOp(DefaultLoc, DefaultLoc,
+                                                      MinLessMax, LBMax, LBMin);
+        if (!MaxLB.isUsable())
+          return nullptr;
+        LBVal = MaxLB.get();
+      }
+    }
+    // UB = TestIsLessOp.getValue() ? max(UB(MinVal), UB(MaxVal)) :
+    // min(UB(MinVal), UB(MaxVal))
+    if (CondDependOnLC) {
+      const LoopIterationSpace &IS =
+          ResultIterSpaces[ResultIterSpaces.size() - 1 -
+                           InitDependOnLC.getValueOr(
+                               CondDependOnLC.getValueOr(0))];
+      if (!IS.MinValue || !IS.MaxValue)
+        return nullptr;
+      // OuterVar = Min
+      ExprResult MinValue =
+          SemaRef.ActOnParenExpr(DefaultLoc, DefaultLoc, IS.MinValue);
+      if (!MinValue.isUsable())
+        return nullptr;
+
+      ExprResult UBMinVal = SemaRef.BuildBinOp(S, DefaultLoc, BO_Assign,
+                                               IS.CounterVar, MinValue.get());
+      if (!UBMinVal.isUsable())
+        return nullptr;
+      // OuterVar = Min, UBVal
+      UBMinVal =
+          SemaRef.BuildBinOp(S, DefaultLoc, BO_Comma, UBMinVal.get(), UBVal);
+      if (!UBMinVal.isUsable())
+        return nullptr;
+      // (OuterVar = Min, UBVal)
+      UBMinVal = SemaRef.ActOnParenExpr(DefaultLoc, DefaultLoc, UBMinVal.get());
+      if (!UBMinVal.isUsable())
+        return nullptr;
+
+      // OuterVar = Max
+      ExprResult MaxValue =
+          SemaRef.ActOnParenExpr(DefaultLoc, DefaultLoc, IS.MaxValue);
+      if (!MaxValue.isUsable())
+        return nullptr;
+
+      ExprResult UBMaxVal = SemaRef.BuildBinOp(S, DefaultLoc, BO_Assign,
+                                               IS.CounterVar, MaxValue.get());
+      if (!UBMaxVal.isUsable())
+        return nullptr;
+      // OuterVar = Max, UBVal
+      UBMaxVal =
+          SemaRef.BuildBinOp(S, DefaultLoc, BO_Comma, UBMaxVal.get(), UBVal);
+      if (!UBMaxVal.isUsable())
+        return nullptr;
+      // (OuterVar = Max, UBVal)
+      UBMaxVal = SemaRef.ActOnParenExpr(DefaultLoc, DefaultLoc, UBMaxVal.get());
+      if (!UBMaxVal.isUsable())
+        return nullptr;
+
+      Expr *UBMin = tryBuildCapture(SemaRef, UBMinVal.get(), Captures).get();
+      Expr *UBMax = tryBuildCapture(SemaRef, UBMaxVal.get(), Captures).get();
+      if (!UBMin || !UBMax)
+        return nullptr;
+      // UB(MinVal) > UB(MaxVal)
+      ExprResult MinGreaterMaxRes =
+          SemaRef.BuildBinOp(S, DefaultLoc, BO_GT, UBMin, UBMax);
+      if (!MinGreaterMaxRes.isUsable())
+        return nullptr;
+      Expr *MinGreaterMax =
+          tryBuildCapture(SemaRef, MinGreaterMaxRes.get(), Captures).get();
+      if (!MinGreaterMax)
+        return nullptr;
+      if (TestIsLessOp.getValue()) {
+        // UB(MinVal) > UB(MaxVal) ? UB(MinVal) : UB(MaxVal) - max(UB(MinVal),
+        // UB(MaxVal))
+        ExprResult MaxUB = SemaRef.ActOnConditionalOp(
+            DefaultLoc, DefaultLoc, MinGreaterMax, UBMin, UBMax);
+        if (!MaxUB.isUsable())
+          return nullptr;
+        UBVal = MaxUB.get();
+      } else {
+        // UB(MinVal) > UB(MaxVal) ? UB(MaxVal) : UB(MinVal) - min(UB(MinVal),
+        // UB(MaxVal))
+        ExprResult MinUB = SemaRef.ActOnConditionalOp(
+            DefaultLoc, DefaultLoc, MinGreaterMax, UBMax, UBMin);
+        if (!MinUB.isUsable())
+          return nullptr;
+        UBVal = MinUB.get();
+      }
+    }
     // Upper - Lower
-    Expr *UBExpr = TestIsLessOp.getValue() ? UB : LB;
-    Expr *LBExpr = TestIsLessOp.getValue() ? LB : UB;
+    Expr *UBExpr = TestIsLessOp.getValue() ? UBVal : LBVal;
+    Expr *LBExpr = TestIsLessOp.getValue() ? LBVal : UBVal;
     Expr *Upper = tryBuildCapture(SemaRef, UBExpr, Captures).get();
     Expr *Lower = tryBuildCapture(SemaRef, LBExpr, Captures).get();
     if (!Upper || !Lower)
@@ -5446,6 +5673,127 @@ Expr *OpenMPIterationSpaceChecker::buildNumIterations(
   return Diff.get();
 }
 
+std::pair<Expr *, Expr *> OpenMPIterationSpaceChecker::buildMinMaxValues(
+    Scope *S, llvm::MapVector<const Expr *, DeclRefExpr *> &Captures) const {
+  // Do not build for iterators, they cannot be used in non-rectangular loop
+  // nests.
+  if (LCDecl->getType()->isRecordType())
+    return std::make_pair(nullptr, nullptr);
+  // If we subtract, the min is in the condition, otherwise the min is in the
+  // init value.
+  Expr *MinExpr = nullptr;
+  Expr *MaxExpr = nullptr;
+  Expr *LBExpr = TestIsLessOp.getValue() ? LB : UB;
+  Expr *UBExpr = TestIsLessOp.getValue() ? UB : LB;
+  bool LBNonRect = TestIsLessOp.getValue() ? InitDependOnLC.hasValue()
+                                           : CondDependOnLC.hasValue();
+  bool UBNonRect = TestIsLessOp.getValue() ? CondDependOnLC.hasValue()
+                                           : InitDependOnLC.hasValue();
+  Expr *Lower =
+      LBNonRect ? LBExpr : tryBuildCapture(SemaRef, LBExpr, Captures).get();
+  Expr *Upper =
+      UBNonRect ? UBExpr : tryBuildCapture(SemaRef, UBExpr, Captures).get();
+  if (!Upper || !Lower)
+    return std::make_pair(nullptr, nullptr);
+
+  if (TestIsLessOp.getValue())
+    MinExpr = Lower;
+  else
+    MaxExpr = Upper;
+
+  // Build minimum/maximum value based on number of iterations.
+  ExprResult Diff;
+  QualType VarType = LCDecl->getType().getNonReferenceType();
+
+  Diff = SemaRef.BuildBinOp(S, DefaultLoc, BO_Sub, Upper, Lower);
+  if (!Diff.isUsable())
+    return std::make_pair(nullptr, nullptr);
+
+  // Upper - Lower [- 1]
+  if (TestIsStrictOp)
+    Diff = SemaRef.BuildBinOp(
+        S, DefaultLoc, BO_Sub, Diff.get(),
+        SemaRef.ActOnIntegerConstant(SourceLocation(), 1).get());
+  if (!Diff.isUsable())
+    return std::make_pair(nullptr, nullptr);
+
+  // Upper - Lower [- 1] + Step
+  ExprResult NewStep = tryBuildCapture(SemaRef, Step, Captures);
+  if (!NewStep.isUsable())
+    return std::make_pair(nullptr, nullptr);
+
+  // Parentheses (for dumping/debugging purposes only).
+  Diff = SemaRef.ActOnParenExpr(DefaultLoc, DefaultLoc, Diff.get());
+  if (!Diff.isUsable())
+    return std::make_pair(nullptr, nullptr);
+
+  // (Upper - Lower [- 1]) / Step
+  Diff = SemaRef.BuildBinOp(S, DefaultLoc, BO_Div, Diff.get(), NewStep.get());
+  if (!Diff.isUsable())
+    return std::make_pair(nullptr, nullptr);
+
+  // ((Upper - Lower [- 1]) / Step) * Step
+  // Parentheses (for dumping/debugging purposes only).
+  Diff = SemaRef.ActOnParenExpr(DefaultLoc, DefaultLoc, Diff.get());
+  if (!Diff.isUsable())
+    return std::make_pair(nullptr, nullptr);
+
+  Diff = SemaRef.BuildBinOp(S, DefaultLoc, BO_Mul, Diff.get(), NewStep.get());
+  if (!Diff.isUsable())
+    return std::make_pair(nullptr, nullptr);
+
+  // Convert to the original type or ptrdiff_t, if original type is pointer.
+  if (!VarType->isAnyPointerType() &&
+      !SemaRef.Context.hasSameType(Diff.get()->getType(), VarType)) {
+    Diff = SemaRef.PerformImplicitConversion(
+        Diff.get(), VarType, Sema::AA_Converting, /*AllowExplicit=*/true);
+  } else if (VarType->isAnyPointerType() &&
+             !SemaRef.Context.hasSameType(
+                 Diff.get()->getType(),
+                 SemaRef.Context.getUnsignedPointerDiffType())) {
+    Diff = SemaRef.PerformImplicitConversion(
+        Diff.get(), SemaRef.Context.getUnsignedPointerDiffType(),
+        Sema::AA_Converting, /*AllowExplicit=*/true);
+  }
+  if (!Diff.isUsable())
+    return std::make_pair(nullptr, nullptr);
+
+  // Parentheses (for dumping/debugging purposes only).
+  Diff = SemaRef.ActOnParenExpr(DefaultLoc, DefaultLoc, Diff.get());
+  if (!Diff.isUsable())
+    return std::make_pair(nullptr, nullptr);
+
+  if (TestIsLessOp.getValue()) {
+    // MinExpr = Lower;
+    // MaxExpr = Lower + (((Upper - Lower [- 1]) / Step) * Step)
+    Diff = SemaRef.BuildBinOp(S, DefaultLoc, BO_Add, Lower, Diff.get());
+    if (!Diff.isUsable())
+      return std::make_pair(nullptr, nullptr);
+    Diff = SemaRef.ActOnFinishFullExpr(Diff.get(), /*DiscardedValue*/ false);
+    if (!Diff.isUsable())
+      return std::make_pair(nullptr, nullptr);
+    MaxExpr = Diff.get();
+  } else {
+    // MaxExpr = Upper;
+    // MinExpr = Upper - (((Upper - Lower [- 1]) / Step) * Step)
+    Diff = SemaRef.BuildBinOp(S, DefaultLoc, BO_Sub, Upper, Diff.get());
+    if (!Diff.isUsable())
+      return std::make_pair(nullptr, nullptr);
+    Diff = SemaRef.ActOnFinishFullExpr(Diff.get(), /*DiscardedValue*/ false);
+    if (!Diff.isUsable())
+      return std::make_pair(nullptr, nullptr);
+    MinExpr = Diff.get();
+  }
+
+  return std::make_pair(MinExpr, MaxExpr);
+}
+
+Expr *OpenMPIterationSpaceChecker::buildFinalCondition(Scope *S) const {
+  if (InitDependOnLC || CondDependOnLC)
+    return Condition;
+  return nullptr;
+}
+
 Expr *OpenMPIterationSpaceChecker::buildPreCond(
     Scope *S, Expr *Cond,
     llvm::MapVector<const Expr *, DeclRefExpr *> &Captures) const {
@@ -5453,8 +5801,10 @@ Expr *OpenMPIterationSpaceChecker::buildPreCond(
   bool Suppress = SemaRef.getDiagnostics().getSuppressAllDiagnostics();
   SemaRef.getDiagnostics().setSuppressAllDiagnostics(/*Val=*/true);
 
-  ExprResult NewLB = tryBuildCapture(SemaRef, LB, Captures);
-  ExprResult NewUB = tryBuildCapture(SemaRef, UB, Captures);
+  ExprResult NewLB =
+      InitDependOnLC ? LB : tryBuildCapture(SemaRef, LB, Captures);
+  ExprResult NewUB =
+      CondDependOnLC ? UB : tryBuildCapture(SemaRef, UB, Captures);
   if (!NewLB.isUsable() || !NewUB.isUsable())
     return nullptr;
 
@@ -5576,36 +5926,6 @@ Expr *OpenMPIterationSpaceChecker::buildOrderedLoopData(
 
   return Diff.get();
 }
-
-/// Iteration space of a single for loop.
-struct LoopIterationSpace final {
-  /// True if the condition operator is the strict compare operator (<, > or
-  /// !=).
-  bool IsStrictCompare = false;
-  /// Condition of the loop.
-  Expr *PreCond = nullptr;
-  /// This expression calculates the number of iterations in the loop.
-  /// It is always possible to calculate it before starting the loop.
-  Expr *NumIterations = nullptr;
-  /// The loop counter variable.
-  Expr *CounterVar = nullptr;
-  /// Private loop counter variable.
-  Expr *PrivateCounterVar = nullptr;
-  /// This is initializer for the initial value of #CounterVar.
-  Expr *CounterInit = nullptr;
-  /// This is step for the #CounterVar used to generate its update:
-  /// #CounterVar = #CounterInit + #CounterStep * CurrentIteration.
-  Expr *CounterStep = nullptr;
-  /// Should step be subtracted?
-  bool Subtract = false;
-  /// Source range of the loop init.
-  SourceRange InitSrcRange;
-  /// Source range of the loop condition.
-  SourceRange CondSrcRange;
-  /// Source range of the loop increment.
-  SourceRange IncSrcRange;
-};
-
 } // namespace
 
 void Sema::ActOnOpenMPLoopInitialization(SourceLocation ForLoc, Stmt *Init) {
@@ -5696,7 +6016,7 @@ static bool checkOpenMPIterationSpace(
     unsigned TotalNestedLoopCount, Expr *CollapseLoopCountExpr,
     Expr *OrderedLoopCountExpr,
     Sema::VarsWithInheritedDSAType &VarsWithImplicitDSA,
-    LoopIterationSpace &ResultIterSpace,
+    llvm::MutableArrayRef<LoopIterationSpace> ResultIterSpaces,
     llvm::MapVector<const Expr *, DeclRefExpr *> &Captures) {
   // OpenMP [2.6, Canonical Loop Form]
   //   for (init-expr; test-expr; incr-expr) structured-block
@@ -5774,37 +6094,57 @@ static bool checkOpenMPIterationSpace(
     return HasErrors;
 
   // Build the loop's iteration space representation.
-  ResultIterSpace.PreCond =
+  ResultIterSpaces[CurrentNestedLoopCount].PreCond =
       ISC.buildPreCond(DSA.getCurScope(), For->getCond(), Captures);
-  ResultIterSpace.NumIterations = ISC.buildNumIterations(
-      DSA.getCurScope(),
-      (isOpenMPWorksharingDirective(DKind) ||
-       isOpenMPTaskLoopDirective(DKind) || isOpenMPDistributeDirective(DKind)),
-      Captures);
-  ResultIterSpace.CounterVar = ISC.buildCounterVar(Captures, DSA);
-  ResultIterSpace.PrivateCounterVar = ISC.buildPrivateCounterVar();
-  ResultIterSpace.CounterInit = ISC.buildCounterInit();
-  ResultIterSpace.CounterStep = ISC.buildCounterStep();
-  ResultIterSpace.InitSrcRange = ISC.getInitSrcRange();
-  ResultIterSpace.CondSrcRange = ISC.getConditionSrcRange();
-  ResultIterSpace.IncSrcRange = ISC.getIncrementSrcRange();
-  ResultIterSpace.Subtract = ISC.shouldSubtractStep();
-  ResultIterSpace.IsStrictCompare = ISC.isStrictTestOp();
-
-  HasErrors |= (ResultIterSpace.PreCond == nullptr ||
-                ResultIterSpace.NumIterations == nullptr ||
-                ResultIterSpace.CounterVar == nullptr ||
-                ResultIterSpace.PrivateCounterVar == nullptr ||
-                ResultIterSpace.CounterInit == nullptr ||
-                ResultIterSpace.CounterStep == nullptr);
+  ResultIterSpaces[CurrentNestedLoopCount].NumIterations =
+      ISC.buildNumIterations(DSA.getCurScope(), ResultIterSpaces,
+                             (isOpenMPWorksharingDirective(DKind) ||
+                              isOpenMPTaskLoopDirective(DKind) ||
+                              isOpenMPDistributeDirective(DKind)),
+                             Captures);
+  ResultIterSpaces[CurrentNestedLoopCount].CounterVar =
+      ISC.buildCounterVar(Captures, DSA);
+  ResultIterSpaces[CurrentNestedLoopCount].PrivateCounterVar =
+      ISC.buildPrivateCounterVar();
+  ResultIterSpaces[CurrentNestedLoopCount].CounterInit = ISC.buildCounterInit();
+  ResultIterSpaces[CurrentNestedLoopCount].CounterStep = ISC.buildCounterStep();
+  ResultIterSpaces[CurrentNestedLoopCount].InitSrcRange = ISC.getInitSrcRange();
+  ResultIterSpaces[CurrentNestedLoopCount].CondSrcRange =
+      ISC.getConditionSrcRange();
+  ResultIterSpaces[CurrentNestedLoopCount].IncSrcRange =
+      ISC.getIncrementSrcRange();
+  ResultIterSpaces[CurrentNestedLoopCount].Subtract = ISC.shouldSubtractStep();
+  ResultIterSpaces[CurrentNestedLoopCount].IsStrictCompare =
+      ISC.isStrictTestOp();
+  std::tie(ResultIterSpaces[CurrentNestedLoopCount].MinValue,
+           ResultIterSpaces[CurrentNestedLoopCount].MaxValue) =
+      ISC.buildMinMaxValues(DSA.getCurScope(), Captures);
+  ResultIterSpaces[CurrentNestedLoopCount].FinalCondition =
+      ISC.buildFinalCondition(DSA.getCurScope());
+  ResultIterSpaces[CurrentNestedLoopCount].IsNonRectangularLB =
+      ISC.doesInitDependOnLC();
+  ResultIterSpaces[CurrentNestedLoopCount].IsNonRectangularUB =
+      ISC.doesCondDependOnLC();
+  ResultIterSpaces[CurrentNestedLoopCount].LoopDependentIdx =
+      ISC.getLoopDependentIdx();
+
+  HasErrors |=
+      (ResultIterSpaces[CurrentNestedLoopCount].PreCond == nullptr ||
+       ResultIterSpaces[CurrentNestedLoopCount].NumIterations == nullptr ||
+       ResultIterSpaces[CurrentNestedLoopCount].CounterVar == nullptr ||
+       ResultIterSpaces[CurrentNestedLoopCount].PrivateCounterVar == nullptr ||
+       ResultIterSpaces[CurrentNestedLoopCount].CounterInit == nullptr ||
+       ResultIterSpaces[CurrentNestedLoopCount].CounterStep == nullptr);
   if (!HasErrors && DSA.isOrderedRegion()) {
     if (DSA.getOrderedRegionParam().second->getNumForLoops()) {
       if (CurrentNestedLoopCount <
           DSA.getOrderedRegionParam().second->getLoopNumIterations().size()) {
         DSA.getOrderedRegionParam().second->setLoopNumIterations(
-            CurrentNestedLoopCount, ResultIterSpace.NumIterations);
+            CurrentNestedLoopCount,
+            ResultIterSpaces[CurrentNestedLoopCount].NumIterations);
         DSA.getOrderedRegionParam().second->setLoopCounter(
-            CurrentNestedLoopCount, ResultIterSpace.CounterVar);
+            CurrentNestedLoopCount,
+            ResultIterSpaces[CurrentNestedLoopCount].CounterVar);
       }
     }
     for (auto &Pair : DSA.getDoacrossDependClauses()) {
@@ -5821,11 +6161,13 @@ static bool checkOpenMPIterationSpace(
       Expr *CntValue;
       if (Pair.first->getDependencyKind() == OMPC_DEPEND_source)
         CntValue = ISC.buildOrderedLoopData(
-            DSA.getCurScope(), ResultIterSpace.CounterVar, Captures,
+            DSA.getCurScope(),
+            ResultIterSpaces[CurrentNestedLoopCount].CounterVar, Captures,
             Pair.first->getDependencyLoc());
       else
         CntValue = ISC.buildOrderedLoopData(
-            DSA.getCurScope(), ResultIterSpace.CounterVar, Captures,
+            DSA.getCurScope(),
+            ResultIterSpaces[CurrentNestedLoopCount].CounterVar, Captures,
             Pair.first->getDependencyLoc(),
             Pair.second[CurrentNestedLoopCount].first,
             Pair.second[CurrentNestedLoopCount].second);
@@ -5839,10 +6181,12 @@ static bool checkOpenMPIterationSpace(
 /// Build 'VarRef = Start.
 static ExprResult
 buildCounterInit(Sema &SemaRef, Scope *S, SourceLocation Loc, ExprResult VarRef,
-                 ExprResult Start,
+                 ExprResult Start, bool IsNonRectangularLB,
                  llvm::MapVector<const Expr *, DeclRefExpr *> &Captures) {
   // Build 'VarRef = Start.
-  ExprResult NewStart = tryBuildCapture(SemaRef, Start.get(), Captures);
+  ExprResult NewStart = IsNonRectangularLB
+                            ? Start.get()
+                            : tryBuildCapture(SemaRef, Start.get(), Captures);
   if (!NewStart.isUsable())
     return ExprError();
   if (!SemaRef.Context.hasSameType(NewStart.get()->getType(),
@@ -5863,6 +6207,7 @@ buildCounterInit(Sema &SemaRef, Scope *S, SourceLocation Loc, ExprResult VarRef,
 static ExprResult buildCounterUpdate(
     Sema &SemaRef, Scope *S, SourceLocation Loc, ExprResult VarRef,
     ExprResult Start, ExprResult Iter, ExprResult Step, bool Subtract,
+    bool IsNonRectangularLB,
     llvm::MapVector<const Expr *, DeclRefExpr *> *Captures = nullptr) {
   // Add parentheses (for debugging purposes only).
   Iter = SemaRef.ActOnParenExpr(Loc, Loc, Iter.get());
@@ -5882,8 +6227,12 @@ static ExprResult buildCounterUpdate(
 
   // Try to build 'VarRef = Start, VarRef (+|-)= Iter * Step' or
   // 'VarRef = Start (+|-) Iter * Step'.
-  ExprResult NewStart = Start;
-  if (Captures)
+  if (!Start.isUsable())
+    return ExprError();
+  ExprResult NewStart = SemaRef.ActOnParenExpr(Loc, Loc, Start.get());
+  if (!NewStart.isUsable())
+    return ExprError();
+  if (Captures && !IsNonRectangularLB)
     NewStart = tryBuildCapture(SemaRef, Start.get(), *Captures);
   if (NewStart.isInvalid())
     return ExprError();
@@ -6054,8 +6403,7 @@ checkOpenMPLoop(OpenMPDirectiveKind DKind, Expr *CollapseLoopCountExpr,
     if (checkOpenMPIterationSpace(
             DKind, CurStmt, SemaRef, DSA, Cnt, NestedLoopCount,
             std::max(OrderedLoopCount, NestedLoopCount), CollapseLoopCountExpr,
-            OrderedLoopCountExpr, VarsWithImplicitDSA, IterSpaces[Cnt],
-            Captures))
+            OrderedLoopCountExpr, VarsWithImplicitDSA, IterSpaces, Captures))
       return 0;
     // Move on to the next nested for loop, or to the loop body.
     // OpenMP [2.8.1, simd construct, Restrictions]
@@ -6068,8 +6416,7 @@ checkOpenMPLoop(OpenMPDirectiveKind DKind, Expr *CollapseLoopCountExpr,
     if (checkOpenMPIterationSpace(
             DKind, CurStmt, SemaRef, DSA, Cnt, NestedLoopCount,
             std::max(OrderedLoopCount, NestedLoopCount), CollapseLoopCountExpr,
-            OrderedLoopCountExpr, VarsWithImplicitDSA, IterSpaces[Cnt],
-            Captures))
+            OrderedLoopCountExpr, VarsWithImplicitDSA, IterSpaces, Captures))
       return 0;
     if (Cnt > 0 && IterSpaces[Cnt].CounterVar) {
       // Handle initialization of captured loop iterator variables.
@@ -6530,6 +6877,9 @@ checkOpenMPLoop(OpenMPDirectiveKind DKind, Expr *CollapseLoopCountExpr,
   Built.Inits.resize(NestedLoopCount);
   Built.Updates.resize(NestedLoopCount);
   Built.Finals.resize(NestedLoopCount);
+  Built.DependentCounters.resize(NestedLoopCount);
+  Built.DependentInits.resize(NestedLoopCount);
+  Built.FinalsConditions.resize(NestedLoopCount);
   {
     // We implement the following algorithm for obtaining the
     // original loop iteration variable values based on the
@@ -6589,24 +6939,26 @@ checkOpenMPLoop(OpenMPDirectiveKind DKind, Expr *CollapseLoopCountExpr,
       DeclRefExpr *CounterVar = buildDeclRefExpr(
           SemaRef, VD, IS.CounterVar->getType(), IS.CounterVar->getExprLoc(),
           /*RefersToCapture=*/true);
-      ExprResult Init = buildCounterInit(SemaRef, CurScope, UpdLoc, CounterVar,
-                                         IS.CounterInit, Captures);
+      ExprResult Init =
+          buildCounterInit(SemaRef, CurScope, UpdLoc, CounterVar,
+                           IS.CounterInit, IS.IsNonRectangularLB, Captures);
       if (!Init.isUsable()) {
         HasErrors = true;
         break;
       }
       ExprResult Update = buildCounterUpdate(
           SemaRef, CurScope, UpdLoc, CounterVar, IS.CounterInit, Iter,
-          IS.CounterStep, IS.Subtract, &Captures);
+          IS.CounterStep, IS.Subtract, IS.IsNonRectangularLB, &Captures);
       if (!Update.isUsable()) {
         HasErrors = true;
         break;
       }
 
       // Build final: IS.CounterVar = IS.Start + IS.NumIters * IS.Step
-      ExprResult Final = buildCounterUpdate(
-          SemaRef, CurScope, UpdLoc, CounterVar, IS.CounterInit,
-          IS.NumIterations, IS.CounterStep, IS.Subtract, &Captures);
+      ExprResult Final =
+          buildCounterUpdate(SemaRef, CurScope, UpdLoc, CounterVar,
+                             IS.CounterInit, IS.NumIterations, IS.CounterStep,
+                             IS.Subtract, IS.IsNonRectangularLB, &Captures);
       if (!Final.isUsable()) {
         HasErrors = true;
         break;
@@ -6622,6 +6974,16 @@ checkOpenMPLoop(OpenMPDirectiveKind DKind, Expr *CollapseLoopCountExpr,
       Built.Inits[Cnt] = Init.get();
       Built.Updates[Cnt] = Update.get();
       Built.Finals[Cnt] = Final.get();
+      Built.DependentCounters[Cnt] = nullptr;
+      Built.DependentInits[Cnt] = nullptr;
+      Built.FinalsConditions[Cnt] = nullptr;
+      if (IS.IsNonRectangularLB) {
+        Built.DependentCounters[Cnt] =
+            Built.Counters[NestedLoopCount - 1 - IS.LoopDependentIdx];
+        Built.DependentInits[Cnt] =
+            Built.Inits[NestedLoopCount - 1 - IS.LoopDependentIdx];
+        Built.FinalsConditions[Cnt] = IS.FinalCondition;
+      }
     }
   }
 
@@ -6634,7 +6996,7 @@ checkOpenMPLoop(OpenMPDirectiveKind DKind, Expr *CollapseLoopCountExpr,
   Built.NumIterations = NumIterations.get();
   Built.CalcLastIteration = SemaRef
                                 .ActOnFinishFullExpr(CalcLastIteration.get(),
-                                                     /*DiscardedValue*/ false)
+                                                     /*DiscardedValue=*/false)
                                 .get();
   Built.PreCond = PreCond.get();
   Built.PreInits = buildPreInits(C, Captures);
@@ -12778,9 +13140,9 @@ static bool FinishOpenMPLinearClause(OMPLinearClause &Clause, DeclRefExpr *IV,
     // Build update: Var = InitExpr + IV * Step
     ExprResult Update;
     if (!Info.first)
-      Update =
-          buildCounterUpdate(SemaRef, S, RefExpr->getExprLoc(), *CurPrivate,
-                             InitExpr, IV, Step, /* Subtract */ false);
+      Update = buildCounterUpdate(
+          SemaRef, S, RefExpr->getExprLoc(), *CurPrivate, InitExpr, IV, Step,
+          /*Subtract=*/false, /*IsNonRectangularLB=*/false);
     else
       Update = *CurPrivate;
     Update = SemaRef.ActOnFinishFullExpr(Update.get(), DE->getBeginLoc(),
@@ -12791,7 +13153,8 @@ static bool FinishOpenMPLinearClause(OMPLinearClause &Clause, DeclRefExpr *IV,
     if (!Info.first)
       Final =
           buildCounterUpdate(SemaRef, S, RefExpr->getExprLoc(), CapturedRef,
-                             InitExpr, NumIterations, Step, /*Subtract=*/false);
+                             InitExpr, NumIterations, Step, /*Subtract=*/false,
+                             /*IsNonRectangularLB=*/false);
     else
       Final = *CurPrivate;
     Final = SemaRef.ActOnFinishFullExpr(Final.get(), DE->getBeginLoc(),
diff --git a/lib/Sema/SemaTemplate.cpp b/lib/Sema/SemaTemplate.cpp
index f11dd22dbdc..dad665d193b 100644
--- a/lib/Sema/SemaTemplate.cpp
+++ b/lib/Sema/SemaTemplate.cpp
@@ -362,13 +362,27 @@ bool Sema::LookupTemplateName(LookupResult &Found,
     // x->B::f, and we are looking into the type of the object.
     assert(!SS.isSet() && "ObjectType and scope specifier cannot coexist");
     LookupCtx = computeDeclContext(ObjectType);
-    IsDependent = !LookupCtx;
+    IsDependent = !LookupCtx && ObjectType->isDependentType();
     assert((IsDependent || !ObjectType->isIncompleteType() ||
             ObjectType->castAs<TagType>()->isBeingDefined()) &&
            "Caller should have completed object type");
 
-    // Template names cannot appear inside an Objective-C class or object type.
-    if (ObjectType->isObjCObjectOrInterfaceType()) {
+    // Template names cannot appear inside an Objective-C class or object type
+    // or a vector type.
+    //
+    // FIXME: This is wrong. For example:
+    //
+    //   template<typename T> using Vec = T __attribute__((ext_vector_type(4)));
+    //   Vec<int> vi;
+    //   vi.Vec<int>::~Vec<int>();
+    //
+    // ... should be accepted but we will not treat 'Vec' as a template name
+    // here. The right thing to do would be to check if the name is a valid
+    // vector component name, and look up a template name if not. And similarly
+    // for lookups into Objective-C class and object types, where the same
+    // problem can arise.
+    if (ObjectType->isObjCObjectOrInterfaceType() ||
+        ObjectType->isVectorType()) {
       Found.clear();
       return false;
     }
diff --git a/lib/Serialization/ASTReaderStmt.cpp b/lib/Serialization/ASTReaderStmt.cpp
index afaaa543bb2..8ab0845d151 100644
--- a/lib/Serialization/ASTReaderStmt.cpp
+++ b/lib/Serialization/ASTReaderStmt.cpp
@@ -2060,6 +2060,18 @@ void ASTStmtReader::VisitOMPLoopDirective(OMPLoopDirective *D) {
   for (unsigned i = 0; i < CollapsedNum; ++i)
     Sub.push_back(Record.readSubExpr());
   D->setFinals(Sub);
+  Sub.clear();
+  for (unsigned i = 0; i < CollapsedNum; ++i)
+    Sub.push_back(Record.readSubExpr());
+  D->setDependentCounters(Sub);
+  Sub.clear();
+  for (unsigned i = 0; i < CollapsedNum; ++i)
+    Sub.push_back(Record.readSubExpr());
+  D->setDependentInits(Sub);
+  Sub.clear();
+  for (unsigned i = 0; i < CollapsedNum; ++i)
+    Sub.push_back(Record.readSubExpr());
+  D->setFinalsConditions(Sub);
 }
 
 void ASTStmtReader::VisitOMPParallelDirective(OMPParallelDirective *D) {
diff --git a/lib/Serialization/ASTWriterStmt.cpp b/lib/Serialization/ASTWriterStmt.cpp
index 4fbcbaabe74..a6927f32c0e 100644
--- a/lib/Serialization/ASTWriterStmt.cpp
+++ b/lib/Serialization/ASTWriterStmt.cpp
@@ -1995,6 +1995,12 @@ void ASTStmtWriter::VisitOMPLoopDirective(OMPLoopDirective *D) {
   for (auto I : D->finals()) {
     Record.AddStmt(I);
   }
+  for (Stmt *S : D->dependent_counters())
+    Record.AddStmt(S);
+  for (Stmt *S : D->dependent_inits())
+    Record.AddStmt(S);
+  for (Stmt *S : D->finals_conditions())
+    Record.AddStmt(S);
 }
 
 void ASTStmtWriter::VisitOMPParallelDirective(OMPParallelDirective *D) {
diff --git a/lib/StaticAnalyzer/Frontend/CheckerRegistration.cpp b/lib/StaticAnalyzer/Frontend/CheckerRegistration.cpp
index 1e45ee96145..27555537c89 100644
--- a/lib/StaticAnalyzer/Frontend/CheckerRegistration.cpp
+++ b/lib/StaticAnalyzer/Frontend/CheckerRegistration.cpp
@@ -74,11 +74,22 @@ void ento::printCheckerConfigList(raw_ostream &OS,
 }
 
 void ento::printAnalyzerConfigList(raw_ostream &out) {
-  out << "OVERVIEW: Clang Static Analyzer -analyzer-config Option List\n\n";
-  out << "USAGE: -analyzer-config <OPTION1=VALUE,OPTION2=VALUE,...>\n\n";
-  out << "       -analyzer-config OPTION1=VALUE, -analyzer-config "
-         "OPTION2=VALUE, ...\n\n";
-  out << "OPTIONS:\n\n";
+  // FIXME: This message sounds scary, should be scary, but incorrectly states
+  // that all configs are super dangerous. In reality, many of them should be
+  // accessible to the user. We should create a user-facing subset of config
+  // options under a different frontend flag.
+  out << R"(
+OVERVIEW: Clang Static Analyzer -analyzer-config Option List
+
+The following list of configurations are meant for development purposes only, as
+some of the variables they define are set to result in the most optimal
+analysis. Setting them to other values may drastically change how the analyzer
+behaves, and may even result in instabilities, crashes!
+
+USAGE: -analyzer-config <OPTION1=VALUE,OPTION2=VALUE,...>
+       -analyzer-config OPTION1=VALUE, -analyzer-config OPTION2=VALUE, ...
+OPTIONS:
+)";
 
   using OptionAndDescriptionTy = std::pair<StringRef, std::string>;
   OptionAndDescriptionTy PrintableOptions[] = {
diff --git a/lib/Tooling/Refactoring/Transformer.cpp b/lib/Tooling/Refactoring/Transformer.cpp
index 24f9b58897e..c27739c5962 100644
--- a/lib/Tooling/Refactoring/Transformer.cpp
+++ b/lib/Tooling/Refactoring/Transformer.cpp
@@ -98,8 +98,10 @@ static std::vector<DynTypedMatcher> taggedMatchers(
   Matchers.reserve(Cases.size());
   for (const auto &Case : Cases) {
     std::string Tag = (TagBase + Twine(Case.first)).str();
-    auto M = Case.second.Matcher.tryBind(Tag);
-    assert(M && "RewriteRule matchers should be bindable.");
+    // HACK: Many matchers are not bindable, so ensure that tryBind will work.
+    DynTypedMatcher BoundMatcher(Case.second.Matcher);
+    BoundMatcher.setAllowBind(true);
+    auto M = BoundMatcher.tryBind(Tag);
     Matchers.push_back(*std::move(M));
   }
   return Matchers;
diff --git a/test/CXX/drs/dr22xx.cpp b/test/CXX/drs/dr22xx.cpp
index 70a26db757c..8896281e9ca 100644
--- a/test/CXX/drs/dr22xx.cpp
+++ b/test/CXX/drs/dr22xx.cpp
@@ -26,3 +26,12 @@ void f() {
 }
 }
 #endif
+
+namespace dr2292 { // dr2292: 9
+#if __cplusplus >= 201103L
+  template<typename T> using id = T;
+  void test(int *p) {
+    p->template id<int>::~id<int>();
+  }
+#endif
+}
diff --git a/test/CXX/drs/dr4xx.cpp b/test/CXX/drs/dr4xx.cpp
index 00393cc2e4c..8eeb7715cad 100644
--- a/test/CXX/drs/dr4xx.cpp
+++ b/test/CXX/drs/dr4xx.cpp
@@ -318,8 +318,8 @@ namespace dr420 { // dr420: yes
     q->~id<int>();
     p->id<int>::~id<int>();
     q->id<int>::~id<int>();
-    p->template id<int>::~id<int>(); // expected-error {{'template' keyword not permitted here}} expected-error {{base type 'int' is not a struct}}
-    q->template id<int>::~id<int>(); // expected-error {{'template' keyword not permitted here}} expected-error {{base type 'int' is not a struct}}
+    p->template id<int>::~id<int>(); // OK since dr2292
+    q->template id<int>::~id<int>(); // OK since dr2292
     p->A::template id<int>::~id<int>();
     q->A::template id<int>::~id<int>();
   }
diff --git a/test/CodeGen/sse2-builtins.c b/test/CodeGen/sse2-builtins.c
index 280640c267c..6e4327f593f 100644
--- a/test/CodeGen/sse2-builtins.c
+++ b/test/CodeGen/sse2-builtins.c
@@ -808,6 +808,23 @@ __m128d test_mm_min_sd(__m128d A, __m128d B) {
   return _mm_min_sd(A, B);
 }
 
+__m64 test_mm_movepi64_pi64(__m128i A)
+{
+  // CHECK-LABEL: test_mm_movepi64_pi64
+  // CHECK: [[EXT:%.*]] = extractelement <2 x i64> %1, i32 0
+  // CHECK: bitcast i64 [[EXT]] to <1 x i64>
+  return _mm_movepi64_pi64(A);
+}
+
+__m128i test_mm_movpi64_epi64(__m64 A)
+{
+  // CHECK-LABEL: test_mm_movpi64_epi64
+  // CHECK: [[CAST:%.*]] = bitcast <1 x i64> %{{.*}} to i64
+  // CHECK: [[INS:%.*]] = insertelement <2 x i64> undef, i64 [[CAST]], i32 0
+  // CHECK: insertelement <2 x i64> [[INS]], i64 0, i32 1
+  return _mm_movpi64_epi64(A);
+}
+
 __m128i test_mm_move_epi64(__m128i A) {
   // CHECK-LABEL: test_mm_move_epi64
   // CHECK: shufflevector <2 x i64> %{{.*}}, <2 x i64> %{{.*}}, <2 x i32> <i32 0, i32 2>
diff --git a/test/CodeGenCXX/pragma-loop-predicate.cpp b/test/CodeGenCXX/pragma-loop-predicate.cpp
index 74aed5d17e6..ec2161d1772 100644
--- a/test/CodeGenCXX/pragma-loop-predicate.cpp
+++ b/test/CodeGenCXX/pragma-loop-predicate.cpp
@@ -27,9 +27,50 @@ void test2(int *List, int Length) {
     List[i] = i * 2;
 }
 
+// vectorize_predicate(enable) implies vectorize(enable)
+void test3(int *List, int Length) {
+// CHECK-LABEL: @{{.*}}test3{{.*}}(
+// CHECK:       br label {{.*}}, !llvm.loop ![[LOOP3:.*]]
+
+  #pragma clang loop vectorize_predicate(enable)
+  for (int i = 0; i < Length; i++)
+    List[i] = i * 2;
+}
+
+// Check that disabling vectorization means a vectorization width of 1, and
+// also that vectorization_predicate isn't enabled.
+void test4(int *List, int Length) {
+// CHECK-LABEL: @{{.*}}test4{{.*}}(
+// CHECK: br label {{.*}}, !llvm.loop ![[LOOP4:.*]]
+
+  #pragma clang loop vectorize(disable)
+  for (int i = 0; i < Length; i++)
+    List[i] = i * 2;
+}
+
+// Check that vectorize and vectorize_predicate are disabled.
+void test5(int *List, int Length) {
+// CHECK-LABEL: @{{.*}}test5{{.*}}(
+// CHECK: br label {{.*}}, !llvm.loop ![[LOOP5:.*]]
+
+  #pragma clang loop vectorize(disable) vectorize_predicate(enable)
+  for (int i = 0; i < Length; i++)
+    List[i] = i * 2;
+}
+
+
 // CHECK:      ![[LOOP0]] = distinct !{![[LOOP0]], !3}
 // CHECK-NEXT: !3 = !{!"llvm.loop.vectorize.enable", i1 true}
+
 // CHECK-NEXT: ![[LOOP1]] = distinct !{![[LOOP1]], !5, !3}
 // CHECK-NEXT: !5 = !{!"llvm.loop.vectorize.predicate.enable", i1 true}
+
 // CHECK-NEXT: ![[LOOP2]] = distinct !{![[LOOP2]], !7, !3}
 // CHECK-NEXT: !7 = !{!"llvm.loop.vectorize.predicate.enable", i1 false}
+
+// CHECK-NEXT: ![[LOOP3]] = distinct !{![[LOOP3]], !5, !3}
+
+// CHECK-NEXT: ![[LOOP4]] = distinct !{![[LOOP4]], !10}
+// CHECK-NEXT: !10 = !{!"llvm.loop.vectorize.width", i32 1}
+
+// CHECK-NEXT: ![[LOOP5]] = distinct !{![[LOOP5]], !10}
diff --git a/test/Driver/amdgpu-mcpu.cl b/test/Driver/amdgpu-mcpu.cl
index 1559c7fafc7..3aa507618bf 100644
--- a/test/Driver/amdgpu-mcpu.cl
+++ b/test/Driver/amdgpu-mcpu.cl
@@ -52,6 +52,7 @@
 // AMDGCN-based processors.
 //
 
+// RUN: %clang -### -target amdgcn %s 2>&1 | FileCheck --check-prefix=GCNDEFAULT %s
 // RUN: %clang -### -target amdgcn -mcpu=gfx600 %s 2>&1 | FileCheck --check-prefix=GFX600 %s
 // RUN: %clang -### -target amdgcn -mcpu=tahiti %s 2>&1 | FileCheck --check-prefix=TAHITI %s
 // RUN: %clang -### -target amdgcn -mcpu=gfx601 %s 2>&1 | FileCheck --check-prefix=GFX601 %s
@@ -90,6 +91,7 @@
 // RUN: %clang -### -target amdgcn -mcpu=gfx1011 %s 2>&1 | FileCheck --check-prefix=GFX1011 %s
 // RUN: %clang -### -target amdgcn -mcpu=gfx1012 %s 2>&1 | FileCheck --check-prefix=GFX1012 %s
 
+// GCNDEFAULT-NOT: -target-cpu
 // GFX600:    "-target-cpu" "gfx600"
 // TAHITI:    "-target-cpu" "tahiti"
 // GFX601:    "-target-cpu" "gfx601"
diff --git a/test/Driver/clang-offload-bundler.c b/test/Driver/clang-offload-bundler.c
index 5411a30f92e..620a5553f47 100644
--- a/test/Driver/clang-offload-bundler.c
+++ b/test/Driver/clang-offload-bundler.c
@@ -1,16 +1,18 @@
 // REQUIRES: x86-registered-target
 // REQUIRES: powerpc-registered-target
+// REQUIRES: shell
+// UNSUPPORTED: ms-sdk
 
 //
 // Generate all the types of files we can bundle.
 //
-// RUN: %clang -O0 -target powerpc64le-ibm-linux-gnu %s -E -o %t.i
-// RUN: %clangxx -O0 -target powerpc64le-ibm-linux-gnu -x c++ %s -E -o %t.ii
-// RUN: %clang -O0 -target powerpc64le-ibm-linux-gnu %s -S -emit-llvm -o %t.ll
-// RUN: %clang -O0 -target powerpc64le-ibm-linux-gnu %s -c -emit-llvm -o %t.bc
-// RUN: %clang -O0 -target powerpc64le-ibm-linux-gnu %s -S -o %t.s
-// RUN: %clang -O0 -target powerpc64le-ibm-linux-gnu %s -c -o %t.o
-// RUN: %clang -O0 -target powerpc64le-ibm-linux-gnu %s -emit-ast -o %t.ast
+// RUN: %clang -O0 -target %itanium_abi_triple %s -E -o %t.i
+// RUN: %clangxx -O0 -target %itanium_abi_triple -x c++ %s -E -o %t.ii
+// RUN: %clang -O0 -target %itanium_abi_triple %s -S -emit-llvm -o %t.ll
+// RUN: %clang -O0 -target %itanium_abi_triple %s -c -emit-llvm -o %t.bc
+// RUN: %clang -O0 -target %itanium_abi_triple %s -S -o %t.s
+// RUN: %clang -O0 -target %itanium_abi_triple %s -c -o %t.o
+// RUN: %clang -O0 -target %itanium_abi_triple %s -emit-ast -o %t.ast
 
 //
 // Generate an empty file to help with the checks of empty files.
@@ -50,27 +52,27 @@
 //
 // Check errors.
 //
-// RUN: not clang-offload-bundler -type=i -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.i,%t.tgt1,%t.tgt2 -outputs=%t.bundle.i -unbundle 2>&1 | FileCheck %s --check-prefix CK-ERR1
+// RUN: not clang-offload-bundler -type=i -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.i,%t.tgt1,%t.tgt2 -outputs=%t.bundle.i -unbundle 2>&1 | FileCheck %s --check-prefix CK-ERR1
 // CK-ERR1: error: only one input file supported in unbundling mode.
 // CK-ERR1: error: number of output files and targets should match in unbundling mode.
 
-// RUN: not clang-offload-bundler -type=i -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu -inputs=%t.i,%t.tgt1,%t.tgt2 -outputs=%t.bundle.i 2>&1 | FileCheck %s --check-prefix CK-ERR2
-// RUN: not clang-offload-bundler -type=i -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.i,%t.tgt1 -outputs=%t.bundle.i 2>&1 | FileCheck %s --check-prefix CK-ERR2
+// RUN: not clang-offload-bundler -type=i -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu -inputs=%t.i,%t.tgt1,%t.tgt2 -outputs=%t.bundle.i 2>&1 | FileCheck %s --check-prefix CK-ERR2
+// RUN: not clang-offload-bundler -type=i -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.i,%t.tgt1 -outputs=%t.bundle.i 2>&1 | FileCheck %s --check-prefix CK-ERR2
 // CK-ERR2: error: number of input files and targets should match in bundling mode.
 
-// RUN: not clang-offload-bundler -type=i -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.i,%t.tgt1,%t.tgt2 -inputs=%t.bundle.i 2>&1 | FileCheck %s --check-prefix CK-ERR3
+// RUN: not clang-offload-bundler -type=i -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.i,%t.tgt1,%t.tgt2 -inputs=%t.bundle.i 2>&1 | FileCheck %s --check-prefix CK-ERR3
 // CK-ERR3: error: only one output file supported in bundling mode.
 // CK-ERR3: error: number of input files and targets should match in bundling mode.
 
-// RUN: not clang-offload-bundler -type=i -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu -outputs=%t.i,%t.tgt1,%t.tgt2 -inputs=%t.bundle.i -unbundle 2>&1 | FileCheck %s --check-prefix CK-ERR4
-// RUN: not clang-offload-bundler -type=i -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.i,%t.tgt1 -inputs=%t.bundle.i -unbundle 2>&1 | FileCheck %s --check-prefix CK-ERR4
+// RUN: not clang-offload-bundler -type=i -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu -outputs=%t.i,%t.tgt1,%t.tgt2 -inputs=%t.bundle.i -unbundle 2>&1 | FileCheck %s --check-prefix CK-ERR4
+// RUN: not clang-offload-bundler -type=i -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.i,%t.tgt1 -inputs=%t.bundle.i -unbundle 2>&1 | FileCheck %s --check-prefix CK-ERR4
 // CK-ERR4: error: number of output files and targets should match in unbundling mode.
 
-// RUN: not clang-offload-bundler -type=i -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.i,%t.tgt1,%t.tgt2.notexist -outputs=%t.bundle.i 2>&1 | FileCheck %s --check-prefix CK-ERR5
-// RUN: not clang-offload-bundler -type=i -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.i,%t.tgt1,%t.tgt2 -inputs=%t.bundle.i.notexist -unbundle 2>&1 | FileCheck %s --check-prefix CK-ERR5
+// RUN: not clang-offload-bundler -type=i -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.i,%t.tgt1,%t.tgt2.notexist -outputs=%t.bundle.i 2>&1 | FileCheck %s --check-prefix CK-ERR5
+// RUN: not clang-offload-bundler -type=i -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.i,%t.tgt1,%t.tgt2 -inputs=%t.bundle.i.notexist -unbundle 2>&1 | FileCheck %s --check-prefix CK-ERR5
 // CK-ERR5: error: Can't open file {{.+}}.notexist: {{N|n}}o such file or directory
 
-// RUN: not clang-offload-bundler -type=invalid -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.i,%t.tgt1,%t.tgt2 -outputs=%t.bundle.i 2>&1 | FileCheck %s --check-prefix CK-ERR6
+// RUN: not clang-offload-bundler -type=invalid -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.i,%t.tgt1,%t.tgt2 -outputs=%t.bundle.i 2>&1 | FileCheck %s --check-prefix CK-ERR6
 // CK-ERR6: error: invalid file type specified.
 
 // RUN: not clang-offload-bundler 2>&1 | FileCheck %s --check-prefix CK-ERR7
@@ -85,28 +87,28 @@
 // CK-ERR8: error: invalid target 'xpenmp-x86_xx-pc-linux-gnu', unknown offloading kind 'xpenmp', unknown target triple 'x86_xx-pc-linux-gnu'.
 
 // RUN: not clang-offload-bundler -type=i -targets=openmp-powerpc64le-linux,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.i,%t.tgt1,%t.tgt2 -outputs=%t.bundle.i 2>&1 | FileCheck %s --check-prefix CK-ERR9A
-// RUN: not clang-offload-bundler -type=i -targets=host-powerpc64le-ibm-linux-gnu,host-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.i,%t.tgt1,%t.tgt2 -outputs=%t.bundle.i 2>&1 | FileCheck %s --check-prefix CK-ERR9B
+// RUN: not clang-offload-bundler -type=i -targets=host-%itanium_abi_triple,host-%itanium_abi_triple,openmp-x86_64-pc-linux-gnu -inputs=%t.i,%t.tgt1,%t.tgt2 -outputs=%t.bundle.i 2>&1 | FileCheck %s --check-prefix CK-ERR9B
 // CK-ERR9A: error: expecting exactly one host target but got 0.
 // CK-ERR9B: error: expecting exactly one host target but got 2.
 
 //
 // Check text bundle. This is a readable format, so we check for the format we expect to find.
 //
-// RUN: clang-offload-bundler -type=i -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.i,%t.tgt1,%t.tgt2 -outputs=%t.bundle3.i
-// RUN: clang-offload-bundler -type=ii -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.ii,%t.tgt1,%t.tgt2 -outputs=%t.bundle3.ii
-// RUN: clang-offload-bundler -type=ll -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.ll,%t.tgt1,%t.tgt2 -outputs=%t.bundle3.ll
-// RUN: clang-offload-bundler -type=s -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.s,%t.tgt1,%t.tgt2 -outputs=%t.bundle3.s
-// RUN: clang-offload-bundler -type=s -targets=openmp-powerpc64le-ibm-linux-gnu,host-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.tgt1,%t.s,%t.tgt2 -outputs=%t.bundle3.unordered.s
+// RUN: clang-offload-bundler -type=i -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.i,%t.tgt1,%t.tgt2 -outputs=%t.bundle3.i
+// RUN: clang-offload-bundler -type=ii -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.ii,%t.tgt1,%t.tgt2 -outputs=%t.bundle3.ii
+// RUN: clang-offload-bundler -type=ll -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.ll,%t.tgt1,%t.tgt2 -outputs=%t.bundle3.ll
+// RUN: clang-offload-bundler -type=s -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.s,%t.tgt1,%t.tgt2 -outputs=%t.bundle3.s
+// RUN: clang-offload-bundler -type=s -targets=openmp-powerpc64le-ibm-linux-gnu,host-%itanium_abi_triple,openmp-x86_64-pc-linux-gnu -inputs=%t.tgt1,%t.s,%t.tgt2 -outputs=%t.bundle3.unordered.s
 // RUN: FileCheck %s --input-file %t.bundle3.i --check-prefix CK-TEXTI
 // RUN: FileCheck %s --input-file %t.bundle3.ii --check-prefix CK-TEXTI
 // RUN: FileCheck %s --input-file %t.bundle3.ll --check-prefix CK-TEXTLL
 // RUN: FileCheck %s --input-file %t.bundle3.s --check-prefix CK-TEXTS
 // RUN: FileCheck %s --input-file %t.bundle3.unordered.s --check-prefix CK-TEXTS-UNORDERED
 
-// CK-TEXTI: // __CLANG_OFFLOAD_BUNDLE____START__ host-powerpc64le-ibm-linux-gnu
+// CK-TEXTI: // __CLANG_OFFLOAD_BUNDLE____START__ host-[[HOST:.+]]
 // CK-TEXTI: int A = 0;
 // CK-TEXTI: test_func(void)
-// CK-TEXTI: // __CLANG_OFFLOAD_BUNDLE____END__ host-powerpc64le-ibm-linux-gnu
+// CK-TEXTI: // __CLANG_OFFLOAD_BUNDLE____END__ host-[[HOST]]
 // CK-TEXTI: // __CLANG_OFFLOAD_BUNDLE____START__ openmp-powerpc64le-ibm-linux-gnu
 // CK-TEXTI: Content of device file 1
 // CK-TEXTI: // __CLANG_OFFLOAD_BUNDLE____END__ openmp-powerpc64le-ibm-linux-gnu
@@ -114,10 +116,10 @@
 // CK-TEXTI: Content of device file 2
 // CK-TEXTI: // __CLANG_OFFLOAD_BUNDLE____END__ openmp-x86_64-pc-linux-gnu
 
-// CK-TEXTLL: ; __CLANG_OFFLOAD_BUNDLE____START__ host-powerpc64le-ibm-linux-gnu
+// CK-TEXTLL: ; __CLANG_OFFLOAD_BUNDLE____START__ host-[[HOST:.+]]
 // CK-TEXTLL: @A = dso_local global i32 0
 // CK-TEXTLL: define {{.*}}@test_func()
-// CK-TEXTLL: ; __CLANG_OFFLOAD_BUNDLE____END__ host-powerpc64le-ibm-linux-gnu
+// CK-TEXTLL: ; __CLANG_OFFLOAD_BUNDLE____END__ host-[[HOST]]
 // CK-TEXTLL: ; __CLANG_OFFLOAD_BUNDLE____START__ openmp-powerpc64le-ibm-linux-gnu
 // CK-TEXTLL: Content of device file 1
 // CK-TEXTLL: ; __CLANG_OFFLOAD_BUNDLE____END__ openmp-powerpc64le-ibm-linux-gnu
@@ -125,10 +127,10 @@
 // CK-TEXTLL: Content of device file 2
 // CK-TEXTLL: ; __CLANG_OFFLOAD_BUNDLE____END__ openmp-x86_64-pc-linux-gnu
 
-// CK-TEXTS: # __CLANG_OFFLOAD_BUNDLE____START__ host-powerpc64le-ibm-linux-gnu
+// CK-TEXTS: # __CLANG_OFFLOAD_BUNDLE____START__ host-[[HOST:.+]]
 // CK-TEXTS: .globl {{.*}}test_func
 // CK-TEXTS: .globl {{.*}}A
-// CK-TEXTS: # __CLANG_OFFLOAD_BUNDLE____END__ host-powerpc64le-ibm-linux-gnu
+// CK-TEXTS: # __CLANG_OFFLOAD_BUNDLE____END__ host-[[HOST]]
 // CK-TEXTS: # __CLANG_OFFLOAD_BUNDLE____START__ openmp-powerpc64le-ibm-linux-gnu
 // CK-TEXTS: Content of device file 1
 // CK-TEXTS: # __CLANG_OFFLOAD_BUNDLE____END__ openmp-powerpc64le-ibm-linux-gnu
@@ -139,10 +141,10 @@
 // CK-TEXTS-UNORDERED: # __CLANG_OFFLOAD_BUNDLE____START__ openmp-powerpc64le-ibm-linux-gnu
 // CK-TEXTS-UNORDERED: Content of device file 1
 // CK-TEXTS-UNORDERED: # __CLANG_OFFLOAD_BUNDLE____END__ openmp-powerpc64le-ibm-linux-gnu
-// CK-TEXTS-UNORDERED: # __CLANG_OFFLOAD_BUNDLE____START__ host-powerpc64le-ibm-linux-gnu
+// CK-TEXTS-UNORDERED: # __CLANG_OFFLOAD_BUNDLE____START__ host-[[HOST:.+]]
 // CK-TEXTS-UNORDERED: .globl {{.*}}test_func
 // CK-TEXTS-UNORDERED: .globl {{.*}}A
-// CK-TEXTS-UNORDERED: # __CLANG_OFFLOAD_BUNDLE____END__ host-powerpc64le-ibm-linux-gnu
+// CK-TEXTS-UNORDERED: # __CLANG_OFFLOAD_BUNDLE____END__ host-[[HOST]]
 // CK-TEXTS-UNORDERED: # __CLANG_OFFLOAD_BUNDLE____START__ openmp-x86_64-pc-linux-gnu
 // CK-TEXTS-UNORDERED: Content of device file 2
 // CK-TEXTS-UNORDERED: # __CLANG_OFFLOAD_BUNDLE____END__ openmp-x86_64-pc-linux-gnu
@@ -150,33 +152,33 @@
 //
 // Check text unbundle. Check if we get the exact same content that we bundled before for each file.
 //
-// RUN: clang-offload-bundler -type=i -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.i,%t.res.tgt1,%t.res.tgt2 -inputs=%t.bundle3.i -unbundle
+// RUN: clang-offload-bundler -type=i -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.i,%t.res.tgt1,%t.res.tgt2 -inputs=%t.bundle3.i -unbundle
 // RUN: diff %t.i %t.res.i
 // RUN: diff %t.tgt1 %t.res.tgt1
 // RUN: diff %t.tgt2 %t.res.tgt2
-// RUN: clang-offload-bundler -type=ii -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.ii,%t.res.tgt1,%t.res.tgt2 -inputs=%t.bundle3.ii -unbundle
+// RUN: clang-offload-bundler -type=ii -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.ii,%t.res.tgt1,%t.res.tgt2 -inputs=%t.bundle3.ii -unbundle
 // RUN: diff %t.ii %t.res.ii
 // RUN: diff %t.tgt1 %t.res.tgt1
 // RUN: diff %t.tgt2 %t.res.tgt2
-// RUN: clang-offload-bundler -type=ll -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.ll,%t.res.tgt1,%t.res.tgt2 -inputs=%t.bundle3.ll -unbundle
+// RUN: clang-offload-bundler -type=ll -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.ll,%t.res.tgt1,%t.res.tgt2 -inputs=%t.bundle3.ll -unbundle
 // RUN: diff %t.ll %t.res.ll
 // RUN: diff %t.tgt1 %t.res.tgt1
 // RUN: diff %t.tgt2 %t.res.tgt2
-// RUN: clang-offload-bundler -type=s -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.s,%t.res.tgt1,%t.res.tgt2 -inputs=%t.bundle3.s -unbundle
+// RUN: clang-offload-bundler -type=s -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.s,%t.res.tgt1,%t.res.tgt2 -inputs=%t.bundle3.s -unbundle
 // RUN: diff %t.s %t.res.s
 // RUN: diff %t.tgt1 %t.res.tgt1
 // RUN: diff %t.tgt2 %t.res.tgt2
-// RUN: clang-offload-bundler -type=s -targets=openmp-powerpc64le-ibm-linux-gnu,host-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.tgt1,%t.res.s,%t.res.tgt2 -inputs=%t.bundle3.s -unbundle
+// RUN: clang-offload-bundler -type=s -targets=openmp-powerpc64le-ibm-linux-gnu,host-%itanium_abi_triple,openmp-x86_64-pc-linux-gnu -outputs=%t.res.tgt1,%t.res.s,%t.res.tgt2 -inputs=%t.bundle3.s -unbundle
 // RUN: diff %t.s %t.res.s
 // RUN: diff %t.tgt1 %t.res.tgt1
 // RUN: diff %t.tgt2 %t.res.tgt2
 
 // Check if we can unbundle a file with no magic strings.
-// RUN: clang-offload-bundler -type=s -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.s,%t.res.tgt1,%t.res.tgt2 -inputs=%t.s -unbundle
+// RUN: clang-offload-bundler -type=s -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.s,%t.res.tgt1,%t.res.tgt2 -inputs=%t.s -unbundle
 // RUN: diff %t.s %t.res.s
 // RUN: diff %t.empty %t.res.tgt1
 // RUN: diff %t.empty %t.res.tgt2
-// RUN: clang-offload-bundler -type=s -targets=openmp-powerpc64le-ibm-linux-gnu,host-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.tgt1,%t.res.s,%t.res.tgt2 -inputs=%t.s -unbundle
+// RUN: clang-offload-bundler -type=s -targets=openmp-powerpc64le-ibm-linux-gnu,host-%itanium_abi_triple,openmp-x86_64-pc-linux-gnu -outputs=%t.res.tgt1,%t.res.s,%t.res.tgt2 -inputs=%t.s -unbundle
 // RUN: diff %t.s %t.res.s
 // RUN: diff %t.empty %t.res.tgt1
 // RUN: diff %t.empty %t.res.tgt2
@@ -184,37 +186,37 @@
 //
 // Check binary bundle/unbundle. The content that we have before bundling must be the same we have after unbundling.
 //
-// RUN: clang-offload-bundler -type=bc -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.bc,%t.tgt1,%t.tgt2 -outputs=%t.bundle3.bc
-// RUN: clang-offload-bundler -type=gch -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.ast,%t.tgt1,%t.tgt2 -outputs=%t.bundle3.gch
-// RUN: clang-offload-bundler -type=ast -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.ast,%t.tgt1,%t.tgt2 -outputs=%t.bundle3.ast
-// RUN: clang-offload-bundler -type=ast -targets=openmp-powerpc64le-ibm-linux-gnu,host-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.tgt1,%t.ast,%t.tgt2 -outputs=%t.bundle3.unordered.ast
-// RUN: clang-offload-bundler -type=bc -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.bc,%t.res.tgt1,%t.res.tgt2 -inputs=%t.bundle3.bc -unbundle
+// RUN: clang-offload-bundler -type=bc -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.bc,%t.tgt1,%t.tgt2 -outputs=%t.bundle3.bc
+// RUN: clang-offload-bundler -type=gch -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.ast,%t.tgt1,%t.tgt2 -outputs=%t.bundle3.gch
+// RUN: clang-offload-bundler -type=ast -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.ast,%t.tgt1,%t.tgt2 -outputs=%t.bundle3.ast
+// RUN: clang-offload-bundler -type=ast -targets=openmp-powerpc64le-ibm-linux-gnu,host-%itanium_abi_triple,openmp-x86_64-pc-linux-gnu -inputs=%t.tgt1,%t.ast,%t.tgt2 -outputs=%t.bundle3.unordered.ast
+// RUN: clang-offload-bundler -type=bc -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.bc,%t.res.tgt1,%t.res.tgt2 -inputs=%t.bundle3.bc -unbundle
 // RUN: diff %t.bc %t.res.bc
 // RUN: diff %t.tgt1 %t.res.tgt1
 // RUN: diff %t.tgt2 %t.res.tgt2
-// RUN: clang-offload-bundler -type=gch -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.gch,%t.res.tgt1,%t.res.tgt2 -inputs=%t.bundle3.gch -unbundle
+// RUN: clang-offload-bundler -type=gch -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.gch,%t.res.tgt1,%t.res.tgt2 -inputs=%t.bundle3.gch -unbundle
 // RUN: diff %t.ast %t.res.gch
 // RUN: diff %t.tgt1 %t.res.tgt1
 // RUN: diff %t.tgt2 %t.res.tgt2
-// RUN: clang-offload-bundler -type=ast -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.ast,%t.res.tgt1,%t.res.tgt2 -inputs=%t.bundle3.ast -unbundle
+// RUN: clang-offload-bundler -type=ast -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.ast,%t.res.tgt1,%t.res.tgt2 -inputs=%t.bundle3.ast -unbundle
 // RUN: diff %t.ast %t.res.ast
 // RUN: diff %t.tgt1 %t.res.tgt1
 // RUN: diff %t.tgt2 %t.res.tgt2
-// RUN: clang-offload-bundler -type=ast -targets=openmp-powerpc64le-ibm-linux-gnu,host-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.tgt1,%t.res.ast,%t.res.tgt2 -inputs=%t.bundle3.ast -unbundle
+// RUN: clang-offload-bundler -type=ast -targets=openmp-powerpc64le-ibm-linux-gnu,host-%itanium_abi_triple,openmp-x86_64-pc-linux-gnu -outputs=%t.res.tgt1,%t.res.ast,%t.res.tgt2 -inputs=%t.bundle3.ast -unbundle
 // RUN: diff %t.ast %t.res.ast
 // RUN: diff %t.tgt1 %t.res.tgt1
 // RUN: diff %t.tgt2 %t.res.tgt2
-// RUN: clang-offload-bundler -type=ast -targets=openmp-powerpc64le-ibm-linux-gnu,host-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.tgt1,%t.res.ast,%t.res.tgt2 -inputs=%t.bundle3.unordered.ast -unbundle
+// RUN: clang-offload-bundler -type=ast -targets=openmp-powerpc64le-ibm-linux-gnu,host-%itanium_abi_triple,openmp-x86_64-pc-linux-gnu -outputs=%t.res.tgt1,%t.res.ast,%t.res.tgt2 -inputs=%t.bundle3.unordered.ast -unbundle
 // RUN: diff %t.ast %t.res.ast
 // RUN: diff %t.tgt1 %t.res.tgt1
 // RUN: diff %t.tgt2 %t.res.tgt2
 
 // Check if we can unbundle a file with no magic strings.
-// RUN: clang-offload-bundler -type=bc -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.bc,%t.res.tgt1,%t.res.tgt2 -inputs=%t.bc -unbundle
+// RUN: clang-offload-bundler -type=bc -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.bc,%t.res.tgt1,%t.res.tgt2 -inputs=%t.bc -unbundle
 // RUN: diff %t.bc %t.res.bc
 // RUN: diff %t.empty %t.res.tgt1
 // RUN: diff %t.empty %t.res.tgt2
-// RUN: clang-offload-bundler -type=bc -targets=openmp-powerpc64le-ibm-linux-gnu,host-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.tgt1,%t.res.bc,%t.res.tgt2 -inputs=%t.bc -unbundle
+// RUN: clang-offload-bundler -type=bc -targets=openmp-powerpc64le-ibm-linux-gnu,host-%itanium_abi_triple,openmp-x86_64-pc-linux-gnu -outputs=%t.res.tgt1,%t.res.bc,%t.res.tgt2 -inputs=%t.bc -unbundle
 // RUN: diff %t.bc %t.res.bc
 // RUN: diff %t.empty %t.res.tgt1
 // RUN: diff %t.empty %t.res.tgt2
@@ -227,28 +229,29 @@
 // tests.
 //
 
-// RUN: clang-offload-bundler -type=o -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.o,%t.tgt1,%t.tgt2 -outputs=%t.bundle3.o -### -dump-temporary-files 2>&1 \
+// RUN: clang-offload-bundler -type=o -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.o,%t.tgt1,%t.tgt2 -outputs=%t.bundle3.o -### -dump-temporary-files 2>&1 \
 // RUN: | FileCheck %s --check-prefix CK-OBJ-CMD
-// CK-OBJ-CMD: private constant [1 x i8] zeroinitializer, section "__CLANG_OFFLOAD_BUNDLE__host-powerpc64le-ibm-linux-gnu"
+// CK-OBJ-CMD: private constant [{{[0-9]+}} x i8] c"{{.+}}", section "__CLANG_OFFLOAD_BUNDLE__host-[[HOST:.+]]"
 // CK-OBJ-CMD: private constant [{{[0-9]+}} x i8] c"Content of device file 1{{.+}}", section "__CLANG_OFFLOAD_BUNDLE__openmp-powerpc64le-ibm-linux-gnu"
 // CK-OBJ-CMD: private constant [{{[0-9]+}} x i8] c"Content of device file 2{{.+}}", section "__CLANG_OFFLOAD_BUNDLE__openmp-x86_64-pc-linux-gnu"
-// CK-OBJ-CMD: clang{{(.exe)?}}" "-r" "-target" "powerpc64le-ibm-linux-gnu" "-o" "{{.+}}.o" "{{.+}}.o" "{{.+}}.bc" "-nostdlib"
+// CK-OBJ-CMD: clang{{(.exe)?}}" "-r" "-target" "[[HOST]]" "-o" "{{.+}}.o" "{{.+}}.o" "{{.+}}.bc" "-nostdlib"
 
-// RUN: clang-offload-bundler -type=o -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.o,%t.res.tgt1,%t.res.tgt2 -inputs=%s.o -unbundle
-// RUN: diff %s.o %t.res.o
+// RUN: clang-offload-bundler -type=o -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -inputs=%t.o,%t.tgt1,%t.tgt2 -outputs=%t.bundle3.o
+// RUN: clang-offload-bundler -type=o -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.o,%t.res.tgt1,%t.res.tgt2 -inputs=%t.bundle3.o -unbundle
+// RUN: diff %t.o %t.res.o
 // RUN: diff %t.tgt1 %t.res.tgt1
 // RUN: diff %t.tgt2 %t.res.tgt2
-// RUN: clang-offload-bundler -type=o -targets=openmp-powerpc64le-ibm-linux-gnu,host-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.tgt1,%t.res.o,%t.res.tgt2 -inputs=%s.o -unbundle
-// RUN: diff %s.o %t.res.o
+// RUN: clang-offload-bundler -type=o -targets=openmp-powerpc64le-ibm-linux-gnu,host-%itanium_abi_triple,openmp-x86_64-pc-linux-gnu -outputs=%t.res.tgt1,%t.res.o,%t.res.tgt2 -inputs=%t.bundle3.o -unbundle
+// RUN: diff %t.o %t.res.o
 // RUN: diff %t.tgt1 %t.res.tgt1
 // RUN: diff %t.tgt2 %t.res.tgt2
 
 // Check if we can unbundle a file with no magic strings.
-// RUN: clang-offload-bundler -type=o -targets=host-powerpc64le-ibm-linux-gnu,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.o,%t.res.tgt1,%t.res.tgt2 -inputs=%t.o -unbundle
+// RUN: clang-offload-bundler -type=o -targets=host-%itanium_abi_triple,openmp-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.o,%t.res.tgt1,%t.res.tgt2 -inputs=%t.o -unbundle
 // RUN: diff %t.o %t.res.o
 // RUN: diff %t.empty %t.res.tgt1
 // RUN: diff %t.empty %t.res.tgt2
-// RUN: clang-offload-bundler -type=o -targets=openmp-powerpc64le-ibm-linux-gnu,host-powerpc64le-ibm-linux-gnu,openmp-x86_64-pc-linux-gnu -outputs=%t.res.tgt1,%t.res.o,%t.res.tgt2 -inputs=%t.o -unbundle
+// RUN: clang-offload-bundler -type=o -targets=openmp-powerpc64le-ibm-linux-gnu,host-%itanium_abi_triple,openmp-x86_64-pc-linux-gnu -outputs=%t.res.tgt1,%t.res.o,%t.res.tgt2 -inputs=%t.o -unbundle
 // RUN: diff %t.o %t.res.o
 // RUN: diff %t.empty %t.res.tgt1
 // RUN: diff %t.empty %t.res.tgt2
diff --git a/test/Driver/clang-offload-bundler.c.o b/test/Driver/clang-offload-bundler.c.o
deleted file mode 100644
index 2cdba14ecb2..00000000000
Binary files a/test/Driver/clang-offload-bundler.c.o and /dev/null differ
diff --git a/test/OpenMP/for_codegen.cpp b/test/OpenMP/for_codegen.cpp
index 47c5be9becf..80afa1c9452 100644
--- a/test/OpenMP/for_codegen.cpp
+++ b/test/OpenMP/for_codegen.cpp
@@ -25,10 +25,176 @@
 
 // CHECK-LABEL: loop_with_counter_collapse
 void loop_with_counter_collapse() {
+  // Captured initializations.
+  // CHECK: store i32 0, i32* [[I_TMP:%.+]],
+  // CHECK: [[VAL:%.+]] = load i32, i32* [[I_TMP]],
+  // CHECK: store i32 [[VAL]], i32* [[J_LB_MIN:%.+]],
+  // CHECK: store i32 3, i32* [[I_TMP]],
+  // CHECK: [[VAL:%.+]] = load i32, i32* [[I_TMP]],
+  // CHECK: store i32 [[VAL]], i32* [[J_LB_MAX:%.+]],
+  // CHECK: [[J_LB_MIN_VAL:%.+]] = load i32, i32* [[J_LB_MIN]],
+  // CHECK: [[J_LB_MAX_VAL:%.+]] = load i32, i32* [[J_LB_MAX]],
+  // CHECK: [[CMP:%.+]] = icmp slt i32 [[J_LB_MIN_VAL]], [[J_LB_MAX_VAL]]
+  // CHECK: [[BOOL:%.+]] = zext i1 [[CMP]] to i8
+  // CHECK: store i8 [[BOOL]], i8* [[J_LB_CMP:%.+]],
+  // CHECK: store i32 0, i32* [[I_TMP]],
+  // CHECK: [[VAL:%.+]] = load i32, i32* [[I_TMP]],
+  // CHECK: [[J_UB_MIN_VAL:%.+]] = add nsw i32 4, [[VAL]]
+  // CHECK: store i32 [[J_UB_MIN_VAL]], i32* [[J_UB_MIN:%.+]],
+  // CHECK: store i32 3, i32* [[I_TMP]],
+  // CHECK: [[VAL:%.+]] = load i32, i32* [[I_TMP]],
+  // CHECK: [[J_UB_MAX_VAL:%.+]] = add nsw i32 4, [[VAL]]
+  // CHECK: store i32 [[J_UB_MAX_VAL]], i32* [[J_UB_MAX:%.+]],
+  // CHECK: [[J_UB_MIN_VAL:%.+]] = load i32, i32* [[J_UB_MIN]],
+  // CHECK: [[J_UB_MAX_VAL:%.+]] = load i32, i32* [[J_UB_MAX]],
+  // CHECK: [[CMP:%.+]] = icmp sgt i32 [[J_UB_MIN_VAL]], [[J_UB_MAX_VAL]]
+  // CHECK: [[BOOL:%.+]] = zext i1 [[CMP]] to i8
+  // CHECK: store i8 [[BOOL]], i8* [[J_UB_CMP:%.+]],
+  // CHECK: [[J_UB_CMP_VAL:%.+]] = load i8, i8* [[J_UB_CMP]],
+  // CHECK: [[BOOL:%.+]] = trunc i8 [[J_UB_CMP_VAL]] to i1
+  // CHECK: br i1 [[BOOL]], label %[[TRUE:[^,]+]], label %[[FALSE:[^,]+]]
+  // CHECK: [[TRUE]]:
+  // CHECK: [[J_UB_MIN_VAL:%.+]] = load i32, i32* [[J_UB_MIN]],
+  // CHECK: br label %[[EXIT:[^,]+]]
+  // CHECK: [[FALSE]]:
+  // CHECK: [[J_UB_MAX_VAL:%.+]] = load i32, i32* [[J_UB_MAX]],
+  // CHECK: br label %[[EXIT]]
+  // CHECK: [[EXIT]]:
+  // CHECK: [[J_UB_VAL:%.+]] = phi i32 [ [[J_UB_MIN_VAL]], %[[TRUE]] ], [ [[J_UB_MAX_VAL]], %[[FALSE]] ]
+  // CHECK: store i32 [[J_UB_VAL]], i32* [[J_UB:%.+]],
+  // CHECK: [[J_LB_CMP_VAL:%.+]] = load i8, i8* [[J_LB_CMP]],
+  // CHECK: [[BOOL:%.+]] = trunc i8 [[J_LB_CMP_VAL]] to i1
+  // CHECK: br i1 [[BOOL]], label %[[TRUE:[^,]+]], label %[[FALSE:[^,]+]]
+  // CHECK: [[TRUE]]:
+  // CHECK: [[J_LB_MIN_VAL:%.+]] = load i32, i32* [[J_LB_MIN]],
+  // CHECK: br label %[[EXIT:[^,]+]]
+  // CHECK: [[FALSE]]:
+  // CHECK: [[J_LB_MAX_VAL:%.+]] = load i32, i32* [[J_LB_MAX]],
+  // CHECK: br label %[[EXIT]]
+  // CHECK: [[EXIT]]:
+  // CHECK: [[J_LB_VAL:%.+]] = phi i32 [ [[J_LB_MIN_VAL]], %[[TRUE]] ], [ [[J_LB_MAX_VAL]], %[[FALSE]] ]
+  // CHECK: store i32 [[J_LB_VAL]], i32* [[J_LB:%.+]],
+  // CHECK: [[J_UB_VAL:%.+]] = load i32, i32* [[J_UB]],
+  // CHECK: [[J_LB_VAL:%.+]] = load i32, i32* [[J_LB]],
+  // CHECK: [[SUB:%.+]] = sub nsw i32 [[J_UB_VAL]], [[J_LB_VAL]]
+  // CHECK: [[SUB_ST:%.+]] = sub nsw i32 [[SUB]], 1
+  // CHECK: [[ADD_ST:%.+]] = add nsw i32 [[SUB_ST]], 1
+  // CHECK: [[DIV_ST:%.+]] = sdiv i32 [[ADD_ST]], 1
+  // CHECK: [[CAST:%.+]] = sext i32 [[DIV_ST]] to i64
+  // CHECK: [[MUL:%.+]] = mul nsw i64 4, [[CAST]]
+  // CHECK: [[NUM_ITERS_VAL:%.+]] = sub nsw i64 [[MUL]], 1
+  // CHECK: store i64 [[NUM_ITERS_VAL]], i64* [[NUM_ITERS:%.+]],
+
+  // Initialization
+  // CHECK: store i32 0, i32* [[I:%.+]],
+  // CHECK: [[I_INIT:%.+]] = load i32, i32* [[I]],
+  // CHECK: store i32 [[I_INIT]], i32* [[J:%.+]],
+
   // LIFETIME: call void @llvm.lifetime.end
   // LIFETIME: call void @llvm.lifetime.end
-  // CHECK: call void @__kmpc_for_static_init_8(%struct.ident_t* @
-  // CHECK: call void @__kmpc_for_static_fini(%struct.ident_t* @
+
+  // Precondition for j counter
+  // CHECK: store i32 0, i32* [[TMP_I:%.+]],
+  // CHECK: [[J_LB_VAL:%.+]] = load i32, i32* [[TMP_I]],
+  // CHECK: [[I_VAL:%.+]] = load i32, i32* [[TMP_I]],
+  // CHECK: [[J_UB_VAL:%.+]] = add nsw i32 4, [[I_VAL]]
+  // CHECK: [[CMP:%.+]] = icmp slt i32 [[J_LB_VAL]], [[J_UB_VAL]]
+  // CHECK: br i1 [[CMP]], label %[[THEN:[^,]+]], label %[[ELSE:[^,]+]]
+
+  // CHECK: [[THEN]]:
+  // CHECK: store i64 0, i64* [[LB:%.+]],
+  // CHECK: [[NUM_ITERS_VAL:%.+]] = load i64, i64* [[NUM_ITERS]],
+  // CHECK: store i64 [[NUM_ITERS_VAL]], i64* [[UB:%.+]],
+  // CHECK: store i64 1, i64* [[STRIDE:%.+]],
+  // CHECK: store i32 0, i32* [[IS_LAST:%.+]],
+  // CHECK: call void @__kmpc_for_static_init_8(%struct.ident_t* @{{.+}}, i32 %{{.+}}, i32 34, i32* [[IS_LAST]], i64* [[LB]], i64* [[UB]], i64* [[STRIDE]], i64 1, i64 1)
+  // CHECK: [[UB_VAL:%.+]] = load i64, i64* [[UB]],
+  // CHECK: [[NUM_ITERS_VAL:%.+]] = load i64, i64* [[NUM_ITERS]],
+  // CHECK: [[CMP:%.+]] = icmp sgt i64 [[UB_VAL]], [[NUM_ITERS_VAL]]
+  // CHECK: br i1 [[CMP]], label %[[TRUE:[^,]+]], label %[[FALSE:[^,]+]]
+  // CHECK: [[TRUE]]:
+  // CHECK: [[NUM_ITERS_VAL:%.+]] = load i64, i64* [[NUM_ITERS]],
+  // CHECK: br label %[[DONE:[^,]+]]
+  // CHECK: [[FALSE]]:
+  // CHECK: [[UB_VAL:%.+]] = load i64, i64* [[UB]],
+  // CHECK: br label %[[DONE]]
+  // CHECK: [[DONE]]:
+  // CHECK: [[TOP:%.+]] = phi i64 [ [[NUM_ITERS_VAL]], %[[TRUE]] ], [ [[UB_VAL]], %[[FALSE]] ]
+  // CHECK: store i64 [[TOP]], i64* [[UB]],
+  // CHECK: [[LB_VAL:%.+]] = load i64, i64* [[LB]],
+  // CHECK: store i64 [[LB_VAL]], i64* [[IV:%.+]],
+  // CHECK: br label %[[COND:[^,]+]]
+  // CHECK: [[COND]]:
+  // CHECK: [[IV_VAL:%.+]] = load i64, i64* [[IV]],
+  // CHECK: [[UB_VAL:%.+]] = load i64, i64* [[UB]],
+  // CHECK: [[CMP:%.+]] = icmp sle i64 [[IV_VAL]], [[UB_VAL]]
+  // CHECK: br i1 [[CMP]], label %[[BODY:[^,]+]], label %[[CLEANUP:[^,]+]]
+  // LIFETIME: [[CLEANUP]]:
+  // LIFETIME: br label %[[CLEANUP:[^,]+]]
+  // CHECK: [[BODY]]:
+  // CHECK: [[IV_VAL:%.+]] = load i64, i64* [[IV]],
+  // CHECK: [[J_UB_VAL:%.+]] = load i32, i32* [[J_UB]],
+  // CHECK: [[J_LB_VAL:%.+]] = load i32, i32* [[J_LB]],
+  // CHECK: [[SUB:%.+]] = sub nsw i32 [[J_UB_VAL]], [[J_LB_VAL]]
+  // CHECK: [[SUB_ST:%.+]] = sub nsw i32 [[SUB]], 1
+  // CHECK: [[ADD_ST:%.+]] = add nsw i32 [[SUB_ST]], 1
+  // CHECK: [[DIV_ST:%.+]] = sdiv i32 [[ADD_ST]], 1
+  // CHECK: [[MUL:%.+]] = mul nsw i32 1, [[DIV_ST]]
+  // CHECK: [[CAST:%.+]] = sext i32 [[MUL]] to i64
+  // CHECK: [[DIV:%.+]] = sdiv i64 [[IV_VAL]], [[CAST]]
+  // CHECK: [[MUL:%.+]] = mul nsw i64 [[DIV]], 1
+  // CHECK: [[ADD:%.+]] = add nsw i64 0, [[MUL]]
+  // CHECK: [[CAST:%.+]] = trunc i64 [[ADD]] to i32
+  // CHECK: store i32 [[CAST]], i32* [[I_PRIV:%.+]],
+  // CHECK: [[I_VAL:%.+]] = load i32, i32* [[I_PRIV]],
+  // CHECK: [[CONV:%.+]] = sext i32 [[I_VAL]] to i64
+  // CHECK: [[IV_VAL:%.+]] = load i64, i64* [[IV]],
+  // CHECK: [[IV_VAL1:%.+]] = load i64, i64* [[IV]],
+  // CHECK: [[J_UB_VAL:%.+]] = load i32, i32* [[J_UB]],
+  // CHECK: [[J_LB_VAL:%.+]] = load i32, i32* [[J_LB]],
+  // CHECK: [[SUB:%.+]] = sub nsw i32 [[J_UB_VAL]], [[J_LB_VAL]]
+  // CHECK: [[SUB_ST:%.+]] = sub nsw i32 [[SUB]], 1
+  // CHECK: [[ADD_ST:%.+]] = add nsw i32 [[SUB_ST]], 1
+  // CHECK: [[DIV_ST:%.+]] = sdiv i32 [[ADD_ST]], 1
+  // CHECK: [[MUL:%.+]] = mul nsw i32 1, [[DIV_ST]]
+  // CHECK: [[CAST:%.+]] = sext i32 [[MUL]] to i64
+  // CHECK: [[DIV:%.+]] = sdiv i64 [[IV_VAL1]], [[CAST]]
+  // CHECK: [[J_UB_VAL:%.+]] = load i32, i32* [[J_UB]],
+  // CHECK: [[J_LB_VAL:%.+]] = load i32, i32* [[J_LB]],
+  // CHECK: [[SUB:%.+]] = sub nsw i32 [[J_UB_VAL]], [[J_LB_VAL]]
+  // CHECK: [[SUB_ST:%.+]] = sub nsw i32 [[SUB]], 1
+  // CHECK: [[ADD_ST:%.+]] = add nsw i32 [[SUB_ST]], 1
+  // CHECK: [[DIV_ST:%.+]] = sdiv i32 [[ADD_ST]], 1
+  // CHECK: [[MUL:%.+]] = mul nsw i32 1, [[DIV_ST]]
+  // CHECK: [[CAST:%.+]] = sext i32 [[MUL]] to i64
+  // CHECK: [[MUL:%.+]] = mul nsw i64 [[DIV]], [[CAST]]
+  // CHECK: [[SUB:%.+]] = sub nsw i64 [[IV_VAL]], [[MUL]]
+  // CHECK: [[MUL:%.+]] = mul nsw i64 [[SUB:%.+]], 1
+  // CHECK: [[ADD:%.+]] = add nsw i64 [[CONV]], [[MUL]]
+  // CHECK: [[CAST:%.+]] = trunc i64 [[ADD]] to i32
+  // CHECK: store i32 [[CAST]], i32* [[J_PRIV:%.+]],
+
+  // Check that the loop variable is not out of its boundaries.
+  // CHECK: [[J_VAL:%.+]] = load i32, i32* [[J_PRIV]],
+  // CHECK: [[I_VAL:%.+]] = load i32, i32* [[I_PRIV]],
+  // CHECK: [[J_COND:%.+]] = add nsw i32 4, [[I_VAL]]
+  // CHECK: [[CMP:%.+]] = icmp slt i32 [[J_VAL]], [[J_COND]]
+  // CHECK: br i1 [[CMP]], label %[[NEXT:[^,]+]], label %[[BODY_CONT:[^,]+]]
+  // CHECK: [[NEXT]]:
+
+  // Main body is empty.
+  // CHECK: br label %[[BODY_CONT]]
+  // CHECK: [[BODY_CONT]]:
+  // CHECK: br label %[[INC:[^,]+]]
+  // CHECK: [[INC]]:
+  // CHECK: [[IV_VAL:%.+]] = load i64, i64* [[IV]],
+  // CHECK: [[ADD:%.+]] = add nsw i64 [[IV_VAL]], 1
+  // CHECK: store i64 [[ADD]], i64* [[IV]],
+  // CHECK: br label %[[COND]]
+  // CHECK: [[CLEANUP]]:
+  // CHECK: br label %[[EXIT:[^,]+]]
+  // CHECK: [[EXIT]]:
+  // CHECK: call void @__kmpc_for_static_fini(%struct.ident_t* @{{.+}}, i32 %{{.+}})
   // LIFETIME: call void @llvm.lifetime.end
   // LIFETIME: call void @llvm.lifetime.end
   // LIFETIME: call void @llvm.lifetime.end
@@ -40,7 +206,7 @@ void loop_with_counter_collapse() {
   // LIFETIME: call void @llvm.lifetime.end
   #pragma omp for collapse(2)
   for (int i = 0; i < 4; i++) {
-    for (int j = i; j < 4; j++) {
+    for (int j = i; j < 4 + i; j++) {
     }
   }
 }
@@ -449,13 +615,14 @@ void for_with_references() {
 // CHECK: [[I:%.+]] = alloca i8,
 // CHECK: [[CNT:%.+]] = alloca i8*,
 // CHECK: [[CNT_PRIV:%.+]] = alloca i8,
-// CHECK: call void @__kmpc_for_static_init_4(
+// CHECK: call void @__kmpc_for_static_init_8(
 // CHECK-NOT: load i8, i8* [[CNT]],
 // CHECK: call void @__kmpc_for_static_fini(
   char i = 0;
   char &cnt = i;
-#pragma omp for
+#pragma omp for collapse(2)
   for (cnt = 0; cnt < 2; ++cnt)
+    for (int j = cnt; j < 4 + cnt; j++)
     k = cnt;
 }
 
@@ -528,13 +695,14 @@ void loop_with_It_plus(It<char> begin, It<char> end) {
 // CHECK: call void @__kmpc_for_static_fini(
 
 void loop_with_stmt_expr() {
-#pragma omp for
+#pragma omp for collapse(2)
   for (int i = __extension__({float b = 0;b; }); i < __extension__({double c = 1;c; }); i += __extension__({char d = 1; d; }))
+    for (int j = i; j < 4 + i; j++)
     ;
 }
 // CHECK-LABEL: loop_with_stmt_expr
 // CHECK: call i32 @__kmpc_global_thread_num(
-// CHECK: call void @__kmpc_for_static_init_4(
+// CHECK: call void @__kmpc_for_static_init_8(
 // CHECK: call void @__kmpc_for_static_fini(
 
 
diff --git a/test/OpenMP/for_loop_messages.cpp b/test/OpenMP/for_loop_messages.cpp
index f5f6d0b7031..37c81123e43 100644
--- a/test/OpenMP/for_loop_messages.cpp
+++ b/test/OpenMP/for_loop_messages.cpp
@@ -651,10 +651,9 @@ class TC {
         ;
 
 #pragma omp parallel
-// expected-error@+6 2 {{expected loop invariant expression or '<invariant1> * ii + <invariant2>' kind of expression}}
-// expected-error@+5 {{expected loop invariant expression or '<invariant1> * TC::ii + <invariant2>' kind of expression}}
 // expected-error@+5 2 {{expected loop invariant expression or '<invariant1> * ii + <invariant2>' kind of expression}}
 // expected-error@+4 {{expected loop invariant expression or '<invariant1> * TC::ii + <invariant2>' kind of expression}}
+// expected-error@+4 {{expected loop invariant expression or '<invariant1> * TC::ii + <invariant2>' kind of expression}}
 #pragma omp for collapse(3)
     for (ii = 10 + 25; ii < 1000; ii += 1)
       for (iii = ii * 10 + 25; iii < ii / ii - 23; iii += 1)
diff --git a/test/OpenMP/parallel_messages.cpp b/test/OpenMP/parallel_messages.cpp
index ac8869fff9d..c9b6dbc98f0 100644
--- a/test/OpenMP/parallel_messages.cpp
+++ b/test/OpenMP/parallel_messages.cpp
@@ -102,6 +102,6 @@ struct h {
 h operator<(h, h);
 void g::j() {
 #pragma omp parallel for default(none) if(a::b)
-  for (auto a = blocks.cbegin; a < blocks; ++a) // expected-error {{invalid operands to binary expression ('f' and 'int')}}
+  for (auto a = blocks.cbegin; a < blocks; ++a) // expected-error 2 {{invalid operands to binary expression ('f' and 'int')}}
     ;
 }
diff --git a/test/Sema/warn-lifetime-analysis-nocfg.cpp b/test/Sema/warn-lifetime-analysis-nocfg.cpp
index ba3cc1fc969..c8016bf55cb 100644
--- a/test/Sema/warn-lifetime-analysis-nocfg.cpp
+++ b/test/Sema/warn-lifetime-analysis-nocfg.cpp
@@ -170,7 +170,15 @@ template<typename T>
 struct optional {
   optional();
   optional(const T&);
-  T &operator*();
+  T &operator*() &;
+  T &&operator*() &&;
+  T &value() &;
+  T &&value() &&;
+};
+
+template<typename T>
+struct stack {
+  T &top();
 };
 }
 
@@ -188,6 +196,16 @@ const char *danglingRawPtrFromLocal() {
   return s.c_str(); // expected-warning {{address of stack memory associated with local variable 's' returned}}
 }
 
+int &danglingRawPtrFromLocal2() {
+  std::optional<int> o;
+  return o.value(); // expected-warning {{reference to stack memory associated with local variable 'o' returned}}
+}
+
+int &danglingRawPtrFromLocal3() {
+  std::optional<int> o;
+  return *o; // expected-warning {{reference to stack memory associated with local variable 'o' returned}}
+}
+
 const char *danglingRawPtrFromTemp() {
   return std::basic_string<char>().c_str(); // expected-warning {{returning address of local temporary object}}
 }
@@ -203,9 +221,10 @@ int *danglingUniquePtrFromTemp2() {
 }
 
 void danglingReferenceFromTempOwner() {
-  int &r = *std::optional<int>(); // expected-warning {{object backing the pointer will be destroyed at the end of the full-expression}}
-  int &r2 = *std::optional<int>(5); // expected-warning {{object backing the pointer will be destroyed at the end of the full-expression}}
-  int &r3 = std::vector<int>().at(3); // expected-warning {{object backing the pointer will be destroyed at the end of the full-expression}}
+  int &&r = *std::optional<int>();          // expected-warning {{object backing the pointer will be destroyed at the end of the full-expression}}
+  int &&r2 = *std::optional<int>(5);        // expected-warning {{object backing the pointer will be destroyed at the end of the full-expression}}
+  int &&r3 = std::optional<int>(5).value(); // expected-warning {{object backing the pointer will be destroyed at the end of the full-expression}}
+  int &r4 = std::vector<int>().at(3);       // expected-warning {{object backing the pointer will be destroyed at the end of the full-expression}}
 }
 
 std::vector<int> getTempVec();
diff --git a/test/SemaCXX/cxx2a-adl-only-template-id.cpp b/test/SemaCXX/cxx2a-adl-only-template-id.cpp
index 4bd9a22f5f4..28ecbade1b1 100644
--- a/test/SemaCXX/cxx2a-adl-only-template-id.cpp
+++ b/test/SemaCXX/cxx2a-adl-only-template-id.cpp
@@ -65,3 +65,11 @@ void xf(g<int> x); // expected-error {{variable has incomplete type 'void'}} exp
 struct B : g<int> { // expected-error {{expected class name}}
   B() : g<int>() {} // expected-error {{expected class member or base class name}}
 };
+
+namespace vector_components {
+  typedef __attribute__((__ext_vector_type__(2))) float vector_float2;
+  bool foo123(vector_float2 &A, vector_float2 &B)
+  {
+    return A.x < B.x && B.y > A.y;
+  }
+}
diff --git a/test/SemaCXX/pseudo-destructors.cpp b/test/SemaCXX/pseudo-destructors.cpp
index 08938bf34a7..fb2d0afdc3f 100644
--- a/test/SemaCXX/pseudo-destructors.cpp
+++ b/test/SemaCXX/pseudo-destructors.cpp
@@ -34,7 +34,7 @@ void f(A* a, Foo *f, int *i, double *d, int ii) {
   g().~Bar(); // expected-error{{non-scalar}}
   
   f->::~Bar();
-  f->N::~Wibble(); // FIXME: technically, Wibble isn't a class-name
+  f->N::~Wibble(); // expected-error{{'N' does not refer to a type}} expected-error{{'Wibble' does not refer to a type}}
   
   f->::~Bar(17, 42); // expected-error{{cannot have any arguments}}
 
@@ -79,7 +79,7 @@ namespace PR11339 {
   template<class T>
   void destroy(T* p) {
     p->~T(); // ok
-    p->~oops(); // expected-error{{expected the class name after '~' to name a destructor}}
+    p->~oops(); // expected-error{{identifier 'oops' in object destruction expression does not name a type}}
   }
 
   template void destroy(int*); // expected-note{{in instantiation of function template specialization}}
diff --git a/test/SemaCXX/vector.cpp b/test/SemaCXX/vector.cpp
index 40dcd35c1bb..295e1e17323 100644
--- a/test/SemaCXX/vector.cpp
+++ b/test/SemaCXX/vector.cpp
@@ -342,3 +342,19 @@ void test_vector_literal(inte4 res) {
   inte4 b = (inte4)(a, a); //expected-error{{C-style cast from vector 'inte2' (vector of 2 'int' values) to vector 'inte4' (vector of 4 'int' values) of different size}} //expected-warning{{expression result unused}}
 }
 
+typedef __attribute__((__ext_vector_type__(4))) float vector_float4;
+typedef __attribute__((__ext_vector_type__(4))) int vector_int4;
+
+namespace swizzle_template_confusion {
+  template<typename T> struct xyzw {};
+  vector_int4 foo123(vector_float4 &A, vector_float4 &B) {
+    return A.xyzw < B.x && B.y > A.y; // OK, not a template-id
+  }
+}
+
+namespace swizzle_typo_correction {
+  template<typename T> struct xyzv {};
+  vector_int4 foo123(vector_float4 &A, vector_float4 &B) {
+    return A.xyzw < B.x && B.y > A.y; // OK, not a typo for 'xyzv'
+  }
+}
diff --git a/tools/clang-offload-bundler/ClangOffloadBundler.cpp b/tools/clang-offload-bundler/ClangOffloadBundler.cpp
index 1c461145cb5..8581fff9b7c 100644
--- a/tools/clang-offload-bundler/ClangOffloadBundler.cpp
+++ b/tools/clang-offload-bundler/ClangOffloadBundler.cpp
@@ -370,13 +370,9 @@ class BinaryFileHandler final : public FileHandler {
 /// designated name.
 ///
 /// In order to bundle we create an IR file with the content of each section and
-/// use incremental linking to produce the resulting object. We also add section
-/// with a single byte to state the name of the component the main object file
-/// (the one we are bundling into) refers to.
+/// use incremental linking to produce the resulting object.
 ///
-/// To unbundle, we use just copy the contents of the designated section. If the
-/// requested bundle refer to the main object file, we just copy it with no
-/// changes.
+/// To unbundle, we just copy the contents of the designated section.
 class ObjectFileHandler final : public FileHandler {
 
   /// The object file we are currently dealing with.
@@ -471,10 +467,7 @@ class ObjectFileHandler final : public FileHandler {
       return;
     }
 
-    if (Content->size() < 2)
-      OS.write(Input.getBufferStart(), Input.getBufferSize());
-    else
-      OS.write(Content->data(), Content->size());
+    OS.write(Content->data(), Content->size());
   }
 
   void WriteHeader(raw_fd_ostream &OS,
@@ -592,22 +585,14 @@ class ObjectFileHandler final : public FileHandler {
     std::string SectionName = OFFLOAD_BUNDLER_MAGIC_STR;
     SectionName += CurrentTriple;
 
-    // Create the constant with the content of the section. For the input we are
-    // bundling into (the host input), this is just a place-holder, so a single
-    // byte is sufficient.
-    assert(HostInputIndex != ~0u && "Host input index undefined??");
-    Constant *Content;
-    if (NumberOfProcessedInputs == HostInputIndex + 1) {
-      uint8_t Byte[] = {0};
-      Content = ConstantDataArray::get(VMContext, Byte);
-    } else
-      Content = ConstantDataArray::get(
-          VMContext, ArrayRef<uint8_t>(reinterpret_cast<const uint8_t *>(
-                                           Input.getBufferStart()),
-                                       Input.getBufferSize()));
-
-    // Create the global in the desired section. We don't want these globals in
-    // the symbol table, so we mark them private.
+    // Create the constant with the content of the section.
+    auto *Content = ConstantDataArray::get(
+        VMContext, ArrayRef<uint8_t>(reinterpret_cast<const uint8_t *>(
+                                         Input.getBufferStart()),
+                                     Input.getBufferSize()));
+
+    // Create the global in the desired section. We don't want these globals
+    // in the symbol table, so we mark them private.
     auto *GV = new GlobalVariable(*M, Content->getType(), /*IsConstant=*/true,
                                   GlobalVariable::PrivateLinkage, Content);
     GV->setSection(SectionName);
diff --git a/www/cxx_dr_status.html b/www/cxx_dr_status.html
index f36db8aae96..a8172eb43d4 100755
--- a/www/cxx_dr_status.html
+++ b/www/cxx_dr_status.html
@@ -4131,11 +4131,11 @@ <h2 id="cxxdr">C++ defect report implementation status</h2>
     <td>Restrictions on declarators with late-specified return types</td>
     <td class="partial" align="center">Partial</td>
   </tr>
-  <tr class="open" id="682">
+  <tr id="682">
     <td><a href="http://wg21.link/cwg682">682</a></td>
-    <td>drafting</td>
+    <td>tentatively ready</td>
     <td>Missing description of lookup of template aliases</td>
-    <td align="center">Not resolved</td>
+    <td class="none" align="center">Unknown</td>
   </tr>
   <tr id="683">
     <td><a href="http://wg21.link/cwg683">683</a></td>
@@ -10847,7 +10847,7 @@ <h2 id="cxxdr">C++ defect report implementation status</h2>
   </tr>
   <tr class="open" id="1839">
     <td><a href="http://wg21.link/cwg1839">1839</a></td>
-    <td>drafting</td>
+    <td>review</td>
     <td>Lookup of block-scope <TT>extern</TT> declarations</td>
     <td align="center">Not resolved</td>
   </tr>
@@ -13053,11 +13053,11 @@ <h2 id="cxxdr">C++ defect report implementation status</h2>
     <td>Composite type of object and function pointers</td>
     <td class="none" align="center">Unknown</td>
   </tr>
-  <tr class="open" id="2207">
+  <tr id="2207">
     <td><a href="http://wg21.link/cwg2207">2207</a></td>
-    <td>drafting</td>
+    <td>tentatively ready</td>
     <td>Alignment of allocation function return value</td>
-    <td align="center">Not resolved</td>
+    <td class="none" align="center">Unknown</td>
   </tr>
   <tr id="2208">
     <td><a href="http://wg21.link/cwg2208">2208</a></td>
@@ -13567,7 +13567,7 @@ <h2 id="cxxdr">C++ defect report implementation status</h2>
     <td><a href="http://wg21.link/cwg2292">2292</a></td>
     <td>DRWP</td>
     <td><I>simple-template-id</I> is ambiguous between <I>class-name</I> and <I>type-name</I></td>
-    <td class="none" align="center">Unknown</td>
+    <td class="svn" align="center">SVN</td>
   </tr>
   <tr id="2293">
     <td><a href="http://wg21.link/cwg2293">2293</a></td>
@@ -13611,11 +13611,11 @@ <h2 id="cxxdr">C++ defect report implementation status</h2>
     <td><TT>constexpr</TT> vararg functions</td>
     <td class="none" align="center">Unknown</td>
   </tr>
-  <tr class="open" id="2300">
+  <tr id="2300">
     <td><a href="http://wg21.link/cwg2300">2300</a></td>
-    <td>drafting</td>
+    <td>tentatively ready</td>
     <td>Lambdas in multiple definitions</td>
-    <td align="center">Not resolved</td>
+    <td class="none" align="center">Unknown</td>
   </tr>
   <tr class="open" id="2301">
     <td><a href="http://wg21.link/cwg2301">2301</a></td>
@@ -13797,11 +13797,11 @@ <h2 id="cxxdr">C++ defect report implementation status</h2>
     <td>Missing references to variable templates</td>
     <td class="none" align="center">Unknown</td>
   </tr>
-  <tr id="2331">
+  <tr class="open" id="2331">
     <td><a href="http://wg21.link/cwg2331">2331</a></td>
-    <td>DR</td>
+    <td>drafting</td>
     <td>Redundancy in description of class scope</td>
-    <td class="none" align="center">Unknown</td>
+    <td align="center">Not resolved</td>
   </tr>
   <tr id="2332">
     <td><a href="http://wg21.link/cwg2332">2332</a></td>
@@ -13895,7 +13895,7 @@ <h2 id="cxxdr">C++ defect report implementation status</h2>
   </tr>
   <tr class="open" id="2347">
     <td><a href="http://wg21.link/cwg2347">2347</a></td>
-    <td>drafting</td>
+    <td>review</td>
     <td>Passing short scoped enumerations to ellipsis</td>
     <td align="center">Not resolved</td>
   </tr>
@@ -14007,11 +14007,11 @@ <h2 id="cxxdr">C++ defect report implementation status</h2>
     <td>Confusing specification for <TT>dynamic_cast</TT></td>
     <td class="none" align="center">Unknown</td>
   </tr>
-  <tr class="open" id="2366">
+  <tr id="2366">
     <td><a href="http://wg21.link/cwg2366">2366</a></td>
-    <td>drafting</td>
+    <td>tentatively ready</td>
     <td>Can default initialization be constant initialization?</td>
-    <td align="center">Not resolved</td>
+    <td class="none" align="center">Unknown</td>
   </tr>
   <tr class="open" id="2367">
     <td><a href="http://wg21.link/cwg2367">2367</a></td>
@@ -14057,7 +14057,7 @@ <h2 id="cxxdr">C++ defect report implementation status</h2>
   </tr>
   <tr class="open" id="2374">
     <td><a href="http://wg21.link/cwg2374">2374</a></td>
-    <td>drafting</td>
+    <td>review</td>
     <td>Overly permissive specification of <TT>enum</TT> direct-list-initialization</td>
     <td align="center">Not resolved</td>
   </tr>
@@ -14067,11 +14067,11 @@ <h2 id="cxxdr">C++ defect report implementation status</h2>
     <td>Multiple redeclarations of <TT>constexpr</TT> static data members</td>
     <td class="none" align="center">Unknown</td>
   </tr>
-  <tr class="open" id="2376">
+  <tr id="2376">
     <td><a href="http://wg21.link/cwg2376">2376</a></td>
-    <td>drafting</td>
+    <td>tentatively ready</td>
     <td>Class template argument deduction with array declarator</td>
-    <td align="center">Not resolved</td>
+    <td class="none" align="center">Unknown</td>
   </tr>
   <tr id="2377">
     <td><a href="http://wg21.link/cwg2377">2377</a></td>
@@ -14139,11 +14139,11 @@ <h2 id="cxxdr">C++ defect report implementation status</h2>
     <td>Linkage of const-qualified variable template</td>
     <td class="svn" align="center">SVN</td>
   </tr>
-  <tr class="open" id="2388">
+  <tr id="2388">
     <td><a href="http://wg21.link/cwg2388">2388</a></td>
-    <td>drafting</td>
+    <td>NAD</td>
     <td>Applicability of <I>contract-attribute-specifier</I>s</td>
-    <td align="center">Not resolved</td>
+    <td class="none" align="center">Unknown</td>
   </tr>
   <tr class="open" id="2389">
     <td><a href="http://wg21.link/cwg2389">2389</a></td>
@@ -14151,11 +14151,11 @@ <h2 id="cxxdr">C++ defect report implementation status</h2>
     <td>Agreement of deduced and explicitly-specified variable types</td>
     <td align="center">Not resolved</td>
   </tr>
-  <tr class="open" id="2390">
+  <tr id="2390">
     <td><a href="http://wg21.link/cwg2390">2390</a></td>
-    <td>drafting</td>
+    <td>tentatively ready</td>
     <td>Is the argument of <TT>__has_cpp_attribute</TT> macro-expanded?</td>
-    <td align="center">Not resolved</td>
+    <td class="none" align="center">Unknown</td>
   </tr>
   <tr id="2391">
     <td><a href="http://wg21.link/cwg2391">2391</a></td>
@@ -14207,15 +14207,15 @@ <h2 id="cxxdr">C++ defect report implementation status</h2>
   </tr>
   <tr class="open" id="2399">
     <td><a href="http://wg21.link/cwg2399">2399</a></td>
-    <td>drafting</td>
+    <td>review</td>
     <td>Unclear referent of &#8220;expression&#8221; in <I>assignment-expression</I></td>
     <td align="center">Not resolved</td>
   </tr>
-  <tr class="open" id="2400">
+  <tr id="2400">
     <td><a href="http://wg21.link/cwg2400">2400</a></td>
-    <td>drafting</td>
+    <td>tentatively ready</td>
     <td>Constexpr virtual functions and temporary objects</td>
-    <td align="center">Not resolved</td>
+    <td class="none" align="center">Unknown</td>
   </tr>
   <tr class="open" id="2401">
     <td><a href="http://wg21.link/cwg2401">2401</a></td>
@@ -14229,6 +14229,108 @@ <h2 id="cxxdr">C++ defect report implementation status</h2>
     <td>When is the restriction to a single <I>c-char</I> in a Unicode literal enforced?</td>
     <td align="center">Not resolved</td>
   </tr>
+  <tr class="open" id="2403">
+    <td><a href="http://wg21.link/cwg2403">2403</a></td>
+    <td>open</td>
+    <td>Temporary materialization and base/member initialization</td>
+    <td align="center">Not resolved</td>
+  </tr>
+  <tr id="2404">
+    <td><a href="http://wg21.link/cwg2404">2404</a></td>
+    <td>tentatively ready</td>
+    <td><TT>[[no_unique_address]]</TT> and allocation order</td>
+    <td class="none" align="center">Unknown</td>
+  </tr>
+  <tr class="open" id="2405">
+    <td><a href="http://wg21.link/cwg2405">2405</a></td>
+    <td>drafting</td>
+    <td>Additional type-dependent expressions</td>
+    <td align="center">Not resolved</td>
+  </tr>
+  <tr id="2406">
+    <td><a href="http://wg21.link/cwg2406">2406</a></td>
+    <td>tentatively ready</td>
+    <td><TT>[[fallthrough]]</TT> attribute and iteration statements</td>
+    <td class="none" align="center">Unknown</td>
+  </tr>
+  <tr class="open" id="2407">
+    <td><a href="http://wg21.link/cwg2407">2407</a></td>
+    <td>review</td>
+    <td>Missing entry in Annex C for defaulted comparison operators</td>
+    <td align="center">Not resolved</td>
+  </tr>
+  <tr class="open" id="2408">
+    <td><a href="http://wg21.link/cwg2408">2408</a></td>
+    <td>open</td>
+    <td>Temporaries and previously-initialized elements in aggregate initialization</td>
+    <td align="center">Not resolved</td>
+  </tr>
+  <tr class="open" id="2409">
+    <td><a href="http://wg21.link/cwg2409">2409</a></td>
+    <td>drafting</td>
+    <td>Explicit specializations of constexpr static data members</td>
+    <td align="center">Not resolved</td>
+  </tr>
+  <tr class="open" id="2410">
+    <td><a href="http://wg21.link/cwg2410">2410</a></td>
+    <td>review</td>
+    <td>Implicit calls of immediate functions</td>
+    <td align="center">Not resolved</td>
+  </tr>
+  <tr class="open" id="2411">
+    <td><a href="http://wg21.link/cwg2411">2411</a></td>
+    <td>open</td>
+    <td>Comparison of pointers to members in template non-type arguments</td>
+    <td align="center">Not resolved</td>
+  </tr>
+  <tr class="open" id="2412">
+    <td><a href="http://wg21.link/cwg2412">2412</a></td>
+    <td>open</td>
+    <td>SFINAE vs undeduced placeholder type</td>
+    <td align="center">Not resolved</td>
+  </tr>
+  <tr class="open" id="2413">
+    <td><a href="http://wg21.link/cwg2413">2413</a></td>
+    <td>drafting</td>
+    <td><TT>typename</TT> in <I>conversion-function-id</I>s</td>
+    <td align="center">Not resolved</td>
+  </tr>
+  <tr class="open" id="2414">
+    <td><a href="http://wg21.link/cwg2414">2414</a></td>
+    <td>drafting</td>
+    <td>Unclear results if both member and friend <TT>operator&lt;=&gt;</TT> are declared</td>
+    <td align="center">Not resolved</td>
+  </tr>
+  <tr id="2415">
+    <td><a href="http://wg21.link/cwg2415">2415</a></td>
+    <td>NAD</td>
+    <td><I>using-declaration</I>s vs copy assignment operators</td>
+    <td class="none" align="center">Unknown</td>
+  </tr>
+  <tr class="open" id="2416">
+    <td><a href="http://wg21.link/cwg2416">2416</a></td>
+    <td>open</td>
+    <td>Explicit specializations vs <TT>constexpr</TT> and <TT>consteval</TT></td>
+    <td align="center">Not resolved</td>
+  </tr>
+  <tr class="open" id="2417">
+    <td><a href="http://wg21.link/cwg2417">2417</a></td>
+    <td>open</td>
+    <td>Explicit instantiation and exception specifications</td>
+    <td align="center">Not resolved</td>
+  </tr>
+  <tr id="2418">
+    <td><a href="http://wg21.link/cwg2418">2418</a></td>
+    <td>tentatively ready</td>
+    <td>Missing cases in definition of &#8220;usable in constant expressions&#8221;</td>
+    <td class="none" align="center">Unknown</td>
+  </tr>
+  <tr class="open" id="2419">
+    <td><a href="http://wg21.link/cwg2419">2419</a></td>
+    <td>open</td>
+    <td>Loss of generality treating pointers to objects as one-element arrays</td>
+    <td align="center">Not resolved</td>
+  </tr>
 </table>
 
 </div>
diff --git a/www/index.html b/www/index.html
index 1ca24973a1d..cb89f44da12 100755
--- a/www/index.html
+++ b/www/index.html
@@ -105,6 +105,7 @@ <h2>Get it and get involved!</h2>
      interested in
      following the development of Clang, signing up for a mailing list is a good
      way to learn about how the project works.</p>
+
 </div>
 </body>
 </html>