[OpenACC] Implement 'worker' clause for combined constructs

This is very similar to 'gang', except with fewer restrictions, and only an
interaction with 'num_workers', plus disallowing 'gang' and 'worker' in
its associated statement.  This patch implements this, the same as how
'gang' implemented it.

Author: erichkeane

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -12766,8 +12766,8 @@ def err_acc_num_arg_conflict
             "construct%select{| associated with a '%3' construct}2 that has a "
             "'%4' clause">;
 def err_acc_num_arg_conflict_reverse
-    : Error<"'num_gangs' clause not allowed on a 'kernels loop' construct that "
-            "has a 'gang' clause with a 'num' argument">;
+    : Error<"'%0' clause not allowed on a 'kernels loop' construct that "
+            "has a '%1' clause with a%select{n| 'num'}2 argument">;
 def err_acc_clause_in_clause_region
     : Error<"loop with a '%0' clause may not exist in the region of a '%1' "
             "clause%select{| on a '%3' construct}2">;

clang/lib/Sema/SemaOpenACC.cpp

@@ -748,7 +748,9 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNumGangsClause(
     for (auto *GC : GangClauses) {
       if (cast<OpenACCGangClause>(GC)->hasExprOfKind(OpenACCGangKind::Num)) {
         SemaRef.Diag(Clause.getBeginLoc(),
-                     diag::err_acc_num_arg_conflict_reverse);
+                     diag::err_acc_num_arg_conflict_reverse)
+            << OpenACCClauseKind::NumGangs << OpenACCClauseKind::Gang
+            << /*Num argument*/ 1;
         SemaRef.Diag(GC->getBeginLoc(), diag::note_acc_previous_clause_here);
         return nullptr;
       }
@@ -768,6 +770,25 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitNumWorkersClause(
   if (checkAlreadyHasClauseOfKind(SemaRef, ExistingClauses, Clause))
     return nullptr;
 
+  // OpenACC 3.3 Section 2.9.2:
+  // An argument is allowed only when the 'num_workers' does not appear on the
+  // kernels construct.
+  if (Clause.getDirectiveKind() == OpenACCDirectiveKind::KernelsLoop) {
+    auto WorkerClauses = llvm::make_filter_range(
+        ExistingClauses, llvm::IsaPred<OpenACCWorkerClause>);
+
+    for (auto *WC : WorkerClauses) {
+      if (cast<OpenACCWorkerClause>(WC)->hasIntExpr()) {
+        SemaRef.Diag(Clause.getBeginLoc(),
+                     diag::err_acc_num_arg_conflict_reverse)
+            << OpenACCClauseKind::NumWorkers << OpenACCClauseKind::Worker
+            << /*num argument*/ 0;
+        SemaRef.Diag(WC->getBeginLoc(), diag::note_acc_previous_clause_here);
+        return nullptr;
+      }
+    }
+  }
+
   assert(Clause.getIntExprs().size() == 1 &&
          "Invalid number of expressions for NumWorkers");
   return OpenACCNumWorkersClause::Create(
@@ -1254,74 +1275,107 @@ OpenACCClause *SemaOpenACCClauseVisitor::VisitWorkerClause(
   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:
-    case OpenACCDirectiveKind::Serial:
-      DiagIntArgInvalid(SemaRef, IntExpr, OpenACCGangKind::Num,
-                        OpenACCClauseKind::Worker, Clause.getDirectiveKind(),
-                        SemaRef.getActiveComputeConstructInfo().Kind);
-      IntExpr = nullptr;
-      break;
-    case OpenACCDirectiveKind::Kernels: {
-      const auto *Itr =
-          llvm::find_if(SemaRef.getActiveComputeConstructInfo().Clauses,
-                        llvm::IsaPred<OpenACCNumWorkersClause>);
-      if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) {
-        SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict)
-            << OpenACCClauseKind::Worker << Clause.getDirectiveKind()
-            << HasAssocKind(Clause.getDirectiveKind(),
-                            SemaRef.getActiveComputeConstructInfo().Kind)
-            << SemaRef.getActiveComputeConstructInfo().Kind
-            << OpenACCClauseKind::NumWorkers;
-        SemaRef.Diag((*Itr)->getBeginLoc(),
-                     diag::note_acc_previous_clause_here);
-
+    if (!isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) {
+      switch (SemaRef.getActiveComputeConstructInfo().Kind) {
+      case OpenACCDirectiveKind::Invalid:
+      case OpenACCDirectiveKind::ParallelLoop:
+      case OpenACCDirectiveKind::SerialLoop:
+      case OpenACCDirectiveKind::Parallel:
+      case OpenACCDirectiveKind::Serial:
+        DiagIntArgInvalid(SemaRef, IntExpr, OpenACCGangKind::Num,
+                          OpenACCClauseKind::Worker, Clause.getDirectiveKind(),
+                          SemaRef.getActiveComputeConstructInfo().Kind);
         IntExpr = nullptr;
+        break;
+      case OpenACCDirectiveKind::KernelsLoop:
+      case OpenACCDirectiveKind::Kernels: {
+        const auto *Itr =
+            llvm::find_if(SemaRef.getActiveComputeConstructInfo().Clauses,
+                          llvm::IsaPred<OpenACCNumWorkersClause>);
+        if (Itr != SemaRef.getActiveComputeConstructInfo().Clauses.end()) {
+          SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict)
+              << OpenACCClauseKind::Worker << Clause.getDirectiveKind()
+              << HasAssocKind(Clause.getDirectiveKind(),
+                              SemaRef.getActiveComputeConstructInfo().Kind)
+              << SemaRef.getActiveComputeConstructInfo().Kind
+              << OpenACCClauseKind::NumWorkers;
+          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::ParallelLoop ||
+          Clause.getDirectiveKind() == OpenACCDirectiveKind::SerialLoop) {
+        DiagIntArgInvalid(SemaRef, IntExpr, OpenACCGangKind::Num,
+                          OpenACCClauseKind::Worker, Clause.getDirectiveKind(),
+                          SemaRef.getActiveComputeConstructInfo().Kind);
+        IntExpr = nullptr;
+      } else {
+        assert(Clause.getDirectiveKind() == OpenACCDirectiveKind::KernelsLoop &&
+               "Unknown combined directive kind?");
+        const auto *Itr = llvm::find_if(ExistingClauses,
+                                        llvm::IsaPred<OpenACCNumWorkersClause>);
+        if (Itr != ExistingClauses.end()) {
+          SemaRef.Diag(IntExpr->getBeginLoc(), diag::err_acc_num_arg_conflict)
+              << OpenACCClauseKind::Worker << Clause.getDirectiveKind()
+              << HasAssocKind(Clause.getDirectiveKind(),
+                              SemaRef.getActiveComputeConstructInfo().Kind)
+              << SemaRef.getActiveComputeConstructInfo().Kind
+              << OpenACCClauseKind::NumWorkers;
+          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.3: The region of a loop with a 'worker' clause may not
-  // contain a loop with a gang or worker clause unless within a nested compute
-  // region.
-  if (SemaRef.LoopWorkerClauseLoc.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::Worker << OpenACCClauseKind::Worker
-        << /*skip kernels construct info*/ 0;
-    SemaRef.Diag(SemaRef.LoopWorkerClauseLoc,
-                 diag::note_acc_previous_clause_here);
-    return nullptr;
-  }
+  if (!isOpenACCCombinedDirectiveKind(Clause.getDirectiveKind())) {
+    // OpenACC 3.3 2.9.3: The region of a loop with a 'worker' clause may not
+    // contain a loop with a gang or worker clause unless within a nested
+    // compute region.
+    if (SemaRef.LoopWorkerClauseLoc.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::Worker << OpenACCClauseKind::Worker
+          << /*skip kernels construct info*/ 0;
+      SemaRef.Diag(SemaRef.LoopWorkerClauseLoc,
+                   diag::note_acc_previous_clause_here);
+      return nullptr;
+    }
 
-  // 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::Worker << OpenACCClauseKind::Vector
-        << /*skip kernels construct info*/ 0;
-    SemaRef.Diag(SemaRef.LoopVectorClauseLoc,
-                 diag::note_acc_previous_clause_here);
-    return nullptr;
+    // 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::Worker << OpenACCClauseKind::Vector
+          << /*skip kernels construct info*/ 0;
+      SemaRef.Diag(SemaRef.LoopVectorClauseLoc,
+                   diag::note_acc_previous_clause_here);
+      return nullptr;
+    }
   }
 
   return OpenACCWorkerClause::Create(Ctx, Clause.getBeginLoc(),

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

@@ -313,4 +313,35 @@ void foo() {
 // CHECK-NEXT: ;
 #pragma acc serial loop gang(static:*)
   for(int i = 0;i<5;++i);
+
+// CHECK: #pragma acc parallel loop worker
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc parallel loop worker
+  for(int i = 0;i<5;++i);
+
+// CHECK-NEXT: #pragma acc parallel loop worker
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc parallel loop worker
+  for(int i = 0;i<5;++i);
+
+// CHECK-NEXT: #pragma acc serial loop worker
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc serial loop worker
+  for(int i = 0;i<5;++i);
+
+// CHECK-NEXT: #pragma acc kernels loop worker(num: 5)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc kernels loop worker(5)
+  for(int i = 0;i<5;++i);
+
+// CHECK-NEXT: #pragma acc kernels loop worker(num: 5)
+// CHECK-NEXT: for (int i = 0; i < 5; ++i)
+// CHECK-NEXT: ;
+#pragma acc kernels loop worker(num:5)
+  for(int i = 0;i<5;++i);
+
 }

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

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

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

@@ -57,7 +57,6 @@ void uses() {
   for(int i = 0; i < 5; ++i);
 #pragma acc kernels loop device_type(*) auto
   for(int i = 0; i < 5; ++i);
-  // expected-warning@+1{{OpenACC clause 'worker' not yet implemented, clause ignored}}
 #pragma acc parallel loop device_type(*) worker
   for(int i = 0; i < 5; ++i);
   // expected-error@+2{{OpenACC clause 'nohost' may not follow a 'device_type' clause in a 'serial loop' construct}}

clang/test/SemaOpenACC/combined-construct-gang-ast.cpp

@@ -1,4 +1,3 @@
-
 // RUN: %clang_cc1 %s -fopenacc -ast-dump | FileCheck %s
 
 // Test this with PCH.