[OpenACC] Implement firstprivate clause for compute constructs

This clause is pretty nearly copy/paste from private, except that it
doesn't support 'loop', and thus 'kernelsloop' for appertainment.

clang/include/clang/AST/OpenACCClause.h

@@ -294,6 +294,25 @@ class OpenACCPrivateClause final
          ArrayRef<Expr *> VarList, SourceLocation EndLoc);
 };
 
+class OpenACCFirstPrivateClause final
+    : public OpenACCClauseWithVarList,
+      public llvm::TrailingObjects<OpenACCFirstPrivateClause, Expr *> {
+
+  OpenACCFirstPrivateClause(SourceLocation BeginLoc, SourceLocation LParenLoc,
+                            ArrayRef<Expr *> VarList, SourceLocation EndLoc)
+      : OpenACCClauseWithVarList(OpenACCClauseKind::FirstPrivate, BeginLoc,
+                                 LParenLoc, EndLoc) {
+    std::uninitialized_copy(VarList.begin(), VarList.end(),
+                            getTrailingObjects<Expr *>());
+    setExprs(MutableArrayRef(getTrailingObjects<Expr *>(), VarList.size()));
+  }
+
+public:
+  static OpenACCFirstPrivateClause *
+  Create(const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc,
+         ArrayRef<Expr *> VarList, SourceLocation EndLoc);
+};
+
 template <class Impl> class OpenACCClauseVisitor {
   Impl &getDerived() { return static_cast<Impl &>(*this); }
 

clang/include/clang/Basic/OpenACCClauses.def

@@ -16,11 +16,12 @@
 // VISIT_CLAUSE(CLAUSE_NAME)
 
 VISIT_CLAUSE(Default)
+VISIT_CLAUSE(FirstPrivate)
 VISIT_CLAUSE(If)
-VISIT_CLAUSE(Self)
 VISIT_CLAUSE(NumGangs)
 VISIT_CLAUSE(NumWorkers)
 VISIT_CLAUSE(Private)
+VISIT_CLAUSE(Self)
 VISIT_CLAUSE(VectorLength)
 
 #undef VISIT_CLAUSE

clang/include/clang/Sema/SemaOpenACC.h

@@ -117,7 +117,8 @@ class SemaOpenACC : public SemaBase {
     }
 
     ArrayRef<Expr *> getVarList() {
-      assert(ClauseKind == OpenACCClauseKind::Private &&
+      assert((ClauseKind == OpenACCClauseKind::Private ||
+              ClauseKind == OpenACCClauseKind::FirstPrivate) &&
              "Parsed clause kind does not have a var-list");
       return std::get<VarListDetails>(Details).VarList;
     }
@@ -165,13 +166,15 @@ class SemaOpenACC : public SemaBase {
     }
 
     void setVarListDetails(ArrayRef<Expr *> VarList) {
-      assert(ClauseKind == OpenACCClauseKind::Private &&
+      assert((ClauseKind == OpenACCClauseKind::Private ||
+              ClauseKind == OpenACCClauseKind::FirstPrivate) &&
              "Parsed clause kind does not have a var-list");
       Details = VarListDetails{{VarList.begin(), VarList.end()}};
     }
 
     void setVarListDetails(llvm::SmallVector<Expr *> &&VarList) {
-      assert(ClauseKind == OpenACCClauseKind::Private &&
+      assert((ClauseKind == OpenACCClauseKind::Private ||
+              ClauseKind == OpenACCClauseKind::FirstPrivate) &&
              "Parsed clause kind does not have a var-list");
       Details = VarListDetails{std::move(VarList)};
     }

clang/lib/AST/OpenACCClause.cpp

@@ -144,6 +144,15 @@ OpenACCPrivateClause *OpenACCPrivateClause::Create(const ASTContext &C,
   return new (Mem) OpenACCPrivateClause(BeginLoc, LParenLoc, VarList, EndLoc);
 }
 
+OpenACCFirstPrivateClause *OpenACCFirstPrivateClause::Create(
+    const ASTContext &C, SourceLocation BeginLoc, SourceLocation LParenLoc,
+    ArrayRef<Expr *> VarList, SourceLocation EndLoc) {
+  void *Mem = C.Allocate(
+      OpenACCFirstPrivateClause::totalSizeToAlloc<Expr *>(VarList.size()));
+  return new (Mem)
+      OpenACCFirstPrivateClause(BeginLoc, LParenLoc, VarList, EndLoc);
+}
+
 //===----------------------------------------------------------------------===//
 //  OpenACC clauses printing methods
 //===----------------------------------------------------------------------===//
@@ -198,3 +207,11 @@ void OpenACCClausePrinter::VisitPrivateClause(const OpenACCPrivateClause &C) {
                         [&](const Expr *E) { printExpr(E); });
   OS << ")";
 }
+
+void OpenACCClausePrinter::VisitFirstPrivateClause(
+    const OpenACCFirstPrivateClause &C) {
+  OS << "firstprivate(";
+  llvm::interleaveComma(C.getVarList(), OS,
+                        [&](const Expr *E) { printExpr(E); });
+  OS << ")";
+}

clang/lib/AST/StmtProfile.cpp

@@ -2515,6 +2515,12 @@ void OpenACCClauseProfiler::VisitPrivateClause(
     Profiler.VisitStmt(E);
 }
 
+void OpenACCClauseProfiler::VisitFirstPrivateClause(
+    const OpenACCFirstPrivateClause &Clause) {
+  for (auto *E : Clause.getVarList())
+    Profiler.VisitStmt(E);
+}
+
 void OpenACCClauseProfiler::VisitVectorLengthClause(
     const OpenACCVectorLengthClause &Clause) {
   assert(Clause.hasIntExpr() &&

clang/lib/AST/TextNodeDumper.cpp

@@ -398,10 +398,11 @@ void TextNodeDumper::Visit(const OpenACCClause *C) {
       OS << '(' << cast<OpenACCDefaultClause>(C)->getDefaultClauseKind() << ')';
       break;
     case OpenACCClauseKind::If:
-    case OpenACCClauseKind::Self:
+    case OpenACCClauseKind::FirstPrivate:
     case OpenACCClauseKind::NumGangs:
     case OpenACCClauseKind::NumWorkers:
     case OpenACCClauseKind::Private:
+    case OpenACCClauseKind::Self:
     case OpenACCClauseKind::VectorLength:
       // The condition expression will be printed as a part of the 'children',
       // but print 'clause' here so it is clear what is happening from the dump.

clang/lib/Parse/ParseOpenACC.cpp

@@ -926,14 +926,14 @@ Parser::OpenACCClauseParseResult Parser::ParseOpenACCClauseParams(
     case OpenACCClauseKind::Device:
     case OpenACCClauseKind::DeviceResident:
     case OpenACCClauseKind::DevicePtr:
-    case OpenACCClauseKind::FirstPrivate:
     case OpenACCClauseKind::Host:
     case OpenACCClauseKind::Link:
     case OpenACCClauseKind::NoCreate:
     case OpenACCClauseKind::Present:
     case OpenACCClauseKind::UseDevice:
       ParseOpenACCVarList();
       break;
+    case OpenACCClauseKind::FirstPrivate:
     case OpenACCClauseKind::Private:
       ParsedClause.setVarListDetails(ParseOpenACCVarList());
       break;

clang/lib/Sema/SemaOpenACC.cpp

@@ -104,6 +104,16 @@ bool doesClauseApplyToDirective(OpenACCDirectiveKind DirectiveKind,
     default:
       return false;
     }
+  case OpenACCClauseKind::FirstPrivate:
+    switch (DirectiveKind) {
+    case OpenACCDirectiveKind::Parallel:
+    case OpenACCDirectiveKind::Serial:
+    case OpenACCDirectiveKind::ParallelLoop:
+    case OpenACCDirectiveKind::SerialLoop:
+      return true;
+    default:
+      return false;
+    }
   case OpenACCClauseKind::Private:
     switch (DirectiveKind) {
     case OpenACCDirectiveKind::Parallel:
@@ -331,6 +341,21 @@ SemaOpenACC::ActOnClause(ArrayRef<const OpenACCClause *> ExistingClauses,
         getASTContext(), Clause.getBeginLoc(), Clause.getLParenLoc(),
         Clause.getVarList(), Clause.getEndLoc());
   }
+  case OpenACCClauseKind::FirstPrivate: {
+    // 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;
+
+    // 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 OpenACCFirstPrivateClause::Create(
+        getASTContext(), Clause.getBeginLoc(), Clause.getLParenLoc(),
+        Clause.getVarList(), Clause.getEndLoc());
+  }
   default:
     break;
   }

clang/lib/Sema/TreeTransform.h

@@ -11112,6 +11112,23 @@ class OpenACCClauseTransform final
   SemaOpenACC::OpenACCParsedClause &ParsedClause;
   OpenACCClause *NewClause = nullptr;
 
+  llvm::SmallVector<Expr *> VisitVarList(ArrayRef<Expr *> VarList) {
+    llvm::SmallVector<Expr *> InstantiatedVarList;
+    for (Expr *CurVar : VarList) {
+      ExprResult Res = Self.TransformExpr(CurVar);
+
+      if (!Res.isUsable())
+        continue;
+
+      Res = Self.getSema().OpenACC().ActOnVar(Res.get());
+
+      if (Res.isUsable())
+        InstantiatedVarList.push_back(Res.get());
+    }
+
+    return InstantiatedVarList;
+  }
+
 public:
   OpenACCClauseTransform(TreeTransform<Derived> &Self,
                          ArrayRef<const OpenACCClause *> ExistingClauses,
@@ -11206,22 +11223,20 @@ void OpenACCClauseTransform<Derived>::VisitNumGangsClause(
 template <typename Derived>
 void OpenACCClauseTransform<Derived>::VisitPrivateClause(
     const OpenACCPrivateClause &C) {
-  llvm::SmallVector<Expr *> InstantiatedVarList;
-
-  for (Expr *CurVar : C.getVarList()) {
-    ExprResult Res = Self.TransformExpr(CurVar);
-
-    if (!Res.isUsable())
-      return;
+  ParsedClause.setVarListDetails(VisitVarList(C.getVarList()));
 
-    Res = Self.getSema().OpenACC().ActOnVar(Res.get());
+  NewClause = OpenACCPrivateClause::Create(
+      Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
+      ParsedClause.getLParenLoc(), ParsedClause.getVarList(),
+      ParsedClause.getEndLoc());
+}
 
-    if (Res.isUsable())
-      InstantiatedVarList.push_back(Res.get());
-  }
-  ParsedClause.setVarListDetails(std::move(InstantiatedVarList));
+template <typename Derived>
+void OpenACCClauseTransform<Derived>::VisitFirstPrivateClause(
+    const OpenACCFirstPrivateClause &C) {
+  ParsedClause.setVarListDetails(VisitVarList(C.getVarList()));
 
-  NewClause = OpenACCPrivateClause::Create(
+  NewClause = OpenACCFirstPrivateClause::Create(
       Self.getSema().getASTContext(), ParsedClause.getBeginLoc(),
       ParsedClause.getLParenLoc(), ParsedClause.getVarList(),
       ParsedClause.getEndLoc());

clang/lib/Serialization/ASTReader.cpp

@@ -11835,6 +11835,12 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
     return OpenACCPrivateClause::Create(getContext(), BeginLoc, LParenLoc,
                                         VarList, EndLoc);
   }
+  case OpenACCClauseKind::FirstPrivate: {
+    SourceLocation LParenLoc = readSourceLocation();
+    llvm::SmallVector<Expr *> VarList = readOpenACCVarList();
+    return OpenACCFirstPrivateClause::Create(getContext(), BeginLoc, LParenLoc,
+                                             VarList, EndLoc);
+  }
   case OpenACCClauseKind::Finalize:
   case OpenACCClauseKind::IfPresent:
   case OpenACCClauseKind::Seq:
@@ -11851,7 +11857,6 @@ OpenACCClause *ASTRecordReader::readOpenACCClause() {
   case OpenACCClauseKind::Device:
   case OpenACCClauseKind::DevicePtr:
   case OpenACCClauseKind::DeviceResident:
-  case OpenACCClauseKind::FirstPrivate:
   case OpenACCClauseKind::Host:
   case OpenACCClauseKind::Link:
   case OpenACCClauseKind::NoCreate:

clang/lib/Serialization/ASTWriter.cpp

@@ -7787,6 +7787,12 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
     writeOpenACCVarList(PC);
     return;
   }
+  case OpenACCClauseKind::FirstPrivate: {
+    const auto *FPC = cast<OpenACCFirstPrivateClause>(C);
+    writeSourceLocation(FPC->getLParenLoc());
+    writeOpenACCVarList(FPC);
+    return;
+  }
   case OpenACCClauseKind::Finalize:
   case OpenACCClauseKind::IfPresent:
   case OpenACCClauseKind::Seq:
@@ -7803,7 +7809,6 @@ void ASTRecordWriter::writeOpenACCClause(const OpenACCClause *C) {
   case OpenACCClauseKind::Device:
   case OpenACCClauseKind::DevicePtr:
   case OpenACCClauseKind::DeviceResident:
-  case OpenACCClauseKind::FirstPrivate:
   case OpenACCClauseKind::Host:
   case OpenACCClauseKind::Link:
   case OpenACCClauseKind::NoCreate:

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

@@ -38,5 +38,9 @@ void foo() {
 // CHECK: #pragma acc parallel private(i, array[1], array, array[1:2])
 #pragma acc parallel private(i, array[1], array, array[1:2])
   while(true);
+
+// CHECK: #pragma acc parallel firstprivate(i, array[1], array, array[1:2])
+#pragma acc parallel firstprivate(i, array[1], array, array[1:2])
+  while(true);
 }
 

clang/test/ParserOpenACC/parse-clauses.c

@@ -598,13 +598,11 @@ void VarListClauses() {
 #pragma acc serial private(s.array[s.value : 5], s.value), seq
   for(;;){}
 
-  // expected-error@+3{{expected ','}}
-  // expected-warning@+2{{OpenACC clause 'firstprivate' not yet implemented, clause ignored}}
+  // expected-error@+2{{expected ','}}
   // expected-warning@+1{{OpenACC clause 'seq' not yet implemented, clause ignored}}
 #pragma acc serial firstprivate(s.array[s.value] s.array[s.value :5] ), seq
   for(;;){}
 
-  // expected-warning@+2{{OpenACC clause 'firstprivate' not yet implemented, clause ignored}}
   // expected-warning@+1{{OpenACC clause 'seq' not yet implemented, clause ignored}}
 #pragma acc serial firstprivate(s.array[s.value : 5], s.value), seq
   for(;;){}

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

@@ -0,0 +1,55 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+typedef struct IsComplete {
+  struct S { int A; } CompositeMember;
+  int ScalarMember;
+  float ArrayMember[5];
+  void *PointerMember;
+} Complete;
+void uses(int IntParam, short *PointerParam, float ArrayParam[5], Complete CompositeParam) {
+  int LocalInt;
+  short *LocalPointer;
+  float LocalArray[5];
+  Complete LocalComposite;
+  // Check Appertainment:
+#pragma acc parallel firstprivate(LocalInt)
+  while(1);
+#pragma acc serial firstprivate(LocalInt)
+  while(1);
+  // expected-error@+1{{OpenACC 'firstprivate' clause is not valid on 'kernels' directive}}
+#pragma acc kernels firstprivate(LocalInt)
+  while(1);
+
+  // Valid cases:
+#pragma acc parallel firstprivate(LocalInt, LocalPointer, LocalArray)
+  while(1);
+#pragma acc parallel firstprivate(LocalArray[2:1])
+  while(1);
+
+#pragma acc parallel firstprivate(LocalComposite.ScalarMember, LocalComposite.ScalarMember)
+  while(1);
+
+  // expected-error@+1{{OpenACC variable is not a valid variable name, sub-array, array element, or composite variable member}}
+#pragma acc parallel firstprivate(1 + IntParam)
+  while(1);
+
+  // expected-error@+1{{OpenACC variable is not a valid variable name, sub-array, array element, or composite variable member}}
+#pragma acc parallel firstprivate(+IntParam)
+  while(1);
+
+  // expected-error@+1{{OpenACC sub-array length is unspecified and cannot be inferred because the subscripted value is not an array}}
+#pragma acc parallel firstprivate(PointerParam[2:])
+  while(1);
+
+  // expected-error@+1{{OpenACC sub-array specified range [2:5] would be out of the range of the subscripted array size of 5}}
+#pragma acc parallel firstprivate(ArrayParam[2:5])
+  while(1);
+
+  // expected-error@+2{{OpenACC sub-array specified range [2:5] would be out of the range of the subscripted array size of 5}}
+  // expected-error@+1{{OpenACC variable is not a valid variable name, sub-array, array element, or composite variable member}}
+#pragma acc parallel firstprivate((float*)ArrayParam[2:5])
+  while(1);
+  // expected-error@+1{{OpenACC variable is not a valid variable name, sub-array, array element, or composite variable member}}
+#pragma acc parallel firstprivate((float)ArrayParam[2])
+  while(1);
+}