[OpenACC] Implement private/firstprivate for combined constructs

This is another pair of clauses where the work is already done from
previous constructs, so this just has to allow them and include tests
for them. This patch adds testing, does a few little cleanup bits on the
clause checking, and enables these.

Author: erichkeane

clang/lib/Sema/SemaOpenACC.cpp

@@ -598,6 +598,7 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitDefaultClause(
 
 OpenACCClause *SemaOpenACCClauseVisitor::VisitTileClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
+  // TODO OpenACC: Remove this when we get combined construct impl for this.
   if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop)
     return isNotImplemented();
 
@@ -699,6 +700,8 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNumGangsClause(
   // 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.
+  // TODO OpenACC:  Remove this check when we have combined constructs for this
+  // clause.
   if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
     return isNotImplemented();
 
@@ -754,6 +757,7 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNumWorkersClause(
   // 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.
+  // TODO: OpenACC: Remove when we get combined constructs.
   if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
     return isNotImplemented();
 
@@ -775,6 +779,7 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorLengthClause(
   // 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.
+  // TODO: OpenACC: Remove when we get combined constructs.
   if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()))
     return isNotImplemented();
 
@@ -815,14 +820,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitAsyncClause(
 
 OpenACCClause *SemaOpenACCClauseVisitor::VisitPrivateClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
-  // Restrictions only properly implemented on 'compute' and 'loop'
-  // constructs, and 'compute'/'loop' constructs are the only construct that
-  // can do anything with this yet, so skip/treat as unimplemented in this
-  // case.
-  if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()) &&
-      Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop)
-    return isNotImplemented();
-
   // 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.
@@ -834,12 +831,6 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitPrivateClause(
 
 OpenACCClause *SemaOpenACCClauseVisitor::VisitFirstPrivateClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
-  // 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()))
-    return isNotImplemented();
-
   // 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.
@@ -1412,6 +1403,7 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitReductionClause(
   // 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.
+  // TODO: OpenACC: Remove check once we get combined constructs for this clause.
   if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()) &&
       Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop)
     return isNotImplemented();
@@ -1502,6 +1494,7 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitReductionClause(
 
 OpenACCClause *SemaOpenACCClauseVisitor::VisitCollapseClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
+  // TODO: Remove this check once we implement this for combined constructs.
   if (!isOpenACCComputeDirectiveKind(Clause.getDirectiveKind()) &&
       Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop)
     return isNotImplemented();

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

@@ -96,4 +96,13 @@ void foo() {
 // CHECK-NEXT: ;
 #pragma acc serial loop default(present)
   for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc parallel loop private(i, array[1], array, array[1:2])
+#pragma acc parallel loop private(i, array[1], array, array[1:2])
+  for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc serial loop firstprivate(i, array[1], array, array[1:2])
+#pragma acc serial loop firstprivate(i, array[1], array, array[1:2])
+  for(int i = 0;i<5;++i);
+
 }

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

@@ -99,8 +99,6 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'device_resident' not yet implemented}}
 #pragma acc parallel loop auto device_resident(VarPtr)
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'firstprivate' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'firstprivate' not yet implemented}}
 #pragma acc parallel loop auto firstprivate(Var)
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'host' not yet implemented}}
@@ -117,7 +115,6 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'present' not yet implemented}}
 #pragma acc parallel loop auto present(Var)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'private' not yet implemented}}
 #pragma acc parallel loop auto private(Var)
   for(unsigned i = 0; i < 5; ++i);
   // TODOexpected-error@+1{{OpenACC 'copyout' clause is not valid on 'parallel loop' directive}}
@@ -267,8 +264,6 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'device_resident' not yet implemented}}
 #pragma acc parallel loop device_resident(VarPtr) auto
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'firstprivate' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'firstprivate' not yet implemented}}
 #pragma acc parallel loop firstprivate(Var) auto
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'host' not yet implemented}}
@@ -285,7 +280,6 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'present' not yet implemented}}
 #pragma acc parallel loop present(Var) auto
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'private' not yet implemented}}
 #pragma acc parallel loop private(Var) auto
   for(unsigned i = 0; i < 5; ++i);
   // TODOexpected-error@+1{{OpenACC 'copyout' clause is not valid on 'parallel loop' directive}}
@@ -436,8 +430,6 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'device_resident' not yet implemented}}
 #pragma acc parallel loop independent device_resident(VarPtr)
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'firstprivate' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'firstprivate' not yet implemented}}
 #pragma acc parallel loop independent firstprivate(Var)
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'host' not yet implemented}}
@@ -454,7 +446,6 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'present' not yet implemented}}
 #pragma acc parallel loop independent present(Var)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'private' not yet implemented}}
 #pragma acc parallel loop independent private(Var)
   for(unsigned i = 0; i < 5; ++i);
   // TODOexpected-error@+1{{OpenACC 'copyout' clause is not valid on 'parallel loop' directive}}
@@ -604,8 +595,6 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'device_resident' not yet implemented}}
 #pragma acc parallel loop device_resident(VarPtr) independent
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'firstprivate' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'firstprivate' not yet implemented}}
 #pragma acc parallel loop firstprivate(Var) independent
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'host' not yet implemented}}
@@ -622,7 +611,6 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'present' not yet implemented}}
 #pragma acc parallel loop present(Var) independent
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'private' not yet implemented}}
 #pragma acc parallel loop private(Var) independent
   for(unsigned i = 0; i < 5; ++i);
   // TODOexpected-error@+1{{OpenACC 'copyout' clause is not valid on 'parallel loop' directive}}
@@ -779,8 +767,6 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'device_resident' not yet implemented}}
 #pragma acc parallel loop seq device_resident(VarPtr)
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'firstprivate' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'firstprivate' not yet implemented}}
 #pragma acc parallel loop seq firstprivate(Var)
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'host' not yet implemented}}
@@ -797,7 +783,6 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'present' not yet implemented}}
 #pragma acc parallel loop seq present(Var)
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'private' not yet implemented}}
 #pragma acc parallel loop seq private(Var)
   for(unsigned i = 0; i < 5; ++i);
   // TODOexpected-error@+1{{OpenACC 'copyout' clause is not valid on 'parallel loop' directive}}
@@ -953,8 +938,6 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'device_resident' not yet implemented}}
 #pragma acc parallel loop device_resident(VarPtr) seq
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+1{{OpenACC 'firstprivate' clause is not valid on 'parallel loop' directive}}
-  // expected-warning@+1{{OpenACC clause 'firstprivate' not yet implemented}}
 #pragma acc parallel loop firstprivate(Var) seq
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'host' not yet implemented}}
@@ -971,7 +954,6 @@ void uses() {
   // expected-warning@+1{{OpenACC clause 'present' not yet implemented}}
 #pragma acc parallel loop present(Var) seq
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'private' not yet implemented}}
 #pragma acc parallel loop private(Var) seq
   for(unsigned i = 0; i < 5; ++i);
   // TODOexpected-error@+1{{OpenACC 'copyout' clause is not valid on 'parallel loop' directive}}

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

@@ -18,7 +18,6 @@ void SingleOnly() {
   #pragma acc kernels loop self default(present) present(i) default(none) copy(i)
   for(int i = 5; i < 10;++i);
 
-  // expected-warning@+4{{OpenACC clause 'private' not yet implemented}}
   // expected-warning@+3{{OpenACC clause 'copy' not yet implemented}}
   // expected-error@+2{{OpenACC 'default' clause cannot appear more than once on a 'parallel loop' directive}}
   // expected-note@+1{{previous clause is here}}

clang/test/SemaOpenACC/combined-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 loop firstprivate(LocalInt)
+  for (int i = 5; i < 10; ++i);
+#pragma acc serial loop firstprivate(LocalInt)
+  for (int i = 5; i < 10; ++i);
+  // expected-error@+1{{OpenACC 'firstprivate' clause is not valid on 'kernels loop' directive}}
+#pragma acc kernels loop firstprivate(LocalInt)
+  for (int i = 5; i < 10; ++i);
+
+  // Valid cases:
+#pragma acc parallel loop firstprivate(LocalInt, LocalPointer, LocalArray)
+  for (int i = 5; i < 10; ++i);
+#pragma acc serial loop firstprivate(LocalArray[2:1])
+  for (int i = 5; i < 10; ++i);
+
+#pragma acc parallel loop firstprivate(LocalComposite.ScalarMember, LocalComposite.ScalarMember)
+  for (int i = 5; i < 10; ++i);
+
+  // expected-error@+1{{OpenACC variable is not a valid variable name, sub-array, array element, member of a composite variable, or composite variable member}}
+#pragma acc serial loop firstprivate(1 + IntParam)
+  for (int i = 5; i < 10; ++i);
+
+  // expected-error@+1{{OpenACC variable is not a valid variable name, sub-array, array element, member of a composite variable, or composite variable member}}
+#pragma acc serial loop firstprivate(+IntParam)
+  for (int i = 5; i < 10; ++i);
+
+  // expected-error@+1{{OpenACC sub-array length is unspecified and cannot be inferred because the subscripted value is not an array}}
+#pragma acc parallel loop firstprivate(PointerParam[2:])
+  for (int i = 5; i < 10; ++i);
+
+  // 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 serial loop firstprivate(ArrayParam[2:5])
+  for (int i = 5; i < 10; ++i);
+
+  // 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, member of a composite variable, or composite variable member}}
+#pragma acc parallel loop firstprivate((float*)ArrayParam[2:5])
+  for (int i = 5; i < 10; ++i);
+  // expected-error@+1{{OpenACC variable is not a valid variable name, sub-array, array element, member of a composite variable, or composite variable member}}
+#pragma acc serial loop firstprivate((float)ArrayParam[2])
+  for (int i = 5; i < 10; ++i);
+}

clang/test/SemaOpenACC/combined-construct-firstprivate-clause.cpp

@@ -0,0 +1,113 @@
+// RUN: %clang_cc1 %s -fopenacc -verify
+
+enum SomeE{};
+typedef struct IsComplete {
+  struct S { int A; } CompositeMember;
+  int ScalarMember;
+  float ArrayMember[5];
+  SomeE EnumMember;
+  char *PointerMember;
+} Complete;
+
+void uses(int IntParam, char *PointerParam, float ArrayParam[5], Complete CompositeParam, int &IntParamRef) {
+  int LocalInt;
+  char *LocalPointer;
+  float LocalArray[5];
+  // Check Appertainment:
+#pragma acc parallel loop firstprivate(LocalInt)
+  for (int i = 5; i < 10; ++i);
+#pragma acc serial loop firstprivate(LocalInt)
+  for (int i = 5; i < 10; ++i);
+  // expected-error@+1{{OpenACC 'firstprivate' clause is not valid on 'kernels loop' directive}}
+#pragma acc kernels loop firstprivate(LocalInt)
+  for (int i = 5; i < 10; ++i);
+
+  // Valid cases:
+#pragma acc parallel loop firstprivate(LocalInt, LocalPointer, LocalArray)
+  for (int i = 5; i < 10; ++i);
+#pragma acc serial loop firstprivate(LocalArray[2:1])
+  for (int i = 5; i < 10; ++i);
+
+  Complete LocalComposite2;
+#pragma acc parallel loop firstprivate(LocalComposite2.ScalarMember, LocalComposite2.ScalarMember)
+  for (int i = 5; i < 10; ++i);
+
+  // expected-error@+1{{OpenACC variable is not a valid variable name, sub-array, array element, member of a composite variable, or composite variable member}}
+#pragma acc parallel loop firstprivate(1 + IntParam)
+  for (int i = 5; i < 10; ++i);
+
+  // expected-error@+1{{OpenACC variable is not a valid variable name, sub-array, array element, member of a composite variable, or composite variable member}}
+#pragma acc serial loop firstprivate(+IntParam)
+  for (int i = 5; i < 10; ++i);
+
+  // expected-error@+1{{OpenACC sub-array length is unspecified and cannot be inferred because the subscripted value is not an array}}
+#pragma acc parallel loop firstprivate(PointerParam[2:])
+  for (int i = 5; i < 10; ++i);
+
+  // 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 loop firstprivate(ArrayParam[2:5])
+  for (int i = 5; i < 10; ++i);
+
+  // 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, member of a composite variable, or composite variable member}}
+#pragma acc serial loop firstprivate((float*)ArrayParam[2:5])
+  for (int i = 5; i < 10; ++i);
+  // expected-error@+1{{OpenACC variable is not a valid variable name, sub-array, array element, member of a composite variable, or composite variable member}}
+#pragma acc parallel loop firstprivate((float)ArrayParam[2])
+  for (int i = 5; i < 10; ++i);
+}
+
+template<typename T, unsigned I, typename V>
+void TemplUses(T t, T (&arrayT)[I], V TemplComp) {
+  // expected-error@+1{{OpenACC variable is not a valid variable name, sub-array, array element, member of a composite variable, or composite variable member}}
+#pragma acc parallel loop firstprivate(+t)
+  for (int i = 5; i < 10; ++i);
+
+  // NTTP's are only valid if it is a reference to something.
+  // expected-error@+2{{OpenACC variable is not a valid variable name, sub-array, array element, member of a composite variable, or composite variable member}}
+  // expected-note@#TEMPL_USES_INST{{in instantiation of}}
+#pragma acc parallel loop firstprivate(I)
+  for (int i = 5; i < 10; ++i);
+
+  // expected-error@+1{{OpenACC variable is not a valid variable name, sub-array, array element, member of a composite variable, or composite variable member}}
+#pragma acc serial loop firstprivate(t, I)
+  for (int i = 5; i < 10; ++i);
+
+#pragma acc parallel loop firstprivate(arrayT)
+  for (int i = 5; i < 10; ++i);
+
+#pragma acc parallel loop firstprivate(TemplComp)
+  for (int i = 5; i < 10; ++i);
+
+#pragma acc parallel loop firstprivate(TemplComp.PointerMember[5])
+  for (int i = 5; i < 10; ++i);
+ int *Pointer;
+#pragma acc serial loop firstprivate(Pointer[:I])
+  for (int i = 5; i < 10; ++i);
+#pragma acc parallel loop firstprivate(Pointer[:t])
+  for (int i = 5; i < 10; ++i);
+  // expected-error@+1{{OpenACC sub-array length is unspecified and cannot be inferred because the subscripted value is not an array}}
+#pragma acc parallel loop firstprivate(Pointer[1:])
+  for (int i = 5; i < 10; ++i);
+}
+
+template<unsigned I, auto &NTTP_REF>
+void NTTP() {
+  // NTTP's are only valid if it is a reference to something.
+  // expected-error@+2{{OpenACC variable is not a valid variable name, sub-array, array element, member of a composite variable, or composite variable member}}
+  // expected-note@#NTTP_INST{{in instantiation of}}
+#pragma acc parallel loop firstprivate(I)
+  for (int i = 5; i < 10; ++i);
+
+#pragma acc serial loop firstprivate(NTTP_REF)
+  for (int i = 5; i < 10; ++i);
+}
+
+void Inst() {
+  static constexpr int NTTP_REFed = 1;
+  int i;
+  int Arr[5];
+  Complete C;
+  TemplUses(i, Arr, C); // #TEMPL_USES_INST
+  NTTP<5, NTTP_REFed>(); // #NTTP_INST
+}