[OpenACC] Implement 'delete' AST/Sema for 'exit data' construct

'delete' is another clause that has very little compile-time
implication, but needs a full AST that takes a var list.  This patch
ipmlements it fully, plus adds sufficient test coverage.

clang/include/clang/AST/OpenACCClause.h

@@ -766,6 +766,29 @@ class OpenACCDetachClause final
          ArrayRef<Expr *> VarList, SourceLocation EndLoc);
 };
 
+class OpenACCDeleteClause final
+    : public OpenACCClauseWithVarList,
+      public llvm::TrailingObjects<OpenACCDeleteClause, Expr *> {
+
+  OpenACCDeleteClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
+                      ArrayRef<Expr *> VarList, SourceLocation EndLoc)
+      : OpenACCClauseWithVarList(OpenACCClauseKind::Delete, BeginLoc, LParenLoc,
+                                 EndLoc) {
+    std::uninitialized_copy(VarList.begin(), VarList.end(),
+                            getTrailingObjects<Expr *>());
+    setExprs(MutableArrayRef(getTrailingObjects<Expr *>(), VarList.size()));
+  }
+
+public:
+  static bool classof(const OpenACCClause *C) {
+    return C->getClauseKind() == OpenACCClauseKind::Delete;
+  }
+  static OpenACCDeleteClause *
+  Create(const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc,
+         ArrayRef<Expr *> VarList, SourceLocation EndLoc);
+};
+
+
 class OpenACCNoCreateClause final
     : public OpenACCClauseWithVarList,
       public llvm::TrailingObjects<OpenACCNoCreateClause, Expr *> {

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(Delete)
 VISIT_CLAUSE(Detach)
 VISIT_CLAUSE(DevicePtr)
 VISIT_CLAUSE(DeviceType)

clang/include/clang/Sema/SemaOpenACC.h

@@ -399,6 +399,7 @@ class SemaOpenACC : public SemaBase {
               ClauseKind == OpenACCClauseKind::PCreate ||
               ClauseKind == OpenACCClauseKind::PresentOrCreate ||
               ClauseKind == OpenACCClauseKind::Attach ||
+              ClauseKind == OpenACCClauseKind::Delete ||
               ClauseKind == OpenACCClauseKind::Detach ||
               ClauseKind == OpenACCClauseKind::DevicePtr ||
               ClauseKind == OpenACCClauseKind::Reduction ||
@@ -536,6 +537,7 @@ class SemaOpenACC : public SemaBase {
               ClauseKind == OpenACCClauseKind::PCreate ||
               ClauseKind == OpenACCClauseKind::PresentOrCreate ||
               ClauseKind == OpenACCClauseKind::Attach ||
+              ClauseKind == OpenACCClauseKind::Delete ||
               ClauseKind == OpenACCClauseKind::Detach ||
               ClauseKind == OpenACCClauseKind::DevicePtr ||
               ClauseKind == OpenACCClauseKind::FirstPrivate) &&
@@ -573,6 +575,7 @@ class SemaOpenACC : public SemaBase {
               ClauseKind == OpenACCClauseKind::PCreate ||
               ClauseKind == OpenACCClauseKind::PresentOrCreate ||
               ClauseKind == OpenACCClauseKind::Attach ||
+              ClauseKind == OpenACCClauseKind::Delete ||
               ClauseKind == OpenACCClauseKind::Detach ||
               ClauseKind == OpenACCClauseKind::DevicePtr ||
               ClauseKind == OpenACCClauseKind::FirstPrivate) &&

clang/lib/AST/OpenACCClause.cpp

@@ -32,7 +32,7 @@ bool OpenACCClauseWithVarList::classof(const OpenACCClause *C) {
   return OpenACCPrivateClause::classof(C) ||
          OpenACCFirstPrivateClause::classof(C) ||
          OpenACCDevicePtrClause::classof(C) ||
-         OpenACCDevicePtrClause::classof(C) ||
+         OpenACCDeleteClause::classof(C) ||
          OpenACCDetachClause::classof(C) || OpenACCAttachClause::classof(C) ||
          OpenACCNoCreateClause::classof(C) ||
          OpenACCPresentClause::classof(C) || OpenACCCopyClause::classof(C) ||
@@ -288,6 +288,16 @@ OpenACCDetachClause *OpenACCDetachClause::Create(const ASTContext &C,
   return new (Mem) OpenACCDetachClause(BeginLoc, LParenLoc, VarList, EndLoc);
 }
 
+OpenACCDeleteClause *OpenACCDeleteClause::Create(const ASTContext &C,
+                                                 SourceLocation BeginLoc,
+                                                 SourceLocation LParenLoc,
+                                                 ArrayRef<Expr *> VarList,
+                                                 SourceLocation EndLoc) {
+  void *Mem =
+      C.Allocate(OpenACCDeleteClause::totalSizeToAlloc<Expr *>(VarList.size()));
+  return new (Mem) OpenACCDeleteClause(BeginLoc, LParenLoc, VarList, EndLoc);
+}
+
 OpenACCDevicePtrClause *OpenACCDevicePtrClause::Create(const ASTContext &C,
                                                        SourceLocation BeginLoc,
                                                        SourceLocation LParenLoc,
@@ -564,6 +574,13 @@ void OpenACCClausePrinter::VisitDetachClause(const OpenACCDetachClause &C) {
   OS << ")";
 }
 
+void OpenACCClausePrinter::VisitDeleteClause(const OpenACCDeleteClause &C) {
+  OS << "delete(";
+  llvm::interleaveComma(C.getVarList(), OS,
+                        [&](const Expr *E) { printExpr(E); });
+  OS << ")";
+}
+
 void OpenACCClausePrinter::VisitDevicePtrClause(
     const OpenACCDevicePtrClause &C) {
   OS << "deviceptr(";

clang/lib/AST/StmtProfile.cpp

@@ -2611,6 +2611,12 @@ void OpenACCClauseProfiler::VisitDetachClause(
     Profiler.VisitStmt(E);
 }
 
+void OpenACCClauseProfiler::VisitDeleteClause(
+    const OpenACCDeleteClause &Clause) {
+  for (auto *E : Clause.getVarList())
+    Profiler.VisitStmt(E);
+}
+
 void OpenACCClauseProfiler::VisitDevicePtrClause(
     const OpenACCDevicePtrClause &Clause) {
   for (auto *E : Clause.getVarList())

clang/lib/AST/TextNodeDumper.cpp

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

clang/lib/Parse/ParseOpenACC.cpp

@@ -998,7 +998,6 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
       // make sure we get the right differentiator.
       assert(DirKind == OpenACCDirectiveKind::Update);
       [[fallthrough]];
-    case OpenACCClauseKind::Delete:
     case OpenACCClauseKind::Device:
     case OpenACCClauseKind::DeviceResident:
     case OpenACCClauseKind::Host:
@@ -1007,6 +1006,7 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
       ParseOpenACCVarList(ClauseKind);
       break;
     case OpenACCClauseKind::Attach:
+    case OpenACCClauseKind::Delete:
     case OpenACCClauseKind::Detach:
     case OpenACCClauseKind::DevicePtr:
       ParsedClause.setVarListDetails(ParseOpenACCVarList(ClauseKind),

clang/lib/Sema/SemaOpenACC.cpp

@@ -425,6 +425,15 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
       return false;
     }
   }
+  case OpenACCClauseKind::Delete: {
+    switch (DirectiveKind) {
+    case OpenACCDirectiveKind::ExitData:
+      return true;
+    default:
+      return false;
+    }
+  }
+
   case OpenACCClauseKind::Detach: {
     switch (DirectiveKind) {
     case OpenACCDirectiveKind::ExitData:
@@ -1066,6 +1075,17 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitDetachClause(
                                      Clause.getEndLoc());
 }
 
+OpenACCClause *SemaOpenACCClauseVisitor::VisitDeleteClause(
+    SemaOpenACC::OpenACCParsedClause &Clause) {
+  // ActOnVar ensured that everything is a valid variable reference, so there
+  // really isn't anything to do here. GCC does some duplicate-finding, though
+  // it isn't apparent in the standard where this is justified.
+  return OpenACCDeleteClause::Create(Ctx, Clause.getBeginLoc(),
+                                     Clause.getLParenLoc(), Clause.getVarList(),
+                                     Clause.getEndLoc());
+}
+
+
 OpenACCClause *SemaOpenACCClauseVisitor::VisitDevicePtrClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
   // Restrictions only properly implemented on 'compute'/'combined'/'data'

clang/lib/Sema/TreeTransform.h

@@ -11777,6 +11777,17 @@ void OpenACCClauseTransform<Derived>::VisitDetachClause(
       ParsedClause.getEndLoc());
 }
 
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitDeleteClause(
+    const OpenACCDeleteClause &C) {
+  ParsedClause.setVarListDetails(VisitVarList(C.getVarList()),
+                                 /*IsReadOnly=*/false, /*IsZero=*/false);
+  NewClause = OpenACCDeleteClause::Create(
+      Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
+      ParsedClause.getLParenLoc(), ParsedClause.getVarList(),
+      ParsedClause.getEndLoc());
+}
+
 template <typename Derived>
 void OpenACCClauseTransform<Derived>::VisitDevicePtrClause(
     const OpenACCDevicePtrClause &C) {

clang/lib/Serialization/ASTReader.cpp

@@ -12435,6 +12435,12 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
     return OpenACCDetachClause::Create(getContext(), BeginLoc, LParenLoc,
                                        VarList, EndLoc);
   }
+  case OpenACCClauseKind::Delete: {
+    SourceLocation LParenLoc = readSourceLocation();
+    llvm::SmallVector<Expr *> VarList = readOpenACCVarList();
+    return OpenACCDeleteClause::Create(getContext(), BeginLoc, LParenLoc,
+                                       VarList, EndLoc);
+  }
   case OpenACCClauseKind::DevicePtr: {
     SourceLocation LParenLoc = readSourceLocation();
     llvm::SmallVector<Expr *> VarList = readOpenACCVarList();
@@ -12578,7 +12584,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
 
   case OpenACCClauseKind::NoHost:
   case OpenACCClauseKind::UseDevice:
-  case OpenACCClauseKind::Delete:
   case OpenACCClauseKind::Device:
   case OpenACCClauseKind::DeviceResident:
   case OpenACCClauseKind::Host:

clang/lib/Serialization/ASTWriter.cpp

@@ -8362,6 +8362,12 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
     writeOpenACCVarList(DC);
     return;
   }
+  case OpenACCClauseKind::Delete: {
+    const auto *DC = cast<OpenACCDeleteClause>(C);
+    writeSourceLocation(DC->getLParenLoc());
+    writeOpenACCVarList(DC);
+    return;
+  }
   case OpenACCClauseKind::DevicePtr: {
     const auto *DPC = cast<OpenACCDevicePtrClause>(C);
     writeSourceLocation(DPC->getLParenLoc());
@@ -8506,7 +8512,6 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
 
   case OpenACCClauseKind::NoHost:
   case OpenACCClauseKind::UseDevice:
-  case OpenACCClauseKind::Delete:
   case OpenACCClauseKind::Device:
   case OpenACCClauseKind::DeviceResident:
   case OpenACCClauseKind::Host:

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

@@ -120,4 +120,10 @@ void foo() {
 // CHECK: #pragma acc exit data copyout(i) detach(iPtr, arrayPtr[0])
 #pragma acc exit data copyout(i) detach(iPtr, arrayPtr[0])
 
+// CHECK: #pragma acc exit data copyout(i) delete(i, array[1], array, array[1:2])
+#pragma acc exit data copyout(i) delete(i, array[1], array, array[1:2])
+  ;
+
+// CHECK: #pragma acc exit data copyout(i) delete(i, array[1], array, array[1:2])
+#pragma acc exit data copyout(i) delete(i, array[1], array, array[1:2])
 }

clang/test/ParserOpenACC/parse-clauses.c

@@ -510,12 +510,10 @@ void VarListClauses() {
 #pragma acc serial firstprivate(s.array[s.value : 5], s.value), self
   for(int i = 0; i < 5;++i) {}
 
-  // expected-error@+2{{expected ','}}
-  // expected-warning@+1{{OpenACC clause 'delete' not yet implemented, clause ignored}}
+  // expected-error@+1{{expected ','}}
 #pragma acc exit data delete(s.array[s.value] s.array[s.value :5] ) async
   for(int i = 0; i < 5;++i) {}
 
-  // expected-warning@+1{{OpenACC clause 'delete' not yet implemented, clause ignored}}
 #pragma acc exit data delete(s.array[s.value : 5], s.value),async
   for(int i = 0; i < 5;++i) {}
 

clang/test/ParserOpenACC/parse-clauses.cpp

@@ -35,8 +35,9 @@ void templ() {
 #pragma acc parallel async
   for(;;){}
 
-  // expected-warning@+1{{OpenACC clause 'delete' not yet implemented, clause ignored}}
-#pragma acc exit data delete(I)
+
+  T t;
+#pragma acc exit data delete(t)
   ;
 }
 

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

@@ -69,7 +69,7 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop auto attach(VarPtr)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'delete' not yet implemented}}
+  // expected-error@+1{{OpenACC 'delete' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop auto delete(Var)
   for(unsigned i = 0; i < 5; ++i);
   // expected-error@+1{{OpenACC 'detach' clause is not valid on 'parallel loop' directive}}
@@ -186,7 +186,7 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop attach(VarPtr) auto
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'delete' not yet implemented}}
+  // expected-error@+1{{OpenACC 'delete' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop delete(Var) auto
   for(unsigned i = 0; i < 5; ++i);
   // expected-error@+1{{OpenACC 'detach' clause is not valid on 'parallel loop' directive}}
@@ -304,7 +304,7 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop independent attach(VarPtr)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'delete' not yet implemented}}
+  // expected-error@+1{{OpenACC 'delete' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop independent delete(Var)
   for(unsigned i = 0; i < 5; ++i);
   // expected-error@+1{{OpenACC 'detach' clause is not valid on 'parallel loop' directive}}
@@ -421,7 +421,7 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop attach(VarPtr) independent
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'delete' not yet implemented}}
+  // expected-error@+1{{OpenACC 'delete' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop delete(Var) independent
   for(unsigned i = 0; i < 5; ++i);
   // expected-error@+1{{OpenACC 'detach' clause is not valid on 'parallel loop' directive}}
@@ -547,7 +547,7 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop seq attach(VarPtr)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'delete' not yet implemented}}
+  // expected-error@+1{{OpenACC 'delete' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop seq delete(Var)
   for(unsigned i = 0; i < 5; ++i);
   // expected-error@+1{{OpenACC 'detach' clause is not valid on 'parallel loop' directive}}
@@ -670,7 +670,7 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop attach(VarPtr) seq
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'delete' not yet implemented}}
+  // expected-error@+1{{OpenACC 'delete' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop delete(Var) seq
   for(unsigned i = 0; i < 5; ++i);
   // expected-error@+1{{OpenACC 'detach' clause is not valid on 'parallel loop' directive}}

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

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

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

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