diff --git a/clang/include/clang/Basic/OpenACCKinds.h b/clang/include/clang/Basic/OpenACCKinds.h
index 8da02b93b2b97..d53b7223b5334 100644
--- a/clang/include/clang/Basic/OpenACCKinds.h
+++ b/clang/include/clang/Basic/OpenACCKinds.h
@@ -33,7 +33,10 @@ enum class OpenACCDirectiveKind {
   Loop,
   // FIXME: 'cache'
 
-  // FIXME: Combined Constructs.
+  // Combined Constructs.
+  ParallelLoop,
+  SerialLoop,
+  KernelsLoop,
 
   // FIXME: atomic Construct variants.
 
diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index e0c8d43955646..fdf0f24daf985 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -22,8 +22,10 @@ using namespace llvm;
 
 namespace {
 
-// Translate single-token string representations to the OpenACC Directive Kind.
-OpenACCDirectiveKind GetOpenACCDirectiveKind(StringRef Name) {
+/// This doesn't completely comprehend 'Compound Constructs' (as it just
+/// identifies the first token) just the first token of each.  So
+/// this should only be used by `ParseOpenACCDirectiveKind`.
+OpenACCDirectiveKind getOpenACCDirectiveKind(StringRef Name) {
   return llvm::StringSwitch<OpenACCDirectiveKind>(Name)
       .Case("parallel", OpenACCDirectiveKind::Parallel)
       .Case("serial", OpenACCDirectiveKind::Serial)
@@ -39,6 +41,42 @@ OpenACCDirectiveKind GetOpenACCDirectiveKind(StringRef Name) {
       .Default(OpenACCDirectiveKind::Invalid);
 }
 
+bool isOpenACCDirectiveKind(OpenACCDirectiveKind Kind, StringRef Tok) {
+  switch (Kind) {
+  case OpenACCDirectiveKind::Parallel:
+    return Tok == "parallel";
+  case OpenACCDirectiveKind::Serial:
+    return Tok == "serial";
+  case OpenACCDirectiveKind::Kernels:
+    return Tok == "kernels";
+  case OpenACCDirectiveKind::Data:
+    return Tok == "data";
+  case OpenACCDirectiveKind::HostData:
+    return Tok == "host_data";
+  case OpenACCDirectiveKind::Loop:
+    return Tok == "loop";
+
+  case OpenACCDirectiveKind::ParallelLoop:
+  case OpenACCDirectiveKind::SerialLoop:
+  case OpenACCDirectiveKind::KernelsLoop:
+    return false;
+
+  case OpenACCDirectiveKind::Declare:
+    return Tok == "declare";
+  case OpenACCDirectiveKind::Init:
+    return Tok == "init";
+  case OpenACCDirectiveKind::Shutdown:
+    return Tok == "shutdown";
+  case OpenACCDirectiveKind::Set:
+    return Tok == "set";
+  case OpenACCDirectiveKind::Update:
+    return Tok == "update";
+  case OpenACCDirectiveKind::Invalid:
+    return false;
+  }
+  llvm_unreachable("Unknown 'Kind' Passed");
+}
+
 // Parse and consume the tokens for OpenACC Directive/Construct kinds.
 OpenACCDirectiveKind ParseOpenACCDirectiveKind(Parser &P) {
   Token FirstTok = P.getCurToken();
@@ -53,11 +91,35 @@ OpenACCDirectiveKind ParseOpenACCDirectiveKind(Parser &P) {
   P.ConsumeToken();
   std::string FirstTokSpelling = P.getPreprocessor().getSpelling(FirstTok);
 
-  OpenACCDirectiveKind DirKind = GetOpenACCDirectiveKind(FirstTokSpelling);
+  OpenACCDirectiveKind DirKind = getOpenACCDirectiveKind(FirstTokSpelling);
 
   if (DirKind == OpenACCDirectiveKind::Invalid)
     P.Diag(FirstTok, diag::err_acc_invalid_directive) << FirstTokSpelling;
 
+  // Combined Constructs allows parallel loop, serial loop, or kernels loop. Any
+  // other attempt at a combined construct will be diagnosed as an invalid
+  // clause.
+  Token SecondTok = P.getCurToken();
+  if (!SecondTok.isAnnotation() &&
+      isOpenACCDirectiveKind(OpenACCDirectiveKind::Loop,
+                             P.getPreprocessor().getSpelling(SecondTok))) {
+    switch (DirKind) {
+    default:
+      // Nothing to do except in the below cases, as they should be diagnosed as
+      // a clause.
+      break;
+    case OpenACCDirectiveKind::Parallel:
+      P.ConsumeToken();
+      return OpenACCDirectiveKind::ParallelLoop;
+    case OpenACCDirectiveKind::Serial:
+      P.ConsumeToken();
+      return OpenACCDirectiveKind::SerialLoop;
+    case OpenACCDirectiveKind::Kernels:
+      P.ConsumeToken();
+      return OpenACCDirectiveKind::KernelsLoop;
+    }
+  }
+
   return DirKind;
 }
 
diff --git a/clang/test/ParserOpenACC/parse-constructs.c b/clang/test/ParserOpenACC/parse-constructs.c
index 8642833691ede..fcad4507e7c1d 100644
--- a/clang/test/ParserOpenACC/parse-constructs.c
+++ b/clang/test/ParserOpenACC/parse-constructs.c
@@ -48,14 +48,24 @@ void func() {
   // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
 #pragma acc parallel loop clause list
   for(;;){}
+
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc parallel loop
+  for(;;){}
   // expected-warning@+2{{OpenACC clause parsing not yet implemented}}
   // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
 #pragma acc serial loop clause list
+  for(;;){}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc serial loop
   for(;;){}
   // expected-warning@+2{{OpenACC clause parsing not yet implemented}}
   // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
 #pragma acc kernels loop clause list
   for(;;){}
+  // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}
+#pragma acc kernels loop
+  for(;;){}
 
   // expected-warning@+2{{OpenACC clause parsing not yet implemented}}
   // expected-warning@+1{{OpenACC directives not yet implemented, pragma ignored}}