diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 654f3cd97c1c..060df967322a 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -694,14 +694,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitSelfClause(
 
 OpenACCClause *SemaOpenACCClauseVisitor::VisitNumGangsClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
-  // Restrictions only properly implemented on 'compute' constructs, and
-  // 'compute' constructs are the only construct that can do anything with
-  // this yet, so skip/treat as unimplemented in this case.
-  // TODO OpenACC:  Remove this check when we have combined constructs for this
-  // clause.
-  if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
-    return isNotImplemented();
-
   // There is no prose in the standard that says duplicates aren't allowed,
   // but this diagnostic is present in other compilers, as well as makes
   // sense.
@@ -730,6 +722,7 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNumGangsClause(
   // OpenACC 3.3 Section 2.5.4:
   // A reduction clause may not appear on a parallel construct with a
   // num_gangs clause that has more than one argument.
+  // TODO: OpenACC: Reduction on Combined Construct needs to do this too.
   if (Clause.getDirectiveKind() == OpenACCDirectiveKind::Parallel &&
       Clause.getIntExprs().size() > 1) {
     auto *Parallel =
@@ -751,13 +744,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNumGangsClause(
 
 OpenACCClause *SemaOpenACCClauseVisitor::VisitNumWorkersClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
-  // Restrictions only properly implemented on 'compute' constructs, and
-  // 'compute' constructs are the only construct that can do anything with
-  // this yet, so skip/treat as unimplemented in this case.
-  // TODO: OpenACC: Remove when we get combined constructs.
-  if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
-    return isNotImplemented();
-
   // There is no prose in the standard that says duplicates aren't allowed,
   // but this diagnostic is present in other compilers, as well as makes
   // sense.
@@ -773,13 +759,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNumWorkersClause(
 
 OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorLengthClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
-  // Restrictions only properly implemented on 'compute' constructs, and
-  // 'compute' constructs are the only construct that can do anything with
-  // this yet, so skip/treat as unimplemented in this case.
-  // TODO: OpenACC: Remove when we get combined constructs.
-  if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
-    return isNotImplemented();
-
   // There is no prose in the standard that says duplicates aren't allowed,
   // but this diagnostic is present in other compilers, as well as makes
   // sense.
diff --git a/clang/test/AST/ast-print-openacc-combined-construct.cpp b/clang/test/AST/ast-print-openacc-combined-construct.cpp
index d16e44670680..435c770c7457 100644
--- a/clang/test/AST/ast-print-openacc-combined-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-combined-construct.cpp
@@ -224,4 +224,22 @@ void foo() {
       for(int i = 0;i<5;++i)
         for(int i = 0;i<5;++i);
 
+// CHECK: #pragma acc parallel loop num_gangs(i, (int)array[2])
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc parallel loop num_gangs(i, (int)array[2])
+  for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc parallel loop num_workers(i)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc parallel loop num_workers(i)
+  for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc parallel loop vector_length((int)array[1])
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc parallel loop vector_length((int)array[1])
+  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 fc5250ce548e..a6f57a63a91d 100644
--- a/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
+++ b/clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c
@@ -134,16 +134,10 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}}
 #pragma acc parallel loop auto bind(Var)
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'vector_length' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'vector_length' not yet implemented}}
 #pragma acc parallel loop auto vector_length(1)
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'num_gangs' not yet implemented}}
 #pragma acc parallel loop auto num_gangs(1)
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'num_workers' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'num_workers' not yet implemented}}
 #pragma acc parallel loop auto num_workers(1)
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'device_num' not yet implemented}}
@@ -261,16 +255,10 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}}
 #pragma acc parallel loop bind(Var) auto
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'vector_length' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'vector_length' not yet implemented}}
 #pragma acc parallel loop vector_length(1) auto
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'num_gangs' not yet implemented}}
 #pragma acc parallel loop num_gangs(1) auto
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'num_workers' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'num_workers' not yet implemented}}
 #pragma acc parallel loop num_workers(1) auto
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'device_num' not yet implemented}}
@@ -389,16 +377,10 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}}
 #pragma acc parallel loop independent bind(Var)
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'vector_length' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'vector_length' not yet implemented}}
 #pragma acc parallel loop independent vector_length(1)
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'num_gangs' not yet implemented}}
 #pragma acc parallel loop independent num_gangs(1)
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'num_workers' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'num_workers' not yet implemented}}
 #pragma acc parallel loop independent num_workers(1)
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'device_num' not yet implemented}}
@@ -516,16 +498,10 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}}
 #pragma acc parallel loop bind(Var) independent
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'vector_length' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'vector_length' not yet implemented}}
 #pragma acc parallel loop vector_length(1) independent
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'num_gangs' not yet implemented}}
 #pragma acc parallel loop num_gangs(1) independent
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'num_workers' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'num_workers' not yet implemented}}
 #pragma acc parallel loop num_workers(1) independent
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'device_num' not yet implemented}}
@@ -650,16 +626,10 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}}
 #pragma acc parallel loop seq bind(Var)
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'vector_length' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'vector_length' not yet implemented}}
 #pragma acc parallel loop seq vector_length(1)
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'num_gangs' not yet implemented}}
 #pragma acc parallel loop seq num_gangs(1)
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'num_workers' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'num_workers' not yet implemented}}
 #pragma acc parallel loop seq num_workers(1)
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'device_num' not yet implemented}}
@@ -783,16 +753,10 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'bind' not yet implemented}}
 #pragma acc parallel loop bind(Var) seq
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'vector_length' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'vector_length' not yet implemented}}
 #pragma acc parallel loop vector_length(1) seq
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'num_gangs' not yet implemented}}
 #pragma acc parallel loop num_gangs(1) seq
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'num_workers' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'num_workers' not yet implemented}}
 #pragma acc parallel loop num_workers(1) seq
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'device_num' not yet implemented}}
diff --git a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
index a5ab39cb12c3..9a60fb4c665e 100644
--- a/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
+++ b/clang/test/SemaOpenACC/combined-construct-device_type-clause.c
@@ -195,7 +195,6 @@ void uses() {
   // expected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'serial loop' directive}}
 #pragma acc serial loop device_type(*) num_gangs(1)
   for(int i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'num_workers' not yet implemented, clause ignored}}
 #pragma acc parallel loop device_type(*) num_workers(1)
   for(int i = 0; i < 5; ++i);
   // expected-error@+2{{OpenACC clause 'device_num' may not follow a 'device_type' clause in a 'serial loop' construct}}
diff --git a/clang/test/SemaOpenACC/combined-construct-num_gangs-ast.cpp b/clang/test/SemaOpenACC/combined-construct-num_gangs-ast.cpp
new file mode 100644
index 000000000000..6e75a0094336
--- /dev/null
+++ b/clang/test/SemaOpenACC/combined-construct-num_gangs-ast.cpp
@@ -0,0 +1,121 @@
+// 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
+int some_int();
+short some_short();
+long some_long();
+struct CorrectConvert {
+  operator int();
+} Convert;
+
+
+void NormalUses() {
+  // CHECK: FunctionDecl{{.*}}NormalUses
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc parallel loop num_gangs(some_int(), some_long(), some_short())
+  for (unsigned i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: num_gangs clause
+  // CHECK-NEXT: CallExpr{{.*}}'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'int (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}}'int ()' lvalue Function{{.*}} 'some_int' 'int ()'
+  // CHECK-NEXT: CallExpr{{.*}}'long'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'long (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}}'long ()' lvalue Function{{.*}} 'some_long' 'long ()'
+  // CHECK-NEXT: CallExpr{{.*}}'short'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'short (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}}'short ()' lvalue Function{{.*}} 'some_short' 'short ()'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+#pragma acc kernels loop num_gangs(some_int())
+  for (unsigned i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+  // CHECK-NEXT: num_gangs clause
+  // CHECK-NEXT: CallExpr{{.*}}'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'int (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}}'int ()' lvalue Function{{.*}} 'some_int' 'int ()'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+}
+
+template<typename T, typename U>
+void TemplUses(T t, U u) {
+  // CHECK-NEXT: FunctionTemplateDecl
+  // CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 0 T
+  // CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 1 U
+  // CHECK-NEXT: FunctionDecl{{.*}} TemplUses 'void (T, U)'
+  // CHECK-NEXT: ParmVarDecl{{.*}} t 'T'
+  // CHECK-NEXT: ParmVarDecl{{.*}} u 'U'
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc kernels loop num_gangs(u)
+  for (unsigned i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+  // CHECK-NEXT: num_gangs clause
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'U' lvalue ParmVar{{.*}} 'u' 'U'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+#pragma acc parallel loop num_gangs(u, U::value)
+  for (unsigned i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: num_gangs clause
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'U' lvalue ParmVar{{.*}} 'u' 'U'
+  // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
+  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'U'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+  // Check the instantiated versions of the above.
+  // CHECK-NEXT: FunctionDecl{{.*}} used TemplUses 'void (CorrectConvert, HasInt)' implicit_instantiation
+  // CHECK-NEXT: TemplateArgument type 'CorrectConvert'
+  // CHECK-NEXT: RecordType{{.*}} 'CorrectConvert'
+  // CHECK-NEXT: CXXRecord{{.*}} 'CorrectConvert'
+  // CHECK-NEXT: TemplateArgument type 'HasInt'
+  // CHECK-NEXT: RecordType{{.*}} 'HasInt'
+  // CHECK-NEXT: CXXRecord{{.*}} 'HasInt'
+  // CHECK-NEXT: ParmVarDecl{{.*}} t 'CorrectConvert'
+  // CHECK-NEXT: ParmVarDecl{{.*}} u 'HasInt'
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+  // CHECK-NEXT: num_gangs clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'char' <UserDefinedConversion>
+  // CHECK-NEXT: CXXMemberCallExpr{{.*}}'char'
+  // CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator char
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'HasInt' lvalue ParmVar
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: num_gangs clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'char' <UserDefinedConversion>
+  // CHECK-NEXT: CXXMemberCallExpr{{.*}}'char'
+  // CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator char
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'HasInt' lvalue ParmVar
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'const int' lvalue Var{{.*}} 'value' 'const int'
+  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'HasInt'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+}
+
+struct HasInt {
+  using IntTy = int;
+  using ShortTy = short;
+  static constexpr int value = 1;
+
+  operator char();
+};
+
+void Inst() {
+  TemplUses<CorrectConvert, HasInt>({}, {});
+}
+#endif // PCH_HELPER
diff --git a/clang/test/SemaOpenACC/combined-construct-num_gangs-clause.c b/clang/test/SemaOpenACC/combined-construct-num_gangs-clause.c
new file mode 100644
index 000000000000..bd035bd4a51a
--- /dev/null
+++ b/clang/test/SemaOpenACC/combined-construct-num_gangs-clause.c
@@ -0,0 +1,45 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+short getS();
+float getF();
+void Test() {
+#pragma acc kernels loop num_gangs(1)
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+1{{OpenACC 'num_gangs' clause is not valid on 'serial loop' directive}}
+#pragma acc serial loop num_gangs(1)
+  for(int i = 5; i < 10;++i);
+
+#pragma acc parallel loop num_gangs(1)
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+1{{OpenACC clause 'num_gangs' requires expression of integer type}}
+#pragma acc parallel loop num_gangs(getF())
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+1{{expected expression}}
+#pragma acc kernels loop num_gangs()
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+1{{expected expression}}
+#pragma acc parallel loop num_gangs()
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+2{{OpenACC 'num_gangs' clause cannot appear more than once on a 'kernels loop' directive}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc kernels loop num_gangs(1) num_gangs(2)
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+2{{OpenACC 'num_gangs' clause cannot appear more than once on a 'parallel loop' directive}}
+  // expected-note@+1{{previous clause is here}}
+#pragma acc parallel loop num_gangs(1) num_gangs(2)
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+1{{too many integer expression arguments provided to OpenACC 'num_gangs' clause: 'kernels loop' directive expects maximum of 1, 2 were provided}}
+#pragma acc kernels loop num_gangs(1, getS())
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+1{{too many integer expression arguments provided to OpenACC 'num_gangs' clause: 'parallel loop' directive expects maximum of 3, 4 were provided}}
+#pragma acc parallel loop num_gangs(getS(), 1, getS(), 1)
+  for(int i = 5; i < 10;++i);
+}
diff --git a/clang/test/SemaOpenACC/combined-construct-num_workers-ast.cpp b/clang/test/SemaOpenACC/combined-construct-num_workers-ast.cpp
new file mode 100644
index 000000000000..8aa361c7b037
--- /dev/null
+++ b/clang/test/SemaOpenACC/combined-construct-num_workers-ast.cpp
@@ -0,0 +1,230 @@
+// 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
+int some_int();
+short some_short();
+long some_long();
+enum E{};
+E some_enum();
+struct CorrectConvert {
+  operator int();
+} Convert;
+
+
+void NormalUses() {
+  // CHECK: FunctionDecl{{.*}}NormalUses
+  // CHECK-NEXT: CompoundStmt
+#pragma acc parallel loop num_workers(some_int())
+  for (unsigned i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: CallExpr{{.*}}'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'int (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}}'int ()' lvalue Function{{.*}} 'some_int' 'int ()'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+#pragma acc kernels loop num_workers(some_short())
+  for (unsigned i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: CallExpr{{.*}}'short'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'short (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}}'short ()' lvalue Function{{.*}} 'some_short' 'short ()'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+#pragma acc parallel loop num_workers(some_long())
+  for (unsigned i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: CallExpr{{.*}}'long'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'long (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}}'long ()' lvalue Function{{.*}} 'some_long' 'long ()'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+#pragma acc parallel loop num_workers(some_enum())
+  for (unsigned i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: CallExpr{{.*}}'E'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'E (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}}'E ()' lvalue Function{{.*}} 'some_enum' 'E ()'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+#pragma acc kernels loop num_workers(Convert)
+  for (unsigned i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <UserDefinedConversion>
+  // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int'
+  // CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator int
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'struct CorrectConvert':'CorrectConvert' lvalue Var
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+}
+
+template<typename T, typename U>
+void TemplUses(T t, U u) {
+  // CHECK-NEXT: FunctionTemplateDecl
+  // CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 0 T
+  // CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 1 U
+  // CHECK-NEXT: FunctionDecl{{.*}} TemplUses 'void (T, U)'
+  // CHECK-NEXT: ParmVarDecl{{.*}} referenced t 'T'
+  // CHECK-NEXT: ParmVarDecl{{.*}} referenced u 'U'
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc parallel loop num_workers(t)
+  for (unsigned i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'T' lvalue ParmVar{{.*}} 't' 'T'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+#pragma acc kernels loop num_workers(u)
+  for (unsigned i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'U' lvalue ParmVar{{.*}} 'u' 'U'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+#pragma acc parallel loop num_workers(U::value)
+  for (unsigned i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
+  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'U'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+#pragma acc kernels loop num_workers(T{})
+  for (unsigned i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}} 'T' 'T' list
+  // CHECK-NEXT: InitListExpr{{.*}} 'void'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+#pragma acc parallel loop num_workers(U{})
+  for (unsigned i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}} 'U' 'U' list
+  // CHECK-NEXT: InitListExpr{{.*}} 'void'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+#pragma acc kernels loop num_workers(typename U::IntTy{})
+  for (unsigned i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}} 'typename U::IntTy' 'typename U::IntTy' list
+  // CHECK-NEXT: InitListExpr{{.*}} 'void'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+#pragma acc parallel loop num_workers(typename U::ShortTy{})
+  for (unsigned i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: CXXUnresolvedConstructExpr{{.*}} 'typename U::ShortTy' 'typename U::ShortTy' list
+  // CHECK-NEXT: InitListExpr{{.*}} 'void'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+  // Check the instantiated versions of the above.
+  // CHECK-NEXT: FunctionDecl{{.*}} used TemplUses 'void (CorrectConvert, HasInt)' implicit_instantiation
+  // CHECK-NEXT: TemplateArgument type 'CorrectConvert'
+  // CHECK-NEXT: RecordType{{.*}} 'CorrectConvert'
+  // CHECK-NEXT: CXXRecord{{.*}} 'CorrectConvert'
+  // CHECK-NEXT: TemplateArgument type 'HasInt'
+  // CHECK-NEXT: RecordType{{.*}} 'HasInt'
+  // CHECK-NEXT: CXXRecord{{.*}} 'HasInt'
+  // CHECK-NEXT: ParmVarDecl{{.*}} used t 'CorrectConvert'
+  // CHECK-NEXT: ParmVarDecl{{.*}} used u 'HasInt'
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <UserDefinedConversion>
+  // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int'
+  // CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator int
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'CorrectConvert' lvalue ParmVar
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'char' <UserDefinedConversion>
+  // CHECK-NEXT: CXXMemberCallExpr{{.*}}'char'
+  // CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator char
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'HasInt' lvalue ParmVar
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'const int' lvalue Var{{.*}} 'value' 'const int'
+  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'HasInt'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <UserDefinedConversion>
+  // CHECK-NEXT: CXXMemberCallExpr{{.*}}'int'
+  // CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator int
+  // CHECK-NEXT: MaterializeTemporaryExpr{{.*}} 'CorrectConvert' lvalue
+  // CHECK-NEXT: CXXFunctionalCastExpr{{.*}} 'CorrectConvert' functional cast to struct CorrectConvert <NoOp>
+  // CHECK-NEXT: InitListExpr{{.*}}'CorrectConvert'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'char' <UserDefinedConversion>
+  // CHECK-NEXT: CXXMemberCallExpr{{.*}}'char'
+  // CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator char
+  // CHECK-NEXT: MaterializeTemporaryExpr{{.*}} 'HasInt' lvalue
+  // CHECK-NEXT: CXXFunctionalCastExpr{{.*}} 'HasInt' functional cast to struct HasInt <NoOp>
+  // CHECK-NEXT: InitListExpr{{.*}}'HasInt'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: CXXFunctionalCastExpr{{.*}} 'typename HasInt::IntTy':'int' functional cast to typename struct HasInt::IntTy <NoOp>
+  // CHECK-NEXT: InitListExpr{{.*}}'typename HasInt::IntTy':'int'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: num_workers clause
+  // CHECK-NEXT: CXXFunctionalCastExpr{{.*}} 'typename HasInt::ShortTy':'short' functional cast to typename struct HasInt::ShortTy <NoOp>
+  // CHECK-NEXT: InitListExpr{{.*}}'typename HasInt::ShortTy':'short'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+}
+struct HasInt {
+  using IntTy = int;
+  using ShortTy = short;
+  static constexpr int value = 1;
+
+  operator char();
+};
+
+void Inst() {
+  TemplUses<CorrectConvert, HasInt>({}, {});
+}
+#endif // PCH_HELPER
diff --git a/clang/test/SemaOpenACC/combined-construct-num_workers-clause.c b/clang/test/SemaOpenACC/combined-construct-num_workers-clause.c
new file mode 100644
index 000000000000..a5891f071bb0
--- /dev/null
+++ b/clang/test/SemaOpenACC/combined-construct-num_workers-clause.c
@@ -0,0 +1,37 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+short getS();
+float getF();
+void Test() {
+#pragma acc kernels loop num_workers(1)
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+1{{OpenACC 'num_workers' clause is not valid on 'serial loop' directive}}
+#pragma acc serial loop num_workers(1)
+  for(int i = 5; i < 10;++i);
+
+#pragma acc parallel loop num_workers(1)
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+1{{OpenACC clause 'num_workers' requires expression of integer type}}
+#pragma acc parallel loop num_workers(getF())
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+1{{expected expression}}
+#pragma acc kernels loop num_workers()
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+1{{expected expression}}
+#pragma acc parallel loop num_workers()
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
+#pragma acc kernels loop num_workers(1, 2)
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
+#pragma acc parallel loop num_workers(1, 2)
+  for(int i = 5; i < 10;++i);
+}
diff --git a/clang/test/SemaOpenACC/combined-construct-vector_length-ast.cpp b/clang/test/SemaOpenACC/combined-construct-vector_length-ast.cpp
new file mode 100644
index 000000000000..6cfc9c6b8b2c
--- /dev/null
+++ b/clang/test/SemaOpenACC/combined-construct-vector_length-ast.cpp
@@ -0,0 +1,98 @@
+// 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
+short some_short();
+
+struct CorrectConvert {
+  operator int();
+} Convert;
+
+
+void NormalUses() {
+  // CHECK: FunctionDecl{{.*}}NormalUses
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc kernels loop vector_length(some_short())
+  for (unsigned i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+  // CHECK-NEXT: vector_length clause
+  // CHECK-NEXT: CallExpr{{.*}}'short'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'short (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}}'short ()' lvalue Function{{.*}} 'some_short' 'short ()'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+}
+template<typename T, typename U>
+void TemplUses(T t, U u) {
+  // CHECK-NEXT: FunctionTemplateDecl
+  // CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 0 T
+  // CHECK-NEXT: TemplateTypeParmDecl{{.*}}typename depth 0 index 1 U
+  // CHECK-NEXT: FunctionDecl{{.*}} TemplUses 'void (T, U)'
+  // CHECK-NEXT: ParmVarDecl{{.*}} t 'T'
+  // CHECK-NEXT: ParmVarDecl{{.*}} u 'U'
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc kernels loop vector_length(u)
+  for (unsigned i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+  // CHECK-NEXT: vector_length clause
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'U' lvalue ParmVar{{.*}} 'u' 'U'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+#pragma acc parallel loop vector_length(U::value)
+  for (unsigned i = 0; i < 5; ++i);
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: vector_length clause
+  // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
+  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'U'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+  // Check the instantiated versions of the above.
+  // CHECK-NEXT: FunctionDecl{{.*}} used TemplUses 'void (CorrectConvert, HasInt)' implicit_instantiation
+  // CHECK-NEXT: TemplateArgument type 'CorrectConvert'
+  // CHECK-NEXT: RecordType{{.*}} 'CorrectConvert'
+  // CHECK-NEXT: CXXRecord{{.*}} 'CorrectConvert'
+  // CHECK-NEXT: TemplateArgument type 'HasInt'
+  // CHECK-NEXT: RecordType{{.*}} 'HasInt'
+  // CHECK-NEXT: CXXRecord{{.*}} 'HasInt'
+  // CHECK-NEXT: ParmVarDecl{{.*}} t 'CorrectConvert'
+  // CHECK-NEXT: ParmVarDecl{{.*}} u 'HasInt'
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} kernels loop
+  // CHECK-NEXT: vector_length clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'char' <UserDefinedConversion>
+  // CHECK-NEXT: CXXMemberCallExpr{{.*}}'char'
+  // CHECK-NEXT: MemberExpr{{.*}} '<bound member function type>' .operator char
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'HasInt' lvalue ParmVar
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+
+  // CHECK-NEXT: OpenACCCombinedConstruct{{.*}} parallel loop
+  // CHECK-NEXT: vector_length clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'const int' lvalue Var{{.*}} 'value' 'const int'
+  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'HasInt'
+  // CHECK-NEXT: ForStmt
+  // CHECK: NullStmt
+}
+
+struct HasInt {
+  using IntTy = int;
+  using ShortTy = short;
+  static constexpr int value = 1;
+
+  operator char();
+};
+
+void Inst() {
+  TemplUses<CorrectConvert, HasInt>({}, {});
+}
+#endif // PCH_HELPER
diff --git a/clang/test/SemaOpenACC/combined-construct-vector_length-clause.c b/clang/test/SemaOpenACC/combined-construct-vector_length-clause.c
new file mode 100644
index 000000000000..8b6dedd9b83b
--- /dev/null
+++ b/clang/test/SemaOpenACC/combined-construct-vector_length-clause.c
@@ -0,0 +1,37 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+short getS();
+float getF();
+void Test() {
+#pragma acc kernels loop vector_length(1)
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+1{{OpenACC 'vector_length' clause is not valid on 'serial loop' directive}}
+#pragma acc serial loop vector_length(1)
+  for(int i = 5; i < 10;++i);
+
+#pragma acc parallel loop vector_length(1)
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+1{{OpenACC clause 'vector_length' requires expression of integer type}}
+#pragma acc parallel loop vector_length(getF())
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+1{{expected expression}}
+#pragma acc kernels loop vector_length()
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+1{{expected expression}}
+#pragma acc parallel loop vector_length()
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
+#pragma acc kernels loop vector_length(1, 2)
+  for(int i = 5; i < 10;++i);
+
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
+#pragma acc parallel loop vector_length(1, 2)
+  for(int i = 5; i < 10;++i);
+}