[OpenACC] Implement 'default_async' sema

A fairly simple one, only valid on the 'set' construct, this clause
takes an int expression.  Most of the work was already done as a part of
parsing, so this patch ends up being a lot of infrastructure.

clang/include/clang/AST/OpenACCClause.h

@@ -633,6 +633,19 @@ class OpenACCDeviceNumClause : public OpenACCClauseWithSingleIntExpr {
                                         SourceLocation EndLoc);
 };
 
+class OpenACCDefaultAsyncClause : public OpenACCClauseWithSingleIntExpr {
+  OpenACCDefaultAsyncClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
+                            Expr *IntExpr, SourceLocation EndLoc);
+
+public:
+  static bool classof(const OpenACCClause *C) {
+    return C->getClauseKind() == OpenACCClauseKind::DefaultAsync;
+  }
+  static OpenACCDefaultAsyncClause *
+  Create(const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc,
+         Expr *IntExpr, SourceLocation EndLoc);
+};
+
 /// Represents a 'collapse' clause on a 'loop' construct. This clause takes an
 /// integer constant expression 'N' that represents how deep to collapse the
 /// construct. It also takes an optional 'force' tag that permits intervening

clang/include/clang/Basic/OpenACCClauses.def

@@ -38,6 +38,7 @@ VISIT_CLAUSE(Create)
 CLAUSE_ALIAS(PCreate, Create, true)
 CLAUSE_ALIAS(PresentOrCreate, Create, true)
 VISIT_CLAUSE(Default)
+VISIT_CLAUSE(DefaultAsync)
 VISIT_CLAUSE(Delete)
 VISIT_CLAUSE(Detach)
 VISIT_CLAUSE(DeviceNum)

clang/include/clang/Sema/SemaOpenACC.h

@@ -297,6 +297,7 @@ class SemaOpenACC : public SemaBase {
               ClauseKind == OpenACCClauseKind::NumWorkers ||
               ClauseKind == OpenACCClauseKind::Async ||
               ClauseKind == OpenACCClauseKind::DeviceNum ||
+              ClauseKind == OpenACCClauseKind::DefaultAsync ||
               ClauseKind == OpenACCClauseKind::Tile ||
               ClauseKind == OpenACCClauseKind::Worker ||
               ClauseKind == OpenACCClauseKind::Vector ||
@@ -349,6 +350,7 @@ class SemaOpenACC : public SemaBase {
               ClauseKind == OpenACCClauseKind::NumWorkers ||
               ClauseKind == OpenACCClauseKind::Async ||
               ClauseKind == OpenACCClauseKind::DeviceNum ||
+              ClauseKind == OpenACCClauseKind::DefaultAsync ||
               ClauseKind == OpenACCClauseKind::Tile ||
               ClauseKind == OpenACCClauseKind::Gang ||
               ClauseKind == OpenACCClauseKind::Worker ||
@@ -486,6 +488,7 @@ class SemaOpenACC : public SemaBase {
               ClauseKind == OpenACCClauseKind::NumWorkers ||
               ClauseKind == OpenACCClauseKind::Async ||
               ClauseKind == OpenACCClauseKind::DeviceNum ||
+              ClauseKind == OpenACCClauseKind::DefaultAsync ||
               ClauseKind == OpenACCClauseKind::Tile ||
               ClauseKind == OpenACCClauseKind::Worker ||
               ClauseKind == OpenACCClauseKind::Vector ||
@@ -498,6 +501,7 @@ class SemaOpenACC : public SemaBase {
               ClauseKind == OpenACCClauseKind::NumWorkers ||
               ClauseKind == OpenACCClauseKind::Async ||
               ClauseKind == OpenACCClauseKind::DeviceNum ||
+              ClauseKind == OpenACCClauseKind::DefaultAsync ||
               ClauseKind == OpenACCClauseKind::Tile ||
               ClauseKind == OpenACCClauseKind::Worker ||
               ClauseKind == OpenACCClauseKind::Vector ||

clang/lib/AST/OpenACCClause.cpp

@@ -47,6 +47,7 @@ bool OpenACCClauseWithSingleIntExpr::classof(const OpenACCClause *C) {
   return OpenACCNumWorkersClause::classof(C) ||
          OpenACCVectorLengthClause::classof(C) ||
          OpenACCDeviceNumClause::classof(C) ||
+         OpenACCDefaultAsyncClause::classof(C) ||
          OpenACCVectorClause::classof(C) || OpenACCWorkerClause::classof(C) ||
          OpenACCCollapseClause::classof(C) || OpenACCAsyncClause::classof(C);
 }
@@ -239,6 +240,27 @@ OpenACCDeviceNumClause *OpenACCDeviceNumClause::Create(const ASTContext &C,
   return new (Mem) OpenACCDeviceNumClause(BeginLoc, LParenLoc, IntExpr, EndLoc);
 }
 
+OpenACCDefaultAsyncClause::OpenACCDefaultAsyncClause(SourceLocation BeginLoc,
+                                                     SourceLocation LParenLoc,
+                                                     Expr *IntExpr,
+                                                     SourceLocation EndLoc)
+    : OpenACCClauseWithSingleIntExpr(OpenACCClauseKind::DefaultAsync, BeginLoc,
+                                     LParenLoc, IntExpr, EndLoc) {
+  assert((IntExpr->isInstantiationDependent() ||
+          IntExpr->getType()->isIntegerType()) &&
+         "default_async expression type not scalar/dependent");
+}
+
+OpenACCDefaultAsyncClause *
+OpenACCDefaultAsyncClause::Create(const ASTContext &C, SourceLocation BeginLoc,
+                                  SourceLocation LParenLoc, Expr *IntExpr,
+                                  SourceLocation EndLoc) {
+  void *Mem = C.Allocate(sizeof(OpenACCDefaultAsyncClause),
+                         alignof(OpenACCDefaultAsyncClause));
+  return new (Mem)
+      OpenACCDefaultAsyncClause(BeginLoc, LParenLoc, IntExpr, EndLoc);
+}
+
 OpenACCWaitClause *OpenACCWaitClause::Create(
     const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc,
     Expr *DevNumExpr, SourceLocation QueuesLoc, ArrayRef<Expr *> QueueIdExprs,
@@ -575,6 +597,13 @@ void OpenACCClausePrinter::VisitDeviceNumClause(
   OS << ")";
 }
 
+void OpenACCClausePrinter::VisitDefaultAsyncClause(
+    const OpenACCDefaultAsyncClause &C) {
+  OS << "default_async(";
+  printExpr(C.getIntExpr());
+  OS << ")";
+}
+
 void OpenACCClausePrinter::VisitAsyncClause(const OpenACCAsyncClause &C) {
   OS << "async";
   if (C.hasIntExpr()) {

clang/lib/AST/StmtProfile.cpp

@@ -2650,6 +2650,11 @@ void OpenACCClauseProfiler::VisitDeviceNumClause(
   Profiler.VisitStmt(Clause.getIntExpr());
 }
 
+void OpenACCClauseProfiler::VisitDefaultAsyncClause(
+    const OpenACCDefaultAsyncClause &Clause) {
+  Profiler.VisitStmt(Clause.getIntExpr());
+}
+
 void OpenACCClauseProfiler::VisitWorkerClause(
     const OpenACCWorkerClause &Clause) {
   if (Clause.hasIntExpr())

clang/lib/AST/TextNodeDumper.cpp

@@ -414,6 +414,7 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
     case OpenACCClauseKind::Detach:
     case OpenACCClauseKind::Delete:
     case OpenACCClauseKind::DeviceNum:
+    case OpenACCClauseKind::DefaultAsync:
     case OpenACCClauseKind::DevicePtr:
     case OpenACCClauseKind::Finalize:
     case OpenACCClauseKind::FirstPrivate:

clang/lib/Parse/ParseOpenACC.cpp

@@ -1082,13 +1082,7 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
         return OpenACCCanContinue();
       }
 
-      // TODO OpenACC: as we implement the 'rest' of the above, this 'if' should
-      // be removed leaving just the 'setIntExprDetails'.
-      if (ClauseKind == OpenACCClauseKind::NumWorkers ||
-          ClauseKind == OpenACCClauseKind::DeviceNum ||
-          ClauseKind == OpenACCClauseKind::VectorLength)
-        ParsedClause.setIntExprDetails(IntExpr.get());
-
+      ParsedClause.setIntExprDetails(IntExpr.get());
       break;
     }
     case OpenACCClauseKind::DType:

clang/lib/Sema/SemaOpenACC.cpp

@@ -463,6 +463,14 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
       return false;
     }
   }
+  case OpenACCClauseKind::DefaultAsync: {
+    switch (DirectiveKind) {
+    case OpenACCDirectiveKind::Set:
+      return true;
+    default:
+      return false;
+    }
+  }
   }
 
   default:
@@ -977,6 +985,20 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitDeviceNumClause(
       Clause.getEndLoc());
 }
 
+OpenACCClause *SemaOpenACCClauseVisitor::VisitDefaultAsyncClause(
+    SemaOpenACC::OpenACCParsedClause &Clause) {
+  // OpenACC 3.3 2.14.3: Two instances of the same clause may not appear on the
+  // same directive.
+  if (checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, Clause))
+    return nullptr;
+
+  assert(Clause.getNumIntExprs() == 1 &&
+         "Invalid number of expressions for default_async");
+  return OpenACCDefaultAsyncClause::Create(
+      Ctx, Clause.getBeginLoc(), Clause.getLParenLoc(), Clause.getIntExprs()[0],
+      Clause.getEndLoc());
+}
+
 OpenACCClause *SemaOpenACCClauseVisitor::VisitPrivateClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
   // ActOnVar ensured that everything is a valid variable reference, so there
@@ -3681,12 +3703,9 @@ bool SemaOpenACC::ActOnStartStmtDirective(
   if (K == OpenACCDirectiveKind::Set &&
       llvm::find_if(
           Clauses,
-          llvm::IsaPred</*OpenACCDefaultAsyncClause,*/ // TODO: ERICH Need to
-                                                       // implement.
-                                                       // DefaultAsyncClause
-                                                       // then enable this.
-                        OpenACCDeviceNumClause, OpenACCDeviceTypeClause,
-                        OpenACCIfClause>) == Clauses.end())
+          llvm::IsaPred<OpenACCDefaultAsyncClause, OpenACCDeviceNumClause,
+                        OpenACCDeviceTypeClause, OpenACCIfClause>) ==
+          Clauses.end())
     return Diag(StartLoc, diag::err_acc_construct_one_clause_of)
            << K
            << GetListOfClauses({OpenACCClauseKind::DefaultAsync,

clang/lib/Sema/TreeTransform.h

@@ -11909,6 +11909,29 @@ void OpenACCClauseTransform<Derived>::VisitDeviceNumClause (
       ParsedClause.getEndLoc());
 }
 
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitDefaultAsyncClause(
+    const OpenACCDefaultAsyncClause &C) {
+  Expr *IntExpr = const_cast<Expr *>(C.getIntExpr());
+  assert(IntExpr && "default_async clause constructed with invalid int expr");
+
+  ExprResult Res = Self.TransformExpr(IntExpr);
+  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 = OpenACCDefaultAsyncClause::Create(
+      Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
+      ParsedClause.getLParenLoc(), ParsedClause.getIntExprs()[0],
+      ParsedClause.getEndLoc());
+}
+
 template <typename Derived>
 void OpenACCClauseTransform<Derived>::VisitVectorLengthClause(
     const OpenACCVectorLengthClause &C) {

clang/lib/Serialization/ASTReader.cpp

@@ -12412,6 +12412,12 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
     return OpenACCDeviceNumClause::Create(getContext(), BeginLoc, LParenLoc,
                                            IntExpr, EndLoc);
   }
+  case OpenACCClauseKind::DefaultAsync: {
+    SourceLocation LParenLoc = readSourceLocation();
+    Expr *IntExpr = readSubExpr();
+    return OpenACCDefaultAsyncClause::Create(getContext(), BeginLoc, LParenLoc,
+                                             IntExpr, EndLoc);
+  }
   case OpenACCClauseKind::VectorLength: {
     SourceLocation LParenLoc = readSourceLocation();
     Expr *IntExpr = readSubExpr();
@@ -12601,7 +12607,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
   case OpenACCClauseKind::Host:
   case OpenACCClauseKind::Link:
   case OpenACCClauseKind::Bind:
-  case OpenACCClauseKind::DefaultAsync:
   case OpenACCClauseKind::Invalid:
     llvm_unreachable("Clause serialization not yet implemented");
   }

clang/lib/Serialization/ASTWriter.cpp

@@ -8340,6 +8340,12 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
     AddStmt(const_cast<Expr*>(DNC->getIntExpr()));
     return;
   }
+  case OpenACCClauseKind::DefaultAsync: {
+    const auto *DAC = cast<OpenACCDefaultAsyncClause>(C);
+    writeSourceLocation(DAC->getLParenLoc());
+    AddStmt(const_cast<Expr *>(DAC->getIntExpr()));
+    return;
+  }
   case OpenACCClauseKind::NumWorkers: {
     const auto *NWC = cast<OpenACCNumWorkersClause>(C);
     writeSourceLocation(NWC->getLParenLoc());
@@ -8536,7 +8542,6 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
   case OpenACCClauseKind::Host:
   case OpenACCClauseKind::Link:
   case OpenACCClauseKind::Bind:
-  case OpenACCClauseKind::DefaultAsync:
   case OpenACCClauseKind::Invalid:
     llvm_unreachable("Clause serialization not yet implemented");
   }

clang/test/AST/ast-print-openacc-set-construct.cpp

@@ -3,19 +3,18 @@
 unsigned Int;
 
 void uses() {
-// CHECK: #pragma acc set if(Int == 5) device_type(I) device_num(Int)
+// CHECK: #pragma acc set default_async(Int) if(Int == 5) device_type(I) device_num(Int)
 #pragma acc set default_async(Int) if (Int == 5) device_type(I) device_num(Int)
-// CHECK: #pragma acc set device_type(I) device_num(Int)
+// CHECK: #pragma acc set default_async(Int) device_type(I) device_num(Int)
 #pragma acc set default_async(Int) device_type(I) device_num(Int)
-// CHECK: #pragma acc set if(Int == 5) device_num(Int)
+// CHECK: #pragma acc set default_async(Int) if(Int == 5) device_num(Int)
 #pragma acc set default_async(Int) if (Int == 5) device_num(Int)
-// CHECK: #pragma acc set if(Int == 5) device_type(I)
+// CHECK: #pragma acc set default_async(Int) if(Int == 5) device_type(I)
 #pragma acc set default_async(Int) if (Int == 5) device_type(I)
 // CHECK: #pragma acc set if(Int == 5) device_type(I) device_num(Int)
 #pragma acc set if (Int == 5) device_type(I) device_num(Int)
-// CHECK-NOT: default_async(Int) 
-// TODO: OpenACC: Enable this when we have default_async implemented.
-// #pragma acc set default_async(Int)
+// CHECK: #pragma acc set default_async(Int)
+#pragma acc set default_async(Int)
 // CHECK: #pragma acc set if(Int == 5)
 #pragma acc set if (Int == 5)
 // CHECK: #pragma acc set device_type(I)

clang/test/ParserOpenACC/parse-clauses.c

@@ -837,12 +837,8 @@ void IntExprParsing() {
   // expected-error@+1{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', 'device_type' or 'if' clause}}
 #pragma acc set default_async(5, 4)
 
-  // expected-error@+2{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', 'device_type' or 'if' clause}}
-  // expected-warning@+1{{OpenACC clause 'default_async' not yet implemented, clause ignored}}
 #pragma acc set default_async(5)
 
-  // expected-error@+2{{OpenACC 'set' construct must have at least one 'default_async', 'device_num', 'device_type' or 'if' clause}}
-  // expected-warning@+1{{OpenACC clause 'default_async' not yet implemented, clause ignored}}
 #pragma acc set default_async(returns_int())
 
 

clang/test/SemaOpenACC/combined-construct-auto_seq_independent-clauses.c

@@ -137,7 +137,7 @@ void uses() {
   // expected-error@+1{{OpenACC 'device_num' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop auto device_num(1)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'default_async' not yet implemented}}
+  // expected-error@+1{{OpenACC 'default_async' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop auto default_async(1)
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop auto device_type(*)
@@ -254,7 +254,7 @@ void uses() {
   // expected-error@+1{{OpenACC 'device_num' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop device_num(1) auto
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'default_async' not yet implemented}}
+  // expected-error@+1{{OpenACC 'default_async' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop default_async(1) auto
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop device_type(*) auto
@@ -372,7 +372,7 @@ void uses() {
   // expected-error@+1{{OpenACC 'device_num' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop independent device_num(1)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'default_async' not yet implemented}}
+  // expected-error@+1{{OpenACC 'default_async' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop independent default_async(1)
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop independent device_type(*)
@@ -489,7 +489,7 @@ void uses() {
   // expected-error@+1{{OpenACC 'device_num' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop device_num(1) independent
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'default_async' not yet implemented}}
+  // expected-error@+1{{OpenACC 'default_async' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop default_async(1) independent
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop device_type(*) independent
@@ -615,7 +615,7 @@ void uses() {
   // expected-error@+1{{OpenACC 'device_num' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop seq device_num(1)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'default_async' not yet implemented}}
+  // expected-error@+1{{OpenACC 'default_async' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop seq default_async(1)
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop seq device_type(*)
@@ -738,7 +738,7 @@ void uses() {
   // expected-error@+1{{OpenACC 'device_num' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop device_num(1) seq
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'default_async' not yet implemented}}
+  // expected-error@+1{{OpenACC 'default_async' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop default_async(1) seq
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop device_type(*) seq

clang/test/SemaOpenACC/combined-construct-device_type-clause.c

@@ -193,8 +193,7 @@ void uses() {
   // expected-error@+1{{OpenACC 'device_num' clause is not valid on 'serial loop' directive}}
 #pragma acc serial loop device_type(*) device_num(1)
   for(int i = 0; i < 5; ++i);
-  // expected-error@+2{{OpenACC clause 'default_async' may not follow a 'device_type' clause in a 'serial loop' construct}}
-  // expected-note@+1{{previous clause is here}}
+  // expected-error@+1{{OpenACC 'default_async' clause is not valid on 'serial loop' directive}}
 #pragma acc serial loop device_type(*) default_async(1)
   for(int i = 0; i < 5; ++i);
 #pragma acc parallel loop device_type(*) async

clang/test/SemaOpenACC/compute-construct-device_type-clause.c

@@ -198,8 +198,7 @@ void uses() {
   // expected-error@+1{{OpenACC 'device_num' clause is not valid on 'kernels' directive}}
 #pragma acc kernels device_type(*) device_num(1)
   while(1);
-  // expected-error@+2{{OpenACC clause 'default_async' may not follow a 'device_type' clause in a 'kernels' construct}}
-  // expected-note@+1{{previous clause is here}}
+  // expected-error@+1{{OpenACC 'default_async' clause is not valid on 'kernels' directive}}
 #pragma acc kernels device_type(*) default_async(1)
   while(1);
 #pragma acc kernels device_type(*) async