diff --git a/clang/include/clang/AST/OpenACCClause.h b/clang/include/clang/AST/OpenACCClause.h
index ec6b4aebcb9f4c..e7b0b411b654fd 100644
--- a/clang/include/clang/AST/OpenACCClause.h
+++ b/clang/include/clang/AST/OpenACCClause.h
@@ -227,7 +227,8 @@ class OpenACCClauseWithSingleIntExpr : public OpenACCClauseWithExprs {
                                  SourceLocation EndLoc)
       : OpenACCClauseWithExprs(K, BeginLoc, LParenLoc, EndLoc),
         IntExpr(IntExpr) {
-    setExprs(MutableArrayRef<Expr *>{&this->IntExpr, 1});
+    if (IntExpr)
+      setExprs(MutableArrayRef<Expr *>{&this->IntExpr, 1});
   }
 
 public:
@@ -260,6 +261,17 @@ class OpenACCVectorLengthClause : public OpenACCClauseWithSingleIntExpr {
          Expr *IntExpr, SourceLocation EndLoc);
 };
 
+class OpenACCAsyncClause : public OpenACCClauseWithSingleIntExpr {
+  OpenACCAsyncClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
+                     Expr *IntExpr, SourceLocation EndLoc);
+
+public:
+  static OpenACCAsyncClause *Create(const ASTContext &C,
+                                    SourceLocation BeginLoc,
+                                    SourceLocation LParenLoc, Expr *IntExpr,
+                                    SourceLocation EndLoc);
+};
+
 /// Represents a clause with one or more 'var' objects, represented as an expr,
 /// as its arguments. Var-list is expected to be stored in trailing storage.
 /// For now, we're just storing the original expression in its entirety, unlike
diff --git a/clang/include/clang/Basic/OpenACCClauses.def b/clang/include/clang/Basic/OpenACCClauses.def
index c92e5eb1e1b634..8933e09b44f9b7 100644
--- a/clang/include/clang/Basic/OpenACCClauses.def
+++ b/clang/include/clang/Basic/OpenACCClauses.def
@@ -21,6 +21,7 @@
 #define CLAUSE_ALIAS(ALIAS_NAME, CLAUSE_NAME)
 #endif
 
+VISIT_CLAUSE(Async)
 VISIT_CLAUSE(Attach)
 VISIT_CLAUSE(Copy)
 CLAUSE_ALIAS(PCopy, Copy)
diff --git a/clang/include/clang/Parse/Parser.h b/clang/include/clang/Parse/Parser.h
index daefd4f28f011a..532b5c125ef599 100644
--- a/clang/include/clang/Parse/Parser.h
+++ b/clang/include/clang/Parse/Parser.h
@@ -3698,7 +3698,9 @@ class Parser : public CodeCompletionHandler {
   bool ParseOpenACCDeviceTypeList();
   /// Parses the 'async-argument', which is an integral value with two
   /// 'special' values that are likely negative (but come from Macros).
-  ExprResult ParseOpenACCAsyncArgument();
+  OpenACCIntExprParseResult ParseOpenACCAsyncArgument(OpenACCDirectiveKind DK,
+                                                      OpenACCClauseKind CK,
+                                                      SourceLocation Loc);
   /// Parses the 'size-expr', which is an integral value, or an asterisk.
   bool ParseOpenACCSizeExpr();
   /// Parses a comma delimited list of 'size-expr's.
diff --git a/clang/include/clang/Sema/SemaOpenACC.h b/clang/include/clang/Sema/SemaOpenACC.h
index 32d94ee8f33fed..2cec2b73e918cf 100644
--- a/clang/include/clang/Sema/SemaOpenACC.h
+++ b/clang/include/clang/Sema/SemaOpenACC.h
@@ -101,16 +101,24 @@ class SemaOpenACC : public SemaBase {
     unsigned getNumIntExprs() const {
       assert((ClauseKind == OpenACCClauseKind::NumGangs ||
               ClauseKind == OpenACCClauseKind::NumWorkers ||
+              ClauseKind == OpenACCClauseKind::Async ||
               ClauseKind == OpenACCClauseKind::VectorLength) &&
              "Parsed clause kind does not have a int exprs");
+      //
+      // 'async' has an optional IntExpr, so be tolerant of that.
+      if (ClauseKind == OpenACCClauseKind::Async &&
+          std::holds_alternative<std::monostate>(Details))
+        return 0;
       return std::get<IntExprDetails>(Details).IntExprs.size();
     }
 
     ArrayRef<Expr *> getIntExprs() {
       assert((ClauseKind == OpenACCClauseKind::NumGangs ||
               ClauseKind == OpenACCClauseKind::NumWorkers ||
+              ClauseKind == OpenACCClauseKind::Async ||
               ClauseKind == OpenACCClauseKind::VectorLength) &&
              "Parsed clause kind does not have a int exprs");
+
       return std::get<IntExprDetails>(Details).IntExprs;
     }
 
@@ -190,6 +198,7 @@ class SemaOpenACC : public SemaBase {
     void setIntExprDetails(ArrayRef<Expr *> IntExprs) {
       assert((ClauseKind == OpenACCClauseKind::NumGangs ||
               ClauseKind == OpenACCClauseKind::NumWorkers ||
+              ClauseKind == OpenACCClauseKind::Async ||
               ClauseKind == OpenACCClauseKind::VectorLength) &&
              "Parsed clause kind does not have a int exprs");
       Details = IntExprDetails{{IntExprs.begin(), IntExprs.end()}};
@@ -197,6 +206,7 @@ class SemaOpenACC : public SemaBase {
     void setIntExprDetails(llvm::SmallVector<Expr *> &&IntExprs) {
       assert((ClauseKind == OpenACCClauseKind::NumGangs ||
               ClauseKind == OpenACCClauseKind::NumWorkers ||
+              ClauseKind == OpenACCClauseKind::Async ||
               ClauseKind == OpenACCClauseKind::VectorLength) &&
              "Parsed clause kind does not have a int exprs");
       Details = IntExprDetails{std::move(IntExprs)};
diff --git a/clang/lib/AST/OpenACCClause.cpp b/clang/lib/AST/OpenACCClause.cpp
index c1affa97b781ca..ffa90884cef5b4 100644
--- a/clang/lib/AST/OpenACCClause.cpp
+++ b/clang/lib/AST/OpenACCClause.cpp
@@ -127,6 +127,26 @@ OpenACCVectorLengthClause::Create(const ASTContext &C, SourceLocation BeginLoc,
       OpenACCVectorLengthClause(BeginLoc, LParenLoc, IntExpr, EndLoc);
 }
 
+OpenACCAsyncClause::OpenACCAsyncClause(SourceLocation BeginLoc,
+                                       SourceLocation LParenLoc, Expr *IntExpr,
+                                       SourceLocation EndLoc)
+    : OpenACCClauseWithSingleIntExpr(OpenACCClauseKind::Async, BeginLoc,
+                                     LParenLoc, IntExpr, EndLoc) {
+  assert((!IntExpr || IntExpr->isInstantiationDependent() ||
+          IntExpr->getType()->isIntegerType()) &&
+         "Condition expression type not scalar/dependent");
+}
+
+OpenACCAsyncClause *OpenACCAsyncClause::Create(const ASTContext &C,
+                                               SourceLocation BeginLoc,
+                                               SourceLocation LParenLoc,
+                                               Expr *IntExpr,
+                                               SourceLocation EndLoc) {
+  void *Mem =
+      C.Allocate(sizeof(OpenACCAsyncClause), alignof(OpenACCAsyncClause));
+  return new (Mem) OpenACCAsyncClause(BeginLoc, LParenLoc, IntExpr, EndLoc);
+}
+
 OpenACCNumGangsClause *OpenACCNumGangsClause::Create(const ASTContext &C,
                                                      SourceLocation BeginLoc,
                                                      SourceLocation LParenLoc,
@@ -287,6 +307,15 @@ void OpenACCClausePrinter::VisitVectorLengthClause(
   OS << ")";
 }
 
+void OpenACCClausePrinter::VisitAsyncClause(const OpenACCAsyncClause &C) {
+  OS << "async";
+  if (C.hasIntExpr()) {
+    OS << "(";
+    printExpr(C.getIntExpr());
+    OS << ")";
+  }
+}
+
 void OpenACCClausePrinter::VisitPrivateClause(const OpenACCPrivateClause &C) {
   OS << "private(";
   llvm::interleaveComma(C.getVarList(), OS,
diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp
index 11d3f3d4cec444..0910471098c9c1 100644
--- a/clang/lib/AST/StmtProfile.cpp
+++ b/clang/lib/AST/StmtProfile.cpp
@@ -2573,6 +2573,11 @@ void OpenACCClauseProfiler::VisitVectorLengthClause(
          "vector_length clause requires a valid int expr");
   Profiler.VisitStmt(Clause.getIntExpr());
 }
+
+void OpenACCClauseProfiler::VisitAsyncClause(const OpenACCAsyncClause &Clause) {
+  if (Clause.hasIntExpr())
+    Profiler.VisitStmt(Clause.getIntExpr());
+}
 } // namespace
 
 void StmtProfiler::VisitOpenACCComputeConstruct(
diff --git a/clang/lib/AST/TextNodeDumper.cpp b/clang/lib/AST/TextNodeDumper.cpp
index 21167ca56e5947..bf02d9545f8449 100644
--- a/clang/lib/AST/TextNodeDumper.cpp
+++ b/clang/lib/AST/TextNodeDumper.cpp
@@ -397,6 +397,7 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
     case OpenACCClauseKind::Default:
       OS << '(' << cast<OpenACCDefaultClause>(C)->getDefaultClauseKind() << ')';
       break;
+    case OpenACCClauseKind::Async:
     case OpenACCClauseKind::Attach:
     case OpenACCClauseKind::Copy:
     case OpenACCClauseKind::PCopy:
diff --git a/clang/lib/Parse/ParseOpenACC.cpp b/clang/lib/Parse/ParseOpenACC.cpp
index b4b81e2ba13ea6..8c8330a5fad75a 100644
--- a/clang/lib/Parse/ParseOpenACC.cpp
+++ b/clang/lib/Parse/ParseOpenACC.cpp
@@ -1081,7 +1081,12 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
         break;
       }
       case OpenACCClauseKind::Async: {
-        ExprResult AsyncArg = ParseOpenACCAsyncArgument();
+        ExprResult AsyncArg =
+            ParseOpenACCAsyncArgument(OpenACCDirectiveKind::Invalid,
+                                      OpenACCClauseKind::Async, ClauseLoc)
+                .first;
+        ParsedClause.setIntExprDetails(AsyncArg.isUsable() ? AsyncArg.get()
+                                                           : nullptr);
         if (AsyncArg.isInvalid()) {
           Parens.skipToEnd();
           return OpenACCCanContinue();
@@ -1120,8 +1125,10 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
 /// defined in the C header file and the Fortran openacc module. The special
 /// values are negative values, so as not to conflict with a user-specified
 /// nonnegative async-argument.
-ExprResult Parser::ParseOpenACCAsyncArgument() {
-  return getActions().CorrectDelayedTyposInExpr(ParseAssignmentExpression());
+Parser::OpenACCIntExprParseResult
+Parser::ParseOpenACCAsyncArgument(OpenACCDirectiveKind DK, OpenACCClauseKind CK,
+                                  SourceLocation Loc) {
+  return ParseOpenACCIntExpr(DK, CK, Loc);
 }
 
 /// OpenACC 3.3, section 2.16:
@@ -1137,14 +1144,12 @@ bool Parser::ParseOpenACCWaitArgument(SourceLocation Loc, bool IsDirective) {
     // Consume colon.
     ConsumeToken();
 
-    ExprResult IntExpr =
-        ParseOpenACCIntExpr(IsDirective ? OpenACCDirectiveKind::Wait
-                                        : OpenACCDirectiveKind::Invalid,
-                            IsDirective ? OpenACCClauseKind::Invalid
-                                        : OpenACCClauseKind::Wait,
-                            Loc)
-            .first;
-    if (IntExpr.isInvalid())
+    OpenACCIntExprParseResult Res = ParseOpenACCIntExpr(
+        IsDirective ? OpenACCDirectiveKind::Wait
+                    : OpenACCDirectiveKind::Invalid,
+        IsDirective ? OpenACCClauseKind::Invalid : OpenACCClauseKind::Wait,
+        Loc);
+    if (Res.first.isInvalid() && Res.second == OpenACCParseCanContinue::Cannot)
       return true;
 
     if (ExpectAndConsume(tok::colon))
@@ -1172,9 +1177,13 @@ bool Parser::ParseOpenACCWaitArgument(SourceLocation Loc, bool IsDirective) {
     }
     FirstArg = false;
 
-    ExprResult CurArg = ParseOpenACCAsyncArgument();
+    OpenACCIntExprParseResult Res = ParseOpenACCAsyncArgument(
+        IsDirective ? OpenACCDirectiveKind::Wait
+                    : OpenACCDirectiveKind::Invalid,
+        IsDirective ? OpenACCClauseKind::Invalid : OpenACCClauseKind::Wait,
+        Loc);
 
-    if (CurArg.isInvalid())
+    if (Res.first.isInvalid() && Res.second == OpenACCParseCanContinue::Cannot)
       return true;
   }
 
diff --git a/clang/lib/Sema/SemaOpenACC.cpp b/clang/lib/Sema/SemaOpenACC.cpp
index 8cf829cf215b54..b1086baa3ae2eb 100644
--- a/clang/lib/Sema/SemaOpenACC.cpp
+++ b/clang/lib/Sema/SemaOpenACC.cpp
@@ -198,6 +198,25 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
     default:
       return false;
     }
+  case OpenACCClauseKind::Async:
+    switch (DirectiveKind) {
+    case OpenACCDirectiveKind::Parallel:
+    case OpenACCDirectiveKind::Serial:
+    case OpenACCDirectiveKind::Kernels:
+    case OpenACCDirectiveKind::Data:
+    case OpenACCDirectiveKind::EnterData:
+    case OpenACCDirectiveKind::ExitData:
+    case OpenACCDirectiveKind::Set:
+    case OpenACCDirectiveKind::Update:
+    case OpenACCDirectiveKind::Wait:
+    case OpenACCDirectiveKind::ParallelLoop:
+    case OpenACCDirectiveKind::SerialLoop:
+    case OpenACCDirectiveKind::KernelsLoop:
+      return true;
+    default:
+      return false;
+    }
+
   default:
     // Do nothing so we can go to the 'unimplemented' diagnostic instead.
     return true;
@@ -398,6 +417,27 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses,
         getASTContext(), Clause.getBeginLoc(), Clause.getLParenLoc(),
         Clause.getIntExprs()[0], Clause.getEndLoc());
   }
+  case OpenACCClauseKind::Async: {
+    // 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.
+    if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
+      break;
+
+    // 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.
+    if (checkAlreadyHasClauseOfKind(*this, ExistingClauses, Clause))
+      return nullptr;
+
+    assert(Clause.getNumIntExprs() < 2 &&
+           "Invalid number of expressions for Async");
+
+    return OpenACCAsyncClause::Create(
+        getASTContext(), Clause.getBeginLoc(), Clause.getLParenLoc(),
+        Clause.getNumIntExprs() != 0 ? Clause.getIntExprs()[0] : nullptr,
+        Clause.getEndLoc());
+  }
   case OpenACCClauseKind::Private: {
     // Restrictions only properly implemented on 'compute' constructs, and
     // 'compute' constructs are the only construct that can do anything with
diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h
index a4ca8b5771a9f2..2d6d6dae680cc9 100644
--- a/clang/lib/Sema/TreeTransform.h
+++ b/clang/lib/Sema/TreeTransform.h
@@ -11401,6 +11401,30 @@ void OpenACCClauseTransform<Derived>::VisitVectorLengthClause(
       ParsedClause.getLParenLoc(), ParsedClause.getIntExprs()[0],
       ParsedClause.getEndLoc());
 }
+
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitAsyncClause(
+    const OpenACCAsyncClause &C) {
+  if (C.hasIntExpr()) {
+    ExprResult Res = Self.TransformExpr(const_cast<Expr *>(C.getIntExpr()));
+    if (!Res.isUsable())
+      return;
+
+    Res = Self.getSema().OpenACC().ActOnIntExpr(OpenACCDirectiveKind::Invalid,
+                                                C.getClauseKind(),
+                                                C.getBeginLoc(), Res.get());
+    if (!Res.isUsable())
+      return;
+    ParsedClause.setIntExprDetails(Res.get());
+  }
+
+  NewClause = OpenACCAsyncClause::Create(
+      Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
+      ParsedClause.getLParenLoc(),
+      ParsedClause.getNumIntExprs() != 0 ? ParsedClause.getIntExprs()[0]
+                                         : nullptr,
+      ParsedClause.getEndLoc());
+}
 } // namespace
 template <typename Derived>
 OpenACCClause *TreeTransform<Derived>::TransformOpenACCClause(
diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp
index 81b78edd9c6cfe..b4b2f999d22597 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -11881,6 +11881,12 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
     return OpenACCCreateClause::Create(getContext(), ClauseKind, BeginLoc,
                                        LParenLoc, IsZero, VarList, EndLoc);
   }
+  case OpenACCClauseKind::Async: {
+    SourceLocation LParenLoc = readSourceLocation();
+    Expr *AsyncExpr = readBool() ? readSubExpr() : nullptr;
+    return OpenACCAsyncClause::Create(getContext(), BeginLoc, LParenLoc,
+                                      AsyncExpr, EndLoc);
+  }
 
   case OpenACCClauseKind::Finalize:
   case OpenACCClauseKind::IfPresent:
@@ -11904,7 +11910,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
   case OpenACCClauseKind::DefaultAsync:
   case OpenACCClauseKind::DeviceType:
   case OpenACCClauseKind::DType:
-  case OpenACCClauseKind::Async:
   case OpenACCClauseKind::Tile:
   case OpenACCClauseKind::Gang:
   case OpenACCClauseKind::Wait:
diff --git a/clang/lib/Serialization/ASTWriter.cpp b/clang/lib/Serialization/ASTWriter.cpp
index 42da50abdc687c..ce2ea4e3d614b2 100644
--- a/clang/lib/Serialization/ASTWriter.cpp
+++ b/clang/lib/Serialization/ASTWriter.cpp
@@ -7908,6 +7908,14 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
     writeOpenACCVarList(CC);
     return;
   }
+  case OpenACCClauseKind::Async: {
+    const auto *AC = cast<OpenACCAsyncClause>(C);
+    writeSourceLocation(AC->getLParenLoc());
+    writeBool(AC->hasIntExpr());
+    if (AC->hasIntExpr())
+      AddStmt(const_cast<Expr *>(AC->getIntExpr()));
+    return;
+  }
 
   case OpenACCClauseKind::Finalize:
   case OpenACCClauseKind::IfPresent:
@@ -7931,7 +7939,6 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
   case OpenACCClauseKind::DefaultAsync:
   case OpenACCClauseKind::DeviceType:
   case OpenACCClauseKind::DType:
-  case OpenACCClauseKind::Async:
   case OpenACCClauseKind::Tile:
   case OpenACCClauseKind::Gang:
   case OpenACCClauseKind::Wait:
diff --git a/clang/test/AST/ast-print-openacc-compute-construct.cpp b/clang/test/AST/ast-print-openacc-compute-construct.cpp
index 1ee1e15bdfc3ac..13597543e9b6d7 100644
--- a/clang/test/AST/ast-print-openacc-compute-construct.cpp
+++ b/clang/test/AST/ast-print-openacc-compute-construct.cpp
@@ -75,5 +75,13 @@ void foo() {
   // CHECK: #pragma acc kernels deviceptr(iPtr, arrayPtr[0])
 #pragma acc kernels deviceptr(iPtr, arrayPtr[0])
   while(true);
+
+  // CHECK: #pragma acc kernels async(*iPtr)
+#pragma acc kernels async(*iPtr)
+  while(true);
+
+  // CHECK: #pragma acc kernels async
+#pragma acc kernels async
+  while(true);
 }
 
diff --git a/clang/test/ParserOpenACC/parse-clauses.c b/clang/test/ParserOpenACC/parse-clauses.c
index 035b7ab4c1f40f..51858b441e9354 100644
--- a/clang/test/ParserOpenACC/parse-clauses.c
+++ b/clang/test/ParserOpenACC/parse-clauses.c
@@ -1233,7 +1233,6 @@ void device_type() {
 
 #define acc_async_sync -1
 void AsyncArgument() {
-  // expected-warning@+1{{OpenACC clause 'async' not yet implemented, clause ignored}}
 #pragma acc parallel async
   {}
 
@@ -1250,15 +1249,12 @@ void AsyncArgument() {
 #pragma acc parallel async(4, 3)
   {}
 
-  // expected-warning@+1{{OpenACC clause 'async' not yet implemented, clause ignored}}
 #pragma acc parallel async(returns_int())
   {}
 
-  // expected-warning@+1{{OpenACC clause 'async' not yet implemented, clause ignored}}
 #pragma acc parallel async(5)
   {}
 
-  // expected-warning@+1{{OpenACC clause 'async' not yet implemented, clause ignored}}
 #pragma acc parallel async(acc_async_sync)
   {}
 }
diff --git a/clang/test/ParserOpenACC/parse-clauses.cpp b/clang/test/ParserOpenACC/parse-clauses.cpp
index 8c1d6437479961..702eb75ca8902d 100644
--- a/clang/test/ParserOpenACC/parse-clauses.cpp
+++ b/clang/test/ParserOpenACC/parse-clauses.cpp
@@ -18,13 +18,14 @@ void templ() {
 #pragma acc parallel vector_length(I)
   for(;;){}
 
-  // expected-warning@+1{{OpenACC clause 'async' not yet implemented, clause ignored}}
 #pragma acc parallel async(T::value)
   for(;;){}
 
-  // expected-warning@+1{{OpenACC clause 'async' not yet implemented, clause ignored}}
 #pragma acc parallel async(I)
   for(;;){}
+
+#pragma acc parallel async
+  for(;;){}
 }
 
 struct S {
diff --git a/clang/test/ParserOpenACC/parse-wait-clause.c b/clang/test/ParserOpenACC/parse-wait-clause.c
index f3e651de458377..64f5b9c8fd7351 100644
--- a/clang/test/ParserOpenACC/parse-wait-clause.c
+++ b/clang/test/ParserOpenACC/parse-wait-clause.c
@@ -84,29 +84,35 @@ void func() {
   #pragma acc parallel wait (devnum: i + j:queues:) clause-list
     {}
 
-  // expected-error@+3{{use of undeclared identifier 'devnum'}}
+  // expected-error@+4{{use of undeclared identifier 'devnum'}}
+  // expected-error@+3{{expected ','}}
   // expected-error@+2{{expected ')'}}
   // expected-note@+1{{to match this '('}}
   #pragma acc parallel wait (queues:devnum: i + j
     {}
 
+  // expected-error@+2{{expected ','}}
   // expected-error@+1{{use of undeclared identifier 'devnum'}}
   #pragma acc parallel wait (queues:devnum: i + j)
     {}
 
+  // expected-error@+3{{expected ','}}
   // expected-error@+2{{use of undeclared identifier 'devnum'}}
   // expected-error@+1{{invalid OpenACC clause 'clause'}}
   #pragma acc parallel wait (queues:devnum: i + j) clause-list
     {}
 
+  // expected-error@+3{{OpenACC clause 'wait' requires expression of integer type ('double' invalid)}}
   // expected-error@+2{{expected ')'}}
   // expected-note@+1{{to match this '('}}
   #pragma acc parallel wait(i, j, 1+1, 3.3
     {}
 
+  // expected-error@+2{{OpenACC clause 'wait' requires expression of integer type ('double' invalid)}}
   // expected-warning@+1{{OpenACC clause 'wait' not yet implemented, clause ignored}}
   #pragma acc parallel wait(i, j, 1+1, 3.3)
     {}
+  // expected-error@+3{{OpenACC clause 'wait' requires expression of integer type ('double' invalid)}}
   // expected-error@+2{{invalid OpenACC clause 'clause'}}
   // expected-warning@+1{{OpenACC clause 'wait' not yet implemented, clause ignored}}
   #pragma acc parallel wait(i, j, 1+1, 3.3) clause-list
@@ -127,45 +133,55 @@ void func() {
   #pragma acc parallel wait(,) clause-list
     {}
 
+  // expected-error@+3{{OpenACC clause 'wait' requires expression of integer type ('double' invalid)}}
   // expected-error@+2{{expected ')'}}
   // expected-note@+1{{to match this '('}}
   #pragma acc parallel wait(queues:i, j, 1+1, 3.3
     {}
 
+  // expected-error@+4{{OpenACC clause 'wait' requires expression of integer type ('double' invalid)}}
   // expected-error@+3{{expected expression}}
   // expected-error@+2{{expected ')'}}
   // expected-note@+1{{to match this '('}}
   #pragma acc parallel wait(queues:i, j, 1+1, 3.3,
     {}
 
+  // expected-error@+2{{OpenACC clause 'wait' requires expression of integer type ('double' invalid)}}
   // expected-warning@+1{{OpenACC clause 'wait' not yet implemented, clause ignored}}
   #pragma acc parallel wait(queues:i, j, 1+1, 3.3)
     {}
 
+  // expected-error@+3{{OpenACC clause 'wait' requires expression of integer type ('double' invalid)}}
   // expected-error@+2{{invalid OpenACC clause 'clause'}}
   // expected-warning@+1{{OpenACC clause 'wait' not yet implemented, clause ignored}}
   #pragma acc parallel wait(queues:i, j, 1+1, 3.3) clause-list
     {}
 
+  // expected-error@+3{{OpenACC clause 'wait' requires expression of integer type ('double' invalid)}}
   // expected-error@+2{{expected ')'}}
   // expected-note@+1{{to match this '('}}
   #pragma acc parallel wait(devnum:3:i, j, 1+1, 3.3
     {}
+  // expected-error@+2{{OpenACC clause 'wait' requires expression of integer type ('double' invalid)}}
   // expected-warning@+1{{OpenACC clause 'wait' not yet implemented, clause ignored}}
   #pragma acc parallel wait(devnum:3:i, j, 1+1, 3.3)
     {}
+  // expected-error@+3{{OpenACC clause 'wait' requires expression of integer type ('double' invalid)}}
   // expected-error@+2{{invalid OpenACC clause 'clause'}}
   // expected-warning@+1{{OpenACC clause 'wait' not yet implemented, clause ignored}}
   #pragma acc parallel wait(devnum:3:i, j, 1+1, 3.3) clause-list
     {}
 
+  // expected-error@+3{{OpenACC clause 'wait' requires expression of integer type ('double' invalid)}}
   // expected-error@+2{{expected ')'}}
   // expected-note@+1{{to match this '('}}
   #pragma acc parallel wait(devnum:3:queues:i, j, 1+1, 3.3
     {}
+  // expected-error@+2{{OpenACC clause 'wait' requires expression of integer type ('double' invalid)}}
   // expected-warning@+1{{OpenACC clause 'wait' not yet implemented, clause ignored}}
   #pragma acc parallel wait(devnum:3:queues:i, j, 1+1, 3.3)
     {}
+  // expected-error@+3{{OpenACC clause 'wait' requires expression of integer type ('double' invalid)}}
   // expected-error@+2{{invalid OpenACC clause 'clause'}}
   // expected-warning@+1{{OpenACC clause 'wait' not yet implemented, clause ignored}}
   #pragma acc parallel wait(devnum:3:queues:i, j, 1+1, 3.3) clause-list
diff --git a/clang/test/ParserOpenACC/parse-wait-construct.c b/clang/test/ParserOpenACC/parse-wait-construct.c
index 30a9fc8c12a474..8f7ea8efd57668 100644
--- a/clang/test/ParserOpenACC/parse-wait-construct.c
+++ b/clang/test/ParserOpenACC/parse-wait-construct.c
@@ -76,28 +76,34 @@ void func() {
   // expected-warning@+1{{OpenACC construct 'wait' not yet implemented, pragma ignored}}
   #pragma acc wait (devnum: i + j:queues:) clause-list
 
-  // expected-error@+4{{use of undeclared identifier 'devnum'}}
+  // expected-error@+5{{use of undeclared identifier 'devnum'}}
+  // expected-error@+4{{expected ','}}
   // expected-error@+3{{expected ')'}}
   // expected-note@+2{{to match this '('}}
   // expected-warning@+1{{OpenACC construct 'wait' not yet implemented, pragma ignored}}
   #pragma acc wait (queues:devnum: i + j
 
-  // expected-error@+2{{use of undeclared identifier 'devnum'}}
+  // expected-error@+3{{use of undeclared identifier 'devnum'}}
+  // expected-error@+2{{expected ','}}
   // expected-warning@+1{{OpenACC construct 'wait' not yet implemented, pragma ignored}}
   #pragma acc wait (queues:devnum: i + j)
 
-  // expected-error@+3{{use of undeclared identifier 'devnum'}}
+  // expected-error@+4{{use of undeclared identifier 'devnum'}}
+  // expected-error@+3{{expected ','}}
   // expected-error@+2{{invalid OpenACC clause 'clause'}}
   // expected-warning@+1{{OpenACC construct 'wait' not yet implemented, pragma ignored}}
   #pragma acc wait (queues:devnum: i + j) clause-list
 
+  // expected-error@+4{{OpenACC directive 'wait' requires expression of integer type ('double' invalid)}}
   // expected-error@+3{{expected ')'}}
   // expected-note@+2{{to match this '('}}
   // expected-warning@+1{{OpenACC construct 'wait' not yet implemented, pragma ignored}}
   #pragma acc wait(i, j, 1+1, 3.3
 
+  // expected-error@+2{{OpenACC directive 'wait' requires expression of integer type ('double' invalid)}}
   // expected-warning@+1{{OpenACC construct 'wait' not yet implemented, pragma ignored}}
   #pragma acc wait(i, j, 1+1, 3.3)
+  // expected-error@+3{{OpenACC directive 'wait' requires expression of integer type ('double' invalid)}}
   // expected-error@+2{{invalid OpenACC clause 'clause'}}
   // expected-warning@+1{{OpenACC construct 'wait' not yet implemented, pragma ignored}}
   #pragma acc wait(i, j, 1+1, 3.3) clause-list
@@ -117,40 +123,50 @@ void func() {
   // expected-warning@+1{{OpenACC construct 'wait' not yet implemented, pragma ignored}}
   #pragma acc wait(,) clause-list
 
+  // expected-error@+4{{OpenACC directive 'wait' requires expression of integer type ('double' invalid)}}
   // expected-error@+3{{expected ')'}}
   // expected-note@+2{{to match this '('}}
   // expected-warning@+1{{OpenACC construct 'wait' not yet implemented, pragma ignored}}
   #pragma acc wait(queues:i, j, 1+1, 3.3
 
+  // expected-error@+5{{OpenACC directive 'wait' requires expression of integer type ('double' invalid)}}
   // expected-error@+4{{expected expression}}
   // expected-error@+3{{expected ')'}}
   // expected-note@+2{{to match this '('}}
   // expected-warning@+1{{OpenACC construct 'wait' not yet implemented, pragma ignored}}
   #pragma acc wait(queues:i, j, 1+1, 3.3,
 
+  // expected-error@+2{{OpenACC directive 'wait' requires expression of integer type ('double' invalid)}}
   // expected-warning@+1{{OpenACC construct 'wait' not yet implemented, pragma ignored}}
   #pragma acc wait(queues:i, j, 1+1, 3.3)
 
+  // expected-error@+3{{OpenACC directive 'wait' requires expression of integer type ('double' invalid)}}
   // expected-error@+2{{invalid OpenACC clause 'clause'}}
   // expected-warning@+1{{OpenACC construct 'wait' not yet implemented, pragma ignored}}
   #pragma acc wait(queues:i, j, 1+1, 3.3) clause-list
 
+  // expected-error@+4{{OpenACC directive 'wait' requires expression of integer type ('double' invalid)}}
   // expected-error@+3{{expected ')'}}
   // expected-note@+2{{to match this '('}}
   // expected-warning@+1{{OpenACC construct 'wait' not yet implemented, pragma ignored}}
   #pragma acc wait(devnum:3:i, j, 1+1, 3.3
+  // expected-error@+2{{OpenACC directive 'wait' requires expression of integer type ('double' invalid)}}
   // expected-warning@+1{{OpenACC construct 'wait' not yet implemented, pragma ignored}}
   #pragma acc wait(devnum:3:i, j, 1+1, 3.3)
+  // expected-error@+3{{OpenACC directive 'wait' requires expression of integer type ('double' invalid)}}
   // expected-error@+2{{invalid OpenACC clause 'clause'}}
   // expected-warning@+1{{OpenACC construct 'wait' not yet implemented, pragma ignored}}
   #pragma acc wait(devnum:3:i, j, 1+1, 3.3) clause-list
 
+  // expected-error@+4{{OpenACC directive 'wait' requires expression of integer type ('double' invalid)}}
   // expected-error@+3{{expected ')'}}
   // expected-note@+2{{to match this '('}}
   // expected-warning@+1{{OpenACC construct 'wait' not yet implemented, pragma ignored}}
   #pragma acc wait(devnum:3:queues:i, j, 1+1, 3.3
+  // expected-error@+2{{OpenACC directive 'wait' requires expression of integer type ('double' invalid)}}
   // expected-warning@+1{{OpenACC construct 'wait' not yet implemented, pragma ignored}}
   #pragma acc wait(devnum:3:queues:i, j, 1+1, 3.3)
+  // expected-error@+3{{OpenACC directive 'wait' requires expression of integer type ('double' invalid)}}
   // expected-error@+2{{invalid OpenACC clause 'clause'}}
   // expected-warning@+1{{OpenACC construct 'wait' not yet implemented, pragma ignored}}
   #pragma acc wait(devnum:3:queues:i, j, 1+1, 3.3) clause-list
diff --git a/clang/test/SemaOpenACC/compute-construct-async-clause.c b/clang/test/SemaOpenACC/compute-construct-async-clause.c
new file mode 100644
index 00000000000000..a8af06bc0afd61
--- /dev/null
+++ b/clang/test/SemaOpenACC/compute-construct-async-clause.c
@@ -0,0 +1,41 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+short getS();
+
+void Test() {
+#pragma acc parallel async
+  while(1);
+#pragma acc parallel async(1)
+  while(1);
+#pragma acc kernels async(1)
+  while(1);
+#pragma acc kernels async(-51)
+  while(1);
+
+#pragma acc serial async(1)
+  while(1);
+
+  // expected-error@+2{{expected ')'}}
+  // expected-note@+1{{to match this '('}}
+#pragma acc serial async(1, 2)
+  while(1);
+
+  struct NotConvertible{} NC;
+  // expected-error@+1{{OpenACC clause 'async' requires expression of integer type ('struct NotConvertible' invalid)}}
+#pragma acc parallel async(NC)
+  while(1);
+
+#pragma acc kernels async(getS())
+  while(1);
+
+  struct Incomplete *SomeIncomplete;
+
+  // expected-error@+1{{OpenACC clause 'async' requires expression of integer type ('struct Incomplete' invalid)}}
+#pragma acc kernels async(*SomeIncomplete)
+  while(1);
+
+  enum E{A} SomeE;
+
+#pragma acc kernels async(SomeE)
+  while(1);
+}
diff --git a/clang/test/SemaOpenACC/compute-construct-async-clause.cpp b/clang/test/SemaOpenACC/compute-construct-async-clause.cpp
new file mode 100644
index 00000000000000..a5da7c8f4e56ec
--- /dev/null
+++ b/clang/test/SemaOpenACC/compute-construct-async-clause.cpp
@@ -0,0 +1,135 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+struct NotConvertible{} NC;
+struct Incomplete *SomeIncomplete; // #INCOMPLETE
+enum E{} SomeE;
+enum class E2{} SomeE2;
+
+struct CorrectConvert {
+  operator int();
+} Convert;
+
+struct ExplicitConvertOnly {
+  explicit operator int() const; // #EXPL_CONV
+} Explicit;
+
+struct AmbiguousConvert{
+  operator int(); // #AMBIG_INT
+  operator short(); // #AMBIG_SHORT
+  operator float();
+} Ambiguous;
+
+void Test() {
+#pragma acc parallel async
+  while(1);
+#pragma acc parallel async(1)
+  while(1);
+#pragma acc kernels async(-51)
+  while(1);
+
+  // expected-error@+1{{OpenACC clause 'async' requires expression of integer type ('struct NotConvertible' invalid}}
+#pragma acc parallel async(NC)
+  while(1);
+
+  // expected-error@+2{{OpenACC integer expression has incomplete class type 'struct Incomplete'}}
+  // expected-note@#INCOMPLETE{{forward declaration of 'Incomplete'}}
+#pragma acc kernels async(*SomeIncomplete)
+  while(1);
+
+#pragma acc parallel async(SomeE)
+  while(1);
+
+  // expected-error@+1{{OpenACC clause 'async' requires expression of integer type ('enum E2' invalid}}
+#pragma acc kernels async(SomeE2)
+  while(1);
+
+#pragma acc parallel async(Convert)
+  while(1);
+
+  // expected-error@+2{{OpenACC integer expression type 'struct ExplicitConvertOnly' requires explicit conversion to 'int'}}
+  // expected-note@#EXPL_CONV{{conversion to integral type 'int'}}
+#pragma acc kernels async(Explicit)
+  while(1);
+
+  // expected-error@+3{{multiple conversions from expression type 'struct AmbiguousConvert' to an integral type}}
+  // expected-note@#AMBIG_INT{{conversion to integral type 'int'}}
+  // expected-note@#AMBIG_SHORT{{conversion to integral type 'short'}}
+#pragma acc parallel async(Ambiguous)
+  while(1);
+}
+
+struct HasInt {
+  using IntTy = int;
+  using ShortTy = short;
+  static constexpr int value = 1;
+  static constexpr AmbiguousConvert ACValue;
+  static constexpr ExplicitConvertOnly EXValue;
+
+  operator char();
+};
+
+template<typename T>
+void TestInst() {
+
+  // expected-error@+1{{no member named 'Invalid' in 'HasInt'}}
+#pragma acc parallel async(HasInt::Invalid)
+  while (1);
+
+  // expected-error@+2{{no member named 'Invalid' in 'HasInt'}}
+  // expected-note@#INST{{in instantiation of function template specialization 'TestInst<HasInt>' requested here}}
+#pragma acc kernels async(T::Invalid)
+  while (1);
+
+  // expected-error@+3{{multiple conversions from expression type 'const AmbiguousConvert' to an integral type}}
+  // expected-note@#AMBIG_INT{{conversion to integral type 'int'}}
+  // expected-note@#AMBIG_SHORT{{conversion to integral type 'short'}}
+#pragma acc parallel async(HasInt::ACValue)
+  while (1);
+
+  // expected-error@+3{{multiple conversions from expression type 'const AmbiguousConvert' to an integral type}}
+  // expected-note@#AMBIG_INT{{conversion to integral type 'int'}}
+  // expected-note@#AMBIG_SHORT{{conversion to integral type 'short'}}
+#pragma acc kernels async(T::ACValue)
+  while (1);
+
+  // expected-error@+2{{OpenACC integer expression type 'const ExplicitConvertOnly' requires explicit conversion to 'int'}}
+  // expected-note@#EXPL_CONV{{conversion to integral type 'int'}}
+#pragma acc parallel async(HasInt::EXValue)
+  while (1);
+
+  // expected-error@+2{{OpenACC integer expression type 'const ExplicitConvertOnly' requires explicit conversion to 'int'}}
+  // expected-note@#EXPL_CONV{{conversion to integral type 'int'}}
+#pragma acc kernels async(T::EXValue)
+  while (1);
+
+#pragma acc parallel async(HasInt::value)
+  while (1);
+
+#pragma acc kernels async(T::value)
+  while (1);
+
+#pragma acc parallel async(HasInt::IntTy{})
+  while (1);
+
+#pragma acc kernels async(typename T::ShortTy{})
+  while (1);
+
+#pragma acc parallel async(HasInt::IntTy{})
+  while (1);
+
+#pragma acc kernels async(typename T::ShortTy{})
+  while (1);
+
+  HasInt HI{};
+  T MyT{};
+
+#pragma acc parallel async(HI)
+  while (1);
+
+#pragma acc kernels async(MyT)
+  while (1);
+}
+
+void Inst() {
+  TestInst<HasInt>(); // #INST
+}
diff --git a/clang/test/SemaOpenACC/compute-construct-intexpr-clause-ast.cpp b/clang/test/SemaOpenACC/compute-construct-intexpr-clause-ast.cpp
index 5a4c9f05ee089e..b85de56c7ae9cc 100644
--- a/clang/test/SemaOpenACC/compute-construct-intexpr-clause-ast.cpp
+++ b/clang/test/SemaOpenACC/compute-construct-intexpr-clause-ast.cpp
@@ -116,8 +116,28 @@ void NormalUses() {
   // CHECK-NEXT: WhileStmt
   // CHECK-NEXT: CXXBoolLiteralExpr
   // CHECK-NEXT: CompoundStmt
+
+#pragma acc kernels async(some_int())
+  while(true){}
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels
+  // CHECK-NEXT: async clause
+  // CHECK-NEXT: CallExpr{{.*}}'int'
+  // CHECK-NEXT: ImplicitCastExpr{{.*}}'int (*)()' <FunctionToPointerDecay>
+  // CHECK-NEXT: DeclRefExpr{{.*}}'int ()' lvalue Function{{.*}} 'some_int' 'int ()'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc kernels async
+  while(true){}
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels
+  // CHECK-NEXT: async clause
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
 }
 
+
 template<typename T, typename U>
 void TemplUses(T t, U u) {
   // CHECK-NEXT: FunctionTemplateDecl
@@ -235,6 +255,33 @@ void TemplUses(T t, U u) {
   // CHECK-NEXT: CXXBoolLiteralExpr
   // CHECK-NEXT: CompoundStmt
 
+#pragma acc kernels async
+  while(true){}
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels
+  // CHECK-NEXT: async clause
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc kernels async(u)
+  while(true){}
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels
+  // CHECK-NEXT: async clause
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'U' lvalue ParmVar{{.*}} 'u' 'U'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+
+#pragma acc parallel async (U::value)
+  while(true){}
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+  // CHECK-NEXT: async clause
+  // CHECK-NEXT: DependentScopeDeclRefExpr{{.*}} '<dependent type>' lvalue
+  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'U'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+
 
   // CHECK-NEXT: DeclStmt
   // CHECK-NEXT: VarDecl{{.*}}EndMarker
@@ -365,6 +412,31 @@ void TemplUses(T t, U u) {
   // CHECK-NEXT: CXXBoolLiteralExpr
   // CHECK-NEXT: CompoundStmt
 
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels
+  // CHECK-NEXT: async clause
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}kernels
+  // CHECK-NEXT: async 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: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: OpenACCComputeConstruct{{.*}}parallel
+  // CHECK-NEXT: async clause
+  // CHECK-NEXT: ImplicitCastExpr{{.*}} 'int' <LValueToRValue>
+  // CHECK-NEXT: DeclRefExpr{{.*}} 'const int' lvalue Var{{.*}} 'value' 'const int'
+  // CHECK-NEXT: NestedNameSpecifier TypeSpec 'HasInt'
+  // CHECK-NEXT: WhileStmt
+  // CHECK-NEXT: CXXBoolLiteralExpr
+  // CHECK-NEXT: CompoundStmt
+
   // CHECK-NEXT: DeclStmt
   // CHECK-NEXT: VarDecl{{.*}}EndMarker
 }
diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp
index 6c07c4d2e30738..b845a381d63bed 100644
--- a/clang/tools/libclang/CIndex.cpp
+++ b/clang/tools/libclang/CIndex.cpp
@@ -2847,6 +2847,10 @@ void OpenACCClauseEnqueue::VisitDevicePtrClause(
     const OpenACCDevicePtrClause &C) {
   VisitVarList(C);
 }
+void OpenACCClauseEnqueue::VisitAsyncClause(const OpenACCAsyncClause &C) {
+  if (C.hasIntExpr())
+    Visitor.AddStmt(C.getIntExpr());
+}
 } // namespace
 
 void EnqueueVisitor::EnqueueChildren(const OpenACCClause *C) {