[OpenACC] Implement 'finalize' clause sema

This is a very simple clause as far as sema is concerned.  It is only
valid on 'exit data', and doesn't have any rules involving it, so it is
simply applied and passed onto the MLIR.

clang/include/clang/AST/OpenACCClause.h

@@ -38,7 +38,7 @@ public:
   SourceLocation getBeginLoc() const { return Location.getBegin(); }
   SourceLocation getEndLoc() const { return Location.getEnd(); }
 
-  static bool classof(const OpenACCClause *) { return false; }
+  static bool classof(const OpenACCClause *) { return true; }
 
   using child_iterator = StmtIterator;
   using const_child_iterator = ConstStmtIterator;
@@ -76,6 +76,28 @@ public:
   }
 };
 
+// Represents the 'finalize' clause.
+class OpenACCFinalizeClause : public OpenACCClause {
+protected:
+  OpenACCFinalizeClause(SourceLocation BeginLoc, SourceLocation EndLoc)
+      : OpenACCClause(OpenACCClauseKind::Finalize, BeginLoc, EndLoc) {}
+
+public:
+  static bool classof(const OpenACCClause *C) {
+    return C->getClauseKind() == OpenACCClauseKind::Finalize;
+  }
+
+  static OpenACCFinalizeClause *
+  Create(const ASTContext &Ctx, SourceLocation BeginLoc, SourceLocation EndLoc);
+
+  child_range children() {
+    return child_range(child_iterator(), child_iterator());
+  }
+  const_child_range children() const {
+    return const_child_range(const_child_iterator(), const_child_iterator());
+  }
+};
+
 // Represents the 'independent' clause.
 class OpenACCIndependentClause : public OpenACCClause {
 protected:

clang/include/clang/Basic/OpenACCClauses.def

@@ -41,6 +41,7 @@ VISIT_CLAUSE(Default)
 VISIT_CLAUSE(DevicePtr)
 VISIT_CLAUSE(DeviceType)
 CLAUSE_ALIAS(DType, DeviceType, false)
+VISIT_CLAUSE(Finalize)
 VISIT_CLAUSE(FirstPrivate)
 VISIT_CLAUSE(Gang)
 VISIT_CLAUSE(If)

clang/lib/AST/OpenACCClause.cpp

@@ -444,6 +444,14 @@ OpenACCVectorClause *OpenACCVectorClause::Create(const ASTContext &C,
   return new (Mem) OpenACCVectorClause(BeginLoc, LParenLoc, IntExpr, EndLoc);
 }
 
+OpenACCFinalizeClause *OpenACCFinalizeClause::Create(const ASTContext &C,
+                                                     SourceLocation BeginLoc,
+                                                     SourceLocation EndLoc) {
+  void *Mem =
+      C.Allocate(sizeof(OpenACCFinalizeClause), alignof(OpenACCFinalizeClause));
+  return new (Mem) OpenACCFinalizeClause(BeginLoc, EndLoc);
+}
+
 //===----------------------------------------------------------------------===//
 //  OpenACC clauses printing methods
 //===----------------------------------------------------------------------===//
@@ -685,3 +693,7 @@ void OpenACCClausePrinter::VisitVectorClause(const OpenACCVectorClause &C) {
     OS << ")";
   }
 }
+
+void OpenACCClausePrinter::VisitFinalizeClause(const OpenACCFinalizeClause &C) {
+  OS << "finalize";
+}

clang/lib/AST/StmtProfile.cpp

@@ -2558,6 +2558,9 @@ void OpenACCClauseProfiler::VisitSelfClause(const OpenACCSelfClause &Clause) {
     Profiler.VisitStmt(Clause.getConditionExpr());
 }
 
+void OpenACCClauseProfiler::VisitFinalizeClause(
+    const OpenACCFinalizeClause &Clause) {}
+
 void OpenACCClauseProfiler::VisitNumGangsClause(
     const OpenACCNumGangsClause &Clause) {
   for (auto *E : Clause.getIntExprs())

clang/lib/AST/TextNodeDumper.cpp

@@ -411,6 +411,7 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
     case OpenACCClauseKind::If:
     case OpenACCClauseKind::Independent:
     case OpenACCClauseKind::DevicePtr:
+    case OpenACCClauseKind::Finalize:
     case OpenACCClauseKind::FirstPrivate:
     case OpenACCClauseKind::NoCreate:
     case OpenACCClauseKind::NumGangs:

clang/lib/Sema/SemaOpenACC.cpp

@@ -408,6 +408,14 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
       return false;
     }
   }
+  case OpenACCClauseKind::Finalize: {
+    switch (DirectiveKind) {
+    case OpenACCDirectiveKind::ExitData:
+      return true;
+    default:
+      return false;
+    }
+  }
   }
 
   default:
@@ -1604,6 +1612,14 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitGangClause(
                                  GangKinds, IntExprs, Clause.getEndLoc());
 }
 
+OpenACCClause *SemaOpenACCClauseVisitor::VisitFinalizeClause(
+    SemaOpenACC::OpenACCParsedClause &Clause) {
+  // There isn't anything to do here, this is only valid on one construct, and
+  // has no associated rules.
+  return OpenACCFinalizeClause::Create(Ctx, Clause.getBeginLoc(),
+                                       Clause.getEndLoc());
+}
+
 OpenACCClause *SemaOpenACCClauseVisitor::VisitSeqClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
   // Restrictions only properly implemented on 'loop' constructs and combined ,

clang/lib/Sema/TreeTransform.h

@@ -11978,6 +11978,13 @@ void OpenACCClauseTransform<Derived>::VisitSeqClause(
                                        ParsedClause.getBeginLoc(),
                                        ParsedClause.getEndLoc());
 }
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitFinalizeClause(
+    const OpenACCFinalizeClause &C) {
+  NewClause = OpenACCFinalizeClause::Create(Self.getSema().getASTContext(),
+                                            ParsedClause.getBeginLoc(),
+                                            ParsedClause.getEndLoc());
+}
 
 template <typename Derived>
 void OpenACCClauseTransform<Derived>::VisitReductionClause(

clang/lib/Serialization/ASTReader.cpp

@@ -12534,6 +12534,8 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
   }
   case OpenACCClauseKind::Seq:
     return OpenACCSeqClause::Create(getContext(), BeginLoc, EndLoc);
+  case OpenACCClauseKind::Finalize:
+    return OpenACCFinalizeClause::Create(getContext(), BeginLoc, EndLoc);
   case OpenACCClauseKind::Independent:
     return OpenACCIndependentClause::Create(getContext(), BeginLoc, EndLoc);
   case OpenACCClauseKind::Auto:
@@ -12579,7 +12581,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
                                        VectorExpr, EndLoc);
   }
 
-  case OpenACCClauseKind::Finalize:
   case OpenACCClauseKind::IfPresent:
   case OpenACCClauseKind::NoHost:
   case OpenACCClauseKind::UseDevice:

clang/lib/Serialization/ASTWriter.cpp

@@ -8451,6 +8451,7 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
   case OpenACCClauseKind::Seq:
   case OpenACCClauseKind::Independent:
   case OpenACCClauseKind::Auto:
+  case OpenACCClauseKind::Finalize:
     // Nothing to do here, there is no additional information beyond the
     // begin/end loc and clause kind.
     return;
@@ -8496,7 +8497,6 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
     return;
   }
 
-  case OpenACCClauseKind::Finalize:
   case OpenACCClauseKind::IfPresent:
   case OpenACCClauseKind::NoHost:
   case OpenACCClauseKind::UseDevice:

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

@@ -110,4 +110,7 @@ void foo() {
 // CHECK: #pragma acc data default(none) attach(iPtr, arrayPtr[0])
 #pragma acc data default(none) attach(iPtr, arrayPtr[0])
   ;
+
+// CHECK: #pragma acc exit data copyout(i) finalize
+#pragma acc exit data copyout(i) finalize
 }

clang/test/ParserOpenACC/parse-clauses.c

@@ -4,23 +4,17 @@
 
 void func() {
 
-  // expected-warning@+1{{OpenACC clause 'finalize' not yet implemented, clause ignored}}
-#pragma acc enter data finalize
+#pragma acc exit data finalize
 
-  // expected-warning@+2{{OpenACC clause 'finalize' not yet implemented, clause ignored}}
-  // expected-warning@+1{{OpenACC clause 'finalize' not yet implemented, clause ignored}}
-#pragma acc enter data finalize finalize
+#pragma acc exit data finalize finalize
 
-  // expected-warning@+2{{OpenACC clause 'finalize' not yet implemented, clause ignored}}
   // expected-error@+1{{invalid OpenACC clause 'invalid'}}
-#pragma acc enter data finalize invalid
+#pragma acc exit data finalize invalid
 
-  // expected-warning@+2{{OpenACC clause 'finalize' not yet implemented, clause ignored}}
   // expected-error@+1{{invalid OpenACC clause 'invalid'}}
-#pragma acc enter data finalize invalid invalid finalize
+#pragma acc exit data finalize invalid invalid finalize
 
-  // expected-warning@+1{{OpenACC clause 'finalize' not yet implemented, clause ignored}}
-#pragma acc enter data wait finalize
+#pragma acc exit data wait finalize
 
   // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented, clause ignored}}
 #pragma acc host_data if_present

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

@@ -37,7 +37,7 @@ void uses() {
   int *VarPtr;
 
   // 'auto' can combine with any other clause.
-  // expected-warning@+1{{OpenACC clause 'finalize' not yet implemented}}
+  // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop auto finalize
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
@@ -154,7 +154,7 @@ void uses() {
 #pragma acc parallel loop auto wait
   for(unsigned i = 0; i < 5; ++i);
 
-  // expected-warning@+1{{OpenACC clause 'finalize' not yet implemented}}
+  // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop finalize auto
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
@@ -272,7 +272,7 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 
   // 'independent' can also be combined with any clauses
-  // expected-warning@+1{{OpenACC clause 'finalize' not yet implemented}}
+  // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop independent finalize
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
@@ -389,7 +389,7 @@ void uses() {
 #pragma acc parallel loop independent wait
   for(unsigned i = 0; i < 5; ++i);
 
-  // expected-warning@+1{{OpenACC clause 'finalize' not yet implemented}}
+  // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop finalize independent
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
@@ -519,7 +519,7 @@ void uses() {
   // expected-note@+1{{previous clause is here}}
 #pragma acc parallel loop seq vector
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'finalize' not yet implemented}}
+  // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop seq finalize
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
@@ -642,7 +642,7 @@ void uses() {
   // expected-note@+1{{previous clause is here}}
 #pragma acc parallel loop vector seq
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'finalize' not yet implemented}}
+  // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'parallel loop' directive}}
 #pragma acc parallel loop finalize seq
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}

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

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

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

@@ -42,8 +42,7 @@ void uses() {
 
   // Only 'async', 'wait', num_gangs', 'num_workers', 'vector_length' allowed after 'device_type'.
 
-  // expected-error@+2{{OpenACC clause 'finalize' may not follow a 'device_type' clause in a 'kernels' construct}}
-  // expected-note@+1{{previous clause is here}}
+  // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'kernels' directive}}
 #pragma acc kernels device_type(*) finalize
   while(1);
   // expected-error@+2{{OpenACC clause 'if_present' may not follow a 'device_type' clause in a 'kernels' construct}}

clang/test/SemaOpenACC/data-construct-finalize-ast.cpp

@@ -0,0 +1,60 @@
+// 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
+
+void Uses() {
+  // CHECK: FunctionDecl{{.*}}Uses
+  // CHECK-NEXT: CompoundStmt
+
+  int I;
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl
+
+#pragma acc exit data copyout(I) finalize
+  // CHECK-NEXT: OpenACCExitDataConstruct{{.*}}exit data
+  // CHECK-NEXT: copyout clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'I' 'int'
+  // CHECK-NEXT: finalize clause
+}
+
+template<typename T>
+void TemplUses() {
+  // CHECK: FunctionTemplateDecl{{.*}}TemplUses
+  // CHECK-NEXT: TemplateTypeParmDecl{{.*}}T
+  // CHECK-NEXT: FunctionDecl{{.*}}TemplUses
+  // CHECK-NEXT: CompoundStmt
+
+  T I;
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl
+
+#pragma acc exit data copyout(I) finalize
+  // CHECK-NEXT: OpenACCExitDataConstruct{{.*}}exit data
+  // CHECK-NEXT: copyout clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'I' 'T'
+  // CHECK-NEXT: finalize clause
+
+  // Instantiations
+  // CHECK-NEXT: FunctionDecl{{.*}} TemplUses 'void ()' implicit_instantiation
+  // CHECK-NEXT: TemplateArgument type 'int'
+  // CHECK-NEXT: BuiltinType{{.*}} 'int'
+  // CHECK-NEXT: CompoundStmt
+
+  // CHECK-NEXT: DeclStmt
+  // CHECK-NEXT: VarDecl
+
+  // CHECK-NEXT: OpenACCExitDataConstruct{{.*}}exit data
+  // CHECK-NEXT: copyout clause
+  // CHECK-NEXT: DeclRefExpr{{.*}}'I' 'int'
+  // CHECK-NEXT: finalize clause
+}
+void Inst() {
+  TemplUses<int>();
+}
+
+
+#endif // PCH_HELPER

clang/test/SemaOpenACC/data-construct-finalize-clause.c

@@ -0,0 +1,20 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+void Test() {
+  int I;
+
+  // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'data' directive}}
+#pragma acc data copyin(I) finalize
+  ;
+  // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'enter data' directive}}
+#pragma acc enter data copyin(I) finalize
+  ;
+
+  // finalize is valid only on exit data, otherwise has no other rules.
+#pragma acc exit data copyout(I) finalize
+  ;
+  // expected-warning@+2{{OpenACC clause 'use_device' not yet implemented}}
+  // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'host_data' directive}}
+#pragma acc host_data use_device(I) finalize
+  ;
+}

clang/test/SemaOpenACC/data-construct.cpp

@@ -84,7 +84,6 @@ void AtLeastOneOf() {
 #pragma acc exit data if(Var)
 #pragma acc exit data async
 #pragma acc exit data wait
-  // expected-warning@+1{{OpenACC clause 'finalize' not yet implemented}}
 #pragma acc exit data finalize
 #pragma acc exit data
 

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

@@ -37,7 +37,7 @@ void uses() {
   int *VarPtr;
 
   // 'auto' can combine with any other clause.
-  // expected-warning@+1{{OpenACC clause 'finalize' not yet implemented}}
+  // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'loop' directive}}
 #pragma acc loop auto finalize
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
@@ -171,7 +171,7 @@ void uses() {
 #pragma acc loop auto wait
   for(unsigned i = 0; i < 5; ++i);
 
-  // expected-warning@+1{{OpenACC clause 'finalize' not yet implemented}}
+  // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'loop' directive}}
 #pragma acc loop finalize auto
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
@@ -306,7 +306,7 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 
   // 'independent' can also be combined with any clauses
-  // expected-warning@+1{{OpenACC clause 'finalize' not yet implemented}}
+  // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'loop' directive}}
 #pragma acc loop independent finalize
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
@@ -440,7 +440,7 @@ void uses() {
 #pragma acc loop independent wait
   for(unsigned i = 0; i < 5; ++i);
 
-  // expected-warning@+1{{OpenACC clause 'finalize' not yet implemented}}
+  // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'loop' directive}}
 #pragma acc loop finalize independent
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
@@ -587,7 +587,7 @@ void uses() {
   // expected-note@+1{{previous clause is here}}
 #pragma acc loop seq vector
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'finalize' not yet implemented}}
+  // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'loop' directive}}
 #pragma acc loop seq finalize
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}
@@ -727,7 +727,7 @@ void uses() {
   // expected-note@+1{{previous clause is here}}
 #pragma acc loop vector seq
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'finalize' not yet implemented}}
+  // expected-error@+1{{OpenACC 'finalize' clause is not valid on 'loop' directive}}
 #pragma acc loop finalize seq
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'if_present' not yet implemented}}

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

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

clang/tools/libclang/CIndex.cpp

@@ -2921,6 +2921,8 @@ void OpenACCClauseEnqueue::VisitAutoClause(const OpenACCAutoClause &C) {}
 void OpenACCClauseEnqueue::VisitIndependentClause(
     const OpenACCIndependentClause &C) {}
 void OpenACCClauseEnqueue::VisitSeqClause(const OpenACCSeqClause &C) {}
+void OpenACCClauseEnqueue::VisitFinalizeClause(const OpenACCFinalizeClause &C) {
+}
 void OpenACCClauseEnqueue::VisitCollapseClause(const OpenACCCollapseClause &C) {
   Visitor.AddStmt(C.getLoopCount());
 }