[OpenACC] 'vector' clause implementation for combined constructs

Similar to 'worker', the 'vector' clause has some rules that needed to
be applied on its argument legality that for combined constructs need to
look at the current construct, not the 'effective' parent construct.
Additionally, it has some interaction with `vector_length` that needed
to be encoded as well.  This patch implements it.

Author: erichkeane

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -12762,9 +12762,9 @@ def err_acc_gang_dim_value
     : Error<"argument to 'gang' clause dimension must be %select{a constant "
             "expression|1, 2, or 3: evaluated to %1}0">;
 def err_acc_num_arg_conflict
-    : Error<"'num' argument to '%0' clause not allowed on a '%1' "
-            "construct%select{| associated with a '%3' construct}2 that has a "
-            "'%4' clause">;
+    : Error<"'%0' argument to '%1' clause not allowed on a '%2' "
+            "construct%select{| associated with a '%4' construct}3 that has a "
+            "'%5' clause">;
 def err_acc_num_arg_conflict_reverse
     : Error<"'%0' clause not allowed on a 'kernels loop' construct that "
             "has a '%1' clause with a%select{n| 'num'}2 argument">;

clang/lib/Sema/SemaOpenACC.cpp

@@ -804,6 +804,25 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorLengthClause(
   if (checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, Clause))
     return nullptr;
 
+  // OpenACC 3.3 Section 2.9.4:
+  // An argument is allowed only when the 'vector_length' does not appear on the
+  // 'kernels' construct.
+  if (Clause.getDirectiveKind() == OpenACCDirectiveKind::KernelsLoop) {
+    auto VectorClauses = llvm::make_filter_range(
+        ExistingClauses, llvm::IsaPred<OpenACCVectorClause>);
+
+    for (auto *VC : VectorClauses) {
+      if (cast<OpenACCVectorClause>(VC)->hasIntExpr()) {
+        SemaRef.Diag(Clause.getBeginLoc(),
+                     diag::err_acc_num_arg_conflict_reverse)
+            << OpenACCClauseKind::VectorLength << OpenACCClauseKind::Vector
+            << /*num argument*/ 0;
+        SemaRef.Diag(VC->getBeginLoc(), diag::note_acc_previous_clause_here);
+        return nullptr;
+      }
+    }
+  }
+
   assert(Clause.getIntExprs().size() == 1 &&
          "Invalid number of expressions for NumWorkers");
   return OpenACCVectorLengthClause::Create(
@@ -1097,6 +1116,14 @@ ExprResult DiagIntArgInvalid(SemaOpenACC &S, Expr *E, OpenACCGangKind GK,
       << HasAssocKind(DK, AssocKind) << AssocKind;
   return ExprError();
 }
+ExprResult DiagIntArgInvalid(SemaOpenACC &S, Expr *E, StringRef TagKind,
+                             OpenACCClauseKind CK, OpenACCDirectiveKind DK,
+                             OpenACCDirectiveKind AssocKind) {
+  S.Diag(E->getBeginLoc(), diag::err_acc_int_arg_invalid)
+      << TagKind << CK << IsOrphanLoop(DK, AssocKind) << DK
+      << HasAssocKind(DK, AssocKind) << AssocKind;
+  return ExprError();
+}
 
 ExprResult CheckGangParallelExpr(SemaOpenACC &S, OpenACCDirectiveKind DK,
                                  OpenACCDirectiveKind AssocKind,
@@ -1172,8 +1199,9 @@ ExprResult CheckGangKernelsExpr(SemaOpenACC &S,
 
     if (Itr != Collection.end()) {
       S.Diag(E->getBeginLoc(), diag::err_acc_num_arg_conflict)
-          << OpenACCClauseKind::Gang << DK << HasAssocKind(DK, AssocKind)
-          << AssocKind << OpenACCClauseKind::NumGangs;
+          << "num" << OpenACCClauseKind::Gang << DK
+          << HasAssocKind(DK, AssocKind) << AssocKind
+          << OpenACCClauseKind::NumGangs;
 
       S.Diag((*Itr)->getBeginLoc(), diag::note_acc_previous_clause_here);
       return ExprError();
@@ -1206,63 +1234,94 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitVectorClause(
     SemaOpenACC::OpenACCParsedClause &Clause) {
   if (DiagIfSeqClause(Clause))
     return nullptr;
-  // Restrictions only properly implemented on 'loop' constructs, and it is
-  // the only construct that can do anything with this, so skip/treat as
-  // unimplemented for the combined constructs.
-  if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop)
+
+  // Restrictions only properly implemented on 'loop'/'combined' constructs, and
+  // it is the only construct that can do anything with this, so skip/treat as
+  // unimplemented for the routine constructs.
+  if (Clause.getDirectiveKind() != OpenACCDirectiveKind::Loop &&
+      !isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind()))
     return isNotImplemented();
 
   Expr *IntExpr =
       Clause.getNumIntExprs() != 0 ? Clause.getIntExprs()[0] : nullptr;
   if (IntExpr) {
-    switch (SemaRef.getActiveComputeConstructInfo().Kind) {
-    case OpenACCDirectiveKind::Invalid:
-    case OpenACCDirectiveKind::Parallel:
-      // No restriction on when 'parallel' can contain an argument.
-      break;
-    case OpenACCDirectiveKind::Serial:
-      // GCC disallows this, and there is no real good reason for us to permit
-      // it, so disallow until we come up with a use case that makes sense.
-      DiagIntArgInvalid(SemaRef, IntExpr, OpenACCGangKind::Num,
-                        OpenACCClauseKind::Vector, Clause.getDirectiveKind(),
-                        SemaRef.getActiveComputeConstructInfo().Kind);
-      IntExpr = nullptr;
-      break;
-    case OpenACCDirectiveKind::Kernels: {
-      const auto *Itr =
-          llvm::find_if(SemaRef.getActiveComputeConstructInfo().Clauses,
-                        llvm::IsaPred<OpenACCVectorLengthClause>);
-      if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) {
-        SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict)
-            << OpenACCClauseKind::Vector << Clause.getDirectiveKind()
-            << HasAssocKind(Clause.getDirectiveKind(),
-                            SemaRef.getActiveComputeConstructInfo().Kind)
-            << SemaRef.getActiveComputeConstructInfo().Kind
-            << OpenACCClauseKind::VectorLength;
-        SemaRef.Diag((*Itr)->getBeginLoc(),
-                     diag::note_acc_previous_clause_here);
+    if (!isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) {
+      switch (SemaRef.getActiveComputeConstructInfo().Kind) {
+      case OpenACCDirectiveKind::Invalid:
+      case OpenACCDirectiveKind::Parallel:
+        // No restriction on when 'parallel' can contain an argument.
+        break;
+      case OpenACCDirectiveKind::Serial:
+        // GCC disallows this, and there is no real good reason for us to permit
+        // it, so disallow until we come up with a use case that makes sense.
+        DiagIntArgInvalid(SemaRef, IntExpr, "length", OpenACCClauseKind::Vector,
+                          Clause.getDirectiveKind(),
+                          SemaRef.getActiveComputeConstructInfo().Kind);
+        IntExpr = nullptr;
+        break;
+      case OpenACCDirectiveKind::Kernels: {
+        const auto *Itr =
+            llvm::find_if(SemaRef.getActiveComputeConstructInfo().Clauses,
+                          llvm::IsaPred<OpenACCVectorLengthClause>);
+        if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) {
+          SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict)
+              << "length" << OpenACCClauseKind::Vector
+              << Clause.getDirectiveKind()
+              << HasAssocKind(Clause.getDirectiveKind(),
+                              SemaRef.getActiveComputeConstructInfo().Kind)
+              << SemaRef.getActiveComputeConstructInfo().Kind
+              << OpenACCClauseKind::VectorLength;
+          SemaRef.Diag((*Itr)->getBeginLoc(),
+                       diag::note_acc_previous_clause_here);
 
+          IntExpr = nullptr;
+        }
+        break;
+      }
+      default:
+        llvm_unreachable("Non compute construct in active compute construct");
+      }
+    } else {
+      if (Clause.getDirectiveKind() == OpenACCDirectiveKind::SerialLoop) {
+        DiagIntArgInvalid(SemaRef, IntExpr, "length", OpenACCClauseKind::Vector,
+                          Clause.getDirectiveKind(),
+                          SemaRef.getActiveComputeConstructInfo().Kind);
         IntExpr = nullptr;
+      } else if (Clause.getDirectiveKind() ==
+                 OpenACCDirectiveKind::KernelsLoop) {
+        const auto *Itr = llvm::find_if(
+            ExistingClauses, llvm::IsaPred<OpenACCVectorLengthClause>);
+        if (Itr != ExistingClauses.end()) {
+          SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict)
+              << "length" << OpenACCClauseKind::Vector
+              << Clause.getDirectiveKind()
+              << HasAssocKind(Clause.getDirectiveKind(),
+                              SemaRef.getActiveComputeConstructInfo().Kind)
+              << SemaRef.getActiveComputeConstructInfo().Kind
+              << OpenACCClauseKind::VectorLength;
+          SemaRef.Diag((*Itr)->getBeginLoc(),
+                       diag::note_acc_previous_clause_here);
+
+          IntExpr = nullptr;
+        }
       }
-      break;
-    }
-    default:
-      llvm_unreachable("Non compute construct in active compute construct");
     }
   }
 
-  // OpenACC 3.3 2.9.4: The region of a loop with a 'vector' clause may not
-  // contain a loop with a gang, worker, or vector clause unless within a nested
-  // compute region.
-  if (SemaRef.LoopVectorClauseLoc.isValid()) {
-    // This handles the 'inner loop' diagnostic, but we cannot set that we're on
-    // one of these until we get to the end of the construct.
-    SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region)
-        << OpenACCClauseKind::Vector << OpenACCClauseKind::Vector
-        << /*skip kernels construct info*/ 0;
-    SemaRef.Diag(SemaRef.LoopVectorClauseLoc,
-                 diag::note_acc_previous_clause_here);
-    return nullptr;
+  if (!isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) {
+    // OpenACC 3.3 2.9.4: The region of a loop with a 'vector' clause may not
+    // contain a loop with a gang, worker, or vector clause unless within a
+    // nested compute region.
+    if (SemaRef.LoopVectorClauseLoc.isValid()) {
+      // This handles the 'inner loop' diagnostic, but we cannot set that we're
+      // on one of these until we get to the end of the construct.
+      SemaRef.Diag(Clause.getBeginLoc(), diag::err_acc_clause_in_clause_region)
+          << OpenACCClauseKind::Vector << OpenACCClauseKind::Vector
+          << /*skip kernels construct info*/ 0;
+      SemaRef.Diag(SemaRef.LoopVectorClauseLoc,
+                   diag::note_acc_previous_clause_here);
+      return nullptr;
+    }
   }
 
   return OpenACCVectorClause::Create(Ctx, Clause.getBeginLoc(),
@@ -1305,7 +1364,7 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitWorkerClause(
                           llvm::IsaPred<OpenACCNumWorkersClause>);
         if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) {
           SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict)
-              << OpenACCClauseKind::Worker << Clause.getDirectiveKind()
+              << "num" << OpenACCClauseKind::Worker << Clause.getDirectiveKind()
               << HasAssocKind(Clause.getDirectiveKind(),
                               SemaRef.getActiveComputeConstructInfo().Kind)
               << SemaRef.getActiveComputeConstructInfo().Kind
@@ -1334,7 +1393,7 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitWorkerClause(
                                         llvm::IsaPred<OpenACCNumWorkersClause>);
         if (Itr != ExistingClauses.end()) {
           SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict)
-              << OpenACCClauseKind::Worker << Clause.getDirectiveKind()
+              << "num" << OpenACCClauseKind::Worker << Clause.getDirectiveKind()
               << HasAssocKind(Clause.getDirectiveKind(),
                               SemaRef.getActiveComputeConstructInfo().Kind)
               << SemaRef.getActiveComputeConstructInfo().Kind

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

@@ -344,4 +344,46 @@ void foo() {
 #pragma acc kernels loop worker(num:5)
   for(int i = 0;i<5;++i);
 
+  // CHECK: #pragma acc parallel loop vector
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc parallel loop vector
+  for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc parallel loop vector(length: 5)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc parallel loop vector(5)
+  for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc parallel loop vector(length: 5)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc parallel loop vector(length:5)
+  for(int i = 0;i<5;++i);
+
+// CHECK-NEXT: #pragma acc kernels loop vector
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc kernels loop vector
+  for(int i = 0;i<5;++i);
+
+// CHECK-NEXT: #pragma acc kernels loop vector(length: 5)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc kernels loop vector(5)
+  for(int i = 0;i<5;++i);
+
+// CHECK-NEXT: #pragma acc kernels loop vector(length: 5)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc kernels loop vector(length:5)
+  for(int i = 0;i<5;++i);
+
+// CHECK-NEXT: #pragma acc serial loop vector
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc serial loop vector
+  for(int i = 0;i<5;++i);
+
 }

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

@@ -48,7 +48,6 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop auto worker
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'vector' not yet implemented}}
 #pragma acc parallel loop auto vector
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'nohost' not yet implemented}}
@@ -167,7 +166,6 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop worker auto
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'vector' not yet implemented}}
 #pragma acc parallel loop vector auto
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'nohost' not yet implemented}}
@@ -287,7 +285,6 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop independent worker
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'vector' not yet implemented}}
 #pragma acc parallel loop independent vector
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'nohost' not yet implemented}}
@@ -406,7 +403,6 @@ void uses() {
   for(unsigned i = 0; i < 5; ++i);
 #pragma acc parallel loop worker independent
   for(unsigned i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'vector' not yet implemented}}
 #pragma acc parallel loop vector independent
   for(unsigned i = 0; i < 5; ++i);
   // expected-warning@+1{{OpenACC clause 'nohost' not yet implemented}}
@@ -650,9 +646,8 @@ void uses() {
   // expected-note@+1{{previous clause is here}}
 #pragma acc parallel loop worker seq
   for(unsigned i = 0; i < 5; ++i);
-  // TODOexpected-error@+2{{OpenACC clause 'seq' may not appear on the same construct as a 'vector' clause on a 'parallel loop' construct}}
-  // TODOexpected-note@+1{{previous clause is here}}
-  // expected-warning@+1{{OpenACC clause 'vector' not yet implemented}}
+  // expected-error@+2{{OpenACC clause 'seq' may not appear on the same construct as a 'vector' clause on a 'parallel loop' construct}}
+  // 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}}

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

@@ -39,7 +39,6 @@ void uses() {
   // 'worker', 'vector', 'seq', 'independent', 'auto', and 'tile'  after
   // 'device_type'.
 
-  //expected-warning@+1{{OpenACC clause 'vector' not yet implemented, clause ignored}}
 #pragma acc parallel loop device_type(*) vector
   for(int i = 0; i < 5; ++i);