diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 8495884dcd058f..8020be6c57bcfb 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -12733,19 +12733,19 @@ def err_acc_size_expr_value
           "OpenACC 'tile' clause size expression must be %select{an asterisk "
           "or a constant expression|positive integer value, evaluated to %1}0">;
 def err_acc_invalid_in_loop
-    : Error<"%select{OpenACC '%2' construct|while loop|do loop}0 cannot appear "
-            "in intervening code of a 'loop' with a '%1' clause">;
+    : Error<"%select{OpenACC '%3' construct|while loop|do loop}0 cannot appear "
+            "in intervening code of a '%1' with a '%2' clause">;
 def note_acc_active_clause_here
     : Note<"active '%0' clause defined here">;
 def err_acc_clause_multiple_loops
-    : Error<"more than one for-loop in a loop associated with OpenACC 'loop' "
-            "construct with a '%select{collapse|tile}0' clause">;
+    : Error<"more than one for-loop in a loop associated with OpenACC '%0' "
+            "construct with a '%1' clause">;
 def err_acc_insufficient_loops
     : Error<"'%0' clause specifies a loop count greater than the number "
             "of available loops">;
 def err_acc_intervening_code
     : Error<"inner loops must be tightly nested inside a '%0' clause on "
-            "a 'loop' construct">;
+            "a '%1' construct">;
 def err_acc_gang_multiple_elt
     : Error<"OpenACC 'gang' clause may have at most one %select{unnamed or "
             "'num'|'dim'|'static'}0 argument">;
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index e8f022362ea168..b968ff779c89b0 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -87,6 +87,10 @@ class SemaOpenACC : public SemaBase {
     /// which allows us to diagnose if the value of 'N' is too large for the
     /// current number of 'for' loops.
     bool CollapseDepthSatisfied = true;
+
+    /// Records the kind of the directive that this clause is attached to, which
+    /// allows us to use it in diagnostics.
+    OpenACCDirectiveKind DirectiveKind = OpenACCDirectiveKind::Invalid;
   } CollapseInfo;
 
   /// The 'tile' clause requires a bit of additional checking as well, so like
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 8775327a25f66a..7585bfca059567 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -1504,10 +1504,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitReductionClause(
 
 OpenACCClause *SemaOpenACCClauseVisitor::VisitCollapseClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
-  // TODO: Remove this check once we implement this for combined constructs.
-  if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()) &&
-      Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop)
-    return isNotImplemented();
   // Duplicates here are not really sensible.  We could possible permit
   // multiples if they all had the same value, but there isn't really a good
   // reason to do so. Also, this simplifies the suppression of duplicates, in
@@ -1701,6 +1697,7 @@ void SemaOpenACC::AssociatedStmtRAII::SetCollapseInfoBeforeAssociatedStmt(
   SemaRef.CollapseInfo.CollapseDepthSatisfied = false;
   SemaRef.CollapseInfo.CurCollapseCount =
       cast<ConstantExpr>(LoopCount)->getResultAsAPSInt();
+  SemaRef.CollapseInfo.DirectiveKind = DirKind;
 }
 
 void SemaOpenACC::AssociatedStmtRAII::SetTileInfoBeforeAssociatedStmt(
@@ -2597,7 +2594,8 @@ void SemaOpenACC::ActOnWhileStmt(SourceLocation WhileLoc) {
 
   if (CollapseInfo.CurCollapseCount && *CollapseInfo.CurCollapseCount > 0) {
     Diag(WhileLoc, diag::err_acc_invalid_in_loop)
-        << /*while loop*/ 1 << OpenACCClauseKind::Collapse;
+        << /*while loop*/ 1 << CollapseInfo.DirectiveKind
+        << OpenACCClauseKind::Collapse;
     assert(CollapseInfo.ActiveCollapse && "Collapse count without object?");
     Diag(CollapseInfo.ActiveCollapse->getBeginLoc(),
          diag::note_acc_active_clause_here)
@@ -2610,7 +2608,8 @@ void SemaOpenACC::ActOnWhileStmt(SourceLocation WhileLoc) {
 
   if (TileInfo.CurTileCount && *TileInfo.CurTileCount > 0) {
     Diag(WhileLoc, diag::err_acc_invalid_in_loop)
-        << /*while loop*/ 1 << OpenACCClauseKind::Tile;
+        << /*while loop*/ 1 << OpenACCDirectiveKind::Loop
+        << OpenACCClauseKind::Tile;
     assert(TileInfo.ActiveTile && "tile count without object?");
     Diag(TileInfo.ActiveTile->getBeginLoc(), diag::note_acc_active_clause_here)
         << OpenACCClauseKind::Tile;
@@ -2630,7 +2629,8 @@ void SemaOpenACC::ActOnDoStmt(SourceLocation DoLoc) {
 
   if (CollapseInfo.CurCollapseCount && *CollapseInfo.CurCollapseCount > 0) {
     Diag(DoLoc, diag::err_acc_invalid_in_loop)
-        << /*do loop*/ 2 << OpenACCClauseKind::Collapse;
+        << /*do loop*/ 2 << CollapseInfo.DirectiveKind
+        << OpenACCClauseKind::Collapse;
     assert(CollapseInfo.ActiveCollapse && "Collapse count without object?");
     Diag(CollapseInfo.ActiveCollapse->getBeginLoc(),
          diag::note_acc_active_clause_here)
@@ -2643,7 +2643,8 @@ void SemaOpenACC::ActOnDoStmt(SourceLocation DoLoc) {
 
   if (TileInfo.CurTileCount && *TileInfo.CurTileCount > 0) {
     Diag(DoLoc, diag::err_acc_invalid_in_loop)
-        << /*do loop*/ 2 << OpenACCClauseKind::Tile;
+        << /*do loop*/ 2 << OpenACCDirectiveKind::Loop
+        << OpenACCClauseKind::Tile;
     assert(TileInfo.ActiveTile && "tile count without object?");
     Diag(TileInfo.ActiveTile->getBeginLoc(), diag::note_acc_active_clause_here)
         << OpenACCClauseKind::Tile;
@@ -2670,7 +2671,8 @@ void SemaOpenACC::ForStmtBeginHelper(SourceLocation ForLoc,
     // This checks for more than 1 loop at the current level, the
     // 'depth'-satisifed checking manages the 'not zero' case.
     if (LoopInfo.CurLevelHasLoopAlready) {
-      Diag(ForLoc, diag::err_acc_clause_multiple_loops) << /*Collapse*/ 0;
+      Diag(ForLoc, diag::err_acc_clause_multiple_loops)
+          << CollapseInfo.DirectiveKind << OpenACCClauseKind::Collapse;
       assert(CollapseInfo.ActiveCollapse && "No collapse object?");
       Diag(CollapseInfo.ActiveCollapse->getBeginLoc(),
            diag::note_acc_active_clause_here)
@@ -2689,7 +2691,8 @@ void SemaOpenACC::ForStmtBeginHelper(SourceLocation ForLoc,
     C.check();
 
     if (LoopInfo.CurLevelHasLoopAlready) {
-      Diag(ForLoc, diag::err_acc_clause_multiple_loops) << /*Tile*/ 1;
+      Diag(ForLoc, diag::err_acc_clause_multiple_loops)
+          << OpenACCDirectiveKind::Loop << OpenACCClauseKind::Tile;
       assert(TileInfo.ActiveTile && "No tile object?");
       Diag(TileInfo.ActiveTile->getBeginLoc(),
            diag::note_acc_active_clause_here)
@@ -3192,7 +3195,7 @@ void SemaOpenACC::ActOnForStmtEnd(SourceLocation ForLoc, StmtResult Body) {
 
     if (OtherStmtLoc.isValid() && IsActiveCollapse) {
       Diag(OtherStmtLoc, diag::err_acc_intervening_code)
-          << OpenACCClauseKind::Collapse;
+          << OpenACCClauseKind::Collapse << CollapseInfo.DirectiveKind;
       Diag(CollapseInfo.ActiveCollapse->getBeginLoc(),
            diag::note_acc_active_clause_here)
           << OpenACCClauseKind::Collapse;
@@ -3200,7 +3203,7 @@ void SemaOpenACC::ActOnForStmtEnd(SourceLocation ForLoc, StmtResult Body) {
 
     if (OtherStmtLoc.isValid() && IsActiveTile) {
       Diag(OtherStmtLoc, diag::err_acc_intervening_code)
-          << OpenACCClauseKind::Tile;
+          << OpenACCClauseKind::Tile << OpenACCDirectiveKind::Loop;
       Diag(TileInfo.ActiveTile->getBeginLoc(),
            diag::note_acc_active_clause_here)
           << OpenACCClauseKind::Tile;
@@ -3220,7 +3223,8 @@ bool SemaOpenACC::ActOnStartStmtDirective(OpenACCDirectiveKind K,
   // ALL constructs are ill-formed if there is an active 'collapse'
   if (CollapseInfo.CurCollapseCount && *CollapseInfo.CurCollapseCount > 0) {
     Diag(StartLoc, diag::err_acc_invalid_in_loop)
-        << /*OpenACC Construct*/ 0 << OpenACCClauseKind::Collapse << K;
+        << /*OpenACC Construct*/ 0 << CollapseInfo.DirectiveKind
+        << OpenACCClauseKind::Collapse << K;
     assert(CollapseInfo.ActiveCollapse && "Collapse count without object?");
     Diag(CollapseInfo.ActiveCollapse->getBeginLoc(),
          diag::note_acc_active_clause_here)
@@ -3228,7 +3232,8 @@ bool SemaOpenACC::ActOnStartStmtDirective(OpenACCDirectiveKind K,
   }
   if (TileInfo.CurTileCount && *TileInfo.CurTileCount > 0) {
     Diag(StartLoc, diag::err_acc_invalid_in_loop)
-        << /*OpenACC Construct*/ 0 << OpenACCClauseKind::Tile << K;
+        << /*OpenACC Construct*/ 0 << OpenACCDirectiveKind::Loop
+        << OpenACCClauseKind::Tile << K;
     assert(TileInfo.ActiveTile && "Tile count without object?");
     Diag(TileInfo.ActiveTile->getBeginLoc(), diag::note_acc_active_clause_here)
         << OpenACCClauseKind::Tile;
diff --git a/clang/test/AST/ast-print-openacc-combined-construct.cpp b/clang/test/AST/ast-print-openacc-combined-construct.cpp
index 114c0bb1c739af..44ec0bb282f00e 100644
--- a/clang/test/AST/ast-print-openacc-combined-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-combined-construct.cpp
@@ -185,4 +185,29 @@ void foo() {
 // CHECK: #pragma acc parallel loop create(i, array[1], array, array[1:2]) pcreate(zero: i, array[1], array, array[1:2]) present_or_create(i, array[1], array, array[1:2])
 #pragma acc parallel loop create(i, array[1], array, array[1:2]) pcreate(zero: i, array[1], array, array[1:2]) present_or_create(i, array[1], array, array[1:2])
   for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc parallel loop collapse(1)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc parallel loop collapse(1)
+  for(int i = 0;i<5;++i);
+// CHECK: #pragma acc serial loop collapse(force:1)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc serial loop collapse(force:1)
+  for(int i = 0;i<5;++i);
+// CHECK: #pragma acc kernels loop collapse(2)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc kernels loop collapse(2)
+  for(int i = 0;i<5;++i)
+    for(int i = 0;i<5;++i);
+// CHECK: #pragma acc parallel loop collapse(force:2)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc parallel loop collapse(force:2)
+  for(int i = 0;i<5;++i)
+    for(int i = 0;i<5;++i);
 }
diff --git a/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c b/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
index c614bbbd1d1769..9a12c1870e7379 100644
--- a/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
+++ b/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
@@ -129,7 +129,6 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'reduction' not yet implemented}}
 #pragma acc parallel loop auto reduction(+:Var)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'collapse' not yet implemented}}
 #pragma acc parallel loop auto collapse(1)
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}}
@@ -258,7 +257,6 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'reduction' not yet implemented}}
 #pragma acc parallel loop reduction(+:Var) auto
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'collapse' not yet implemented}}
 #pragma acc parallel loop collapse(1) auto
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}}
@@ -388,7 +386,6 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'reduction' not yet implemented}}
 #pragma acc parallel loop independent reduction(+:Var)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'collapse' not yet implemented}}
 #pragma acc parallel loop independent collapse(1)
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}}
@@ -517,7 +514,6 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'reduction' not yet implemented}}
 #pragma acc parallel loop reduction(+:Var) independent
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'collapse' not yet implemented}}
 #pragma acc parallel loop collapse(1) independent
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}}
@@ -653,7 +649,6 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'reduction' not yet implemented}}
 #pragma acc parallel loop seq reduction(+:Var)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'collapse' not yet implemented}}
 #pragma acc parallel loop seq collapse(1)
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}}
@@ -788,7 +783,6 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'reduction' not yet implemented}}
 #pragma acc parallel loop reduction(+:Var) seq
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'collapse' not yet implemented}}
 #pragma acc parallel loop collapse(1) seq
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}}
diff --git a/clang/test/SemaOpenACC/combined-construct-collapse-ast.cpp b/clang/test/SemaOpenACC/combined-construct-collapse-ast.cpp
new file mode 100644
index 00000000000000..89ba7cb24814f6
--- /dev/null
+++ b/clang/test/SemaOpenACC/combined-construct-collapse-ast.cpp
@@ -0,0 +1,219 @@
+
+// RUN: %clang_cc1 %s -fopenacc -ast-dump | FileCheck %s
+
+// Test this with PCH.
+// RUN: %clang_cc1 %s -fopenacc -emit-pch -o %t %s
+// RUN: %clang_cc1 %s -fopenacc -include-pch %t -ast-dump-all | FileCheck %s
+#ifndef PCH_HELPER
+#define PCH_HELPER
+
+struct S {
+  constexpr S(){};
+  constexpr operator auto() {return 1;}
+};
+
+void NormalUses() {
+  // CHECK: FunctionDecl{{.*}}NormalUses
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc parallel loop collapse(1)
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: collapse clause
+  // CHECK-NEXT: ConstantExpr{{.*}}'int'
+  // CHECK-NEXT: value: Int 1
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 1
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} i 'int'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: BinaryOperator{{.*}}'<'
+  // CHECK-NEXT: ImplicitCastExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+  // CHECK-NEXT: UnaryOperator{{.*}}++
+  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
+  // CHECK-NEXT: NullStmt
+
+#pragma acc serial loop collapse(force:S{})
+  for(int i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} serial loop
+  // CHECK-NEXT: collapse clause
+  // CHECK-NEXT: ConstantExpr{{.*}}'int'
+  // CHECK-NEXT: value: Int 1
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <UserDefinedConversion>
+  // CHECK-NEXT: CXXMemberCallExpr{{.*}} 'int'
+  // CHECK-NEXT: MemberExpr{{.*}} .operator auto
+  // CHECK-NEXT: MaterializeTemporaryExpr{{.*}}'S' lvalue
+  // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}}'S' 'void ()' list
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} i 'int'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: BinaryOperator{{.*}}'<'
+  // CHECK-NEXT: ImplicitCastExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+  // CHECK-NEXT: UnaryOperator{{.*}}++
+  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
+  // CHECK-NEXT: NullStmt
+}
+
+template<typename T, unsigned Value>
+void TemplUses() {
+  // CHECK: FunctionTemplateDecl{{.*}}TemplUses
+  // CHECK-NEXT: TemplateTypeParmDecl {{.*}} referenced typename depth 0 index 0 T
+  // CHECK-NEXT: NonTypeTemplateParmDecl{{.*}} referenced 'unsigned int' depth 0 index 1 Value
+  // CHECK-NEXT: FunctionDecl{{.*}} TemplUses 'void ()'
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc parallel loop collapse(Value)
+  for(int i = 0; i < 5; ++i)
+    for(int j = 0; j < 5; ++j);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: collapse clause
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'unsigned int' NonTypeTemplateParm{{.*}} 'Value'
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} i 'int'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: BinaryOperator{{.*}}'<'
+  // CHECK-NEXT: ImplicitCastExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+  // CHECK-NEXT: UnaryOperator{{.*}}++
+  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} j 'int'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: BinaryOperator{{.*}}'<'
+  // CHECK-NEXT: ImplicitCastExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}}'j' 'int'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+  // CHECK-NEXT: UnaryOperator{{.*}}++
+  // CHECK-NEXT: DeclRefExpr{{.*}}'j' 'int'
+  // CHECK-NEXT: NullStmt
+
+#pragma acc kernels loop collapse(force:T{} + S{})
+  for(int i = 0; i < 5; ++i)
+    for(int j = 0; j < 5; ++j);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+  // CHECK-NEXT: collapse clause
+  // CHECK-NEXT: BinaryOperator {{.*}}'+'
+  // CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}}'T' 'T' list
+  // CHECK-NEXT: InitListExpr
+  // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}}'S' 'void ()' list
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} i 'int'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: BinaryOperator{{.*}}'<'
+  // CHECK-NEXT: ImplicitCastExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+  // CHECK-NEXT: UnaryOperator{{.*}}++
+  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} j 'int'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: BinaryOperator{{.*}}'<'
+  // CHECK-NEXT: ImplicitCastExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}}'j' 'int'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+  // CHECK-NEXT: UnaryOperator{{.*}}++
+  // CHECK-NEXT: DeclRefExpr{{.*}}'j' 'int'
+  // CHECK-NEXT: NullStmt
+
+  // Instantiation:
+  // CHECK-NEXT: FunctionDecl{{.*}} used TemplUses 'void ()' implicit_instantiation
+  // CHECK-NEXT: TemplateArgument type 'S'
+  // CHECK-NEXT: RecordType{{.*}} 'S'
+  // CHECK-NEXT: CXXRecord{{.*}} 'S'
+  // CHECK-NEXT: TemplateArgument integral '2U'
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: collapse clause
+  // CHECK-NEXT: ConstantExpr{{.*}}'unsigned int'
+  // CHECK-NEXT: value: Int 2
+  // CHECK-NEXT: SubstNonTypeTemplateParmExpr{{.*}} 'unsigned int'
+  // CHECK-NEXT: NonTypeTemplateParmDecl
+  // CHECK-NEXT: IntegerLiteral {{.*}} 'unsigned int' 2
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} i 'int'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: BinaryOperator{{.*}}'<'
+  // CHECK-NEXT: ImplicitCastExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+  // CHECK-NEXT: UnaryOperator{{.*}}++
+  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} j 'int'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: BinaryOperator{{.*}}'<'
+  // CHECK-NEXT: ImplicitCastExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}}'j' 'int'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+  // CHECK-NEXT: UnaryOperator{{.*}}++
+  // CHECK-NEXT: DeclRefExpr{{.*}}'j' 'int'
+  // CHECK-NEXT: NullStmt
+
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+  // CHECK-NEXT: collapse clause
+  // CHECK-NEXT: ConstantExpr{{.*}}'int'
+  // CHECK-NEXT: value: Int 2
+  // CHECK-NEXT: BinaryOperator {{.*}}'+'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <UserDefinedConversion>
+  // CHECK-NEXT: CXXMemberCallExpr{{.*}} 'int'
+  // CHECK-NEXT: MemberExpr{{.*}} .operator auto
+  // CHECK-NEXT: MaterializeTemporaryExpr{{.*}}'S' lvalue
+  // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}}'S' 'void ()' list
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <UserDefinedConversion>
+  // CHECK-NEXT: CXXMemberCallExpr{{.*}} 'int'
+  // CHECK-NEXT: MemberExpr{{.*}} .operator auto
+  // CHECK-NEXT: MaterializeTemporaryExpr{{.*}}'S' lvalue
+  // CHECK-NEXT: CXXTemporaryObjectExpr{{.*}}'S' 'void ()' list
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} i 'int'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: BinaryOperator{{.*}}'<'
+  // CHECK-NEXT: ImplicitCastExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+  // CHECK-NEXT: UnaryOperator{{.*}}++
+  // CHECK-NEXT: DeclRefExpr{{.*}}'i' 'int'
+  // CHECK-NEXT: ForStmt
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl{{.*}} j 'int'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 0
+  // CHECK-NEXT: <<<NULL>>>
+  // CHECK-NEXT: BinaryOperator{{.*}}'<'
+  // CHECK-NEXT: ImplicitCastExpr
+  // CHECK-NEXT: DeclRefExpr{{.*}}'j' 'int'
+  // CHECK-NEXT: IntegerLiteral{{.*}} 'int' 5
+  // CHECK-NEXT: UnaryOperator{{.*}}++
+  // CHECK-NEXT: DeclRefExpr{{.*}}'j' 'int'
+  // CHECK-NEXT: NullStmt
+
+}
+
+void Inst() {
+  TemplUses<S, 2>();
+}
+
+#endif // PCH_HELPER
diff --git a/clang/test/SemaOpenACC/combined-construct-collapse-clause.cpp b/clang/test/SemaOpenACC/combined-construct-collapse-clause.cpp
new file mode 100644
index 00000000000000..c7db9669a9879b
--- /dev/null
+++ b/clang/test/SemaOpenACC/combined-construct-collapse-clause.cpp
@@ -0,0 +1,390 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+template<unsigned Val>
+void depth_too_high_templ() {
+  // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}}
+  // expected-note@+1{{active 'collapse' clause defined here}}
+#pragma acc parallel loop collapse(Val)
+  for(unsigned i = 0; i < 5; ++i)
+    for(unsigned j = 0; j < 5; ++j);
+}
+constexpr int three() { return 3; }
+struct ConvertsThree{
+  constexpr ConvertsThree(){};
+
+  constexpr operator int(){ return 3; }
+};
+
+void depth_too_high() {
+  depth_too_high_templ<3>(); // expected-note{{in instantiation of function template specialization}}
+
+  // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}}
+  // expected-note@+1{{active 'collapse' clause defined here}}
+#pragma acc serial loop collapse(3)
+  for(unsigned i = 0; i < 5; ++i)
+    for(unsigned j = 0; j < 5; ++j);
+
+  // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}}
+  // expected-note@+1{{active 'collapse' clause defined here}}
+#pragma acc kernels loop collapse(three())
+  for(unsigned i = 0; i < 5; ++i)
+    for(unsigned j = 0; j < 5; ++j);
+
+  // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}}
+  // expected-note@+1{{active 'collapse' clause defined here}}
+#pragma acc parallel loop collapse(ConvertsThree{})
+  for(unsigned i = 0; i < 5; ++i)
+    for(unsigned j = 0; j < 5; ++j);
+}
+
+template<typename T, unsigned Three>
+void not_single_loop_templ() {
+  T Arr[5];
+  // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}}
+  // expected-note@+1 2{{active 'collapse' clause defined here}}
+#pragma acc parallel loop collapse(3)
+  for(auto x : Arr) {
+    for(auto y : Arr){
+      do{}while(true); // expected-error{{do loop cannot appear in intervening code of a 'parallel loop' with a 'collapse' clause}}
+    }
+  }
+
+  // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}}
+  // expected-note@+1 2{{active 'collapse' clause defined here}}
+#pragma acc serial loop collapse(Three)
+  for(unsigned i = 0; i < 5; ++i) {
+    for(unsigned j = 0; j < 5; ++j) {
+      do{}while(true); // expected-error{{do loop cannot appear in intervening code of a 'serial loop' with a 'collapse' clause}}
+    }
+  }
+
+#pragma acc kernels loop collapse(Three)
+  for(unsigned i = 0; i < 5; ++i) {
+    for(unsigned j = 0; j < 5; ++j) {
+      for(unsigned k = 0; k < 5;++k) {
+        do{}while(true);
+      }
+    }
+  }
+  // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}}
+  // expected-note@+1 2{{active 'collapse' clause defined here}}
+#pragma acc parallel loop collapse(Three)
+  for(auto x : Arr) {
+    for(auto y: Arr) {
+      do{}while(true); // expected-error{{do loop cannot appear in intervening code of a 'parallel loop' with a 'collapse' clause}}
+    }
+  }
+
+#pragma acc serial loop collapse(Three)
+  for(auto x : Arr) {
+    for(auto y: Arr) {
+      for(auto z: Arr) {
+        do{}while(true);
+      }
+    }
+  }
+}
+
+void not_single_loop() {
+  not_single_loop_templ<int, 3>(); // expected-note{{in instantiation of function template}}
+
+  // expected-note@+1{{active 'collapse' clause defined here}}
+#pragma acc kernels loop collapse(3)
+  for(unsigned i = 0; i < 5; ++i) {
+    for(unsigned j = 0; j < 5; ++j) {
+      for(unsigned k = 0; k < 5;++k);
+    }
+    while(true); // expected-error{{while loop cannot appear in intervening code of a 'kernels loop' with a 'collapse' clause}}
+  }
+
+  // expected-note@+1{{active 'collapse' clause defined here}}
+#pragma acc parallel loop collapse(3)
+  for(unsigned i = 0; i < 5; ++i) {
+    for(unsigned j = 0; j < 5; ++j) {
+      for(unsigned k = 0; k < 5;++k);
+    }
+    do{}while(true); // expected-error{{do loop cannot appear in intervening code of a 'parallel loop' with a 'collapse' clause}}
+  }
+
+  // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}}
+  // expected-note@+1 2{{active 'collapse' clause defined here}}
+#pragma acc serial loop collapse(3)
+  for(unsigned i = 0; i < 5; ++i) {
+    for(unsigned j = 0; j < 5; ++j) {
+      while(true); // expected-error{{while loop cannot appear in intervening code of a 'serial loop' with a 'collapse' clause}}
+    }
+  }
+  // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}}
+  // expected-note@+1 2{{active 'collapse' clause defined here}}
+#pragma acc kernels loop collapse(3)
+  for(unsigned i = 0; i < 5; ++i) {
+    for(unsigned j = 0; j < 5; ++j) {
+      do{}while(true); // expected-error{{do loop cannot appear in intervening code of a 'kernels loop' with a 'collapse' clause}}
+    }
+  }
+
+#pragma acc parallel loop collapse(2)
+  for(unsigned i = 0; i < 5; ++i) {
+    for(unsigned j = 0; j < 5; ++j) {
+      do{}while(true);
+    }
+  }
+#pragma acc serial loop collapse(2)
+  for(unsigned i = 0; i < 5; ++i) {
+    for(unsigned j = 0; j < 5; ++j) {
+      while(true);
+    }
+  }
+
+  int Arr[5];
+  // expected-error@+2{{'collapse' clause specifies a loop count greater than the number of available loops}}
+  // expected-note@+1 2{{active 'collapse' clause defined here}}
+#pragma acc kernels loop collapse(3)
+  for(auto x : Arr) {
+    for(auto y : Arr){
+      do{}while(true); // expected-error{{do loop cannot appear in intervening code of a 'kernels loop' with a 'collapse' clause}}
+    }
+  }
+
+  // expected-note@+1 {{active 'collapse' clause defined here}}
+#pragma acc parallel loop collapse(3)
+  for(unsigned i = 0; i < 5; ++i) {
+    for(unsigned j = 0; j < 5; ++j) {
+      for(unsigned k = 0; k < 5;++k);
+    }
+    // expected-error@+1{{more than one for-loop in a loop associated with OpenACC 'parallel loop' construct with a 'collapse' clause}}
+      for(unsigned k = 0; k < 5;++k);
+  }
+
+  // expected-note@+1 {{active 'collapse' clause defined here}}
+#pragma acc serial loop collapse(3)
+  for(unsigned i = 0; i < 5; ++i) {
+    for(unsigned j = 0; j < 5; ++j) {
+      for(unsigned k = 0; k < 5;++k);
+    // expected-error@+1{{more than one for-loop in a loop associated with OpenACC 'serial loop' construct with a 'collapse' clause}}
+      for(unsigned k = 0; k < 5;++k);
+    }
+  }
+
+  for(unsigned k = 0; k < 5;++k);
+#pragma acc kernels loop collapse(3)
+  for(unsigned i = 0; i < 5; ++i) {
+    for(unsigned j = 0; j < 5; ++j) {
+      for(unsigned k = 0; k < 5;++k);
+    }
+  }
+}
+
+template<unsigned Two, unsigned Three>
+void no_other_directives() {
+#pragma acc parallel loop collapse(Two)
+  for(unsigned i = 0; i < 5; ++i) {
+    for(unsigned j = 0; j < 5; ++j) {// last loop associated with the top level.
+    // expected-error@+1{{'collapse' clause specifies a loop count greater than the number of available loops}}
+#pragma acc serial loop collapse(Three) // expected-note 2{{active 'collapse' clause defined here}}
+      for(unsigned k = 0; k < 6;++k) {
+        for(unsigned l = 0; l < 5; ++l) {
+    // expected-error@+1{{OpenACC 'serial' construct cannot appear in intervening code of a 'serial loop' with a 'collapse' clause}}
+#pragma acc serial
+          ;
+        }
+      }
+    }
+  }
+#pragma acc kernels loop collapse(Two)// expected-note{{active 'collapse' clause defined here}}
+  for(unsigned i = 0; i < 5; ++i) {
+    for(unsigned j = 0; j < 5; ++j) {// last loop associated with the top level.
+#pragma acc parallel loop collapse(Three)
+      for(unsigned k = 0; k < 6;++k) {
+        for(unsigned l = 0; l < 5; ++l) {
+          for(unsigned m = 0; m < 5; ++m);
+        }
+      }
+    }
+    // expected-error@+1{{OpenACC 'serial' construct cannot appear in intervening code of a 'kernels loop' with a 'collapse' clause}}
+#pragma acc serial
+          ;
+  }
+}
+
+void no_other_directives() {
+  no_other_directives<2,3>(); // expected-note{{in instantiation of function template specialization}}
+
+  // Ok, not inside the intervening list
+#pragma acc serial loop collapse(2)
+  for(unsigned i = 0; i < 5; ++i) {
+    for(unsigned j = 0; j < 5; ++j) {
+#pragma acc data // expected-warning{{OpenACC construct 'data' not yet implemented}}
+    }
+  }
+  // expected-note@+1{{active 'collapse' clause defined here}}
+#pragma acc kernels loop collapse(2)
+  for(unsigned i = 0; i < 5; ++i) {
+    // expected-error@+1{{OpenACC 'data' construct cannot appear in intervening code of a 'kernels loop' with a 'collapse' clause}}
+#pragma acc data // expected-warning{{OpenACC construct 'data' not yet implemented}}
+    for(unsigned j = 0; j < 5; ++j) {
+    }
+  }
+}
+
+void call();
+
+template<unsigned Two>
+void intervening_without_force_templ() {
+  // expected-note@+1{{active 'collapse' clause defined here}}
+#pragma acc parallel loop collapse(2)
+  for(unsigned i = 0; i < 5; ++i) {
+    // expected-error@+1{{inner loops must be tightly nested inside a 'collapse' clause on a 'parallel loop' construct}}
+    call();
+    for(unsigned j = 0; j < 5; ++j);
+  }
+
+  // expected-note@+1{{active 'collapse' clause defined here}}
+#pragma acc serial loop collapse(Two)
+  for(unsigned i = 0; i < 5; ++i) {
+    // expected-error@+1{{inner loops must be tightly nested inside a 'collapse' clause on a 'serial loop' construct}}
+    call();
+    for(unsigned j = 0; j < 5; ++j);
+  }
+
+  // expected-note@+1{{active 'collapse' clause defined here}}
+#pragma acc kernels loop collapse(2)
+  for(unsigned i = 0; i < 5; ++i) {
+    for(unsigned j = 0; j < 5; ++j);
+    // expected-error@+1{{inner loops must be tightly nested inside a 'collapse' clause on a 'kernels loop' construct}}
+    call();
+  }
+
+#pragma acc parallel loop collapse(force:2)
+  for(unsigned i = 0; i < 5; ++i) {
+    call();
+    for(unsigned j = 0; j < 5; ++j);
+  }
+
+#pragma acc parallel loop collapse(force:Two)
+  for(unsigned i = 0; i < 5; ++i) {
+    call();
+    for(unsigned j = 0; j < 5; ++j);
+  }
+
+
+#pragma acc parallel loop collapse(force:2)
+  for(unsigned i = 0; i < 5; ++i) {
+    for(unsigned j = 0; j < 5; ++j);
+    call();
+  }
+
+#pragma acc parallel loop collapse(force:Two)
+  for(unsigned i = 0; i < 5; ++i) {
+    for(unsigned j = 0; j < 5; ++j);
+    call();
+  }
+
+#pragma acc parallel loop collapse(Two)
+  for(unsigned i = 0; i < 5; ++i) {
+    for(unsigned j = 0; j < 5; ++j) {
+    call();
+    }
+  }
+
+#pragma acc parallel loop collapse(Two)
+  for(unsigned i = 0; i < 5; ++i) {
+    {
+      {
+        for(unsigned j = 0; j < 5; ++j) {
+          call();
+        }
+      }
+    }
+  }
+
+#pragma acc parallel loop collapse(force:Two)
+  for(unsigned i = 0; i < 5; ++i) {
+    for(unsigned j = 0; j < 5; ++j) {
+    call();
+    }
+  }
+
+  // expected-note@+1{{active 'collapse' clause defined here}}
+#pragma acc parallel loop collapse(Two)
+  for(unsigned i = 0; i < 5; ++i) {
+    for(unsigned j = 0; j < 5; ++j);
+    // expected-error@+1{{inner loops must be tightly nested inside a 'collapse' clause on a 'parallel loop' construct}}
+    call();
+  }
+
+#pragma acc parallel loop collapse(2)
+  // expected-error@+2{{OpenACC 'parallel loop' construct must have a terminating condition}}
+  // expected-note@-2{{'parallel loop' construct is here}}
+  for(int i = 0;;++i)
+  // expected-error@+2{{OpenACC 'parallel loop' construct must have a terminating condition}}
+  // expected-note@-5{{'parallel loop' construct is here}}
+    for(int j = 0;;++j)
+      for(;;);
+}
+
+void intervening_without_force() {
+  intervening_without_force_templ<2>(); // expected-note{{in instantiation of function template specialization}}
+  // expected-note@+1{{active 'collapse' clause defined here}}
+#pragma acc parallel loop collapse(2)
+  for(unsigned i = 0; i < 5; ++i) {
+    // expected-error@+1{{inner loops must be tightly nested inside a 'collapse' clause on a 'parallel loop' construct}}
+    call();
+    for(unsigned j = 0; j < 5; ++j);
+  }
+
+  // expected-note@+1{{active 'collapse' clause defined here}}
+#pragma acc parallel loop collapse(2)
+  for(unsigned i = 0; i < 5; ++i) {
+    for(unsigned j = 0; j < 5; ++j);
+    // expected-error@+1{{inner loops must be tightly nested inside a 'collapse' clause on a 'parallel loop' construct}}
+    call();
+  }
+
+  // The below two are fine, as they use the 'force' tag.
+#pragma acc parallel loop collapse(force:2)
+  for(unsigned i = 0; i < 5; ++i) {
+    call();
+    for(unsigned j = 0; j < 5; ++j);
+  }
+
+#pragma acc parallel loop collapse(force:2)
+  for(unsigned i = 0; i < 5; ++i) {
+    for(unsigned j = 0; j < 5; ++j);
+    call();
+  }
+
+#pragma acc parallel loop collapse(2)
+  for(unsigned i = 0; i < 5; ++i) {
+    for(unsigned j = 0; j < 5; ++j) {
+    call();
+    }
+  }
+#pragma acc parallel loop collapse(2)
+  for(unsigned i = 0; i < 5; ++i) {
+    {
+      {
+        for(unsigned j = 0; j < 5; ++j) {
+          call();
+        }
+      }
+    }
+  }
+
+#pragma acc parallel loop collapse(force:2)
+  for(unsigned i = 0; i < 5; ++i) {
+    for(unsigned j = 0; j < 5; ++j) {
+    call();
+    }
+  }
+
+#pragma acc parallel loop collapse(2)
+  // expected-error@+2{{OpenACC 'parallel loop' construct must have a terminating condition}}
+  // expected-note@-2{{'parallel loop' construct is here}}
+  for(int i = 0;;++i)
+  // expected-error@+2{{OpenACC 'parallel loop' construct must have a terminating condition}}
+  // expected-note@-5{{'parallel loop' construct is here}}
+    for(int j = 0;;++j)
+      for(;;);
+}
+
diff --git a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
index 58b3cdc2820bcf..8933112a0063de 100644
--- a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
@@ -183,7 +183,6 @@ void uses() {
   // expected-note@+1{{previous clause is here}}
 #pragma acc serial loop device_type(*) reduction(+:Var)
   for(int i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'collapse' not yet implemented, clause ignored}}
 #pragma acc serial loop device_type(*) collapse(1)
   for(int i = 0; i < 5; ++i);
   // expected-error@+2{{OpenACC clause 'bind' may not follow a 'device_type' clause in a 'parallel loop' construct}}